OmniSciDB  c0231cc57d
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
NativeCodegen.cpp
Go to the documentation of this file.
1 /*
2  * Copyright 2022 HEAVY.AI, 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 "QueryEngine/Execute.h"
18 
19 #if LLVM_VERSION_MAJOR < 9
20 static_assert(false, "LLVM Version >= 9 is required.");
21 #endif
22 
23 #include <llvm/Analysis/ScopedNoAliasAA.h>
24 #include <llvm/Analysis/TypeBasedAliasAnalysis.h>
25 #include <llvm/Bitcode/BitcodeReader.h>
26 #include <llvm/Bitcode/BitcodeWriter.h>
27 #include <llvm/ExecutionEngine/MCJIT.h>
28 #include <llvm/IR/Attributes.h>
29 #include <llvm/IR/GlobalValue.h>
30 #include <llvm/IR/InstIterator.h>
31 #include <llvm/IR/IntrinsicInst.h>
32 #include <llvm/IR/Intrinsics.h>
33 #include <llvm/IR/LegacyPassManager.h>
34 #include <llvm/IR/Verifier.h>
35 #include <llvm/IRReader/IRReader.h>
36 #if 14 <= LLVM_VERSION_MAJOR
37 #include <llvm/MC/TargetRegistry.h>
38 #else
39 #include <llvm/Support/TargetRegistry.h>
40 #endif
41 #include <llvm/Support/Casting.h>
42 #include <llvm/Support/FileSystem.h>
43 #include <llvm/Support/FormattedStream.h>
44 #include <llvm/Support/MemoryBuffer.h>
45 #include <llvm/Support/SourceMgr.h>
46 #include <llvm/Support/TargetSelect.h>
47 #include <llvm/Support/raw_os_ostream.h>
48 #include <llvm/Support/raw_ostream.h>
49 #include <llvm/Transforms/IPO.h>
50 #include <llvm/Transforms/IPO/AlwaysInliner.h>
51 #include <llvm/Transforms/IPO/InferFunctionAttrs.h>
52 #include <llvm/Transforms/IPO/PassManagerBuilder.h>
53 #include <llvm/Transforms/InstCombine/InstCombine.h>
54 #include <llvm/Transforms/Instrumentation.h>
55 #include <llvm/Transforms/Scalar.h>
56 #include <llvm/Transforms/Scalar/GVN.h>
57 #include <llvm/Transforms/Scalar/InstSimplifyPass.h>
58 #include <llvm/Transforms/Utils.h>
59 #include <llvm/Transforms/Utils/BasicBlockUtils.h>
60 #include <llvm/Transforms/Utils/Cloning.h>
61 
62 #if LLVM_VERSION_MAJOR >= 11
63 #include <llvm/Support/Host.h>
64 #endif
65 
66 #include "CudaMgr/CudaMgr.h"
76 #include "Shared/MathUtils.h"
77 #include "StreamingTopN.h"
78 
80 
81 static llvm::sys::Mutex g_ee_create_mutex;
82 
83 #ifdef ENABLE_GEOS
84 
85 #include <llvm/Support/DynamicLibrary.h>
86 
87 #ifndef GEOS_LIBRARY_FILENAME
88 #error Configuration should include GEOS library file name
89 #endif
90 std::unique_ptr<std::string> g_libgeos_so_filename(
91  new std::string(GEOS_LIBRARY_FILENAME));
92 static llvm::sys::DynamicLibrary geos_dynamic_library;
93 static std::mutex geos_init_mutex;
94 
95 namespace {
96 
97 void load_geos_dynamic_library() {
98  std::lock_guard<std::mutex> guard(geos_init_mutex);
99 
100  if (!geos_dynamic_library.isValid()) {
101  if (!g_libgeos_so_filename || g_libgeos_so_filename->empty()) {
102  LOG(WARNING) << "Misconfigured GEOS library file name, trying 'libgeos_c.so'";
103  g_libgeos_so_filename.reset(new std::string("libgeos_c.so"));
104  }
105  auto filename = *g_libgeos_so_filename;
106  std::string error_message;
107  geos_dynamic_library =
108  llvm::sys::DynamicLibrary::getPermanentLibrary(filename.c_str(), &error_message);
109  if (!geos_dynamic_library.isValid()) {
110  LOG(ERROR) << "Failed to load GEOS library '" + filename + "'";
111  std::string exception_message = "Failed to load GEOS library: " + error_message;
112  throw std::runtime_error(exception_message.c_str());
113  } else {
114  LOG(INFO) << "Loaded GEOS library '" + filename + "'";
115  }
116  }
117 }
118 
119 } // namespace
120 #endif
121 
122 namespace {
123 
124 void throw_parseIR_error(const llvm::SMDiagnostic& parse_error,
125  std::string src = "",
126  const bool is_gpu = false) {
127  std::string excname = (is_gpu ? "NVVM IR ParseError: " : "LLVM IR ParseError: ");
128  llvm::raw_string_ostream ss(excname);
129  parse_error.print(src.c_str(), ss, false, false);
130  throw ParseIRError(ss.str());
131 }
132 
133 /* SHOW_DEFINED(<llvm::Module instance>) prints the function names
134  that are defined in the given LLVM Module instance.
135 
136  SHOW_FUNCTIONS(<llvm::Module instance>) prints the function names
137  of all used functions in the given LLVM Module
138  instance. Declarations are marked with `[decl]` as a name suffix.
139 
140  Useful for debugging.
141 */
142 
143 #define SHOW_DEFINED(MODULE) \
144  { \
145  std::cout << __func__ << "#" << __LINE__ << ": " #MODULE << " "; \
146  ::show_defined(MODULE); \
147  }
148 
149 #define SHOW_FUNCTIONS(MODULE) \
150  { \
151  std::cout << __func__ << "#" << __LINE__ << ": " #MODULE << " "; \
152  ::show_functions(MODULE); \
153  }
154 
155 template <typename T = void>
156 void show_defined(llvm::Module& llvm_module) {
157  std::cout << "defines: ";
158  for (auto& f : llvm_module.getFunctionList()) {
159  if (!f.isDeclaration()) {
160  std::cout << f.getName().str() << ", ";
161  }
162  }
163  std::cout << std::endl;
164 }
165 
166 template <typename T = void>
167 void show_defined(llvm::Module* llvm_module) {
168  if (llvm_module == nullptr) {
169  std::cout << "is null" << std::endl;
170  } else {
171  show_defined(*llvm_module);
172  }
173 }
174 
175 template <typename T = void>
176 void show_defined(std::unique_ptr<llvm::Module>& llvm_module) {
177  show_defined(llvm_module.get());
178 }
179 
180 /*
181  scan_function_calls(module, defined, undefined, ignored) computes
182  defined and undefined sets of function names:
183 
184  - defined functions are those that are defined in the given module
185 
186  - undefined functions are those that are called by defined functions
187  but that are not defined in the given module
188 
189  - ignored functions are functions that may be undefined but will not
190  be listed in the set of undefined functions.
191 
192  Useful for debugging.
193 */
194 template <typename T = void>
195 void scan_function_calls(llvm::Function& F,
196  std::unordered_set<std::string>& defined,
197  std::unordered_set<std::string>& undefined,
198  const std::unordered_set<std::string>& ignored) {
199  for (llvm::inst_iterator I = llvm::inst_begin(F), E = llvm::inst_end(F); I != E; ++I) {
200  if (auto* CI = llvm::dyn_cast<llvm::CallInst>(&*I)) {
201  auto* F2 = CI->getCalledFunction();
202  if (F2 != nullptr) {
203  auto F2name = F2->getName().str();
204  if (F2->isDeclaration()) {
205  if (F2name.rfind("__", 0) !=
206  0 // assume symbols with double underscore are defined
207  && F2name.rfind("llvm.", 0) !=
208  0 // TODO: this may give false positive for NVVM intrinsics
209  && ignored.find(F2name) == ignored.end() // not in ignored list
210  ) {
211  undefined.emplace(F2name);
212  }
213  } else {
214  if (defined.find(F2name) == defined.end()) {
215  defined.emplace(F2name);
216  scan_function_calls<T>(*F2, defined, undefined, ignored);
217  }
218  }
219  }
220  }
221  }
222 }
223 
224 template <typename T = void>
225 void scan_function_calls(llvm::Module& llvm_module,
226  std::unordered_set<std::string>& defined,
227  std::unordered_set<std::string>& undefined,
228  const std::unordered_set<std::string>& ignored) {
229  for (auto& F : llvm_module) {
230  if (!F.isDeclaration()) {
231  scan_function_calls(F, defined, undefined, ignored);
232  }
233  }
234 }
235 
236 template <typename T = void>
237 std::tuple<std::unordered_set<std::string>, std::unordered_set<std::string>>
238 scan_function_calls(llvm::Module& llvm_module,
239  const std::unordered_set<std::string>& ignored = {}) {
240  std::unordered_set<std::string> defined, undefined;
241  scan_function_calls(llvm_module, defined, undefined, ignored);
242  return std::make_tuple(defined, undefined);
243 }
244 
245 #if defined(HAVE_CUDA) || !defined(WITH_JIT_DEBUG)
247  llvm::Module& M,
248  const std::unordered_set<llvm::Function*>& live_funcs) {
249  std::vector<llvm::Function*> dead_funcs;
250  for (auto& F : M) {
251  bool bAlive = false;
252  if (live_funcs.count(&F)) {
253  continue;
254  }
255  for (auto U : F.users()) {
256  auto* C = llvm::dyn_cast<const llvm::CallInst>(U);
257  if (!C || C->getParent()->getParent() != &F) {
258  bAlive = true;
259  break;
260  }
261  }
262  if (!bAlive) {
263  dead_funcs.push_back(&F);
264  }
265  }
266  for (auto pFn : dead_funcs) {
267  pFn->eraseFromParent();
268  }
269 }
270 
271 #ifdef HAVE_CUDA
272 
273 // check if linking with libdevice is required
274 // libdevice functions have a __nv_* prefix
275 bool check_module_requires_libdevice(llvm::Module* llvm_module) {
276  auto timer = DEBUG_TIMER(__func__);
277  for (llvm::Function& F : *llvm_module) {
278  if (F.hasName() && F.getName().startswith("__nv_")) {
279  LOG(INFO) << "Module requires linking with libdevice: " << std::string(F.getName());
280  return true;
281  }
282  }
283  LOG(DEBUG1) << "module does not require linking against libdevice";
284  return false;
285 }
286 
287 // Adds the missing intrinsics declarations to the given module
288 void add_intrinsics_to_module(llvm::Module* llvm_module) {
289  for (llvm::Function& F : *llvm_module) {
290  for (llvm::Instruction& I : instructions(F)) {
291  if (llvm::IntrinsicInst* ii = llvm::dyn_cast<llvm::IntrinsicInst>(&I)) {
292  if (llvm::Intrinsic::isOverloaded(ii->getIntrinsicID())) {
293  llvm::Type* Tys[] = {ii->getFunctionType()->getReturnType()};
294  llvm::Function& decl_fn =
295  *llvm::Intrinsic::getDeclaration(llvm_module, ii->getIntrinsicID(), Tys);
296  ii->setCalledFunction(&decl_fn);
297  } else {
298  // inserts the declaration into the module if not present
299  llvm::Intrinsic::getDeclaration(llvm_module, ii->getIntrinsicID());
300  }
301  }
302  }
303  }
304 }
305 
306 #endif
307 
308 void optimize_ir(llvm::Function* query_func,
309  llvm::Module* llvm_module,
310  llvm::legacy::PassManager& pass_manager,
311  const std::unordered_set<llvm::Function*>& live_funcs,
312  const bool is_gpu_smem_used,
313  const CompilationOptions& co) {
314  auto timer = DEBUG_TIMER(__func__);
315  // the always inliner legacy pass must always run first
316  pass_manager.add(llvm::createVerifierPass());
317  pass_manager.add(llvm::createAlwaysInlinerLegacyPass());
318 
319  pass_manager.add(new AnnotateInternalFunctionsPass());
320 
321  pass_manager.add(llvm::createSROAPass());
322  // mem ssa drops unused load and store instructions, e.g. passing variables directly
323  // where possible
324  pass_manager.add(
325  llvm::createEarlyCSEPass(/*enable_mem_ssa=*/true)); // Catch trivial redundancies
326 
327  if (!is_gpu_smem_used) {
328  // thread jumps can change the execution order around SMEM sections guarded by
329  // `__syncthreads()`, which results in race conditions. For now, disable jump
330  // threading for shared memory queries. In the future, consider handling shared memory
331  // aggregations with a separate kernel launch
332  pass_manager.add(llvm::createJumpThreadingPass()); // Thread jumps.
333  }
334  pass_manager.add(llvm::createCFGSimplificationPass());
335 
336  // remove load/stores in PHIs if instructions can be accessed directly post thread jumps
337  pass_manager.add(llvm::createNewGVNPass());
338 
339  pass_manager.add(llvm::createDeadStoreEliminationPass());
340  pass_manager.add(llvm::createLICMPass());
341 
342  pass_manager.add(llvm::createInstructionCombiningPass());
343 
344  // module passes
345  pass_manager.add(llvm::createPromoteMemoryToRegisterPass());
346  pass_manager.add(llvm::createGlobalOptimizerPass());
347 
348  pass_manager.add(llvm::createCFGSimplificationPass()); // cleanup after everything
349 
350  pass_manager.run(*llvm_module);
351 
352  eliminate_dead_self_recursive_funcs(*llvm_module, live_funcs);
353 }
354 #endif
355 
356 } // namespace
357 
359 
360 ExecutionEngineWrapper::ExecutionEngineWrapper(llvm::ExecutionEngine* execution_engine)
361  : execution_engine_(execution_engine) {}
362 
363 ExecutionEngineWrapper::ExecutionEngineWrapper(llvm::ExecutionEngine* execution_engine,
364  const CompilationOptions& co)
365  : execution_engine_(execution_engine) {
366  if (execution_engine_) {
368 #ifdef ENABLE_INTEL_JIT_LISTENER
369  intel_jit_listener_.reset(llvm::JITEventListener::createIntelJITEventListener());
371  execution_engine_->RegisterJITEventListener(intel_jit_listener_.get());
372  LOG(INFO) << "Registered IntelJITEventListener";
373 #else
374  LOG(WARNING) << "This build is not Intel JIT Listener enabled. Ignoring Intel JIT "
375  "listener configuration parameter.";
376 #endif // ENABLE_INTEL_JIT_LISTENER
377  }
378  }
379 }
380 
382  llvm::ExecutionEngine* execution_engine) {
383  execution_engine_.reset(execution_engine);
384  intel_jit_listener_ = nullptr;
385  return *this;
386 }
387 
388 void verify_function_ir(const llvm::Function* func) {
389  std::stringstream err_ss;
390  llvm::raw_os_ostream err_os(err_ss);
391  err_os << "\n-----\n";
392  if (llvm::verifyFunction(*func, &err_os)) {
393  err_os << "\n-----\n";
394  func->print(err_os, nullptr);
395  err_os << "\n-----\n";
396  LOG(FATAL) << err_ss.str();
397  }
398 }
399 
400 namespace {
401 
402 std::string assemblyForCPU(ExecutionEngineWrapper& execution_engine,
403  llvm::Module* llvm_module) {
404  llvm::legacy::PassManager pass_manager;
405  auto cpu_target_machine = execution_engine->getTargetMachine();
406  CHECK(cpu_target_machine);
407  llvm::SmallString<256> code_str;
408  llvm::raw_svector_ostream os(code_str);
409 #if LLVM_VERSION_MAJOR >= 10
410  cpu_target_machine->addPassesToEmitFile(
411  pass_manager, os, nullptr, llvm::CGFT_AssemblyFile);
412 #else
413  cpu_target_machine->addPassesToEmitFile(
414  pass_manager, os, nullptr, llvm::TargetMachine::CGFT_AssemblyFile);
415 #endif
416  pass_manager.run(*llvm_module);
417  return "Assembly for the CPU:\n" + std::string(code_str.str()) + "\nEnd of assembly";
418 }
419 
421  llvm::EngineBuilder& eb,
422  const CompilationOptions& co) {
423  auto timer = DEBUG_TIMER(__func__);
424  // Avoids data race in
425  // llvm::sys::DynamicLibrary::getPermanentLibrary and
426  // GDBJITRegistrationListener::notifyObjectLoaded while creating a
427  // new ExecutionEngine instance. Unfortunately we have to use global
428  // mutex here.
429  std::lock_guard<llvm::sys::Mutex> lock(g_ee_create_mutex);
430  ExecutionEngineWrapper execution_engine(eb.create(), co);
431  CHECK(execution_engine.get());
432  // Force the module data layout to match the layout for the selected target
433  llvm_module->setDataLayout(execution_engine->getDataLayout());
434 
435  LOG(ASM) << assemblyForCPU(execution_engine, llvm_module);
436 
437  execution_engine->finalizeObject();
438  return execution_engine;
439 }
440 
441 } // namespace
442 
444  llvm::Function* func,
445  const std::unordered_set<llvm::Function*>& live_funcs,
446  const CompilationOptions& co) {
447  auto timer = DEBUG_TIMER(__func__);
448  llvm::Module* llvm_module = func->getParent();
449  // run optimizations
450 #ifndef WITH_JIT_DEBUG
451  llvm::legacy::PassManager pass_manager;
452  optimize_ir(
453  func, llvm_module, pass_manager, live_funcs, /*is_gpu_smem_used=*/false, co);
454 #endif // WITH_JIT_DEBUG
455 
456  auto init_err = llvm::InitializeNativeTarget();
457  CHECK(!init_err);
458 
459  llvm::InitializeAllTargetMCs();
460  llvm::InitializeNativeTargetAsmPrinter();
461  llvm::InitializeNativeTargetAsmParser();
462 
463  std::string err_str;
464  std::unique_ptr<llvm::Module> owner(llvm_module);
465  llvm::EngineBuilder eb(std::move(owner));
466  eb.setErrorStr(&err_str);
467  eb.setEngineKind(llvm::EngineKind::JIT);
468  llvm::TargetOptions to;
469  to.EnableFastISel = true;
470  eb.setTargetOptions(to);
472  eb.setOptLevel(llvm::CodeGenOpt::None);
473  }
474 
475  return create_execution_engine(llvm_module, eb, co);
476 }
477 
478 std::shared_ptr<CompilationContext> Executor::optimizeAndCodegenCPU(
479  llvm::Function* query_func,
480  llvm::Function* multifrag_query_func,
481  const std::unordered_set<llvm::Function*>& live_funcs,
482  const CompilationOptions& co) {
483  CodeCacheKey key{serialize_llvm_object(query_func),
484  serialize_llvm_object(cgen_state_->row_func_)};
485 
486  llvm::Module* M = query_func->getParent();
487  auto* flag = llvm::mdconst::extract_or_null<llvm::ConstantInt>(
488  M->getModuleFlag("manage_memory_buffer"));
489  if (flag and flag->getZExtValue() == 1 and M->getFunction("allocate_varlen_buffer") and
490  M->getFunction("register_buffer_with_executor_rsm")) {
491  LOG(INFO) << "including executor addr to cache key\n";
492  key.push_back(std::to_string(reinterpret_cast<int64_t>(this)));
493  }
494  if (cgen_state_->filter_func_) {
495  key.push_back(serialize_llvm_object(cgen_state_->filter_func_));
496  }
497  for (const auto helper : cgen_state_->helper_functions_) {
498  key.push_back(serialize_llvm_object(helper));
499  }
500  auto cached_code = QueryEngine::getInstance()->cpu_code_accessor->get_value(key);
501  if (cached_code) {
502  return cached_code;
503  }
504 
505  if (cgen_state_->needs_geos_) {
506 #ifdef ENABLE_GEOS
507  auto llvm_module = multifrag_query_func->getParent();
508  load_geos_dynamic_library();
509 
510  // Read geos runtime module and bind GEOS API function references to GEOS library
511  auto rt_geos_module_copy = llvm::CloneModule(
512  *get_geos_module(), cgen_state_->vmap_, [](const llvm::GlobalValue* gv) {
513  auto func = llvm::dyn_cast<llvm::Function>(gv);
514  if (!func) {
515  return true;
516  }
517  return (func->getLinkage() == llvm::GlobalValue::LinkageTypes::PrivateLinkage ||
518  func->getLinkage() ==
519  llvm::GlobalValue::LinkageTypes::InternalLinkage ||
520  func->getLinkage() == llvm::GlobalValue::LinkageTypes::ExternalLinkage);
521  });
522  CodeGenerator::link_udf_module(rt_geos_module_copy,
523  *llvm_module,
524  cgen_state_.get(),
525  llvm::Linker::Flags::LinkOnlyNeeded);
526 #else
527  throw std::runtime_error("GEOS is disabled in this build");
528 #endif
529  }
530 
531  auto execution_engine =
532  CodeGenerator::generateNativeCPUCode(query_func, live_funcs, co);
533  auto cpu_compilation_context =
534  std::make_shared<CpuCompilationContext>(std::move(execution_engine));
535  cpu_compilation_context->setFunctionPointer(multifrag_query_func);
536  QueryEngine::getInstance()->cpu_code_accessor->put(key, cpu_compilation_context);
537  return std::dynamic_pointer_cast<CompilationContext>(cpu_compilation_context);
538 }
539 
540 void CodeGenerator::link_udf_module(const std::unique_ptr<llvm::Module>& udf_module,
541  llvm::Module& llvm_module,
542  CgenState* cgen_state,
543  llvm::Linker::Flags flags) {
544  auto timer = DEBUG_TIMER(__func__);
545  // throw a runtime error if the target module contains functions
546  // with the same name as in module of UDF functions.
547  for (auto& f : *udf_module) {
548  auto func = llvm_module.getFunction(f.getName());
549  if (!(func == nullptr) && !f.isDeclaration() && flags == llvm::Linker::Flags::None) {
550  LOG(ERROR) << " Attempt to overwrite " << f.getName().str() << " in "
551  << llvm_module.getModuleIdentifier() << " from `"
552  << udf_module->getModuleIdentifier() << "`" << std::endl;
553  throw std::runtime_error(
554  "link_udf_module: *** attempt to overwrite a runtime function with a UDF "
555  "function ***");
556  } else {
557  VLOG(1) << " Adding " << f.getName().str() << " to "
558  << llvm_module.getModuleIdentifier() << " from `"
559  << udf_module->getModuleIdentifier() << "`" << std::endl;
560  }
561  }
562 
563  auto udf_module_copy = llvm::CloneModule(*udf_module, cgen_state->vmap_);
564 
565  udf_module_copy->setDataLayout(llvm_module.getDataLayout());
566  udf_module_copy->setTargetTriple(llvm_module.getTargetTriple());
567 
568  // Initialize linker with module for RuntimeFunctions.bc
569  llvm::Linker ld(llvm_module);
570  bool link_error = false;
571 
572  link_error = ld.linkInModule(std::move(udf_module_copy), flags);
573 
574  if (link_error) {
575  throw std::runtime_error("link_udf_module: *** error linking module ***");
576  }
577 }
578 
579 namespace {
580 
581 std::string cpp_to_llvm_name(const std::string& s) {
582  if (s == "int8_t") {
583  return "i8";
584  }
585  if (s == "int16_t") {
586  return "i16";
587  }
588  if (s == "int32_t") {
589  return "i32";
590  }
591  if (s == "int64_t") {
592  return "i64";
593  }
594  CHECK(s == "float" || s == "double");
595  return s;
596 }
597 
598 std::string gen_array_any_all_sigs() {
599  std::string result;
600  for (const std::string any_or_all : {"any", "all"}) {
601  for (const std::string elem_type :
602  {"int8_t", "int16_t", "int32_t", "int64_t", "float", "double"}) {
603  for (const std::string needle_type :
604  {"int8_t", "int16_t", "int32_t", "int64_t", "float", "double"}) {
605  for (const std::string op_name : {"eq", "ne", "lt", "le", "gt", "ge"}) {
606  result += ("declare i1 @array_" + any_or_all + "_" + op_name + "_" + elem_type +
607  "_" + needle_type + "(i8*, i64, " + cpp_to_llvm_name(needle_type) +
608  ", " + cpp_to_llvm_name(elem_type) + ");\n");
609  }
610  }
611  }
612  }
613  return result;
614 }
615 
617  std::string result;
618  for (const std::string key_type : {"int8_t", "int16_t", "int32_t", "int64_t"}) {
619  const auto key_llvm_type = cpp_to_llvm_name(key_type);
620  result += "declare i64 @translate_null_key_" + key_type + "(" + key_llvm_type + ", " +
621  key_llvm_type + ", i64);\n";
622  }
623  return result;
624 }
625 
626 const std::string cuda_rt_decls =
627  R"( declare void @llvm.dbg.declare(metadata, metadata, metadata) declare void @llvm.dbg.value(metadata, metadata, metadata) declare double @llvm.fmuladd.f64(double, double, double) declare void @llvm.lifetime.start(i64, i8* nocapture) nounwind declare void @llvm.lifetime.end(i64, i8* nocapture) nounwind declare void @llvm.lifetime.start.p0i8(i64, i8* nocapture) nounwind declare void @llvm.lifetime.end.p0i8(i64, i8* nocapture) nounwind declare i64 @get_thread_index(); declare i64 @get_block_index(); declare i32 @pos_start_impl(i32*); declare i32 @group_buff_idx_impl(); declare i32 @pos_step_impl(); declare i8 @thread_warp_idx(i8); declare i64* @init_shared_mem(i64*, i32); declare i64* @init_shared_mem_nop(i64*, i32); declare i64* @declare_dynamic_shared_memory(); declare void @write_back_nop(i64*, i64*, i32); declare void @write_back_non_grouped_agg(i64*, i64*, i32); declare void @init_group_by_buffer_gpu(i64*, i64*, i32, i32, i32, i1, i8); declare i64* @get_group_value(i64*, i32, i64*, i32, i32, i32); declare i64* @get_group_value_with_watchdog(i64*, i32, i64*, i32, i32, i32); declare i32 @get_group_value_columnar_slot(i64*, i32, i64*, i32, i32); declare i32 @get_group_value_columnar_slot_with_watchdog(i64*, i32, i64*, i32, i32); declare i64* @get_group_value_fast(i64*, i64, i64, i64, i32); declare i64* @get_group_value_fast_with_original_key(i64*, i64, i64, i64, i64, i32); declare i32 @get_columnar_group_bin_offset(i64*, i64, i64, i64); declare i64 @baseline_hash_join_idx_32(i8*, i8*, i64, i64); declare i64 @baseline_hash_join_idx_64(i8*, i8*, i64, i64); declare i64 @get_composite_key_index_32(i32*, i64, i32*, i64); declare i64 @get_composite_key_index_64(i64*, i64, i64*, i64); declare i64 @get_bucket_key_for_range_compressed(i8*, i64, double); declare i64 @get_bucket_key_for_range_double(i8*, i64, double); declare i32 @get_num_buckets_for_bounds(i8*, i32, double, double); declare i64 @get_candidate_rows(i32*, i32, i8*, i32, double, double, i32, i64, i64*, i64, i64, i64); declare i64 @agg_count_shared(i64*, i64); declare i64 @agg_count_skip_val_shared(i64*, i64, i64); declare i32 @agg_count_int32_shared(i32*, i32); declare i32 @agg_count_int32_skip_val_shared(i32*, i32, i32); declare i64 @agg_count_double_shared(i64*, double); declare i64 @agg_count_double_skip_val_shared(i64*, double, double); declare i32 @agg_count_float_shared(i32*, float); declare i32 @agg_count_float_skip_val_shared(i32*, float, float); declare i64 @agg_sum_shared(i64*, i64); declare i64 @agg_sum_skip_val_shared(i64*, i64, i64); declare i32 @agg_sum_int32_shared(i32*, i32); declare i32 @agg_sum_int32_skip_val_shared(i32*, i32, i32); declare void @agg_sum_double_shared(i64*, double); declare void @agg_sum_double_skip_val_shared(i64*, double, double); declare void @agg_sum_float_shared(i32*, float); declare void @agg_sum_float_skip_val_shared(i32*, float, float); declare void @agg_max_shared(i64*, i64); declare void @agg_max_skip_val_shared(i64*, i64, i64); declare void @agg_max_int32_shared(i32*, i32); declare void @agg_max_int32_skip_val_shared(i32*, i32, i32); declare void @agg_max_int16_shared(i16*, i16); declare void @agg_max_int16_skip_val_shared(i16*, i16, i16); declare void @agg_max_int8_shared(i8*, i8); declare void @agg_max_int8_skip_val_shared(i8*, i8, i8); declare void @agg_max_double_shared(i64*, double); declare void @agg_max_double_skip_val_shared(i64*, double, double); declare void @agg_max_float_shared(i32*, float); declare void @agg_max_float_skip_val_shared(i32*, float, float); declare void @agg_min_shared(i64*, i64); declare void @agg_min_skip_val_shared(i64*, i64, i64); declare void @agg_min_int32_shared(i32*, i32); declare void @agg_min_int32_skip_val_shared(i32*, i32, i32); declare void @agg_min_int16_shared(i16*, i16); declare void @agg_min_int16_skip_val_shared(i16*, i16, i16); declare void @agg_min_int8_shared(i8*, i8); declare void @agg_min_int8_skip_val_shared(i8*, i8, i8); declare void @agg_min_double_shared(i64*, double); declare void @agg_min_double_skip_val_shared(i64*, double, double); declare void @agg_min_float_shared(i32*, float); declare void @agg_min_float_skip_val_shared(i32*, float, float); declare void @agg_id_shared(i64*, i64); declare i8* @agg_id_varlen_shared(i8*, i64, i8*, i64); declare void @agg_id_int32_shared(i32*, i32); declare void @agg_id_int16_shared(i16*, i16); declare void @agg_id_int8_shared(i8*, i8); declare void @agg_id_double_shared(i64*, double); declare void @agg_id_double_shared_slow(i64*, double*); declare void @agg_id_float_shared(i32*, float); declare i32 @checked_single_agg_id_shared(i64*, i64, i64); declare i32 @checked_single_agg_id_double_shared(i64*, double, double); declare i32 @checked_single_agg_id_double_shared_slow(i64*, double*, double); declare i32 @checked_single_agg_id_float_shared(i32*, float, float); declare i1 @slotEmptyKeyCAS(i64*, i64, i64); declare i1 @slotEmptyKeyCAS_int32(i32*, i32, i32); declare i1 @slotEmptyKeyCAS_int16(i16*, i16, i16); declare i1 @slotEmptyKeyCAS_int8(i8*, i8, i8); declare i64 @datetrunc_century(i64); declare i64 @datetrunc_day(i64); declare i64 @datetrunc_decade(i64); declare i64 @datetrunc_hour(i64); declare i64 @datetrunc_millennium(i64); declare i64 @datetrunc_minute(i64); declare i64 @datetrunc_month(i64); declare i64 @datetrunc_quarter(i64); declare i64 @datetrunc_quarterday(i64); declare i64 @datetrunc_week_monday(i64); declare i64 @datetrunc_week_sunday(i64); declare i64 @datetrunc_week_saturday(i64); declare i64 @datetrunc_year(i64); declare i64 @extract_epoch(i64); declare i64 @extract_dateepoch(i64); declare i64 @extract_quarterday(i64); declare i64 @extract_hour(i64); declare i64 @extract_minute(i64); declare i64 @extract_second(i64); declare i64 @extract_millisecond(i64); declare i64 @extract_microsecond(i64); declare i64 @extract_nanosecond(i64); declare i64 @extract_dow(i64); declare i64 @extract_isodow(i64); declare i64 @extract_day(i64); declare i64 @extract_week_monday(i64); declare i64 @extract_week_sunday(i64); declare i64 @extract_week_saturday(i64); declare i64 @extract_day_of_year(i64); declare i64 @extract_month(i64); declare i64 @extract_quarter(i64); declare i64 @extract_year(i64); declare i64 @ExtractTimeFromHPTimestamp(i64,i64); declare i64 @ExtractTimeFromHPTimestampNullable(i64,i64,i64); declare i64 @ExtractTimeFromLPTimestamp(i64); declare i64 @ExtractTimeFromLPTimestampNullable(i64,i64); declare i64 @DateTruncateHighPrecisionToDate(i64, i64); declare i64 @DateTruncateHighPrecisionToDateNullable(i64, i64, i64); declare i64 @DateDiff(i32, i64, i64); declare i64 @DateDiffNullable(i32, i64, i64, i64); declare i64 @DateDiffHighPrecision(i32, i64, i64, i32, i32); declare i64 @DateDiffHighPrecisionNullable(i32, i64, i64, i32, i32, i64); declare i64 @DateAdd(i32, i64, i64); declare i64 @DateAddNullable(i32, i64, i64, i64); declare i64 @DateAddHighPrecision(i32, i64, i64, i32); declare i64 @DateAddHighPrecisionNullable(i32, i64, i64, i32, i64); declare i64 @string_decode(i8*, i64); declare i32 @array_size(i8*, i64, i32); declare i32 @array_size_nullable(i8*, i64, i32, i32); declare i32 @fast_fixlen_array_size(i8*, i32); declare i1 @array_is_null(i8*, i64); declare i1 @point_coord_array_is_null(i8*, i64); declare i8* @array_buff(i8*, i64); declare i8* @fast_fixlen_array_buff(i8*, i64); declare i8 @array_at_int8_t(i8*, i64, i32); declare i16 @array_at_int16_t(i8*, i64, i32); declare i32 @array_at_int32_t(i8*, i64, i32); declare i64 @array_at_int64_t(i8*, i64, i32); declare float @array_at_float(i8*, i64, i32); declare double @array_at_double(i8*, i64, i32); declare i8 @varlen_array_at_int8_t(i8*, i64, i32); declare i16 @varlen_array_at_int16_t(i8*, i64, i32); declare i32 @varlen_array_at_int32_t(i8*, i64, i32); declare i64 @varlen_array_at_int64_t(i8*, i64, i32); declare float @varlen_array_at_float(i8*, i64, i32); declare double @varlen_array_at_double(i8*, i64, i32); declare i8 @varlen_notnull_array_at_int8_t(i8*, i64, i32); declare i16 @varlen_notnull_array_at_int16_t(i8*, i64, i32); declare i32 @varlen_notnull_array_at_int32_t(i8*, i64, i32); declare i64 @varlen_notnull_array_at_int64_t(i8*, i64, i32); declare float @varlen_notnull_array_at_float(i8*, i64, i32); declare double @varlen_notnull_array_at_double(i8*, i64, i32); declare i8 @array_at_int8_t_checked(i8*, i64, i64, i8); declare i16 @array_at_int16_t_checked(i8*, i64, i64, i16); declare i32 @array_at_int32_t_checked(i8*, i64, i64, i32); declare i64 @array_at_int64_t_checked(i8*, i64, i64, i64); declare float @array_at_float_checked(i8*, i64, i64, float); declare double @array_at_double_checked(i8*, i64, i64, double); declare i32 @char_length(i8*, i32); declare i32 @char_length_nullable(i8*, i32, i32); declare i32 @char_length_encoded(i8*, i32); declare i32 @char_length_encoded_nullable(i8*, i32, i32); declare i32 @key_for_string_encoded(i32); declare i1 @sample_ratio(double, i64); declare double @width_bucket(double, double, double, double, i32); declare double @width_bucket_reverse(double, double, double, double, i32); declare double @width_bucket_nullable(double, double, double, double, i32, double); declare double @width_bucket_reversed_nullable(double, double, double, double, i32, double); declare double @width_bucket_no_oob_check(double, double, double); declare double @width_bucket_reverse_no_oob_check(double, double, double); declare double @width_bucket_expr(double, i1, double, double, i32); declare double @width_bucket_expr_nullable(double, i1, double, double, i32, double); declare double @width_bucket_expr_no_oob_check(double, i1, double, double, i32); declare i1 @string_like(i8*, i32, i8*, i32, i8); declare i1 @string_ilike(i8*, i32, i8*, i32, i8); declare i8 @string_like_nullable(i8*, i32, i8*, i32, i8, i8); declare i8 @string_ilike_nullable(i8*, i32, i8*, i32, i8, i8); declare i1 @string_like_simple(i8*, i32, i8*, i32); declare i1 @string_ilike_simple(i8*, i32, i8*, i32); declare i8 @string_like_simple_nullable(i8*, i32, i8*, i32, i8); declare i8 @string_ilike_simple_nullable(i8*, i32, i8*, i32, i8); declare i1 @string_lt(i8*, i32, i8*, i32); declare i1 @string_le(i8*, i32, i8*, i32); declare i1 @string_gt(i8*, i32, i8*, i32); declare i1 @string_ge(i8*, i32, i8*, i32); declare i1 @string_eq(i8*, i32, i8*, i32); declare i1 @string_ne(i8*, i32, i8*, i32); declare i8 @string_lt_nullable(i8*, i32, i8*, i32, i8); declare i8 @string_le_nullable(i8*, i32, i8*, i32, i8); declare i8 @string_gt_nullable(i8*, i32, i8*, i32, i8); declare i8 @string_ge_nullable(i8*, i32, i8*, i32, i8); declare i8 @string_eq_nullable(i8*, i32, i8*, i32, i8); declare i8 @string_ne_nullable(i8*, i32, i8*, i32, i8); declare i1 @regexp_like(i8*, i32, i8*, i32, i8); declare i8 @regexp_like_nullable(i8*, i32, i8*, i32, i8, i8); declare void @linear_probabilistic_count(i8*, i32, i8*, i32); declare void @agg_count_distinct_bitmap_gpu(i64*, i64, i64, i64, i64, i64, i64); declare void @agg_count_distinct_bitmap_skip_val_gpu(i64*, i64, i64, i64, i64, i64, i64, i64); declare void @agg_approximate_count_distinct_gpu(i64*, i64, i32, i64, i64); declare void @record_error_code(i32, i32*); declare i32 @get_error_code(i32*); declare i1 @dynamic_watchdog(); declare i1 @check_interrupt(); declare void @force_sync(); declare void @sync_warp(); declare void @sync_warp_protected(i64, i64); declare void @sync_threadblock(); declare i64* @get_bin_from_k_heap_int32_t(i64*, i32, i32, i32, i1, i1, i1, i32, i32); declare i64* @get_bin_from_k_heap_int64_t(i64*, i32, i32, i32, i1, i1, i1, i64, i64); declare i64* @get_bin_from_k_heap_float(i64*, i32, i32, i32, i1, i1, i1, float, float); declare i64* @get_bin_from_k_heap_double(i64*, i32, i32, i32, i1, i1, i1, double, double); declare double @decompress_x_coord_geoint(i32); declare double @decompress_y_coord_geoint(i32); declare i32 @compress_x_coord_geoint(double); declare i32 @compress_y_coord_geoint(double); )" + gen_array_any_all_sigs() +
629 
630 #ifdef HAVE_CUDA
631 std::string extension_function_decls(const std::unordered_set<std::string>& udf_decls) {
632  const auto decls =
633  ExtensionFunctionsWhitelist::getLLVMDeclarations(udf_decls, /*is_gpu=*/true);
634  return boost::algorithm::join(decls, "\n");
635 }
636 
637 void legalize_nvvm_ir(llvm::Function* query_func) {
638  // optimizations might add attributes to the function
639  // and NVPTX doesn't understand all of them; play it
640  // safe and clear all attributes
641  clear_function_attributes(query_func);
642  verify_function_ir(query_func);
643 
644  std::vector<llvm::Instruction*> stackrestore_intrinsics;
645  std::vector<llvm::Instruction*> stacksave_intrinsics;
646  std::vector<llvm::Instruction*> lifetime;
647  for (auto& BB : *query_func) {
648  for (llvm::Instruction& I : BB) {
649  if (const llvm::IntrinsicInst* II = llvm::dyn_cast<llvm::IntrinsicInst>(&I)) {
650  if (II->getIntrinsicID() == llvm::Intrinsic::stacksave) {
651  stacksave_intrinsics.push_back(&I);
652  } else if (II->getIntrinsicID() == llvm::Intrinsic::stackrestore) {
653  stackrestore_intrinsics.push_back(&I);
654  } else if (II->getIntrinsicID() == llvm::Intrinsic::lifetime_start ||
655  II->getIntrinsicID() == llvm::Intrinsic::lifetime_end) {
656  lifetime.push_back(&I);
657  }
658  }
659  }
660  }
661 
662  // stacksave and stackrestore intrinsics appear together, and
663  // stackrestore uses stacksaved result as its argument
664  // so it should be removed first.
665  for (auto& II : stackrestore_intrinsics) {
666  II->eraseFromParent();
667  }
668  for (auto& II : stacksave_intrinsics) {
669  II->eraseFromParent();
670  }
671  // Remove lifetime intrinsics as well. NVPTX don't like them
672  for (auto& II : lifetime) {
673  II->eraseFromParent();
674  }
675 }
676 #endif // HAVE_CUDA
677 
678 } // namespace
679 
680 llvm::StringRef get_gpu_target_triple_string() {
681  return llvm::StringRef("nvptx64-nvidia-cuda");
682 }
683 
684 llvm::StringRef get_gpu_data_layout() {
685  return llvm::StringRef(
686  "e-p:64:64:64-i1:8:8-i8:8:8-"
687  "i16:16:16-i32:32:32-i64:64:64-"
688  "f32:32:32-f64:64:64-v16:16:16-"
689  "v32:32:32-v64:64:64-v128:128:128-n16:32:64");
690 }
691 
692 std::map<std::string, std::string> get_device_parameters(bool cpu_only) {
693  std::map<std::string, std::string> result;
694 
695  result.insert(std::make_pair("cpu_name", llvm::sys::getHostCPUName()));
696  result.insert(std::make_pair("cpu_triple", llvm::sys::getProcessTriple()));
697  result.insert(
698  std::make_pair("cpu_cores", std::to_string(llvm::sys::getHostNumPhysicalCores())));
699  result.insert(std::make_pair("cpu_threads", std::to_string(cpu_threads())));
700 
701  // https://en.cppreference.com/w/cpp/language/types
702  std::string sizeof_types;
703  sizeof_types += "bool:" + std::to_string(sizeof(bool)) + ";";
704  sizeof_types += "size_t:" + std::to_string(sizeof(size_t)) + ";";
705  sizeof_types += "ssize_t:" + std::to_string(sizeof(ssize_t)) + ";";
706  sizeof_types += "char:" + std::to_string(sizeof(char)) + ";";
707  sizeof_types += "uchar:" + std::to_string(sizeof(unsigned char)) + ";";
708  sizeof_types += "short:" + std::to_string(sizeof(short)) + ";";
709  sizeof_types += "ushort:" + std::to_string(sizeof(unsigned short int)) + ";";
710  sizeof_types += "int:" + std::to_string(sizeof(int)) + ";";
711  sizeof_types += "uint:" + std::to_string(sizeof(unsigned int)) + ";";
712  sizeof_types += "long:" + std::to_string(sizeof(long int)) + ";";
713  sizeof_types += "ulong:" + std::to_string(sizeof(unsigned long int)) + ";";
714  sizeof_types += "longlong:" + std::to_string(sizeof(long long int)) + ";";
715  sizeof_types += "ulonglong:" + std::to_string(sizeof(unsigned long long int)) + ";";
716  sizeof_types += "float:" + std::to_string(sizeof(float)) + ";";
717  sizeof_types += "double:" + std::to_string(sizeof(double)) + ";";
718  sizeof_types += "longdouble:" + std::to_string(sizeof(long double)) + ";";
719  sizeof_types += "voidptr:" + std::to_string(sizeof(void*)) + ";";
720 
721  result.insert(std::make_pair("type_sizeof", sizeof_types));
722 
723  std::string null_values;
724  null_values += "boolean1:" + std::to_string(serialized_null_value<bool>()) + ";";
725  null_values += "boolean8:" + std::to_string(serialized_null_value<int8_t>()) + ";";
726  null_values += "int8:" + std::to_string(serialized_null_value<int8_t>()) + ";";
727  null_values += "int16:" + std::to_string(serialized_null_value<int16_t>()) + ";";
728  null_values += "int32:" + std::to_string(serialized_null_value<int32_t>()) + ";";
729  null_values += "int64:" + std::to_string(serialized_null_value<int64_t>()) + ";";
730  null_values += "uint8:" + std::to_string(serialized_null_value<uint8_t>()) + ";";
731  null_values += "uint16:" + std::to_string(serialized_null_value<uint16_t>()) + ";";
732  null_values += "uint32:" + std::to_string(serialized_null_value<uint32_t>()) + ";";
733  null_values += "uint64:" + std::to_string(serialized_null_value<uint64_t>()) + ";";
734  null_values += "float32:" + std::to_string(serialized_null_value<float>()) + ";";
735  null_values += "float64:" + std::to_string(serialized_null_value<double>()) + ";";
736  null_values +=
737  "Array<boolean8>:" + std::to_string(serialized_null_value<int8_t, true>()) + ";";
738  null_values +=
739  "Array<int8>:" + std::to_string(serialized_null_value<int8_t, true>()) + ";";
740  null_values +=
741  "Array<int16>:" + std::to_string(serialized_null_value<int16_t, true>()) + ";";
742  null_values +=
743  "Array<int32>:" + std::to_string(serialized_null_value<int32_t, true>()) + ";";
744  null_values +=
745  "Array<int64>:" + std::to_string(serialized_null_value<int64_t, true>()) + ";";
746  null_values +=
747  "Array<float32>:" + std::to_string(serialized_null_value<float, true>()) + ";";
748  null_values +=
749  "Array<float64>:" + std::to_string(serialized_null_value<double, true>()) + ";";
750 
751  result.insert(std::make_pair("null_values", null_values));
752 
753  llvm::StringMap<bool> cpu_features;
754  if (llvm::sys::getHostCPUFeatures(cpu_features)) {
755  std::string features_str = "";
756  for (auto it = cpu_features.begin(); it != cpu_features.end(); ++it) {
757  features_str += (it->getValue() ? " +" : " -");
758  features_str += it->getKey().str();
759  }
760  result.insert(std::make_pair("cpu_features", features_str));
761  }
762 
763  result.insert(std::make_pair("llvm_version",
764  std::to_string(LLVM_VERSION_MAJOR) + "." +
765  std::to_string(LLVM_VERSION_MINOR) + "." +
766  std::to_string(LLVM_VERSION_PATCH)));
767 
768 #ifdef HAVE_CUDA
769  if (!cpu_only) {
770  int device_count = 0;
771  checkCudaErrors(cuDeviceGetCount(&device_count));
772  if (device_count) {
773  CUdevice device{};
774  char device_name[256];
775  int major = 0, minor = 0;
776  int driver_version;
777  checkCudaErrors(cuDeviceGet(&device, 0)); // assuming homogeneous multi-GPU system
778  checkCudaErrors(cuDeviceGetName(device_name, 256, device));
779  checkCudaErrors(cuDeviceGetAttribute(
780  &major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device));
781  checkCudaErrors(cuDeviceGetAttribute(
782  &minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device));
783  checkCudaErrors(cuDriverGetVersion(&driver_version));
784 
785  result.insert(std::make_pair("gpu_name", device_name));
786  result.insert(std::make_pair("gpu_count", std::to_string(device_count)));
787  result.insert(std::make_pair("gpu_compute_capability",
788  std::to_string(major) + "." + std::to_string(minor)));
789  result.insert(std::make_pair("gpu_triple", get_gpu_target_triple_string()));
790  result.insert(std::make_pair("gpu_datalayout", get_gpu_data_layout()));
791  result.insert(std::make_pair("gpu_driver",
792  "CUDA " + std::to_string(driver_version / 1000) + "." +
793  std::to_string((driver_version % 1000) / 10)));
794  }
795  }
796 #endif
797 
798  return result;
799 }
800 
801 namespace {
802 
803 #ifdef HAVE_CUDA
804 std::unordered_set<llvm::Function*> findAliveRuntimeFuncs(
805  llvm::Module& llvm_module,
806  const std::vector<llvm::Function*>& roots) {
807  std::queue<llvm::Function*> queue;
808  std::unordered_set<llvm::Function*> visited;
809  for (llvm::Function* F : roots) {
810  queue.push(F);
811  }
812 
813  while (!queue.empty()) {
814  llvm::Function* F = queue.front();
815  queue.pop();
816  if (visited.find(F) != visited.end()) {
817  continue;
818  }
819  visited.insert(F);
820 
821  for (llvm::inst_iterator I = inst_begin(F), E = inst_end(F); I != E; ++I) {
822  if (llvm::CallInst* CI = llvm::dyn_cast<llvm::CallInst>(&*I)) {
823  if (CI->isInlineAsm()) // libdevice calls inline assembly code
824  continue;
825  llvm::Function* called = CI->getCalledFunction();
826  if (!called || visited.find(called) != visited.end()) {
827  continue;
828  }
829  queue.push(called);
830  }
831  }
832  }
833  return visited;
834 }
835 #endif
836 
837 } // namespace
838 
840  Executor* executor,
841  llvm::Module& llvm_module,
842  llvm::PassManagerBuilder& pass_manager_builder,
843  const GPUTarget& gpu_target) {
844 #ifdef HAVE_CUDA
845  auto timer = DEBUG_TIMER(__func__);
846 
847  if (!executor->has_libdevice_module()) {
848  // raise error
849  throw std::runtime_error(
850  "libdevice library is not available but required by the UDF module");
851  }
852 
853  // Saves functions \in module
854  std::vector<llvm::Function*> roots;
855  for (llvm::Function& fn : llvm_module) {
856  if (!fn.isDeclaration())
857  roots.emplace_back(&fn);
858  }
859 
860  // Bind libdevice to the current module
861  CodeGenerator::link_udf_module(executor->get_libdevice_module(),
862  llvm_module,
863  gpu_target.cgen_state,
864  llvm::Linker::Flags::OverrideFromSrc);
865 
866  std::unordered_set<llvm::Function*> live_funcs =
867  findAliveRuntimeFuncs(llvm_module, roots);
868 
869  std::vector<llvm::Function*> funcs_to_delete;
870  for (llvm::Function& fn : llvm_module) {
871  if (!live_funcs.count(&fn)) {
872  // deleting the function were would invalidate the iterator
873  funcs_to_delete.emplace_back(&fn);
874  }
875  }
876 
877  for (llvm::Function* f : funcs_to_delete) {
878  f->eraseFromParent();
879  }
880 
881  // activate nvvm-reflect-ftz flag on the module
882 #if LLVM_VERSION_MAJOR >= 11
883  llvm::LLVMContext& ctx = llvm_module.getContext();
884  llvm_module.setModuleFlag(llvm::Module::Override,
885  "nvvm-reflect-ftz",
886  llvm::ConstantAsMetadata::get(llvm::ConstantInt::get(
887  llvm::Type::getInt32Ty(ctx), uint32_t(1))));
888 #else
889  llvm_module.addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz", uint32_t(1));
890 #endif
891  for (llvm::Function& fn : llvm_module) {
892  fn.addFnAttr("nvptx-f32ftz", "true");
893  }
894 
895  // add nvvm reflect pass replacing any NVVM conditionals with constants
896  gpu_target.nvptx_target_machine->adjustPassManager(pass_manager_builder);
897  llvm::legacy::FunctionPassManager FPM(&llvm_module);
898  pass_manager_builder.populateFunctionPassManager(FPM);
899 
900  // Run the NVVMReflectPass here rather than inside optimize_ir
901  FPM.doInitialization();
902  for (auto& F : llvm_module) {
903  FPM.run(F);
904  }
905  FPM.doFinalization();
906 #endif
907 }
908 
909 std::shared_ptr<GpuCompilationContext> CodeGenerator::generateNativeGPUCode(
910  Executor* executor,
911  llvm::Function* func,
912  llvm::Function* wrapper_func,
913  const std::unordered_set<llvm::Function*>& live_funcs,
914  const bool is_gpu_smem_used,
915  const CompilationOptions& co,
916  const GPUTarget& gpu_target) {
917 #ifdef HAVE_CUDA
918  auto timer = DEBUG_TIMER(__func__);
919  auto llvm_module = func->getParent();
920  /*
921  `func` is one of the following generated functions:
922  - `call_table_function(i8** %input_col_buffers, i64*
923  %input_row_count, i64** %output_buffers, i64* %output_row_count)`
924  that wraps the user-defined table function.
925  - `multifrag_query`
926  - `multifrag_query_hoisted_literals`
927  - ...
928 
929  `wrapper_func` is table_func_kernel(i32*, i8**, i64*, i64**,
930  i64*) that wraps `call_table_function`.
931 
932  `llvm_module` is from `build/QueryEngine/RuntimeFunctions.bc` and it
933  contains `func` and `wrapper_func`. `llvm_module` should also contain
934  the definitions of user-defined table functions.
935 
936  `live_funcs` contains table_func_kernel and call_table_function
937 
938  `gpu_target.cgen_state->module_` appears to be the same as `llvm_module`
939  */
940  CHECK(gpu_target.cgen_state->module_ == llvm_module);
941  CHECK(func->getParent() == wrapper_func->getParent());
942  llvm_module->setDataLayout(
943  "e-p:64:64:64-i1:8:8-i8:8:8-"
944  "i16:16:16-i32:32:32-i64:64:64-"
945  "f32:32:32-f64:64:64-v16:16:16-"
946  "v32:32:32-v64:64:64-v128:128:128-n16:32:64");
947  llvm_module->setTargetTriple("nvptx64-nvidia-cuda");
948  CHECK(gpu_target.nvptx_target_machine);
949  llvm::PassManagerBuilder pass_manager_builder = llvm::PassManagerBuilder();
950 
951  pass_manager_builder.OptLevel = 0;
952  llvm::legacy::PassManager module_pass_manager;
953  pass_manager_builder.populateModulePassManager(module_pass_manager);
954 
955  bool requires_libdevice = check_module_requires_libdevice(llvm_module);
956 
957  if (requires_libdevice) {
958  linkModuleWithLibdevice(executor, *llvm_module, pass_manager_builder, gpu_target);
959  }
960 
961  // run optimizations
962  optimize_ir(func, llvm_module, module_pass_manager, live_funcs, is_gpu_smem_used, co);
963  legalize_nvvm_ir(func);
964 
965  std::stringstream ss;
966  llvm::raw_os_ostream os(ss);
967 
968  llvm::LLVMContext& ctx = llvm_module->getContext();
969  // Get "nvvm.annotations" metadata node
970  llvm::NamedMDNode* md = llvm_module->getOrInsertNamedMetadata("nvvm.annotations");
971 
972  llvm::Metadata* md_vals[] = {llvm::ConstantAsMetadata::get(wrapper_func),
973  llvm::MDString::get(ctx, "kernel"),
974  llvm::ConstantAsMetadata::get(llvm::ConstantInt::get(
975  llvm::Type::getInt32Ty(ctx), 1))};
976 
977  // Append metadata to nvvm.annotations
978  md->addOperand(llvm::MDNode::get(ctx, md_vals));
979 
980  std::unordered_set<llvm::Function*> roots{wrapper_func, func};
981  if (gpu_target.row_func_not_inlined) {
982  clear_function_attributes(gpu_target.cgen_state->row_func_);
983  roots.insert(gpu_target.cgen_state->row_func_);
984  if (gpu_target.cgen_state->filter_func_) {
985  roots.insert(gpu_target.cgen_state->filter_func_);
986  }
987  }
988 
989  // prevent helper functions from being removed
990  for (auto f : gpu_target.cgen_state->helper_functions_) {
991  roots.insert(f);
992  }
993 
994  if (requires_libdevice) {
995  for (llvm::Function& F : *llvm_module) {
996  // Some libdevice functions calls another functions that starts with "__internal_"
997  // prefix.
998  // __internal_trig_reduction_slowpathd
999  // __internal_accurate_pow
1000  // __internal_lgamma_pos
1001  // Those functions have a "noinline" attribute which prevents the optimizer from
1002  // inlining them into the body of @query_func
1003  if (F.hasName() && F.getName().startswith("__internal") && !F.isDeclaration()) {
1004  roots.insert(&F);
1005  }
1006  legalize_nvvm_ir(&F);
1007  }
1008  }
1009 
1010  // Prevent the udf function(s) from being removed the way the runtime functions are
1011  std::unordered_set<std::string> udf_declarations;
1012 
1013  if (executor->has_udf_module(/*is_gpu=*/true)) {
1014  for (auto& f : executor->get_udf_module(/*is_gpu=*/true)->getFunctionList()) {
1015  llvm::Function* udf_function = llvm_module->getFunction(f.getName());
1016 
1017  if (udf_function) {
1018  legalize_nvvm_ir(udf_function);
1019  roots.insert(udf_function);
1020 
1021  // If we have a udf that declares a external function
1022  // note it so we can avoid duplicate declarations
1023  if (f.isDeclaration()) {
1024  udf_declarations.insert(f.getName().str());
1025  }
1026  }
1027  }
1028  }
1029 
1030  if (executor->has_rt_udf_module(/*is_gpu=*/true)) {
1031  for (auto& f : executor->get_rt_udf_module(/*is_gpu=*/true)->getFunctionList()) {
1032  llvm::Function* udf_function = llvm_module->getFunction(f.getName());
1033  if (udf_function) {
1034  legalize_nvvm_ir(udf_function);
1035  roots.insert(udf_function);
1036 
1037  // If we have a udf that declares a external function
1038  // note it so we can avoid duplicate declarations
1039  if (f.isDeclaration()) {
1040  udf_declarations.insert(f.getName().str());
1041  }
1042  }
1043  }
1044  }
1045 
1046  std::vector<llvm::Function*> rt_funcs;
1047  for (auto& Fn : *llvm_module) {
1048  if (roots.count(&Fn)) {
1049  continue;
1050  }
1051  rt_funcs.push_back(&Fn);
1052  }
1053  for (auto& pFn : rt_funcs) {
1054  pFn->removeFromParent();
1055  }
1056 
1057  if (requires_libdevice) {
1058  add_intrinsics_to_module(llvm_module);
1059  }
1060 
1061  llvm_module->print(os, nullptr);
1062  os.flush();
1063 
1064  for (auto& pFn : rt_funcs) {
1065  llvm_module->getFunctionList().push_back(pFn);
1066  }
1067  llvm_module->eraseNamedMetadata(md);
1068 
1069  auto cuda_llir = ss.str() + cuda_rt_decls + extension_function_decls(udf_declarations);
1070  std::string ptx;
1071  try {
1072  ptx = generatePTX(
1073  cuda_llir, gpu_target.nvptx_target_machine, gpu_target.cgen_state->context_);
1074  } catch (ParseIRError& e) {
1075  LOG(WARNING) << "Failed to generate PTX: " << e.what()
1076  << ". Switching to CPU execution target.";
1077  throw QueryMustRunOnCpu();
1078  }
1079  LOG(PTX) << "PTX for the GPU:\n" << ptx << "\nEnd of PTX";
1080 
1081  auto cubin_result = ptx_to_cubin(ptx, gpu_target.block_size, gpu_target.cuda_mgr);
1082  auto& option_keys = cubin_result.option_keys;
1083  auto& option_values = cubin_result.option_values;
1084  auto cubin = cubin_result.cubin;
1085  auto link_state = cubin_result.link_state;
1086  const auto num_options = option_keys.size();
1087 
1088  auto func_name = wrapper_func->getName().str();
1089  auto gpu_compilation_context = std::make_shared<GpuCompilationContext>();
1090  for (int device_id = 0; device_id < gpu_target.cuda_mgr->getDeviceCount();
1091  ++device_id) {
1092  gpu_compilation_context->addDeviceCode(
1093  std::make_unique<GpuDeviceCompilationContext>(cubin,
1094  func_name,
1095  device_id,
1096  gpu_target.cuda_mgr,
1097  num_options,
1098  &option_keys[0],
1099  &option_values[0]));
1100  }
1101 
1102  checkCudaErrors(cuLinkDestroy(link_state));
1103  return gpu_compilation_context;
1104 #else
1105  return {};
1106 #endif
1107 }
1108 
1109 std::shared_ptr<CompilationContext> Executor::optimizeAndCodegenGPU(
1110  llvm::Function* query_func,
1111  llvm::Function* multifrag_query_func,
1112  std::unordered_set<llvm::Function*>& live_funcs,
1113  const bool no_inline,
1114  const CudaMgr_Namespace::CudaMgr* cuda_mgr,
1115  const bool is_gpu_smem_used,
1116  const CompilationOptions& co) {
1117 #ifdef HAVE_CUDA
1118  auto timer = DEBUG_TIMER(__func__);
1119 
1120  CHECK(cuda_mgr);
1121  CodeCacheKey key{serialize_llvm_object(query_func),
1122  serialize_llvm_object(cgen_state_->row_func_)};
1123  if (cgen_state_->filter_func_) {
1124  key.push_back(serialize_llvm_object(cgen_state_->filter_func_));
1125  }
1126  for (const auto helper : cgen_state_->helper_functions_) {
1127  key.push_back(serialize_llvm_object(helper));
1128  }
1129  auto cached_code = QueryEngine::getInstance()->gpu_code_accessor->get_value(key);
1130  if (cached_code) {
1131  return cached_code;
1132  }
1133 
1134  bool row_func_not_inlined = false;
1135  if (no_inline) {
1136  for (auto it = llvm::inst_begin(cgen_state_->row_func_),
1137  e = llvm::inst_end(cgen_state_->row_func_);
1138  it != e;
1139  ++it) {
1140  if (llvm::isa<llvm::CallInst>(*it)) {
1141  auto& get_gv_call = llvm::cast<llvm::CallInst>(*it);
1142  if (get_gv_call.getCalledFunction()->getName() == "array_size" ||
1143  get_gv_call.getCalledFunction()->getName() == "linear_probabilistic_count") {
1144  mark_function_never_inline(cgen_state_->row_func_);
1145  row_func_not_inlined = true;
1146  break;
1147  }
1148  }
1149  }
1150  }
1151 
1152  initializeNVPTXBackend();
1153  CodeGenerator::GPUTarget gpu_target{nvptx_target_machine_.get(),
1154  cuda_mgr,
1155  blockSize(),
1156  cgen_state_.get(),
1157  row_func_not_inlined};
1158  std::shared_ptr<GpuCompilationContext> compilation_context;
1159 
1160  try {
1161  compilation_context = CodeGenerator::generateNativeGPUCode(this,
1162  query_func,
1163  multifrag_query_func,
1164  live_funcs,
1165  is_gpu_smem_used,
1166  co,
1167  gpu_target);
1168  } catch (CudaMgr_Namespace::CudaErrorException& cuda_error) {
1169  if (cuda_error.getStatus() == CUDA_ERROR_OUT_OF_MEMORY) {
1170  // Thrown if memory not able to be allocated on gpu
1171  // Retry once after evicting portion of code cache
1172  LOG(WARNING) << "Failed to allocate GPU memory for generated code. Evicting "
1174  << "% of GPU code cache and re-trying.";
1175  QueryEngine::getInstance()->gpu_code_accessor->evictFractionEntries(
1177  compilation_context = CodeGenerator::generateNativeGPUCode(this,
1178  query_func,
1179  multifrag_query_func,
1180  live_funcs,
1181  is_gpu_smem_used,
1182  co,
1183  gpu_target);
1184  } else {
1185  throw;
1186  }
1187  }
1188  QueryEngine::getInstance()->gpu_code_accessor->put(key, compilation_context);
1189  return std::dynamic_pointer_cast<CompilationContext>(compilation_context);
1190 #else
1191  return nullptr;
1192 #endif
1193 }
1194 
1195 std::string CodeGenerator::generatePTX(const std::string& cuda_llir,
1196  llvm::TargetMachine* nvptx_target_machine,
1197  llvm::LLVMContext& context) {
1198  auto timer = DEBUG_TIMER(__func__);
1199  auto mem_buff = llvm::MemoryBuffer::getMemBuffer(cuda_llir, "", false);
1200 
1201  llvm::SMDiagnostic parse_error;
1202 
1203  auto llvm_module = llvm::parseIR(mem_buff->getMemBufferRef(), parse_error, context);
1204  if (!llvm_module) {
1205  LOG(IR) << "CodeGenerator::generatePTX:NVVM IR:\n" << cuda_llir << "\nEnd of NNVM IR";
1206  throw_parseIR_error(parse_error, "generatePTX", /* is_gpu= */ true);
1207  }
1208 
1209  llvm::SmallString<256> code_str;
1210  llvm::raw_svector_ostream formatted_os(code_str);
1211  CHECK(nvptx_target_machine);
1212  {
1213  llvm::legacy::PassManager ptxgen_pm;
1214  llvm_module->setDataLayout(nvptx_target_machine->createDataLayout());
1215 
1216 #if LLVM_VERSION_MAJOR >= 10
1217  nvptx_target_machine->addPassesToEmitFile(
1218  ptxgen_pm, formatted_os, nullptr, llvm::CGFT_AssemblyFile);
1219 #else
1220  nvptx_target_machine->addPassesToEmitFile(
1221  ptxgen_pm, formatted_os, nullptr, llvm::TargetMachine::CGFT_AssemblyFile);
1222 #endif
1223  ptxgen_pm.run(*llvm_module);
1224  }
1225 
1226 #if LLVM_VERSION_MAJOR >= 11
1227  return std::string(code_str);
1228 #else
1229  return code_str.str();
1230 #endif
1231 }
1232 
1233 std::unique_ptr<llvm::TargetMachine> CodeGenerator::initializeNVPTXBackend(
1235  auto timer = DEBUG_TIMER(__func__);
1236  llvm::InitializeAllTargets();
1237  llvm::InitializeAllTargetMCs();
1238  llvm::InitializeAllAsmPrinters();
1239  std::string err;
1240  auto target = llvm::TargetRegistry::lookupTarget("nvptx64", err);
1241  if (!target) {
1242  LOG(FATAL) << err;
1243  }
1244  return std::unique_ptr<llvm::TargetMachine>(
1245  target->createTargetMachine("nvptx64-nvidia-cuda",
1247  "",
1248  llvm::TargetOptions(),
1249  llvm::Reloc::Static));
1250 }
1251 
1252 std::string Executor::generatePTX(const std::string& cuda_llir) const {
1254  cuda_llir, nvptx_target_machine_.get(), cgen_state_->context_);
1255 }
1256 
1257 void Executor::initializeNVPTXBackend() const {
1258  if (nvptx_target_machine_) {
1259  return;
1260  }
1261  const auto arch = cudaMgr()->getDeviceArch();
1262  nvptx_target_machine_ = CodeGenerator::initializeNVPTXBackend(arch);
1263 }
1264 
1265 // A small number of runtime functions don't get through CgenState::emitCall. List them
1266 // explicitly here and always clone their implementation from the runtime module.
1267 bool CodeGenerator::alwaysCloneRuntimeFunction(const llvm::Function* func) {
1268  return func->getName() == "query_stub_hoisted_literals" ||
1269  func->getName() == "multifrag_query_hoisted_literals" ||
1270  func->getName() == "query_stub" || func->getName() == "multifrag_query" ||
1271  func->getName() == "fixed_width_int_decode" ||
1272  func->getName() == "fixed_width_unsigned_decode" ||
1273  func->getName() == "diff_fixed_width_int_decode" ||
1274  func->getName() == "fixed_width_double_decode" ||
1275  func->getName() == "fixed_width_float_decode" ||
1276  func->getName() == "fixed_width_small_date_decode" ||
1277  func->getName() == "fixed_width_date_encode" ||
1278  func->getName() == "record_error_code" || func->getName() == "get_error_code" ||
1279  func->getName() == "pos_start_impl" || func->getName() == "pos_step_impl" ||
1280  func->getName() == "group_buff_idx_impl" ||
1281  func->getName() == "init_shared_mem" ||
1282  func->getName() == "init_shared_mem_nop" || func->getName() == "write_back_nop";
1283 }
1284 
1285 std::unique_ptr<llvm::Module> read_llvm_module_from_bc_file(
1286  const std::string& bc_filename,
1287  llvm::LLVMContext& context) {
1288  llvm::SMDiagnostic err;
1289 
1290  auto buffer_or_error = llvm::MemoryBuffer::getFile(bc_filename);
1291  CHECK(!buffer_or_error.getError()) << "bc_filename=" << bc_filename;
1292  llvm::MemoryBuffer* buffer = buffer_or_error.get().get();
1293 
1294  auto owner = llvm::parseBitcodeFile(buffer->getMemBufferRef(), context);
1295  CHECK(!owner.takeError());
1296  CHECK(owner->get());
1297  return std::move(owner.get());
1298 }
1299 
1300 std::unique_ptr<llvm::Module> read_llvm_module_from_ir_file(
1301  const std::string& udf_ir_filename,
1302  llvm::LLVMContext& ctx,
1303  bool is_gpu = false) {
1304  llvm::SMDiagnostic parse_error;
1305 
1306  llvm::StringRef file_name_arg(udf_ir_filename);
1307 
1308  auto owner = llvm::parseIRFile(file_name_arg, parse_error, ctx);
1309  if (!owner) {
1310  throw_parseIR_error(parse_error, udf_ir_filename, is_gpu);
1311  }
1312 
1313  if (is_gpu) {
1314  llvm::Triple gpu_triple(owner->getTargetTriple());
1315  if (!gpu_triple.isNVPTX()) {
1316  LOG(WARNING)
1317  << "Expected triple nvptx64-nvidia-cuda for NVVM IR of loadtime UDFs but got "
1318  << gpu_triple.str() << ". Disabling the NVVM IR module.";
1319  return std::unique_ptr<llvm::Module>();
1320  }
1321  }
1322  return owner;
1323 }
1324 
1325 std::unique_ptr<llvm::Module> read_llvm_module_from_ir_string(
1326  const std::string& udf_ir_string,
1327  llvm::LLVMContext& ctx,
1328  bool is_gpu = false) {
1329  llvm::SMDiagnostic parse_error;
1330 
1331  auto buf = std::make_unique<llvm::MemoryBufferRef>(udf_ir_string,
1332  "Runtime UDF/UDTF LLVM/NVVM IR");
1333 
1334  auto owner = llvm::parseIR(*buf, parse_error, ctx);
1335  if (!owner) {
1336  LOG(IR) << "read_llvm_module_from_ir_string:\n"
1337  << udf_ir_string << "\nEnd of LLVM/NVVM IR";
1338  throw_parseIR_error(parse_error, "", /* is_gpu= */ is_gpu);
1339  }
1340 
1341  if (is_gpu) {
1342  llvm::Triple gpu_triple(owner->getTargetTriple());
1343  if (!gpu_triple.isNVPTX()) {
1344  LOG(IR) << "read_llvm_module_from_ir_string:\n"
1345  << udf_ir_string << "\nEnd of NNVM IR";
1346  LOG(WARNING) << "Expected triple nvptx64-nvidia-cuda for NVVM IR but got "
1347  << gpu_triple.str()
1348  << ". Executing runtime UDF/UDTFs on GPU will be disabled.";
1349  return std::unique_ptr<llvm::Module>();
1350  ;
1351  }
1352  }
1353  return owner;
1354 }
1355 
1356 namespace {
1357 
1358 void bind_pos_placeholders(const std::string& pos_fn_name,
1359  const bool use_resume_param,
1360  llvm::Function* query_func,
1361  llvm::Module* llvm_module) {
1362  for (auto it = llvm::inst_begin(query_func), e = llvm::inst_end(query_func); it != e;
1363  ++it) {
1364  if (!llvm::isa<llvm::CallInst>(*it)) {
1365  continue;
1366  }
1367  auto& pos_call = llvm::cast<llvm::CallInst>(*it);
1368  if (std::string(pos_call.getCalledFunction()->getName()) == pos_fn_name) {
1369  if (use_resume_param) {
1370  const auto error_code_arg = get_arg_by_name(query_func, "error_code");
1371  llvm::ReplaceInstWithInst(
1372  &pos_call,
1373  llvm::CallInst::Create(llvm_module->getFunction(pos_fn_name + "_impl"),
1374  error_code_arg));
1375  } else {
1376  llvm::ReplaceInstWithInst(
1377  &pos_call,
1378  llvm::CallInst::Create(llvm_module->getFunction(pos_fn_name + "_impl")));
1379  }
1380  break;
1381  }
1382  }
1383 }
1384 
1385 void set_row_func_argnames(llvm::Function* row_func,
1386  const size_t in_col_count,
1387  const size_t agg_col_count,
1388  const bool hoist_literals) {
1389  auto arg_it = row_func->arg_begin();
1390 
1391  if (agg_col_count) {
1392  for (size_t i = 0; i < agg_col_count; ++i) {
1393  arg_it->setName("out");
1394  ++arg_it;
1395  }
1396  } else {
1397  arg_it->setName("group_by_buff");
1398  ++arg_it;
1399  arg_it->setName("varlen_output_buff");
1400  ++arg_it;
1401  arg_it->setName("crt_matched");
1402  ++arg_it;
1403  arg_it->setName("total_matched");
1404  ++arg_it;
1405  arg_it->setName("old_total_matched");
1406  ++arg_it;
1407  arg_it->setName("max_matched");
1408  ++arg_it;
1409  }
1410 
1411  arg_it->setName("agg_init_val");
1412  ++arg_it;
1413 
1414  arg_it->setName("pos");
1415  ++arg_it;
1416 
1417  arg_it->setName("frag_row_off");
1418  ++arg_it;
1419 
1420  arg_it->setName("num_rows_per_scan");
1421  ++arg_it;
1422 
1423  if (hoist_literals) {
1424  arg_it->setName("literals");
1425  ++arg_it;
1426  }
1427 
1428  for (size_t i = 0; i < in_col_count; ++i) {
1429  arg_it->setName("col_buf" + std::to_string(i));
1430  ++arg_it;
1431  }
1432 
1433  arg_it->setName("join_hash_tables");
1434 }
1435 
1436 llvm::Function* create_row_function(const size_t in_col_count,
1437  const size_t agg_col_count,
1438  const bool hoist_literals,
1439  llvm::Module* llvm_module,
1440  llvm::LLVMContext& context) {
1441  std::vector<llvm::Type*> row_process_arg_types;
1442 
1443  if (agg_col_count) {
1444  // output (aggregate) arguments
1445  for (size_t i = 0; i < agg_col_count; ++i) {
1446  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1447  }
1448  } else {
1449  // group by buffer
1450  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1451  // varlen output buffer
1452  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1453  // current match count
1454  row_process_arg_types.push_back(llvm::Type::getInt32PtrTy(context));
1455  // total match count passed from the caller
1456  row_process_arg_types.push_back(llvm::Type::getInt32PtrTy(context));
1457  // old total match count returned to the caller
1458  row_process_arg_types.push_back(llvm::Type::getInt32PtrTy(context));
1459  // max matched (total number of slots in the output buffer)
1460  row_process_arg_types.push_back(llvm::Type::getInt32PtrTy(context));
1461  }
1462 
1463  // aggregate init values
1464  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1465 
1466  // position argument
1467  row_process_arg_types.push_back(llvm::Type::getInt64Ty(context));
1468 
1469  // fragment row offset argument
1470  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1471 
1472  // number of rows for each scan
1473  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1474 
1475  // literals buffer argument
1476  if (hoist_literals) {
1477  row_process_arg_types.push_back(llvm::Type::getInt8PtrTy(context));
1478  }
1479 
1480  // column buffer arguments
1481  for (size_t i = 0; i < in_col_count; ++i) {
1482  row_process_arg_types.emplace_back(llvm::Type::getInt8PtrTy(context));
1483  }
1484 
1485  // join hash table argument
1486  row_process_arg_types.push_back(llvm::Type::getInt64PtrTy(context));
1487 
1488  // generate the function
1489  auto ft =
1490  llvm::FunctionType::get(get_int_type(32, context), row_process_arg_types, false);
1491 
1492  auto row_func = llvm::Function::Create(
1493  ft, llvm::Function::ExternalLinkage, "row_func", llvm_module);
1494 
1495  // set the row function argument names; for debugging purposes only
1496  set_row_func_argnames(row_func, in_col_count, agg_col_count, hoist_literals);
1497 
1498  return row_func;
1499 }
1500 
1501 // Iterate through multifrag_query_func, replacing calls to query_fname with query_func.
1502 void bind_query(llvm::Function* query_func,
1503  const std::string& query_fname,
1504  llvm::Function* multifrag_query_func,
1505  llvm::Module* llvm_module) {
1506  std::vector<llvm::CallInst*> query_stubs;
1507  for (auto it = llvm::inst_begin(multifrag_query_func),
1508  e = llvm::inst_end(multifrag_query_func);
1509  it != e;
1510  ++it) {
1511  if (!llvm::isa<llvm::CallInst>(*it)) {
1512  continue;
1513  }
1514  auto& query_call = llvm::cast<llvm::CallInst>(*it);
1515  if (std::string(query_call.getCalledFunction()->getName()) == query_fname) {
1516  query_stubs.push_back(&query_call);
1517  }
1518  }
1519  for (auto& S : query_stubs) {
1520  std::vector<llvm::Value*> args;
1521  for (size_t i = 0; i < S->getNumOperands() - 1; ++i) {
1522  args.push_back(S->getArgOperand(i));
1523  }
1524  llvm::ReplaceInstWithInst(S, llvm::CallInst::Create(query_func, args, ""));
1525  }
1527 
1528 std::vector<std::string> get_agg_fnames(const std::vector<Analyzer::Expr*>& target_exprs,
1529  const bool is_group_by) {
1530  std::vector<std::string> result;
1531  for (size_t target_idx = 0, agg_col_idx = 0; target_idx < target_exprs.size();
1532  ++target_idx, ++agg_col_idx) {
1533  const auto target_expr = target_exprs[target_idx];
1534  CHECK(target_expr);
1535  const auto target_type_info = target_expr->get_type_info();
1536  const auto agg_expr = dynamic_cast<Analyzer::AggExpr*>(target_expr);
1537  const bool is_varlen =
1538  (target_type_info.is_string() &&
1539  target_type_info.get_compression() == kENCODING_NONE) ||
1540  target_type_info.is_array(); // TODO: should it use is_varlen_array() ?
1541  if (!agg_expr || agg_expr->get_aggtype() == kSAMPLE) {
1542  result.emplace_back(target_type_info.is_fp() ? "agg_id_double" : "agg_id");
1543  if (is_varlen) {
1544  result.emplace_back("agg_id");
1545  }
1546  if (target_type_info.is_geometry()) {
1547  result.emplace_back("agg_id");
1548  for (auto i = 2; i < 2 * target_type_info.get_physical_coord_cols(); ++i) {
1549  result.emplace_back("agg_id");
1550  }
1551  }
1552  continue;
1553  }
1554  const auto agg_type = agg_expr->get_aggtype();
1555  const auto& agg_type_info =
1556  agg_type != kCOUNT ? agg_expr->get_arg()->get_type_info() : target_type_info;
1557  switch (agg_type) {
1558  case kAVG: {
1559  if (!agg_type_info.is_integer() && !agg_type_info.is_decimal() &&
1560  !agg_type_info.is_fp()) {
1561  throw std::runtime_error("AVG is only valid on integer and floating point");
1562  }
1563  result.emplace_back((agg_type_info.is_integer() || agg_type_info.is_time())
1564  ? "agg_sum"
1565  : "agg_sum_double");
1566  result.emplace_back((agg_type_info.is_integer() || agg_type_info.is_time())
1567  ? "agg_count"
1568  : "agg_count_double");
1569  break;
1570  }
1571  case kMIN: {
1572  if (agg_type_info.is_string() || agg_type_info.is_array() ||
1573  agg_type_info.is_geometry()) {
1574  throw std::runtime_error(
1575  "MIN on strings, arrays or geospatial types not supported yet");
1576  }
1577  result.emplace_back((agg_type_info.is_integer() || agg_type_info.is_time())
1578  ? "agg_min"
1579  : "agg_min_double");
1580  break;
1581  }
1582  case kMAX: {
1583  if (agg_type_info.is_string() || agg_type_info.is_array() ||
1584  agg_type_info.is_geometry()) {
1585  throw std::runtime_error(
1586  "MAX on strings, arrays or geospatial types not supported yet");
1587  }
1588  result.emplace_back((agg_type_info.is_integer() || agg_type_info.is_time())
1589  ? "agg_max"
1590  : "agg_max_double");
1591  break;
1592  }
1593  case kSUM: {
1594  if (!agg_type_info.is_integer() && !agg_type_info.is_decimal() &&
1595  !agg_type_info.is_fp()) {
1596  throw std::runtime_error("SUM is only valid on integer and floating point");
1597  }
1598  result.emplace_back((agg_type_info.is_integer() || agg_type_info.is_time())
1599  ? "agg_sum"
1600  : "agg_sum_double");
1601  break;
1602  }
1603  case kCOUNT:
1604  result.emplace_back(agg_expr->get_is_distinct() ? "agg_count_distinct"
1605  : "agg_count");
1606  break;
1607  case kSINGLE_VALUE: {
1608  result.emplace_back(agg_type_info.is_fp() ? "agg_id_double" : "agg_id");
1609  break;
1610  }
1611  case kSAMPLE: {
1612  // Note that varlen SAMPLE arguments are handled separately above
1613  result.emplace_back(agg_type_info.is_fp() ? "agg_id_double" : "agg_id");
1614  break;
1615  }
1617  result.emplace_back("agg_approximate_count_distinct");
1618  break;
1619  case kAPPROX_QUANTILE:
1620  result.emplace_back("agg_approx_quantile");
1621  break;
1622  default:
1623  CHECK(false);
1624  }
1625  }
1626  return result;
1627 }
1628 
1629 } // namespace
1630 
1631 void Executor::addUdfIrToModule(const std::string& udf_ir_filename,
1632  const bool is_cuda_ir) {
1636  udf_ir_filename;
1637 }
1638 
1639 std::unordered_set<llvm::Function*> CodeGenerator::markDeadRuntimeFuncs(
1640  llvm::Module& llvm_module,
1641  const std::vector<llvm::Function*>& roots,
1642  const std::vector<llvm::Function*>& leaves) {
1643  auto timer = DEBUG_TIMER(__func__);
1644  std::unordered_set<llvm::Function*> live_funcs;
1645  live_funcs.insert(roots.begin(), roots.end());
1646  live_funcs.insert(leaves.begin(), leaves.end());
1647 
1648  if (auto F = llvm_module.getFunction("init_shared_mem_nop")) {
1649  live_funcs.insert(F);
1650  }
1651  if (auto F = llvm_module.getFunction("write_back_nop")) {
1652  live_funcs.insert(F);
1653  }
1654 
1655  for (const llvm::Function* F : roots) {
1656  for (const llvm::BasicBlock& BB : *F) {
1657  for (const llvm::Instruction& I : BB) {
1658  if (const llvm::CallInst* CI = llvm::dyn_cast<const llvm::CallInst>(&I)) {
1659  live_funcs.insert(CI->getCalledFunction());
1660  }
1661  }
1662  }
1663  }
1664 
1665  for (llvm::Function& F : llvm_module) {
1666  if (!live_funcs.count(&F) && !F.isDeclaration()) {
1667  F.setLinkage(llvm::GlobalValue::InternalLinkage);
1668  }
1669  }
1670 
1671  return live_funcs;
1672 }
1673 
1674 namespace {
1675 // searches for a particular variable within a specific basic block (or all if bb_name is
1676 // empty)
1677 template <typename InstType>
1678 llvm::Value* find_variable_in_basic_block(llvm::Function* func,
1679  std::string bb_name,
1680  std::string variable_name) {
1681  llvm::Value* result = nullptr;
1682  if (func == nullptr || variable_name.empty()) {
1683  return result;
1684  }
1685  bool is_found = false;
1686  for (auto bb_it = func->begin(); bb_it != func->end() && !is_found; ++bb_it) {
1687  if (!bb_name.empty() && bb_it->getName() != bb_name) {
1688  continue;
1689  }
1690  for (auto inst_it = bb_it->begin(); inst_it != bb_it->end(); inst_it++) {
1691  if (llvm::isa<InstType>(*inst_it)) {
1692  if (inst_it->getName() == variable_name) {
1693  result = &*inst_it;
1694  is_found = true;
1695  break;
1696  }
1697  }
1698  }
1699  }
1700  return result;
1701 }
1702 }; // namespace
1703 
1705  llvm::Function* query_func,
1706  bool run_with_dynamic_watchdog,
1707  bool run_with_allowing_runtime_interrupt,
1708  const std::vector<JoinLoop>& join_loops,
1709  ExecutorDeviceType device_type,
1710  const std::vector<InputTableInfo>& input_table_infos) {
1711  AUTOMATIC_IR_METADATA(cgen_state_.get());
1712 
1713  // check whether the row processing was successful; currently, it can
1714  // fail by running out of group by buffer slots
1715 
1716  if (run_with_dynamic_watchdog && run_with_allowing_runtime_interrupt) {
1717  // when both dynamic watchdog and runtime interrupt turns on
1718  // we use dynamic watchdog
1719  run_with_allowing_runtime_interrupt = false;
1720  }
1721 
1722  {
1723  // disable injecting query interrupt checker if the session info is invalid
1725  executor_session_mutex_);
1726  if (current_query_session_.empty()) {
1727  run_with_allowing_runtime_interrupt = false;
1728  }
1729  }
1730 
1731  llvm::Value* row_count = nullptr;
1732  if ((run_with_dynamic_watchdog || run_with_allowing_runtime_interrupt) &&
1733  device_type == ExecutorDeviceType::GPU) {
1734  row_count =
1735  find_variable_in_basic_block<llvm::LoadInst>(query_func, ".entry", "row_count");
1736  }
1737 
1738  bool done_splitting = false;
1739  for (auto bb_it = query_func->begin(); bb_it != query_func->end() && !done_splitting;
1740  ++bb_it) {
1741  llvm::Value* pos = nullptr;
1742  for (auto inst_it = bb_it->begin(); inst_it != bb_it->end(); ++inst_it) {
1743  if ((run_with_dynamic_watchdog || run_with_allowing_runtime_interrupt) &&
1744  llvm::isa<llvm::PHINode>(*inst_it)) {
1745  if (inst_it->getName() == "pos") {
1746  pos = &*inst_it;
1747  }
1748  continue;
1749  }
1750  if (!llvm::isa<llvm::CallInst>(*inst_it)) {
1751  continue;
1752  }
1753  auto& row_func_call = llvm::cast<llvm::CallInst>(*inst_it);
1754  if (std::string(row_func_call.getCalledFunction()->getName()) == "row_process") {
1755  auto next_inst_it = inst_it;
1756  ++next_inst_it;
1757  auto new_bb = bb_it->splitBasicBlock(next_inst_it);
1758  auto& br_instr = bb_it->back();
1759  llvm::IRBuilder<> ir_builder(&br_instr);
1760  llvm::Value* err_lv = &*inst_it;
1761  llvm::Value* err_lv_returned_from_row_func = nullptr;
1762  if (run_with_dynamic_watchdog) {
1763  CHECK(pos);
1764  llvm::Value* call_watchdog_lv = nullptr;
1765  if (device_type == ExecutorDeviceType::GPU) {
1766  // In order to make sure all threads within a block see the same barrier,
1767  // only those blocks whose none of their threads have experienced the critical
1768  // edge will go through the dynamic watchdog computation
1769  CHECK(row_count);
1770  auto crit_edge_rem =
1771  (blockSize() & (blockSize() - 1))
1772  ? ir_builder.CreateSRem(
1773  row_count,
1774  cgen_state_->llInt(static_cast<int64_t>(blockSize())))
1775  : ir_builder.CreateAnd(
1776  row_count,
1777  cgen_state_->llInt(static_cast<int64_t>(blockSize() - 1)));
1778  auto crit_edge_threshold = ir_builder.CreateSub(row_count, crit_edge_rem);
1779  crit_edge_threshold->setName("crit_edge_threshold");
1780 
1781  // only those threads where pos < crit_edge_threshold go through dynamic
1782  // watchdog call
1783  call_watchdog_lv =
1784  ir_builder.CreateICmp(llvm::ICmpInst::ICMP_SLT, pos, crit_edge_threshold);
1785  } else {
1786  // CPU path: run watchdog for every 64th row
1787  auto dw_predicate = ir_builder.CreateAnd(pos, uint64_t(0x3f));
1788  call_watchdog_lv = ir_builder.CreateICmp(
1789  llvm::ICmpInst::ICMP_EQ, dw_predicate, cgen_state_->llInt(int64_t(0LL)));
1790  }
1791  CHECK(call_watchdog_lv);
1792  auto error_check_bb = bb_it->splitBasicBlock(
1793  llvm::BasicBlock::iterator(br_instr), ".error_check");
1794  auto& watchdog_br_instr = bb_it->back();
1795 
1796  auto watchdog_check_bb = llvm::BasicBlock::Create(
1797  cgen_state_->context_, ".watchdog_check", query_func, error_check_bb);
1798  llvm::IRBuilder<> watchdog_ir_builder(watchdog_check_bb);
1799  auto detected_timeout = watchdog_ir_builder.CreateCall(
1800  cgen_state_->module_->getFunction("dynamic_watchdog"), {});
1801  auto timeout_err_lv = watchdog_ir_builder.CreateSelect(
1802  detected_timeout, cgen_state_->llInt(Executor::ERR_OUT_OF_TIME), err_lv);
1803  watchdog_ir_builder.CreateBr(error_check_bb);
1804 
1805  llvm::ReplaceInstWithInst(
1806  &watchdog_br_instr,
1807  llvm::BranchInst::Create(
1808  watchdog_check_bb, error_check_bb, call_watchdog_lv));
1809  ir_builder.SetInsertPoint(&br_instr);
1810  auto unified_err_lv = ir_builder.CreatePHI(err_lv->getType(), 2);
1811 
1812  unified_err_lv->addIncoming(timeout_err_lv, watchdog_check_bb);
1813  unified_err_lv->addIncoming(err_lv, &*bb_it);
1814  err_lv = unified_err_lv;
1815  } else if (run_with_allowing_runtime_interrupt) {
1816  CHECK(pos);
1817  llvm::Value* call_check_interrupt_lv{nullptr};
1818  llvm::Value* interrupt_err_lv{nullptr};
1819  llvm::BasicBlock* error_check_bb{nullptr};
1820  llvm::BasicBlock* interrupt_check_bb{nullptr};
1821  llvm::Instruction* check_interrupt_br_instr{nullptr};
1822 
1823  auto has_loop_join = std::any_of(
1824  join_loops.begin(), join_loops.end(), [](const JoinLoop& join_loop) {
1825  return join_loop.isNestedLoopJoin();
1826  });
1827  auto codegen_interrupt_checker = [&]() {
1828  error_check_bb = bb_it->splitBasicBlock(llvm::BasicBlock::iterator(br_instr),
1829  ".error_check");
1830  check_interrupt_br_instr = &bb_it->back();
1831 
1832  interrupt_check_bb = llvm::BasicBlock::Create(
1833  cgen_state_->context_, ".interrupt_check", query_func, error_check_bb);
1834  llvm::IRBuilder<> interrupt_checker_ir_builder(interrupt_check_bb);
1835  auto detected_interrupt = interrupt_checker_ir_builder.CreateCall(
1836  cgen_state_->module_->getFunction("check_interrupt"), {});
1837  interrupt_err_lv = interrupt_checker_ir_builder.CreateSelect(
1838  detected_interrupt,
1839  cgen_state_->llInt(Executor::ERR_INTERRUPTED),
1840  err_lv);
1841  interrupt_checker_ir_builder.CreateBr(error_check_bb);
1842  };
1843  if (has_loop_join) {
1844  codegen_interrupt_checker();
1845  CHECK(interrupt_check_bb);
1846  CHECK(check_interrupt_br_instr);
1847  llvm::ReplaceInstWithInst(check_interrupt_br_instr,
1848  llvm::BranchInst::Create(interrupt_check_bb));
1849  ir_builder.SetInsertPoint(&br_instr);
1850  err_lv = interrupt_err_lv;
1851  } else {
1852  if (device_type == ExecutorDeviceType::GPU) {
1853  // approximate how many times the %pos variable
1854  // is increased --> the number of iteration
1855  // here we calculate the # bit shift by considering grid/block/fragment
1856  // sizes since if we use the fixed one (i.e., per 64-th increment) some CUDA
1857  // threads cannot enter the interrupt checking block depending on the
1858  // fragment size --> a thread may not take care of 64 threads if an outer
1859  // table is not sufficiently large, and so cannot be interrupted
1860  int32_t num_shift_by_gridDim = shared::getExpOfTwo(gridSize());
1861  int32_t num_shift_by_blockDim = shared::getExpOfTwo(blockSize());
1862  int64_t total_num_shift = num_shift_by_gridDim + num_shift_by_blockDim;
1863  uint64_t interrupt_checking_freq = 32;
1864  auto freq_control_knob = g_running_query_interrupt_freq;
1865  CHECK_GT(freq_control_knob, 0);
1866  CHECK_LE(freq_control_knob, 1.0);
1867  if (!input_table_infos.empty()) {
1868  const auto& outer_table_info = *input_table_infos.begin();
1869  auto num_outer_table_tuples =
1870  outer_table_info.info.getFragmentNumTuplesUpperBound();
1871  if (num_outer_table_tuples > 0) {
1872  // gridSize * blockSize --> pos_step (idx of the next row per thread)
1873  // we additionally multiply two to pos_step since the number of
1874  // dispatched blocks are double of the gridSize
1875  // # tuples (of fragment) / pos_step --> maximum # increment (K)
1876  // also we multiply 1 / freq_control_knob to K to control the frequency
1877  // So, needs to check the interrupt status more frequently? make K
1878  // smaller
1879  auto max_inc = uint64_t(
1880  floor(num_outer_table_tuples / (gridSize() * blockSize() * 2)));
1881  if (max_inc < 2) {
1882  // too small `max_inc`, so this correction is necessary to make
1883  // `interrupt_checking_freq` be valid (i.e., larger than zero)
1884  max_inc = 2;
1885  }
1886  auto calibrated_inc =
1887  uint64_t(floor(max_inc * (1 - freq_control_knob)));
1888  interrupt_checking_freq =
1889  uint64_t(pow(2, shared::getExpOfTwo(calibrated_inc)));
1890  // add the coverage when interrupt_checking_freq > K
1891  // if so, some threads still cannot be branched to the interrupt checker
1892  // so we manually use smaller but close to the max_inc as freq
1893  if (interrupt_checking_freq > max_inc) {
1894  interrupt_checking_freq = max_inc / 2;
1895  }
1896  if (interrupt_checking_freq < 8) {
1897  // such small freq incurs too frequent interrupt status checking,
1898  // so we fixup to the minimum freq value at some reasonable degree
1899  interrupt_checking_freq = 8;
1900  }
1901  }
1902  }
1903  VLOG(1) << "Set the running query interrupt checking frequency: "
1904  << interrupt_checking_freq;
1905  // check the interrupt flag for every interrupt_checking_freq-th iteration
1906  llvm::Value* pos_shifted_per_iteration =
1907  ir_builder.CreateLShr(pos, cgen_state_->llInt(total_num_shift));
1908  auto interrupt_predicate = ir_builder.CreateAnd(pos_shifted_per_iteration,
1909  interrupt_checking_freq);
1910  call_check_interrupt_lv =
1911  ir_builder.CreateICmp(llvm::ICmpInst::ICMP_EQ,
1912  interrupt_predicate,
1913  cgen_state_->llInt(int64_t(0LL)));
1914  } else {
1915  // CPU path: run interrupt checker for every 64th row
1916  auto interrupt_predicate = ir_builder.CreateAnd(pos, uint64_t(0x3f));
1917  call_check_interrupt_lv =
1918  ir_builder.CreateICmp(llvm::ICmpInst::ICMP_EQ,
1919  interrupt_predicate,
1920  cgen_state_->llInt(int64_t(0LL)));
1921  }
1922  codegen_interrupt_checker();
1923  CHECK(call_check_interrupt_lv);
1924  CHECK(interrupt_err_lv);
1925  CHECK(interrupt_check_bb);
1926  CHECK(error_check_bb);
1927  CHECK(check_interrupt_br_instr);
1928  llvm::ReplaceInstWithInst(
1929  check_interrupt_br_instr,
1930  llvm::BranchInst::Create(
1931  interrupt_check_bb, error_check_bb, call_check_interrupt_lv));
1932  ir_builder.SetInsertPoint(&br_instr);
1933  auto unified_err_lv = ir_builder.CreatePHI(err_lv->getType(), 2);
1934 
1935  unified_err_lv->addIncoming(interrupt_err_lv, interrupt_check_bb);
1936  unified_err_lv->addIncoming(err_lv, &*bb_it);
1937  err_lv = unified_err_lv;
1938  }
1939  }
1940  if (!err_lv_returned_from_row_func) {
1941  err_lv_returned_from_row_func = err_lv;
1942  }
1943  if (device_type == ExecutorDeviceType::GPU && g_enable_dynamic_watchdog) {
1944  // let kernel execution finish as expected, regardless of the observed error,
1945  // unless it is from the dynamic watchdog where all threads within that block
1946  // return together.
1947  err_lv = ir_builder.CreateICmp(llvm::ICmpInst::ICMP_EQ,
1948  err_lv,
1949  cgen_state_->llInt(Executor::ERR_OUT_OF_TIME));
1950  } else {
1951  err_lv = ir_builder.CreateICmp(llvm::ICmpInst::ICMP_NE,
1952  err_lv,
1953  cgen_state_->llInt(static_cast<int32_t>(0)));
1954  }
1955  auto error_bb = llvm::BasicBlock::Create(
1956  cgen_state_->context_, ".error_exit", query_func, new_bb);
1957  const auto error_code_arg = get_arg_by_name(query_func, "error_code");
1958  llvm::CallInst::Create(
1959  cgen_state_->module_->getFunction("record_error_code"),
1960  std::vector<llvm::Value*>{err_lv_returned_from_row_func, error_code_arg},
1961  "",
1962  error_bb);
1963  llvm::ReturnInst::Create(cgen_state_->context_, error_bb);
1964  llvm::ReplaceInstWithInst(&br_instr,
1965  llvm::BranchInst::Create(error_bb, new_bb, err_lv));
1966  done_splitting = true;
1967  break;
1968  }
1969  }
1970  }
1971  CHECK(done_splitting);
1972 }
1973 
1975  llvm::Module* M = cgen_state_->module_;
1976  if (M->getFunction("allocate_varlen_buffer") == nullptr)
1977  return;
1978 
1979  // read metadata
1980  bool should_track = false;
1981  auto* flag = M->getModuleFlag("manage_memory_buffer");
1982  if (auto* cnt = llvm::mdconst::extract_or_null<llvm::ConstantInt>(flag)) {
1983  if (cnt->getZExtValue() == 1) {
1984  should_track = true;
1985  }
1986  }
1987 
1988  if (!should_track) {
1989  // metadata is not present
1990  return;
1991  }
1992 
1993  LOG(INFO) << "Found 'manage_memory_buffer' metadata.";
1994  llvm::SmallVector<llvm::CallInst*, 4> calls_to_analyze;
1995 
1996  for (llvm::Function& F : *M) {
1997  for (llvm::BasicBlock& BB : F) {
1998  for (llvm::Instruction& I : BB) {
1999  if (llvm::CallInst* CI = llvm::dyn_cast<llvm::CallInst>(&I)) {
2000  // Keep track of calls to "allocate_varlen_buffer" for later processing
2001  llvm::Function* called = CI->getCalledFunction();
2002  if (called) {
2003  if (called->getName() == "allocate_varlen_buffer") {
2004  calls_to_analyze.push_back(CI);
2005  }
2006  }
2007  }
2008  }
2009  }
2010  }
2011 
2012  // for each call to "allocate_varlen_buffer", check if there's a corresponding
2013  // call to "register_buffer_with_executor_rsm". If not, add a call to it
2014  llvm::IRBuilder<> Builder(cgen_state_->context_);
2015  auto i64 = get_int_type(64, cgen_state_->context_);
2016  auto i8p = get_int_ptr_type(8, cgen_state_->context_);
2017  auto void_ = llvm::Type::getVoidTy(cgen_state_->context_);
2018  llvm::FunctionType* fnty = llvm::FunctionType::get(void_, {i64, i8p}, false);
2019  llvm::FunctionCallee register_buffer_fn =
2020  M->getOrInsertFunction("register_buffer_with_executor_rsm", fnty, {});
2021 
2022  int64_t executor_addr = reinterpret_cast<int64_t>(this);
2023  for (llvm::CallInst* CI : calls_to_analyze) {
2024  bool found = false;
2025  // for each user of the function, check if its a callinst
2026  // and if the callinst is calling "register_buffer_with_executor_rsm"
2027  // if no such instruction exist, add one registering the buffer
2028  for (llvm::User* U : CI->users()) {
2029  if (llvm::CallInst* call = llvm::dyn_cast<llvm::CallInst>(U)) {
2030  if (call->getCalledFunction() and
2031  call->getCalledFunction()->getName() == "register_buffer_with_executor_rsm") {
2032  found = true;
2033  break;
2034  }
2035  }
2036  }
2037  if (!found) {
2038  Builder.SetInsertPoint(CI->getNextNode());
2039  Builder.CreateCall(register_buffer_fn,
2040  {ll_int(executor_addr, cgen_state_->context_), CI});
2041  }
2042  }
2043 }
2044 
2045 std::vector<llvm::Value*> Executor::inlineHoistedLiterals() {
2046  AUTOMATIC_IR_METADATA(cgen_state_.get());
2047 
2048  std::vector<llvm::Value*> hoisted_literals;
2049 
2050  // row_func_ is using literals whose defs have been hoisted up to the query_func_,
2051  // extend row_func_ signature to include extra args to pass these literal values.
2052  std::vector<llvm::Type*> row_process_arg_types;
2053 
2054  for (llvm::Function::arg_iterator I = cgen_state_->row_func_->arg_begin(),
2055  E = cgen_state_->row_func_->arg_end();
2056  I != E;
2057  ++I) {
2058  row_process_arg_types.push_back(I->getType());
2059  }
2060 
2061  for (auto& element : cgen_state_->query_func_literal_loads_) {
2062  for (auto value : element.second) {
2063  row_process_arg_types.push_back(value->getType());
2064  }
2065  }
2066 
2067  auto ft = llvm::FunctionType::get(
2068  get_int_type(32, cgen_state_->context_), row_process_arg_types, false);
2069  auto row_func_with_hoisted_literals =
2070  llvm::Function::Create(ft,
2071  llvm::Function::ExternalLinkage,
2072  "row_func_hoisted_literals",
2073  cgen_state_->row_func_->getParent());
2074 
2075  auto row_func_arg_it = row_func_with_hoisted_literals->arg_begin();
2076  for (llvm::Function::arg_iterator I = cgen_state_->row_func_->arg_begin(),
2077  E = cgen_state_->row_func_->arg_end();
2078  I != E;
2079  ++I) {
2080  if (I->hasName()) {
2081  row_func_arg_it->setName(I->getName());
2082  }
2083  ++row_func_arg_it;
2084  }
2085 
2086  decltype(row_func_with_hoisted_literals) filter_func_with_hoisted_literals{nullptr};
2087  decltype(row_func_arg_it) filter_func_arg_it{nullptr};
2088  if (cgen_state_->filter_func_) {
2089  // filter_func_ is using literals whose defs have been hoisted up to the row_func_,
2090  // extend filter_func_ signature to include extra args to pass these literal values.
2091  std::vector<llvm::Type*> filter_func_arg_types;
2092 
2093  for (llvm::Function::arg_iterator I = cgen_state_->filter_func_->arg_begin(),
2094  E = cgen_state_->filter_func_->arg_end();
2095  I != E;
2096  ++I) {
2097  filter_func_arg_types.push_back(I->getType());
2098  }
2099 
2100  for (auto& element : cgen_state_->query_func_literal_loads_) {
2101  for (auto value : element.second) {
2102  filter_func_arg_types.push_back(value->getType());
2103  }
2104  }
2105 
2106  auto ft2 = llvm::FunctionType::get(
2107  get_int_type(32, cgen_state_->context_), filter_func_arg_types, false);
2108  filter_func_with_hoisted_literals =
2109  llvm::Function::Create(ft2,
2110  llvm::Function::ExternalLinkage,
2111  "filter_func_hoisted_literals",
2112  cgen_state_->filter_func_->getParent());
2113 
2114  filter_func_arg_it = filter_func_with_hoisted_literals->arg_begin();
2115  for (llvm::Function::arg_iterator I = cgen_state_->filter_func_->arg_begin(),
2116  E = cgen_state_->filter_func_->arg_end();
2117  I != E;
2118  ++I) {
2119  if (I->hasName()) {
2120  filter_func_arg_it->setName(I->getName());
2121  }
2122  ++filter_func_arg_it;
2123  }
2124  }
2125 
2126  std::unordered_map<int, std::vector<llvm::Value*>>
2127  query_func_literal_loads_function_arguments,
2128  query_func_literal_loads_function_arguments2;
2129 
2130  for (auto& element : cgen_state_->query_func_literal_loads_) {
2131  std::vector<llvm::Value*> argument_values, argument_values2;
2132 
2133  for (auto value : element.second) {
2134  hoisted_literals.push_back(value);
2135  argument_values.push_back(&*row_func_arg_it);
2136  if (cgen_state_->filter_func_) {
2137  argument_values2.push_back(&*filter_func_arg_it);
2138  cgen_state_->filter_func_args_[&*row_func_arg_it] = &*filter_func_arg_it;
2139  }
2140  if (value->hasName()) {
2141  row_func_arg_it->setName("arg_" + value->getName());
2142  if (cgen_state_->filter_func_) {
2143  filter_func_arg_it->getContext();
2144  filter_func_arg_it->setName("arg_" + value->getName());
2145  }
2146  }
2147  ++row_func_arg_it;
2148  ++filter_func_arg_it;
2149  }
2150 
2151  query_func_literal_loads_function_arguments[element.first] = argument_values;
2152  query_func_literal_loads_function_arguments2[element.first] = argument_values2;
2153  }
2154 
2155  // copy the row_func function body over
2156  // see
2157  // https://stackoverflow.com/questions/12864106/move-function-body-avoiding-full-cloning/18751365
2158  row_func_with_hoisted_literals->getBasicBlockList().splice(
2159  row_func_with_hoisted_literals->begin(),
2160  cgen_state_->row_func_->getBasicBlockList());
2161 
2162  // also replace row_func arguments with the arguments from row_func_hoisted_literals
2163  for (llvm::Function::arg_iterator I = cgen_state_->row_func_->arg_begin(),
2164  E = cgen_state_->row_func_->arg_end(),
2165  I2 = row_func_with_hoisted_literals->arg_begin();
2166  I != E;
2167  ++I) {
2168  I->replaceAllUsesWith(&*I2);
2169  I2->takeName(&*I);
2170  cgen_state_->filter_func_args_.replace(&*I, &*I2);
2171  ++I2;
2172  }
2173 
2174  cgen_state_->row_func_ = row_func_with_hoisted_literals;
2175 
2176  // and finally replace literal placeholders
2177  std::vector<llvm::Instruction*> placeholders;
2178  std::string prefix("__placeholder__literal_");
2179  for (auto it = llvm::inst_begin(row_func_with_hoisted_literals),
2180  e = llvm::inst_end(row_func_with_hoisted_literals);
2181  it != e;
2182  ++it) {
2183  if (it->hasName() && it->getName().startswith(prefix)) {
2184  auto offset_and_index_entry =
2185  cgen_state_->row_func_hoisted_literals_.find(llvm::dyn_cast<llvm::Value>(&*it));
2186  CHECK(offset_and_index_entry != cgen_state_->row_func_hoisted_literals_.end());
2187 
2188  int lit_off = offset_and_index_entry->second.offset_in_literal_buffer;
2189  int lit_idx = offset_and_index_entry->second.index_of_literal_load;
2190 
2191  it->replaceAllUsesWith(
2192  query_func_literal_loads_function_arguments[lit_off][lit_idx]);
2193  placeholders.push_back(&*it);
2194  }
2195  }
2196  for (auto placeholder : placeholders) {
2197  placeholder->removeFromParent();
2198  }
2199 
2200  if (cgen_state_->filter_func_) {
2201  // copy the filter_func function body over
2202  // see
2203  // https://stackoverflow.com/questions/12864106/move-function-body-avoiding-full-cloning/18751365
2204  filter_func_with_hoisted_literals->getBasicBlockList().splice(
2205  filter_func_with_hoisted_literals->begin(),
2206  cgen_state_->filter_func_->getBasicBlockList());
2207 
2208  // also replace filter_func arguments with the arguments from
2209  // filter_func_hoisted_literals
2210  for (llvm::Function::arg_iterator I = cgen_state_->filter_func_->arg_begin(),
2211  E = cgen_state_->filter_func_->arg_end(),
2212  I2 = filter_func_with_hoisted_literals->arg_begin();
2213  I != E;
2214  ++I) {
2215  I->replaceAllUsesWith(&*I2);
2216  I2->takeName(&*I);
2217  ++I2;
2218  }
2219 
2220  cgen_state_->filter_func_ = filter_func_with_hoisted_literals;
2221 
2222  // and finally replace literal placeholders
2223  std::vector<llvm::Instruction*> placeholders;
2224  std::string prefix("__placeholder__literal_");
2225  for (auto it = llvm::inst_begin(filter_func_with_hoisted_literals),
2226  e = llvm::inst_end(filter_func_with_hoisted_literals);
2227  it != e;
2228  ++it) {
2229  if (it->hasName() && it->getName().startswith(prefix)) {
2230  auto offset_and_index_entry = cgen_state_->row_func_hoisted_literals_.find(
2231  llvm::dyn_cast<llvm::Value>(&*it));
2232  CHECK(offset_and_index_entry != cgen_state_->row_func_hoisted_literals_.end());
2233 
2234  int lit_off = offset_and_index_entry->second.offset_in_literal_buffer;
2235  int lit_idx = offset_and_index_entry->second.index_of_literal_load;
2236 
2237  it->replaceAllUsesWith(
2238  query_func_literal_loads_function_arguments2[lit_off][lit_idx]);
2239  placeholders.push_back(&*it);
2240  }
2241  }
2242  for (auto placeholder : placeholders) {
2243  placeholder->removeFromParent();
2244  }
2245  }
2246 
2247  return hoisted_literals;
2248 }
2249 
2250 namespace {
2251 
2252 size_t get_shared_memory_size(const bool shared_mem_used,
2253  const QueryMemoryDescriptor* query_mem_desc_ptr) {
2254  return shared_mem_used
2255  ? (query_mem_desc_ptr->getRowSize() * query_mem_desc_ptr->getEntryCount())
2256  : 0;
2257 }
2258 
2259 bool is_gpu_shared_mem_supported(const QueryMemoryDescriptor* query_mem_desc_ptr,
2260  const RelAlgExecutionUnit& ra_exe_unit,
2261  const CudaMgr_Namespace::CudaMgr* cuda_mgr,
2262  const ExecutorDeviceType device_type,
2263  const unsigned gpu_blocksize,
2264  const unsigned num_blocks_per_mp) {
2265  if (device_type == ExecutorDeviceType::CPU) {
2266  return false;
2267  }
2268  if (query_mem_desc_ptr->didOutputColumnar()) {
2269  return false;
2270  }
2271  CHECK(query_mem_desc_ptr);
2272  CHECK(cuda_mgr);
2273  /*
2274  * We only use shared memory strategy if GPU hardware provides native shared
2275  * memory atomics support. From CUDA Toolkit documentation:
2276  * https://docs.nvidia.com/cuda/pascal-tuning-guide/index.html#atomic-ops "Like
2277  * Maxwell, Pascal [and Volta] provides native shared memory atomic operations
2278  * for 32-bit integer arithmetic, along with native 32 or 64-bit compare-and-swap
2279  * (CAS)."
2280  *
2281  */
2282  if (!cuda_mgr->isArchMaxwellOrLaterForAll()) {
2283  return false;
2284  }
2285 
2286  if (query_mem_desc_ptr->getQueryDescriptionType() ==
2289  query_mem_desc_ptr->countDistinctDescriptorsLogicallyEmpty()) {
2290  // TODO: relax this, if necessary
2291  if (gpu_blocksize < query_mem_desc_ptr->getEntryCount()) {
2292  return false;
2293  }
2294  // skip shared memory usage when dealing with 1) variable length targets, 2)
2295  // not a COUNT aggregate
2296  const auto target_infos =
2297  target_exprs_to_infos(ra_exe_unit.target_exprs, *query_mem_desc_ptr);
2298  std::unordered_set<SQLAgg> supported_aggs{kCOUNT};
2299  if (std::find_if(target_infos.begin(),
2300  target_infos.end(),
2301  [&supported_aggs](const TargetInfo& ti) {
2302  if (ti.sql_type.is_varlen() ||
2303  !supported_aggs.count(ti.agg_kind)) {
2304  return true;
2305  } else {
2306  return false;
2307  }
2308  }) == target_infos.end()) {
2309  return true;
2310  }
2311  }
2312  if (query_mem_desc_ptr->getQueryDescriptionType() ==
2323  if (gpu_blocksize < query_mem_desc_ptr->getEntryCount()) {
2324  return false;
2325  }
2326 
2327  // Fundamentally, we should use shared memory whenever the output buffer
2328  // is small enough so that we can fit it in the shared memory and yet expect
2329  // good occupancy.
2330  // For now, we allow keyless, row-wise layout, and only for perfect hash
2331  // group by operations.
2332  if (query_mem_desc_ptr->hasKeylessHash() &&
2333  query_mem_desc_ptr->countDistinctDescriptorsLogicallyEmpty() &&
2334  !query_mem_desc_ptr->useStreamingTopN()) {
2335  const size_t shared_memory_threshold_bytes = std::min(
2336  g_gpu_smem_threshold == 0 ? SIZE_MAX : g_gpu_smem_threshold,
2337  cuda_mgr->getMinSharedMemoryPerBlockForAllDevices() / num_blocks_per_mp);
2338  const auto output_buffer_size =
2339  query_mem_desc_ptr->getRowSize() * query_mem_desc_ptr->getEntryCount();
2340  if (output_buffer_size > shared_memory_threshold_bytes) {
2341  return false;
2342  }
2343 
2344  // skip shared memory usage when dealing with 1) variable length targets, 2)
2345  // non-basic aggregates (COUNT, SUM, MIN, MAX, AVG)
2346  // TODO: relax this if necessary
2347  const auto target_infos =
2348  target_exprs_to_infos(ra_exe_unit.target_exprs, *query_mem_desc_ptr);
2349  std::unordered_set<SQLAgg> supported_aggs{kCOUNT};
2351  supported_aggs = {kCOUNT, kMIN, kMAX, kSUM, kAVG};
2352  }
2353  if (std::find_if(target_infos.begin(),
2354  target_infos.end(),
2355  [&supported_aggs](const TargetInfo& ti) {
2356  if (ti.sql_type.is_varlen() ||
2357  !supported_aggs.count(ti.agg_kind)) {
2358  return true;
2359  } else {
2360  return false;
2361  }
2362  }) == target_infos.end()) {
2363  return true;
2364  }
2365  }
2366  }
2367  return false;
2368 }
2369 
2370 #ifndef NDEBUG
2371 std::string serialize_llvm_metadata_footnotes(llvm::Function* query_func,
2372  CgenState* cgen_state) {
2373  std::string llvm_ir;
2374  std::unordered_set<llvm::MDNode*> md;
2375 
2376  // Loop over all instructions in the query function.
2377  for (auto bb_it = query_func->begin(); bb_it != query_func->end(); ++bb_it) {
2378  for (auto instr_it = bb_it->begin(); instr_it != bb_it->end(); ++instr_it) {
2379  llvm::SmallVector<std::pair<unsigned, llvm::MDNode*>, 100> imd;
2380  instr_it->getAllMetadata(imd);
2381  for (auto [kind, node] : imd) {
2382  md.insert(node);
2383  }
2384  }
2385  }
2386 
2387  // Loop over all instructions in the row function.
2388  for (auto bb_it = cgen_state->row_func_->begin(); bb_it != cgen_state->row_func_->end();
2389  ++bb_it) {
2390  for (auto instr_it = bb_it->begin(); instr_it != bb_it->end(); ++instr_it) {
2391  llvm::SmallVector<std::pair<unsigned, llvm::MDNode*>, 100> imd;
2392  instr_it->getAllMetadata(imd);
2393  for (auto [kind, node] : imd) {
2394  md.insert(node);
2395  }
2396  }
2397  }
2398 
2399  // Loop over all instructions in the filter function.
2400  if (cgen_state->filter_func_) {
2401  for (auto bb_it = cgen_state->filter_func_->begin();
2402  bb_it != cgen_state->filter_func_->end();
2403  ++bb_it) {
2404  for (auto instr_it = bb_it->begin(); instr_it != bb_it->end(); ++instr_it) {
2405  llvm::SmallVector<std::pair<unsigned, llvm::MDNode*>, 100> imd;
2406  instr_it->getAllMetadata(imd);
2407  for (auto [kind, node] : imd) {
2408  md.insert(node);
2409  }
2410  }
2411  }
2412  }
2413 
2414  // Sort the metadata by canonical number and convert to text.
2415  if (!md.empty()) {
2416  std::map<size_t, std::string> sorted_strings;
2417  for (auto p : md) {
2418  std::string str;
2419  llvm::raw_string_ostream os(str);
2420  p->print(os, cgen_state->module_, true);
2421  os.flush();
2422  auto fields = split(str, {}, 1);
2423  if (fields.empty() || fields[0].empty()) {
2424  continue;
2425  }
2426  sorted_strings.emplace(std::stoul(fields[0].substr(1)), str);
2427  }
2428  llvm_ir += "\n";
2429  for (auto [id, text] : sorted_strings) {
2430  llvm_ir += text;
2431  llvm_ir += "\n";
2432  }
2433  }
2434 
2435  return llvm_ir;
2436 }
2437 #endif // NDEBUG
2438 
2439 } // namespace
2440 
2441 std::tuple<CompilationResult, std::unique_ptr<QueryMemoryDescriptor>>
2442 Executor::compileWorkUnit(const std::vector<InputTableInfo>& query_infos,
2443  const PlanState::DeletedColumnsMap& deleted_cols_map,
2444  const RelAlgExecutionUnit& ra_exe_unit,
2445  const CompilationOptions& co,
2446  const ExecutionOptions& eo,
2447  const CudaMgr_Namespace::CudaMgr* cuda_mgr,
2448  const bool allow_lazy_fetch,
2449  std::shared_ptr<RowSetMemoryOwner> row_set_mem_owner,
2450  const size_t max_groups_buffer_entry_guess,
2451  const int8_t crt_min_byte_width,
2452  const bool has_cardinality_estimation,
2453  ColumnCacheMap& column_cache,
2454  RenderInfo* render_info) {
2455  auto timer = DEBUG_TIMER(__func__);
2456 
2458  const auto cuda_mgr = data_mgr_->getCudaMgr();
2459  if (!cuda_mgr) {
2460  throw QueryMustRunOnCpu();
2461  }
2462  }
2463 
2464 #ifndef NDEBUG
2465  static std::uint64_t counter = 0;
2466  ++counter;
2467  VLOG(1) << "CODEGEN #" << counter << ":";
2468  LOG(IR) << "CODEGEN #" << counter << ":";
2469  LOG(PTX) << "CODEGEN #" << counter << ":";
2470  LOG(ASM) << "CODEGEN #" << counter << ":";
2471 #endif
2472 
2473  // cgenstate_manager uses RAII pattern to manage the live time of
2474  // CgenState instances.
2475  Executor::CgenStateManager cgenstate_manager(*this,
2476  allow_lazy_fetch,
2477  query_infos,
2478  deleted_cols_map,
2479  &ra_exe_unit); // locks compilation_mutex
2480 
2481  addTransientStringLiterals(ra_exe_unit, row_set_mem_owner);
2482 
2483  GroupByAndAggregate group_by_and_aggregate(
2484  this,
2486  ra_exe_unit,
2487  query_infos,
2488  row_set_mem_owner,
2489  has_cardinality_estimation ? std::optional<int64_t>(max_groups_buffer_entry_guess)
2490  : std::nullopt);
2491  auto query_mem_desc =
2492  group_by_and_aggregate.initQueryMemoryDescriptor(eo.allow_multifrag,
2493  max_groups_buffer_entry_guess,
2494  crt_min_byte_width,
2495  render_info,
2497 
2498  if (query_mem_desc->getQueryDescriptionType() ==
2500  !has_cardinality_estimation && (!render_info || !render_info->isInSitu()) &&
2501  !eo.just_explain) {
2502  const auto col_range_info = group_by_and_aggregate.getColRangeInfo();
2503  throw CardinalityEstimationRequired(col_range_info.max - col_range_info.min);
2504  }
2505 
2506  const bool output_columnar = query_mem_desc->didOutputColumnar();
2507  const bool gpu_shared_mem_optimization =
2509  ra_exe_unit,
2510  cuda_mgr,
2511  co.device_type,
2512  cuda_mgr ? this->blockSize() : 1,
2513  cuda_mgr ? this->numBlocksPerMP() : 1);
2514  if (gpu_shared_mem_optimization) {
2515  // disable interleaved bins optimization on the GPU
2516  query_mem_desc->setHasInterleavedBinsOnGpu(false);
2517  LOG(DEBUG1) << "GPU shared memory is used for the " +
2518  query_mem_desc->queryDescTypeToString() + " query(" +
2519  std::to_string(get_shared_memory_size(gpu_shared_mem_optimization,
2520  query_mem_desc.get())) +
2521  " out of " + std::to_string(g_gpu_smem_threshold) + " bytes).";
2522  }
2523 
2524  const GpuSharedMemoryContext gpu_smem_context(
2525  get_shared_memory_size(gpu_shared_mem_optimization, query_mem_desc.get()));
2526 
2528  const size_t num_count_distinct_descs =
2529  query_mem_desc->getCountDistinctDescriptorsSize();
2530  for (size_t i = 0; i < num_count_distinct_descs; i++) {
2531  const auto& count_distinct_descriptor =
2532  query_mem_desc->getCountDistinctDescriptor(i);
2533  if (count_distinct_descriptor.impl_type_ == CountDistinctImplType::UnorderedSet ||
2534  (count_distinct_descriptor.impl_type_ != CountDistinctImplType::Invalid &&
2535  !co.hoist_literals)) {
2536  throw QueryMustRunOnCpu();
2537  }
2538  }
2539 
2540  // we currently do not support varlen projection based on baseline groupby when
2541  // 1) target table is multi-fragmented and 2) multiple gpus are involved for query
2542  // processing in this case, we punt the query to cpu to avoid server crash
2543  for (const auto expr : ra_exe_unit.target_exprs) {
2544  if (auto gby_expr = dynamic_cast<Analyzer::AggExpr*>(expr)) {
2545  bool has_multiple_gpus = cuda_mgr ? cuda_mgr->getDeviceCount() > 1 : false;
2546  if (gby_expr->get_aggtype() == SQLAgg::kSAMPLE && has_multiple_gpus &&
2547  !g_leaf_count) {
2548  std::set<const Analyzer::ColumnVar*,
2549  bool (*)(const Analyzer::ColumnVar*, const Analyzer::ColumnVar*)>
2551  gby_expr->collect_column_var(colvar_set, true);
2552  for (const auto cv : colvar_set) {
2553  if (cv->get_type_info().is_varlen()) {
2554  const auto tbl_id = cv->get_table_id();
2555  std::for_each(query_infos.begin(),
2556  query_infos.end(),
2557  [tbl_id](const InputTableInfo& input_table_info) {
2558  if (input_table_info.table_id == tbl_id &&
2559  input_table_info.info.fragments.size() > 1) {
2560  throw QueryMustRunOnCpu();
2561  }
2562  });
2563  }
2564  }
2565  }
2566  }
2567  }
2568  }
2569 
2570  // Read the module template and target either CPU or GPU
2571  // by binding the stream position functions to the right implementation:
2572  // stride access for GPU, contiguous for CPU
2573  CHECK(cgen_state_->module_ == nullptr);
2574  cgen_state_->set_module_shallow_copy(get_rt_module(), /*always_clone=*/true);
2575 
2576  auto is_gpu = co.device_type == ExecutorDeviceType::GPU;
2577  if (is_gpu) {
2578  cgen_state_->module_->setDataLayout(get_gpu_data_layout());
2579  cgen_state_->module_->setTargetTriple(get_gpu_target_triple_string());
2580  }
2581  if (has_udf_module(/*is_gpu=*/is_gpu)) {
2583  get_udf_module(/*is_gpu=*/is_gpu), *cgen_state_->module_, cgen_state_.get());
2584  }
2585  if (has_rt_udf_module(/*is_gpu=*/is_gpu)) {
2587  get_rt_udf_module(/*is_gpu=*/is_gpu), *cgen_state_->module_, cgen_state_.get());
2588  }
2589 
2590  AUTOMATIC_IR_METADATA(cgen_state_.get());
2591 
2592  auto agg_fnames =
2593  get_agg_fnames(ra_exe_unit.target_exprs, !ra_exe_unit.groupby_exprs.empty());
2594 
2595  const auto agg_slot_count = ra_exe_unit.estimator ? size_t(1) : agg_fnames.size();
2596 
2597  const bool is_group_by{query_mem_desc->isGroupBy()};
2598  auto [query_func, row_func_call] = is_group_by
2599  ? query_group_by_template(cgen_state_->module_,
2600  co.hoist_literals,
2601  *query_mem_desc,
2602  co.device_type,
2603  ra_exe_unit.scan_limit,
2604  gpu_smem_context)
2605  : query_template(cgen_state_->module_,
2606  agg_slot_count,
2607  co.hoist_literals,
2608  !!ra_exe_unit.estimator,
2609  gpu_smem_context);
2610  bind_pos_placeholders("pos_start", true, query_func, cgen_state_->module_);
2611  bind_pos_placeholders("group_buff_idx", false, query_func, cgen_state_->module_);
2612  bind_pos_placeholders("pos_step", false, query_func, cgen_state_->module_);
2613 
2614  cgen_state_->query_func_ = query_func;
2615  cgen_state_->row_func_call_ = row_func_call;
2616  cgen_state_->query_func_entry_ir_builder_.SetInsertPoint(
2617  &query_func->getEntryBlock().front());
2618 
2619  // Generate the function signature and column head fetches s.t.
2620  // double indirection isn't needed in the inner loop
2621  auto& fetch_bb = query_func->front();
2622  llvm::IRBuilder<> fetch_ir_builder(&fetch_bb);
2623  fetch_ir_builder.SetInsertPoint(&*fetch_bb.begin());
2624  auto col_heads = generate_column_heads_load(ra_exe_unit.input_col_descs.size(),
2625  query_func->args().begin(),
2626  fetch_ir_builder,
2627  cgen_state_->context_);
2628  CHECK_EQ(ra_exe_unit.input_col_descs.size(), col_heads.size());
2629 
2630  cgen_state_->row_func_ = create_row_function(ra_exe_unit.input_col_descs.size(),
2631  is_group_by ? 0 : agg_slot_count,
2632  co.hoist_literals,
2633  cgen_state_->module_,
2634  cgen_state_->context_);
2635  CHECK(cgen_state_->row_func_);
2636  cgen_state_->row_func_bb_ =
2637  llvm::BasicBlock::Create(cgen_state_->context_, "entry", cgen_state_->row_func_);
2638 
2640  auto filter_func_ft =
2641  llvm::FunctionType::get(get_int_type(32, cgen_state_->context_), {}, false);
2642  cgen_state_->filter_func_ = llvm::Function::Create(filter_func_ft,
2643  llvm::Function::ExternalLinkage,
2644  "filter_func",
2645  cgen_state_->module_);
2646  CHECK(cgen_state_->filter_func_);
2647  cgen_state_->filter_func_bb_ = llvm::BasicBlock::Create(
2648  cgen_state_->context_, "entry", cgen_state_->filter_func_);
2649  }
2650 
2651  cgen_state_->current_func_ = cgen_state_->row_func_;
2652  cgen_state_->ir_builder_.SetInsertPoint(cgen_state_->row_func_bb_);
2653 
2654  preloadFragOffsets(ra_exe_unit.input_descs, query_infos);
2655  RelAlgExecutionUnit body_execution_unit = ra_exe_unit;
2656  const auto join_loops =
2657  buildJoinLoops(body_execution_unit, co, eo, query_infos, column_cache);
2658 
2659  plan_state_->allocateLocalColumnIds(ra_exe_unit.input_col_descs);
2660  for (auto& simple_qual : ra_exe_unit.simple_quals) {
2661  plan_state_->addSimpleQual(simple_qual);
2662  }
2663  const auto is_not_deleted_bb = codegenSkipDeletedOuterTableRow(ra_exe_unit, co);
2664  if (is_not_deleted_bb) {
2665  cgen_state_->row_func_bb_ = is_not_deleted_bb;
2666  }
2667  if (!join_loops.empty()) {
2668  codegenJoinLoops(join_loops,
2669  body_execution_unit,
2670  group_by_and_aggregate,
2671  query_func,
2672  cgen_state_->row_func_bb_,
2673  *(query_mem_desc.get()),
2674  co,
2675  eo);
2676  } else {
2677  const bool can_return_error = compileBody(
2678  ra_exe_unit, group_by_and_aggregate, *query_mem_desc, co, gpu_smem_context);
2679  if (can_return_error || cgen_state_->needs_error_check_ || eo.with_dynamic_watchdog ||
2681  createErrorCheckControlFlow(query_func,
2684  join_loops,
2685  co.device_type,
2686  group_by_and_aggregate.query_infos_);
2687  }
2688  }
2689  std::vector<llvm::Value*> hoisted_literals;
2690 
2691  if (co.hoist_literals) {
2692  VLOG(1) << "number of hoisted literals: "
2693  << cgen_state_->query_func_literal_loads_.size()
2694  << " / literal buffer usage: " << cgen_state_->getLiteralBufferUsage(0)
2695  << " bytes";
2696  }
2697 
2698  if (co.hoist_literals && !cgen_state_->query_func_literal_loads_.empty()) {
2699  // we have some hoisted literals...
2700  hoisted_literals = inlineHoistedLiterals();
2701  }
2702 
2703  // replace the row func placeholder call with the call to the actual row func
2704  std::vector<llvm::Value*> row_func_args;
2705  for (size_t i = 0; i < cgen_state_->row_func_call_->getNumOperands() - 1; ++i) {
2706  row_func_args.push_back(cgen_state_->row_func_call_->getArgOperand(i));
2707  }
2708  row_func_args.insert(row_func_args.end(), col_heads.begin(), col_heads.end());
2709  row_func_args.push_back(get_arg_by_name(query_func, "join_hash_tables"));
2710  // push hoisted literals arguments, if any
2711  row_func_args.insert(
2712  row_func_args.end(), hoisted_literals.begin(), hoisted_literals.end());
2713  llvm::ReplaceInstWithInst(
2714  cgen_state_->row_func_call_,
2715  llvm::CallInst::Create(cgen_state_->row_func_, row_func_args, ""));
2716 
2717  // replace the filter func placeholder call with the call to the actual filter func
2718  if (cgen_state_->filter_func_) {
2719  std::vector<llvm::Value*> filter_func_args;
2720  for (auto arg_it = cgen_state_->filter_func_args_.begin();
2721  arg_it != cgen_state_->filter_func_args_.end();
2722  ++arg_it) {
2723  filter_func_args.push_back(arg_it->first);
2724  }
2725  llvm::ReplaceInstWithInst(
2726  cgen_state_->filter_func_call_,
2727  llvm::CallInst::Create(cgen_state_->filter_func_, filter_func_args, ""));
2728  }
2729 
2730  // Aggregate
2731  plan_state_->init_agg_vals_ =
2732  init_agg_val_vec(ra_exe_unit.target_exprs, ra_exe_unit.quals, *query_mem_desc);
2733 
2734  /*
2735  * If we have decided to use GPU shared memory (decision is not made here), then
2736  * we generate proper code for extra components that it needs (buffer initialization and
2737  * gpu reduction from shared memory to global memory). We then replace these functions
2738  * into the already compiled query_func (replacing two placeholders, write_back_nop and
2739  * init_smem_nop). The rest of the code should be as before (row_func, etc.).
2740  */
2741  if (gpu_smem_context.isSharedMemoryUsed()) {
2742  if (query_mem_desc->getQueryDescriptionType() ==
2744  GpuSharedMemCodeBuilder gpu_smem_code(
2745  cgen_state_->module_,
2746  cgen_state_->context_,
2747  *query_mem_desc,
2749  plan_state_->init_agg_vals_,
2750  executor_id_);
2751  gpu_smem_code.codegen();
2752  gpu_smem_code.injectFunctionsInto(query_func);
2753 
2754  // helper functions are used for caching purposes later
2755  cgen_state_->helper_functions_.push_back(gpu_smem_code.getReductionFunction());
2756  cgen_state_->helper_functions_.push_back(gpu_smem_code.getInitFunction());
2757  LOG(IR) << gpu_smem_code.toString();
2758  }
2759  }
2760 
2761  auto multifrag_query_func = cgen_state_->module_->getFunction(
2762  "multifrag_query" + std::string(co.hoist_literals ? "_hoisted_literals" : ""));
2763  CHECK(multifrag_query_func);
2764 
2766  insertErrorCodeChecker(
2767  multifrag_query_func, co.hoist_literals, eo.allow_runtime_query_interrupt);
2768  }
2769 
2770  bind_query(query_func,
2771  "query_stub" + std::string(co.hoist_literals ? "_hoisted_literals" : ""),
2772  multifrag_query_func,
2773  cgen_state_->module_);
2774 
2775  std::vector<llvm::Function*> root_funcs{query_func, cgen_state_->row_func_};
2776  if (cgen_state_->filter_func_) {
2777  root_funcs.push_back(cgen_state_->filter_func_);
2778  }
2779  auto live_funcs = CodeGenerator::markDeadRuntimeFuncs(
2780  *cgen_state_->module_, root_funcs, {multifrag_query_func});
2781 
2782  // Always inline the row function and the filter function.
2783  // We don't want register spills in the inner loops.
2784  // LLVM seems to correctly free up alloca instructions
2785  // in these functions even when they are inlined.
2786  mark_function_always_inline(cgen_state_->row_func_);
2787  if (cgen_state_->filter_func_) {
2788  mark_function_always_inline(cgen_state_->filter_func_);
2789  }
2790 
2791 #ifndef NDEBUG
2792  // Add helpful metadata to the LLVM IR for debugging.
2794 #endif
2795 
2796  // Serialize the important LLVM IR functions to text for SQL EXPLAIN.
2797  std::string llvm_ir;
2798  if (eo.just_explain) {
2800 #ifdef WITH_JIT_DEBUG
2801  throw std::runtime_error(
2802  "Explain optimized not available when JIT runtime debug symbols are enabled");
2803 #else
2804  // Note that we don't run the NVVM reflect pass here. Use LOG(IR) to get the
2805  // optimized IR after NVVM reflect
2806  llvm::legacy::PassManager pass_manager;
2807  optimize_ir(query_func,
2808  cgen_state_->module_,
2809  pass_manager,
2810  live_funcs,
2811  gpu_smem_context.isSharedMemoryUsed(),
2812  co);
2813 #endif // WITH_JIT_DEBUG
2814  }
2815  llvm_ir =
2816  serialize_llvm_object(multifrag_query_func) + serialize_llvm_object(query_func) +
2817  serialize_llvm_object(cgen_state_->row_func_) +
2818  (cgen_state_->filter_func_ ? serialize_llvm_object(cgen_state_->filter_func_)
2819  : "");
2820 
2821 #ifndef NDEBUG
2822  llvm_ir += serialize_llvm_metadata_footnotes(query_func, cgen_state_.get());
2823 #endif
2824  }
2825 
2826  LOG(IR) << "\n\n" << query_mem_desc->toString() << "\n";
2827  LOG(IR) << "IR for the "
2828  << (co.device_type == ExecutorDeviceType::CPU ? "CPU:\n" : "GPU:\n");
2829 #ifdef NDEBUG
2830  LOG(IR) << serialize_llvm_object(query_func)
2831  << serialize_llvm_object(cgen_state_->row_func_)
2832  << (cgen_state_->filter_func_ ? serialize_llvm_object(cgen_state_->filter_func_)
2833  : "")
2834  << "\nEnd of IR";
2835 #else
2836  LOG(IR) << serialize_llvm_object(cgen_state_->module_) << "\nEnd of IR";
2837 #endif
2838 
2839  // Insert calls to "register_buffer_with_executor_rsm" for allocations
2840  // in runtime functions (i.e. from RBC) without it
2841  AutoTrackBuffersInRuntimeIR();
2842 
2843  // Run some basic validation checks on the LLVM IR before code is generated below.
2844  verify_function_ir(cgen_state_->row_func_);
2845  if (cgen_state_->filter_func_) {
2846  verify_function_ir(cgen_state_->filter_func_);
2847  }
2848 
2849  // Generate final native code from the LLVM IR.
2850  return std::make_tuple(
2853  ? optimizeAndCodegenCPU(query_func, multifrag_query_func, live_funcs, co)
2854  : optimizeAndCodegenGPU(query_func,
2855  multifrag_query_func,
2856  live_funcs,
2857  is_group_by || ra_exe_unit.estimator,
2858  cuda_mgr,
2859  gpu_smem_context.isSharedMemoryUsed(),
2860  co),
2861  cgen_state_->getLiterals(),
2862  output_columnar,
2863  llvm_ir,
2864  std::move(gpu_smem_context)},
2865  std::move(query_mem_desc));
2866 }
2867 
2868 void Executor::insertErrorCodeChecker(llvm::Function* query_func,
2869  bool hoist_literals,
2870  bool allow_runtime_query_interrupt) {
2871  auto query_stub_func_name =
2872  "query_stub" + std::string(hoist_literals ? "_hoisted_literals" : "");
2873  for (auto bb_it = query_func->begin(); bb_it != query_func->end(); ++bb_it) {
2874  for (auto inst_it = bb_it->begin(); inst_it != bb_it->end(); ++inst_it) {
2875  if (!llvm::isa<llvm::CallInst>(*inst_it)) {
2876  continue;
2877  }
2878  auto& row_func_call = llvm::cast<llvm::CallInst>(*inst_it);
2879  if (std::string(row_func_call.getCalledFunction()->getName()) ==
2880  query_stub_func_name) {
2881  auto next_inst_it = inst_it;
2882  ++next_inst_it;
2883  auto new_bb = bb_it->splitBasicBlock(next_inst_it);
2884  auto& br_instr = bb_it->back();
2885  llvm::IRBuilder<> ir_builder(&br_instr);
2886  llvm::Value* err_lv = &*inst_it;
2887  auto error_check_bb =
2888  bb_it->splitBasicBlock(llvm::BasicBlock::iterator(br_instr), ".error_check");
2889  llvm::Value* error_code_arg = nullptr;
2890  auto arg_cnt = 0;
2891  for (auto arg_it = query_func->arg_begin(); arg_it != query_func->arg_end();
2892  arg_it++, ++arg_cnt) {
2893  // since multi_frag_* func has anonymous arguments so we use arg_offset
2894  // explicitly to capture "error_code" argument in the func's argument list
2895  if (hoist_literals) {
2896  if (arg_cnt == 9) {
2897  error_code_arg = &*arg_it;
2898  break;
2899  }
2900  } else {
2901  if (arg_cnt == 8) {
2902  error_code_arg = &*arg_it;
2903  break;
2904  }
2905  }
2906  }
2907  CHECK(error_code_arg);
2908  llvm::Value* err_code = nullptr;
2909  if (allow_runtime_query_interrupt) {
2910  // decide the final error code with a consideration of interrupt status
2911  auto& check_interrupt_br_instr = bb_it->back();
2912  auto interrupt_check_bb = llvm::BasicBlock::Create(
2913  cgen_state_->context_, ".interrupt_check", query_func, error_check_bb);
2914  llvm::IRBuilder<> interrupt_checker_ir_builder(interrupt_check_bb);
2915  auto detected_interrupt = interrupt_checker_ir_builder.CreateCall(
2916  cgen_state_->module_->getFunction("check_interrupt"), {});
2917  auto detected_error = interrupt_checker_ir_builder.CreateCall(
2918  cgen_state_->module_->getFunction("get_error_code"),
2919  std::vector<llvm::Value*>{error_code_arg});
2920  err_code = interrupt_checker_ir_builder.CreateSelect(
2921  detected_interrupt,
2922  cgen_state_->llInt(Executor::ERR_INTERRUPTED),
2923  detected_error);
2924  interrupt_checker_ir_builder.CreateBr(error_check_bb);
2925  llvm::ReplaceInstWithInst(&check_interrupt_br_instr,
2926  llvm::BranchInst::Create(interrupt_check_bb));
2927  ir_builder.SetInsertPoint(&br_instr);
2928  } else {
2929  // uses error code returned from row_func and skip to check interrupt status
2930  ir_builder.SetInsertPoint(&br_instr);
2931  err_code =
2932  ir_builder.CreateCall(cgen_state_->module_->getFunction("get_error_code"),
2933  std::vector<llvm::Value*>{error_code_arg});
2934  }
2935  err_lv = ir_builder.CreateICmp(
2936  llvm::ICmpInst::ICMP_NE, err_code, cgen_state_->llInt(0));
2937  auto error_bb = llvm::BasicBlock::Create(
2938  cgen_state_->context_, ".error_exit", query_func, new_bb);
2939  llvm::CallInst::Create(cgen_state_->module_->getFunction("record_error_code"),
2940  std::vector<llvm::Value*>{err_code, error_code_arg},
2941  "",
2942  error_bb);
2943  llvm::ReturnInst::Create(cgen_state_->context_, error_bb);
2944  llvm::ReplaceInstWithInst(&br_instr,
2945  llvm::BranchInst::Create(error_bb, new_bb, err_lv));
2946  break;
2947  }
2948  }
2949  }
2950 }
2951 
2953  const RelAlgExecutionUnit& ra_exe_unit,
2954  const CompilationOptions& co) {
2955  AUTOMATIC_IR_METADATA(cgen_state_.get());
2956  if (!co.filter_on_deleted_column) {
2957  return nullptr;
2958  }
2959  CHECK(!ra_exe_unit.input_descs.empty());
2960  const auto& outer_input_desc = ra_exe_unit.input_descs[0];
2961  if (outer_input_desc.getSourceType() != InputSourceType::TABLE) {
2962  return nullptr;
2963  }
2964  const auto deleted_cd =
2965  plan_state_->getDeletedColForTable(outer_input_desc.getTableId());
2966  if (!deleted_cd) {
2967  return nullptr;
2968  }
2969  CHECK(deleted_cd->columnType.is_boolean());
2970  const auto deleted_expr =
2971  makeExpr<Analyzer::ColumnVar>(deleted_cd->columnType,
2972  outer_input_desc.getTableId(),
2973  deleted_cd->columnId,
2974  outer_input_desc.getNestLevel());
2975  CodeGenerator code_generator(this);
2976  const auto is_deleted =
2977  code_generator.toBool(code_generator.codegen(deleted_expr.get(), true, co).front());
2978  const auto is_deleted_bb = llvm::BasicBlock::Create(
2979  cgen_state_->context_, "is_deleted", cgen_state_->row_func_);
2980  llvm::BasicBlock* bb = llvm::BasicBlock::Create(
2981  cgen_state_->context_, "is_not_deleted", cgen_state_->row_func_);
2982  cgen_state_->ir_builder_.CreateCondBr(is_deleted, is_deleted_bb, bb);
2983  cgen_state_->ir_builder_.SetInsertPoint(is_deleted_bb);
2984  cgen_state_->ir_builder_.CreateRet(cgen_state_->llInt<int32_t>(0));
2985  cgen_state_->ir_builder_.SetInsertPoint(bb);
2986  return bb;
2987 }
2988 
2989 bool Executor::compileBody(const RelAlgExecutionUnit& ra_exe_unit,
2990  GroupByAndAggregate& group_by_and_aggregate,
2992  const CompilationOptions& co,
2993  const GpuSharedMemoryContext& gpu_smem_context) {
2994  AUTOMATIC_IR_METADATA(cgen_state_.get());
2995 
2996  // Switch the code generation into a separate filter function if enabled.
2997  // Note that accesses to function arguments are still codegenned from the
2998  // row function's arguments, then later automatically forwarded and
2999  // remapped into filter function arguments by redeclareFilterFunction().
3000  cgen_state_->row_func_bb_ = cgen_state_->ir_builder_.GetInsertBlock();
3001  llvm::Value* loop_done{nullptr};
3002  std::unique_ptr<Executor::FetchCacheAnchor> fetch_cache_anchor;
3003  if (cgen_state_->filter_func_) {
3004  if (cgen_state_->row_func_bb_->getName() == "loop_body") {
3005  auto row_func_entry_bb = &cgen_state_->row_func_->getEntryBlock();
3006  cgen_state_->ir_builder_.SetInsertPoint(row_func_entry_bb,
3007  row_func_entry_bb->begin());
3008  loop_done = cgen_state_->ir_builder_.CreateAlloca(
3009  get_int_type(1, cgen_state_->context_), nullptr, "loop_done");
3010  cgen_state_->ir_builder_.SetInsertPoint(cgen_state_->row_func_bb_);
3011  cgen_state_->ir_builder_.CreateStore(cgen_state_->llBool(true), loop_done);
3012  }
3013  cgen_state_->ir_builder_.SetInsertPoint(cgen_state_->filter_func_bb_);
3014  cgen_state_->current_func_ = cgen_state_->filter_func_;
3015  fetch_cache_anchor = std::make_unique<Executor::FetchCacheAnchor>(cgen_state_.get());
3016  }
3017 
3018  // generate the code for the filter
3019  std::vector<Analyzer::Expr*> primary_quals;
3020  std::vector<Analyzer::Expr*> deferred_quals;
3021  bool short_circuited = CodeGenerator::prioritizeQuals(
3022  ra_exe_unit, primary_quals, deferred_quals, plan_state_->hoisted_filters_);
3023  if (short_circuited) {
3024  VLOG(1) << "Prioritized " << std::to_string(primary_quals.size()) << " quals, "
3025  << "short-circuited and deferred " << std::to_string(deferred_quals.size())
3026  << " quals";
3027  }
3028  llvm::Value* filter_lv = cgen_state_->llBool(true);
3029  CodeGenerator code_generator(this);
3030  for (auto expr : primary_quals) {
3031  // Generate the filter for primary quals
3032  auto cond = code_generator.toBool(code_generator.codegen(expr, true, co).front());
3033  filter_lv = cgen_state_->ir_builder_.CreateAnd(filter_lv, cond);
3034  }
3035  CHECK(filter_lv->getType()->isIntegerTy(1));
3036  llvm::BasicBlock* sc_false{nullptr};
3037  if (!deferred_quals.empty()) {
3038  auto sc_true = llvm::BasicBlock::Create(
3039  cgen_state_->context_, "sc_true", cgen_state_->current_func_);
3040  sc_false = llvm::BasicBlock::Create(
3041  cgen_state_->context_, "sc_false", cgen_state_->current_func_);
3042  cgen_state_->ir_builder_.CreateCondBr(filter_lv, sc_true, sc_false);
3043  cgen_state_->ir_builder_.SetInsertPoint(sc_false);
3044  if (ra_exe_unit.join_quals.empty()) {
3045  cgen_state_->ir_builder_.CreateRet(cgen_state_->llInt(int32_t(0)));
3046  }
3047  cgen_state_->ir_builder_.SetInsertPoint(sc_true);
3048  filter_lv = cgen_state_->llBool(true);
3049  }
3050  for (auto expr : deferred_quals) {
3051  filter_lv = cgen_state_->ir_builder_.CreateAnd(
3052  filter_lv, code_generator.toBool(code_generator.codegen(expr, true, co).front()));
3053  }
3054 
3055  CHECK(filter_lv->getType()->isIntegerTy(1));
3056  auto ret = group_by_and_aggregate.codegen(
3057  filter_lv, sc_false, query_mem_desc, co, gpu_smem_context);
3058 
3059  // Switch the code generation back to the row function if a filter
3060  // function was enabled.
3061  if (cgen_state_->filter_func_) {
3062  if (cgen_state_->row_func_bb_->getName() == "loop_body") {
3063  cgen_state_->ir_builder_.CreateStore(cgen_state_->llBool(false), loop_done);
3064  cgen_state_->ir_builder_.CreateRet(cgen_state_->llInt<int32_t>(0));
3065  }
3066 
3067  cgen_state_->ir_builder_.SetInsertPoint(cgen_state_->row_func_bb_);
3068  cgen_state_->current_func_ = cgen_state_->row_func_;
3069  cgen_state_->filter_func_call_ =
3070  cgen_state_->ir_builder_.CreateCall(cgen_state_->filter_func_, {});
3071 
3072  // Create real filter function declaration after placeholder call
3073  // is emitted.
3074  redeclareFilterFunction();
3075 
3076  if (cgen_state_->row_func_bb_->getName() == "loop_body") {
3077  auto loop_done_true = llvm::BasicBlock::Create(
3078  cgen_state_->context_, "loop_done_true", cgen_state_->row_func_);
3079  auto loop_done_false = llvm::BasicBlock::Create(
3080  cgen_state_->context_, "loop_done_false", cgen_state_->row_func_);
3081  auto loop_done_flag = cgen_state_->ir_builder_.CreateLoad(
3082  loop_done->getType()->getPointerElementType(), loop_done);
3083  cgen_state_->ir_builder_.CreateCondBr(
3084  loop_done_flag, loop_done_true, loop_done_false);
3085  cgen_state_->ir_builder_.SetInsertPoint(loop_done_true);
3086  cgen_state_->ir_builder_.CreateRet(cgen_state_->filter_func_call_);
3087  cgen_state_->ir_builder_.SetInsertPoint(loop_done_false);
3088  } else {
3089  cgen_state_->ir_builder_.CreateRet(cgen_state_->filter_func_call_);
3090  }
3091  }
3092  return ret;
3093 }
3095 std::vector<llvm::Value*> generate_column_heads_load(const int num_columns,
3096  llvm::Value* byte_stream_arg,
3097  llvm::IRBuilder<>& ir_builder,
3098  llvm::LLVMContext& ctx) {
3099  CHECK(byte_stream_arg);
3100  const auto max_col_local_id = num_columns - 1;
3101 
3102  std::vector<llvm::Value*> col_heads;
3103  for (int col_id = 0; col_id <= max_col_local_id; ++col_id) {
3104  auto* gep = ir_builder.CreateGEP(
3105  byte_stream_arg->getType()->getScalarType()->getPointerElementType(),
3106  byte_stream_arg,
3107  llvm::ConstantInt::get(llvm::Type::getInt32Ty(ctx), col_id));
3108  col_heads.emplace_back(
3109  ir_builder.CreateLoad(gep->getType()->getPointerElementType(), gep));
3110  }
3111  return col_heads;
3112 }
3113 
void createErrorCheckControlFlow(llvm::Function *query_func, bool run_with_dynamic_watchdog, bool run_with_allowing_runtime_interrupt, const std::vector< JoinLoop > &join_loops, ExecutorDeviceType device_type, const std::vector< InputTableInfo > &input_table_infos)
std::vector< Analyzer::Expr * > target_exprs
#define CHECK_EQ(x, y)
Definition: Logger.h:230
double g_running_query_interrupt_freq
Definition: Execute.cpp:129
llvm::Value * find_variable_in_basic_block(llvm::Function *func, std::string bb_name, std::string variable_name)
bool g_enable_smem_group_by
bool countDistinctDescriptorsLogicallyEmpty() const
static const int32_t ERR_INTERRUPTED
Definition: Execute.h:1378
static bool colvar_comp(const ColumnVar *l, const ColumnVar *r)
Definition: Analyzer.h:216
void mark_function_never_inline(llvm::Function *func)
bool codegen(llvm::Value *filter_result, llvm::BasicBlock *sc_false, QueryMemoryDescriptor &query_mem_desc, const CompilationOptions &co, const GpuSharedMemoryContext &gpu_smem_context)
void collect_column_var(std::set< const ColumnVar *, bool(*)(const ColumnVar *, const ColumnVar *)> &colvar_set, bool include_agg) const override
Definition: Analyzer.h:221
ExecutorDeviceType
void optimize_ir(llvm::Function *query_func, llvm::Module *llvm_module, llvm::legacy::PassManager &pass_manager, const std::unordered_set< llvm::Function * > &live_funcs, const bool is_gpu_smem_used, const CompilationOptions &co)
Streaming Top N algorithm.
#define LOG(tag)
Definition: Logger.h:216
void eliminate_dead_self_recursive_funcs(llvm::Module &M, const std::unordered_set< llvm::Function * > &live_funcs)
void AutoTrackBuffersInRuntimeIR()
void checkCudaErrors(CUresult err)
Definition: sample.cpp:38
void mark_function_always_inline(llvm::Function *func)
llvm::StringRef get_gpu_data_layout()
llvm::ConstantInt * ll_int(const T v, llvm::LLVMContext &context)
std::string assemblyForCPU(ExecutionEngineWrapper &execution_engine, llvm::Module *llvm_module)
std::string join(T const &container, std::string const &delim)
std::vector< InputDescriptor > input_descs
std::string serialize_llvm_metadata_footnotes(llvm::Function *query_func, CgenState *cgen_state)
std::unique_ptr< llvm::Module > read_llvm_module_from_ir_string(const std::string &udf_ir_string, llvm::LLVMContext &ctx, bool is_gpu=false)
std::tuple< llvm::Function *, llvm::CallInst * > query_template(llvm::Module *mod, const size_t aggr_col_count, const bool hoist_literals, const bool is_estimate_query, const GpuSharedMemoryContext &gpu_smem_context)
std::vector< std::string > CodeCacheKey
Definition: CodeCache.h:25
ExecutorOptLevel opt_level
bool g_enable_dynamic_watchdog
Definition: Execute.cpp:80
static ExecutionEngineWrapper generateNativeCPUCode(llvm::Function *func, const std::unordered_set< llvm::Function * > &live_funcs, const CompilationOptions &co)
const std::list< std::shared_ptr< Analyzer::Expr > > groupby_exprs
llvm::Type * get_int_type(const int width, llvm::LLVMContext &context)
static std::string generatePTX(const std::string &cuda_llir, llvm::TargetMachine *nvptx_target_machine, llvm::LLVMContext &context)
#define CHECK_GT(x, y)
Definition: Logger.h:234
ExecutionEngineWrapper & operator=(const ExecutionEngineWrapper &other)=delete
std::tuple< llvm::Function *, llvm::CallInst * > query_group_by_template(llvm::Module *mod, const bool hoist_literals, const QueryMemoryDescriptor &query_mem_desc, const ExecutorDeviceType device_type, const bool check_scan_limit, const GpuSharedMemoryContext &gpu_smem_context)
constexpr double f
Definition: Utm.h:31
std::vector< std::string > get_agg_fnames(const std::vector< Analyzer::Expr * > &target_exprs, const bool is_group_by)
std::string to_string(char const *&&v)
std::vector< std::string > split(std::string_view str, std::string_view delim, std::optional< size_t > maxsplit)
split apart a string into a vector of substrings
void throw_parseIR_error(const llvm::SMDiagnostic &parse_error, std::string src="", const bool is_gpu=false)
llvm::Function * row_func_
Definition: CgenState.h:431
bool g_enable_smem_non_grouped_agg
Definition: Execute.cpp:138
Definition: sqldefs.h:74
std::shared_lock< T > shared_lock
unsigned getExpOfTwo(unsigned n)
Definition: MathUtils.cpp:23
bool is_gpu_shared_mem_supported(const QueryMemoryDescriptor *query_mem_desc_ptr, const RelAlgExecutionUnit &ra_exe_unit, const CudaMgr_Namespace::CudaMgr *cuda_mgr, const ExecutorDeviceType device_type, const unsigned gpu_blocksize, const unsigned num_blocks_per_mp)
llvm::StringRef get_gpu_target_triple_string()
llvm::Module * module_
Definition: CgenState.h:430
Supported runtime functions management and retrieval.
std::tuple< CompilationResult, std::unique_ptr< QueryMemoryDescriptor > > compileWorkUnit(const std::vector< InputTableInfo > &query_infos, const PlanState::DeletedColumnsMap &deleted_cols_map, const RelAlgExecutionUnit &ra_exe_unit, const CompilationOptions &co, const ExecutionOptions &eo, const CudaMgr_Namespace::CudaMgr *cuda_mgr, const bool allow_lazy_fetch, std::shared_ptr< RowSetMemoryOwner >, const size_t max_groups_buffer_entry_count, const int8_t crt_min_byte_width, const bool has_cardinality_estimation, ColumnCacheMap &column_cache, RenderInfo *render_info=nullptr)
void scan_function_calls(llvm::Function &F, std::unordered_set< std::string > &defined, std::unordered_set< std::string > &undefined, const std::unordered_set< std::string > &ignored)
void verify_function_ir(const llvm::Function *func)
bool compileBody(const RelAlgExecutionUnit &ra_exe_unit, GroupByAndAggregate &group_by_and_aggregate, QueryMemoryDescriptor &query_mem_desc, const CompilationOptions &co, const GpuSharedMemoryContext &gpu_smem_context={})
llvm::Value * get_arg_by_name(llvm::Function *func, const std::string &name)
Definition: Execute.h:166
static std::unordered_set< llvm::Function * > markDeadRuntimeFuncs(llvm::Module &module, const std::vector< llvm::Function * > &roots, const std::vector< llvm::Function * > &leaves)
std::string generatePTX(const std::string &) const
ExecutionEngineWrapper create_execution_engine(llvm::Module *llvm_module, llvm::EngineBuilder &eb, const CompilationOptions &co)
std::unique_ptr< llvm::JITEventListener > intel_jit_listener_
std::unordered_map< TableId, const ColumnDescriptor * > DeletedColumnsMap
Definition: PlanState.h:44
const JoinQualsPerNestingLevel join_quals
std::unique_ptr< llvm::Module > read_llvm_module_from_ir_file(const std::string &udf_ir_filename, llvm::LLVMContext &ctx, bool is_gpu=false)
ExecutorExplainType explain_type
std::shared_ptr< CompilationContext > optimizeAndCodegenCPU(llvm::Function *, llvm::Function *, const std::unordered_set< llvm::Function * > &, const CompilationOptions &)
void insertErrorCodeChecker(llvm::Function *query_func, bool hoist_literals, bool allow_runtime_query_interrupt)
static const int32_t ERR_OUT_OF_TIME
Definition: Execute.h:1377
void initializeNVPTXBackend() const
Definition: sqldefs.h:76
size_t getMinSharedMemoryPerBlockForAllDevices() const
Definition: CudaMgr.h:122
static void link_udf_module(const std::unique_ptr< llvm::Module > &udf_module, llvm::Module &module, CgenState *cgen_state, llvm::Linker::Flags flags=llvm::Linker::Flags::None)
const std::shared_ptr< Analyzer::Estimator > estimator
#define AUTOMATIC_IR_METADATA(CGENSTATE)
this
Definition: Execute.cpp:253
QueryDescriptionType getQueryDescriptionType() const
std::map< std::string, std::string > get_device_parameters(bool cpu_only)
static std::string deviceArchToSM(const NvidiaDeviceArch arch)
Definition: CudaMgr.h:156
#define AUTOMATIC_IR_METADATA_DONE()
llvm::Function * create_row_function(const size_t in_col_count, const size_t agg_col_count, const bool hoist_literals, llvm::Module *llvm_module, llvm::LLVMContext &context)
ExecutorDeviceType device_type
void bind_pos_placeholders(const std::string &pos_fn_name, const bool use_resume_param, llvm::Function *query_func, llvm::Module *llvm_module)
std::unordered_map< int, std::unordered_map< int, std::shared_ptr< const ColumnarResults >>> ColumnCacheMap
llvm::Function * filter_func_
Definition: CgenState.h:432
std::unique_ptr< llvm::ExecutionEngine > execution_engine_
static void addUdfIrToModule(const std::string &udf_ir_filename, const bool is_cuda_ir)
bool isArchMaxwellOrLaterForAll() const
Definition: CudaMgr.cpp:331
llvm::BasicBlock * codegenSkipDeletedOuterTableRow(const RelAlgExecutionUnit &ra_exe_unit, const CompilationOptions &co)
void bind_query(llvm::Function *query_func, const std::string &query_fname, llvm::Function *multifrag_query_func, llvm::Module *llvm_module)
#define CHECK_LE(x, y)
Definition: Logger.h:233
void set_row_func_argnames(llvm::Function *row_func, const size_t in_col_count, const size_t agg_col_count, const bool hoist_literals)
std::string cpp_to_llvm_name(const std::string &s)
std::string serialize_llvm_object(const T *llvm_obj)
void clear_function_attributes(llvm::Function *func)
std::shared_ptr< CompilationContext > optimizeAndCodegenGPU(llvm::Function *, llvm::Function *, std::unordered_set< llvm::Function * > &, const bool no_inline, const CudaMgr_Namespace::CudaMgr *cuda_mgr, const bool is_gpu_smem_used, const CompilationOptions &)
static std::shared_ptr< GpuCompilationContext > generateNativeGPUCode(Executor *executor, llvm::Function *func, llvm::Function *wrapper_func, const std::unordered_set< llvm::Function * > &live_funcs, const bool is_gpu_smem_used, const CompilationOptions &co, const GPUTarget &gpu_target)
bool g_enable_smem_grouped_non_count_agg
Definition: Execute.cpp:135
Definition: sqldefs.h:77
static bool alwaysCloneRuntimeFunction(const llvm::Function *func)
std::vector< llvm::Value * > generate_column_heads_load(const int num_columns, llvm::Value *byte_stream_arg, llvm::IRBuilder<> &ir_builder, llvm::LLVMContext &ctx)
static std::map< ExtModuleKinds, std::string > extension_module_sources
Definition: Execute.h:479
void show_defined(llvm::Module &llvm_module)
int CUdevice
Definition: nocuda.h:20
bool g_enable_filter_function
Definition: Execute.cpp:84
static void linkModuleWithLibdevice(Executor *executor, llvm::Module &module, llvm::PassManagerBuilder &pass_manager_builder, const GPUTarget &gpu_target)
float g_fraction_code_cache_to_evict
static bool prioritizeQuals(const RelAlgExecutionUnit &ra_exe_unit, std::vector< Analyzer::Expr * > &primary_quals, std::vector< Analyzer::Expr * > &deferred_quals, const PlanState::HoistedFiltersSet &hoisted_quals)
Definition: LogicalIR.cpp:157
data_mgr_(data_mgr)
SQLAgg get_aggtype() const
Definition: Analyzer.h:1202
std::list< std::shared_ptr< Analyzer::Expr > > quals
#define CHECK(condition)
Definition: Logger.h:222
#define DEBUG_TIMER(name)
Definition: Logger.h:371
llvm::ValueToValueMapTy vmap_
Definition: CgenState.h:440
std::vector< llvm::Value * > inlineHoistedLiterals()
static std::shared_ptr< QueryEngine > getInstance()
Definition: QueryEngine.h:81
std::vector< TargetInfo > target_exprs_to_infos(const std::vector< Analyzer::Expr * > &targets, const QueryMemoryDescriptor &query_mem_desc)
std::list< std::shared_ptr< const InputColDescriptor > > input_col_descs
size_t g_leaf_count
Definition: ParserNode.cpp:76
Definition: sqldefs.h:75
static llvm::sys::Mutex g_ee_create_mutex
int cpu_threads()
Definition: thread_count.h:25
static std::vector< std::string > getLLVMDeclarations(const std::unordered_set< std::string > &udf_decls, const bool is_gpu=false)
std::vector< int64_t > init_agg_val_vec(const std::vector< TargetInfo > &targets, const QueryMemoryDescriptor &query_mem_desc)
CubinResult ptx_to_cubin(const std::string &ptx, const unsigned block_size, const CudaMgr_Namespace::CudaMgr *cuda_mgr)
Definition: sqldefs.h:73
llvm::Type * get_int_ptr_type(const int width, llvm::LLVMContext &context)
#define VLOG(n)
Definition: Logger.h:316
size_t get_shared_memory_size(const bool shared_mem_used, const QueryMemoryDescriptor *query_mem_desc_ptr)
std::list< std::shared_ptr< Analyzer::Expr > > simple_quals
std::unique_ptr< llvm::Module > read_llvm_module_from_bc_file(const std::string &udf_ir_filename, llvm::LLVMContext &ctx)
static std::unique_ptr< llvm::TargetMachine > initializeNVPTXBackend(const CudaMgr_Namespace::NvidiaDeviceArch arch)
size_t g_gpu_smem_threshold
Definition: Execute.cpp:130