OmniSciDB  5ade3759e0
CodeGeneratorTest.cpp
Go to the documentation of this file.
1 /*
2  * Copyright 2019 OmniSci, Inc.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
17 #include <gtest/gtest.h>
18 #include <llvm/Bitcode/BitcodeReader.h>
19 #include <llvm/IR/Function.h>
20 #include <llvm/IR/Module.h>
21 #include <llvm/IR/Verifier.h>
22 #include <llvm/Support/SourceMgr.h>
23 #include <llvm/Support/raw_os_ostream.h>
24 
25 #include "../Analyzer/Analyzer.h"
26 #include "../QueryEngine/CodeGenerator.h"
27 #include "../QueryEngine/Execute.h"
28 #include "../QueryEngine/IRCodegenUtils.h"
29 #include "../QueryEngine/LLVMGlobalContext.h"
30 #include "../Shared/mapdpath.h"
31 #include "TestHelpers.h"
32 
33 TEST(CodeGeneratorTest, IntegerConstant) {
34  auto& ctx = getGlobalLLVMContext();
35  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
36  ScalarCodeGenerator code_generator(std::move(module));
38 
39  Datum d;
40  d.intval = 42;
41  auto constant = makeExpr<Analyzer::Constant>(kINT, false, d);
42  const auto compiled_expr = code_generator.compile(constant.get(), true, co);
43  verify_function_ir(compiled_expr.func);
44  ASSERT_TRUE(compiled_expr.inputs.empty());
45 
46  using FuncPtr = int (*)(int*);
47  auto func_ptr = reinterpret_cast<FuncPtr>(
48  code_generator.generateNativeCode(compiled_expr, co).front());
49  CHECK(func_ptr);
50  int out;
51  int err = func_ptr(&out);
52  ASSERT_EQ(err, 0);
53  ASSERT_EQ(out, d.intval);
54 }
55 
56 TEST(CodeGeneratorTest, IntegerAdd) {
57  auto& ctx = getGlobalLLVMContext();
58  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
59  ScalarCodeGenerator code_generator(std::move(module));
61 
62  Datum d;
63  d.intval = 42;
64  auto lhs = makeExpr<Analyzer::Constant>(kINT, false, d);
65  auto rhs = makeExpr<Analyzer::Constant>(kINT, false, d);
66  auto plus = makeExpr<Analyzer::BinOper>(kINT, kPLUS, kONE, lhs, rhs);
67  const auto compiled_expr = code_generator.compile(plus.get(), true, co);
68  verify_function_ir(compiled_expr.func);
69  ASSERT_TRUE(compiled_expr.inputs.empty());
70 
71  using FuncPtr = int (*)(int*);
72  auto func_ptr = reinterpret_cast<FuncPtr>(
73  code_generator.generateNativeCode(compiled_expr, co).front());
74  CHECK(func_ptr);
75  int out;
76  int err = func_ptr(&out);
77  ASSERT_EQ(err, 0);
78  ASSERT_EQ(out, d.intval + d.intval);
79 }
80 
81 TEST(CodeGeneratorTest, IntegerColumn) {
82  auto& ctx = getGlobalLLVMContext();
83  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
84  ScalarCodeGenerator code_generator(std::move(module));
86 
87  SQLTypeInfo ti(kINT, false);
88  int table_id = 1;
89  int column_id = 5;
90  int rte_idx = 0;
91  auto col = makeExpr<Analyzer::ColumnVar>(ti, table_id, column_id, rte_idx);
92  const auto compiled_expr = code_generator.compile(col.get(), true, co);
93  verify_function_ir(compiled_expr.func);
94  ASSERT_EQ(compiled_expr.inputs.size(), size_t(1));
95  ASSERT_TRUE(*compiled_expr.inputs.front() == *col);
96 
97  using FuncPtr = int (*)(int*, int);
98  auto func_ptr = reinterpret_cast<FuncPtr>(
99  code_generator.generateNativeCode(compiled_expr, co).front());
100  CHECK(func_ptr);
101  int out;
102  int err = func_ptr(&out, 17);
103  ASSERT_EQ(err, 0);
104  ASSERT_EQ(out, 17);
105 }
106 
107 TEST(CodeGeneratorTest, IntegerExpr) {
108  auto& ctx = getGlobalLLVMContext();
109  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
110  ScalarCodeGenerator code_generator(std::move(module));
112 
113  SQLTypeInfo ti(kINT, false);
114  int table_id = 1;
115  int column_id = 5;
116  int rte_idx = 0;
117  auto lhs = makeExpr<Analyzer::ColumnVar>(ti, table_id, column_id, rte_idx);
118  Datum d;
119  d.intval = 42;
120  auto rhs = makeExpr<Analyzer::Constant>(kINT, false, d);
121  auto plus = makeExpr<Analyzer::BinOper>(kINT, kPLUS, kONE, lhs, rhs);
122  const auto compiled_expr = code_generator.compile(plus.get(), true, co);
123  verify_function_ir(compiled_expr.func);
124  ASSERT_EQ(compiled_expr.inputs.size(), size_t(1));
125  ASSERT_TRUE(*compiled_expr.inputs.front() == *lhs);
126 
127  using FuncPtr = int (*)(int*, int);
128  auto func_ptr = reinterpret_cast<FuncPtr>(
129  code_generator.generateNativeCode(compiled_expr, co).front());
130  CHECK(func_ptr);
131  int out;
132  int err = func_ptr(&out, 58);
133  ASSERT_EQ(err, 0);
134  ASSERT_EQ(out, 100);
135 }
136 
137 #ifdef HAVE_CUDA
138 void free_param_pointers(const std::vector<void*>& param_ptrs,
139  CudaMgr_Namespace::CudaMgr* cuda_mgr) {
140  for (const auto param_ptr : param_ptrs) {
141  const auto device_ptr =
142  reinterpret_cast<int8_t*>(*reinterpret_cast<CUdeviceptr*>(param_ptr));
143  cuda_mgr->freeDeviceMem(device_ptr);
144  }
145 }
146 
147 TEST(CodeGeneratorTest, IntegerConstantGPU) {
148  auto& ctx = getGlobalLLVMContext();
149  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
150  ScalarCodeGenerator code_generator(std::move(module));
152 
153  Datum d;
154  d.intval = 42;
155  auto constant = makeExpr<Analyzer::Constant>(kINT, false, d);
156  const auto compiled_expr = code_generator.compile(constant.get(), true, co);
157  verify_function_ir(compiled_expr.func);
158  ASSERT_TRUE(compiled_expr.inputs.empty());
159 
160  const auto native_function_pointers =
161  code_generator.generateNativeCode(compiled_expr, co);
162 
163  for (size_t gpu_idx = 0; gpu_idx < native_function_pointers.size(); ++gpu_idx) {
164  const auto native_function_pointer = native_function_pointers[gpu_idx];
165  auto func_ptr = reinterpret_cast<CUfunction>(native_function_pointer);
166 
167  std::vector<void*> param_ptrs;
168  CUdeviceptr err = reinterpret_cast<CUdeviceptr>(
169  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
170  CUdeviceptr out = reinterpret_cast<CUdeviceptr>(
171  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
172  param_ptrs.push_back(&err);
173  param_ptrs.push_back(&out);
174  cuLaunchKernel(func_ptr, 1, 1, 1, 1, 1, 1, 0, nullptr, &param_ptrs[0], nullptr);
175  int32_t host_err;
176  int32_t host_out;
177  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_err),
178  reinterpret_cast<const int8_t*>(err),
179  4,
180  gpu_idx);
181  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_out),
182  reinterpret_cast<const int8_t*>(out),
183  4,
184  gpu_idx);
185 
186  ASSERT_EQ(host_err, 0);
187  ASSERT_EQ(host_out, d.intval);
188  free_param_pointers(param_ptrs, code_generator.getCudaMgr());
189  }
190 }
191 
192 TEST(CodeGeneratorTest, IntegerAddGPU) {
193  auto& ctx = getGlobalLLVMContext();
194  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
195  ScalarCodeGenerator code_generator(std::move(module));
197 
198  Datum d;
199  d.intval = 42;
200  auto lhs = makeExpr<Analyzer::Constant>(kINT, false, d);
201  auto rhs = makeExpr<Analyzer::Constant>(kINT, false, d);
202  auto plus = makeExpr<Analyzer::BinOper>(kINT, kPLUS, kONE, lhs, rhs);
203  const auto compiled_expr = code_generator.compile(plus.get(), true, co);
204  verify_function_ir(compiled_expr.func);
205  ASSERT_TRUE(compiled_expr.inputs.empty());
206 
207  const auto native_function_pointers =
208  code_generator.generateNativeCode(compiled_expr, co);
209 
210  for (size_t gpu_idx = 0; gpu_idx < native_function_pointers.size(); ++gpu_idx) {
211  const auto native_function_pointer = native_function_pointers[gpu_idx];
212  auto func_ptr = reinterpret_cast<CUfunction>(native_function_pointer);
213 
214  std::vector<void*> param_ptrs;
215  CUdeviceptr err = reinterpret_cast<CUdeviceptr>(
216  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
217  CUdeviceptr out = reinterpret_cast<CUdeviceptr>(
218  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
219  param_ptrs.push_back(&err);
220  param_ptrs.push_back(&out);
221  cuLaunchKernel(func_ptr, 1, 1, 1, 1, 1, 1, 0, nullptr, &param_ptrs[0], nullptr);
222  int32_t host_err;
223  int32_t host_out;
224  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_err),
225  reinterpret_cast<const int8_t*>(err),
226  4,
227  gpu_idx);
228  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_out),
229  reinterpret_cast<const int8_t*>(out),
230  4,
231  gpu_idx);
232 
233  ASSERT_EQ(host_err, 0);
234  ASSERT_EQ(host_out, d.intval + d.intval);
235  free_param_pointers(param_ptrs, code_generator.getCudaMgr());
236  }
237 }
238 
239 TEST(CodeGeneratorTest, IntegerColumnGPU) {
240  auto& ctx = getGlobalLLVMContext();
241  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
242  ScalarCodeGenerator code_generator(std::move(module));
244 
245  SQLTypeInfo ti(kINT, false);
246  int table_id = 1;
247  int column_id = 5;
248  int rte_idx = 0;
249  auto col = makeExpr<Analyzer::ColumnVar>(ti, table_id, column_id, rte_idx);
250  const auto compiled_expr = code_generator.compile(col.get(), true, co);
251  verify_function_ir(compiled_expr.func);
252  ASSERT_EQ(compiled_expr.inputs.size(), size_t(1));
253  ASSERT_TRUE(*compiled_expr.inputs.front() == *col);
254 
255  const auto native_function_pointers =
256  code_generator.generateNativeCode(compiled_expr, co);
257 
258  for (size_t gpu_idx = 0; gpu_idx < native_function_pointers.size(); ++gpu_idx) {
259  const auto native_function_pointer = native_function_pointers[gpu_idx];
260  auto func_ptr = reinterpret_cast<CUfunction>(native_function_pointer);
261 
262  std::vector<void*> param_ptrs;
263  CUdeviceptr err = reinterpret_cast<CUdeviceptr>(
264  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
265  CUdeviceptr out = reinterpret_cast<CUdeviceptr>(
266  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
267  CUdeviceptr in = reinterpret_cast<CUdeviceptr>(
268  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
269  int host_in = 17;
270  code_generator.getCudaMgr()->copyHostToDevice(
271  reinterpret_cast<int8_t*>(in),
272  reinterpret_cast<const int8_t*>(&host_in),
273  4,
274  gpu_idx);
275  param_ptrs.push_back(&err);
276  param_ptrs.push_back(&out);
277  param_ptrs.push_back(&in);
278  cuLaunchKernel(func_ptr, 1, 1, 1, 1, 1, 1, 0, nullptr, &param_ptrs[0], nullptr);
279  int32_t host_err;
280  int32_t host_out;
281  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_err),
282  reinterpret_cast<const int8_t*>(err),
283  4,
284  gpu_idx);
285  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_out),
286  reinterpret_cast<const int8_t*>(out),
287  4,
288  gpu_idx);
289 
290  ASSERT_EQ(host_err, 0);
291  ASSERT_EQ(host_out, 17);
292  free_param_pointers(param_ptrs, code_generator.getCudaMgr());
293  }
294 }
295 
296 TEST(CodeGeneratorTest, IntegerExprGPU) {
297  auto& ctx = getGlobalLLVMContext();
298  std::unique_ptr<llvm::Module> module(read_template_module(ctx));
299  ScalarCodeGenerator code_generator(std::move(module));
301 
302  SQLTypeInfo ti(kINT, false);
303  int table_id = 1;
304  int column_id = 5;
305  int rte_idx = 0;
306  auto lhs = makeExpr<Analyzer::ColumnVar>(ti, table_id, column_id, rte_idx);
307  Datum d;
308  d.intval = 42;
309  auto rhs = makeExpr<Analyzer::Constant>(kINT, false, d);
310  auto plus = makeExpr<Analyzer::BinOper>(kINT, kPLUS, kONE, lhs, rhs);
311  const auto compiled_expr = code_generator.compile(plus.get(), true, co);
312  verify_function_ir(compiled_expr.func);
313  ASSERT_EQ(compiled_expr.inputs.size(), size_t(1));
314  ASSERT_TRUE(*compiled_expr.inputs.front() == *lhs);
315 
316  const auto native_function_pointers =
317  code_generator.generateNativeCode(compiled_expr, co);
318 
319  for (size_t gpu_idx = 0; gpu_idx < native_function_pointers.size(); ++gpu_idx) {
320  const auto native_function_pointer = native_function_pointers[gpu_idx];
321  auto func_ptr = reinterpret_cast<CUfunction>(native_function_pointer);
322 
323  std::vector<void*> param_ptrs;
324  CUdeviceptr err = reinterpret_cast<CUdeviceptr>(
325  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
326  CUdeviceptr out = reinterpret_cast<CUdeviceptr>(
327  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
328  CUdeviceptr in = reinterpret_cast<CUdeviceptr>(
329  code_generator.getCudaMgr()->allocateDeviceMem(4, gpu_idx));
330  int host_in = 58;
331  code_generator.getCudaMgr()->copyHostToDevice(
332  reinterpret_cast<int8_t*>(in),
333  reinterpret_cast<const int8_t*>(&host_in),
334  4,
335  gpu_idx);
336  param_ptrs.push_back(&err);
337  param_ptrs.push_back(&out);
338  param_ptrs.push_back(&in);
339  cuLaunchKernel(func_ptr, 1, 1, 1, 1, 1, 1, 0, nullptr, &param_ptrs[0], nullptr);
340  int32_t host_err;
341  int32_t host_out;
342  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_err),
343  reinterpret_cast<const int8_t*>(err),
344  4,
345  gpu_idx);
346  code_generator.getCudaMgr()->copyDeviceToHost(reinterpret_cast<int8_t*>(&host_out),
347  reinterpret_cast<const int8_t*>(out),
348  4,
349  gpu_idx);
350 
351  ASSERT_EQ(host_err, 0);
352  ASSERT_EQ(host_out, 100);
353  free_param_pointers(param_ptrs, code_generator.getCudaMgr());
354  }
355 }
356 #endif // HAVE_CUDA
357 
358 int main(int argc, char** argv) {
360  testing::InitGoogleTest(&argc, argv);
361  int err = RUN_ALL_TESTS();
362  return err;
363 }
CompiledExpression compile(const Analyzer::Expr *expr, const bool fetch_columns, const CompilationOptions &co)
TEST(CodeGeneratorTest, IntegerConstant)
void d(const SQLTypes expected_type, const std::string &str)
Definition: ImportTest.cpp:268
std::vector< void * > generateNativeCode(const CompiledExpression &compiled_expression, const CompilationOptions &co)
void copyDeviceToHost(int8_t *host_ptr, const int8_t *device_ptr, const size_t num_bytes, const int device_num)
Definition: CudaMgr.cpp:80
CudaMgr_Namespace::CudaMgr * getCudaMgr() const
unsigned long long CUdeviceptr
Definition: nocuda.h:27
Definition: sqldefs.h:41
int32_t intval
Definition: sqltypes.h:125
void freeDeviceMem(int8_t *device_ptr)
Definition: CudaMgr.cpp:215
void verify_function_ir(const llvm::Function *func)
int main(int argc, char **argv)
void * CUfunction
Definition: nocuda.h:24
void copyHostToDevice(int8_t *device_ptr, const int8_t *host_ptr, const size_t num_bytes, const int device_num)
Definition: CudaMgr.cpp:71
Definition: sqldefs.h:69
llvm::Module * read_template_module(llvm::LLVMContext &context)
llvm::LLVMContext & getGlobalLLVMContext()
#define CHECK(condition)
Definition: Logger.h:187
void init_logger_stderr_only(int argc, char const *const *argv)
Definition: TestHelpers.h:194
Definition: sqltypes.h:47
const int8_t const int64_t const uint64_t const int32_t const int64_t int64_t ** out
int8_t * allocateDeviceMem(const size_t num_bytes, const int device_num)
Definition: CudaMgr.cpp:204