First phase of udf array support

Handle gpu calling conventions for array udfs

Add array unit tests

Add support for GPU array abstraction calling convention

clang-format formatting

Support all array types

Inital code to deal with mangled struct names for array udfs

Merge with runtime udf work and use mangled names for udf arrays, add more unit tests

Add some assert checking for null pointers

If the udf gpu code is compiled for cpu, throw exception to force execution on cpu

Only throw exception if udf module us present

Use CHECK instead of assert

Remove debug functions and clang-formating

QueryEngine/CodeGenerator.h

@@ -405,6 +405,16 @@ class CodeGenerator {
   llvm::Value* codegenFunctionOperNullArg(const Analyzer::FunctionOper*,
                                           const std::vector<llvm::Value*>&);
 
+  llvm::StructType* createArrayStructType(const std::string& udf_func_name,
+                                          size_t param_num);
+
+  void codegenArrayArgs(const std::string& udf_func_name,
+                        size_t param_num,
+                        llvm::Value* array_buf,
+                        llvm::Value* array_size,
+                        llvm::Value* array_is_null,
+                        std::vector<llvm::Value*>& output_args);
+
   std::vector<llvm::Value*> codegenFunctionOperCastArgs(
       const Analyzer::FunctionOper*,
       const ExtensionFunction*,

QueryEngine/ExtensionFunctionsBinding.cpp

@@ -153,13 +153,24 @@ static int match_arguments(const SQLTypeInfo& arg_type,
       break;
     case kLINESTRING:
     case kPOINT:
+      if ((stype == ExtArgumentType::PInt8 || stype == ExtArgumentType::PInt16 ||
+           stype == ExtArgumentType::PInt32 || stype == ExtArgumentType::PInt64 ||
+           stype == ExtArgumentType::PFloat || stype == ExtArgumentType::PDouble) &&
+          sig_pos < max_pos && sig_types[sig_pos + 1] == ExtArgumentType::Int64) {
+        penalty_score += 1000;
+        return 2;
+      }
+
     case kARRAY:
       if ((stype == ExtArgumentType::PInt8 || stype == ExtArgumentType::PInt16 ||
            stype == ExtArgumentType::PInt32 || stype == ExtArgumentType::PInt64 ||
            stype == ExtArgumentType::PFloat || stype == ExtArgumentType::PDouble) &&
           sig_pos < max_pos && sig_types[sig_pos + 1] == ExtArgumentType::Int64) {
         penalty_score += 1000;
         return 2;
+      } else if (is_ext_arg_type_array(stype)) {
+        penalty_score += 1000;
+        return 1;
       }
       break;
     case kPOLYGON:
@@ -327,3 +338,18 @@ ExtensionFunction bind_function(const Analyzer::FunctionOper* function_oper) {
   }
   return bind_function(name, func_args);
 }
+
+bool is_ext_arg_type_array(const ExtArgumentType ext_arg_type) {
+  switch (ext_arg_type) {
+    case ExtArgumentType::ArrayInt8:
+    case ExtArgumentType::ArrayInt16:
+    case ExtArgumentType::ArrayInt32:
+    case ExtArgumentType::ArrayInt64:
+    case ExtArgumentType::ArrayFloat:
+    case ExtArgumentType::ArrayDouble:
+      return true;
+
+    default:
+      return false;
+  }
+}

QueryEngine/ExtensionFunctionsBinding.h

@@ -41,4 +41,5 @@ ExtensionFunction bind_function(std::string name,
 
 ExtensionFunction bind_function(const Analyzer::FunctionOper* function_oper);
 
+bool is_ext_arg_type_array(const ExtArgumentType ext_arg_type);
 #endif  // QUERYENGINE_EXTENSIONFUNCTIONSBINDING_H

QueryEngine/ExtensionFunctionsWhitelist.cpp

@@ -141,6 +141,18 @@ std::string serialize_type(const ExtArgumentType type) {
       return "float*";
     case ExtArgumentType::PDouble:
       return "double*";
+    case ExtArgumentType::ArrayInt8:
+      return "array_i8";
+    case ExtArgumentType::ArrayInt16:
+      return "array_i16";
+    case ExtArgumentType::ArrayInt32:
+      return "array_i32";
+    case ExtArgumentType::ArrayInt64:
+      return "array_i64";
+    case ExtArgumentType::ArrayFloat:
+      return "array_float";
+    case ExtArgumentType::ArrayDouble:
+      return "array_double";
     default:
       CHECK(false);
   }
@@ -279,6 +291,25 @@ ExtArgumentType deserialize_type(const std::string& type_name) {
   if (type_name == "double*") {
     return ExtArgumentType::PDouble;
   }
+  if (type_name == "array_i8") {
+    return ExtArgumentType::ArrayInt8;
+  }
+  if (type_name == "array_i16") {
+    return ExtArgumentType::ArrayInt16;
+  }
+  if (type_name == "array_i32") {
+    return ExtArgumentType::ArrayInt32;
+  }
+  if (type_name == "array_i64") {
+    return ExtArgumentType::ArrayInt64;
+  }
+  if (type_name == "array_float") {
+    return ExtArgumentType::ArrayFloat;
+  }
+  if (type_name == "array_double") {
+    return ExtArgumentType::ArrayDouble;
+  }
+
   CHECK(false);
   return ExtArgumentType::Int16;
 }

QueryEngine/ExtensionFunctionsWhitelist.h

@@ -44,7 +44,13 @@ enum class ExtArgumentType {
   PInt64,
   PFloat,
   PDouble,
-  Bool
+  Bool,
+  ArrayInt8,
+  ArrayInt16,
+  ArrayInt32,
+  ArrayInt64,
+  ArrayFloat,
+  ArrayDouble
 };
 
 SQLTypeInfo ext_arg_type_to_type_info(const ExtArgumentType ext_arg_type);

QueryEngine/ExtensionsIR.cpp

@@ -20,6 +20,9 @@
 #include "ExtensionFunctionsBinding.h"
 #include "ExtensionFunctionsWhitelist.h"
 
+extern std::unique_ptr<llvm::Module> udf_gpu_module;
+extern std::unique_ptr<llvm::Module> udf_cpu_module;
+
 namespace {
 
 llvm::Type* ext_arg_type_to_llvm_type(const ExtArgumentType ext_arg_type,
@@ -171,6 +174,7 @@ llvm::Value* CodeGenerator::codegenFunctionOper(
 
   auto ext_call_nullcheck = endArgsNullcheck(bbs, ext_call, function_oper);
   cgen_state_->ext_call_cache_.push_back({function_oper, ext_call_nullcheck});
+
   return ext_call_nullcheck;
 }
 
@@ -328,6 +332,59 @@ llvm::Value* CodeGenerator::codegenFunctionOperNullArg(
   return one_arg_null;
 }
 
+llvm::StructType* CodeGenerator::createArrayStructType(const std::string& udf_func_name,
+                                                       size_t param_num) {
+  llvm::Function* udf_func = cgen_state_->module_->getFunction(udf_func_name);
+  llvm::Module* module_for_lookup = cgen_state_->module_;
+
+  CHECK(udf_func);
+
+  llvm::FunctionType* udf_func_type = udf_func->getFunctionType();
+  CHECK(param_num < udf_func_type->getNumParams());
+  llvm::Type* param_type = udf_func_type->getParamType(param_num);
+  CHECK(param_type->isPointerTy());
+  llvm::Type* struct_type = param_type->getPointerElementType();
+  CHECK(struct_type->isStructTy());
+  CHECK(struct_type->getStructNumElements() == 3);
+
+  llvm::StringRef struct_name = struct_type->getStructName();
+
+  llvm::StructType* array_type = module_for_lookup->getTypeByName(struct_name);
+  CHECK(array_type);
+
+  return (array_type);
+}
+
+void CodeGenerator::codegenArrayArgs(const std::string& udf_func_name,
+                                     size_t param_num,
+                                     llvm::Value* array_buf,
+                                     llvm::Value* array_size,
+                                     llvm::Value* array_null,
+                                     std::vector<llvm::Value*>& output_args) {
+  CHECK(array_buf);
+  CHECK(array_size);
+  CHECK(array_null);
+
+  auto array_abstraction = createArrayStructType(udf_func_name, param_num);
+  auto alloc_mem = cgen_state_->ir_builder_.CreateAlloca(array_abstraction, nullptr);
+
+  auto array_buf_ptr =
+      cgen_state_->ir_builder_.CreateStructGEP(array_abstraction, alloc_mem, 0);
+  cgen_state_->ir_builder_.CreateStore(array_buf, array_buf_ptr);
+
+  auto array_size_ptr =
+      cgen_state_->ir_builder_.CreateStructGEP(array_abstraction, alloc_mem, 1);
+  cgen_state_->ir_builder_.CreateStore(array_size, array_size_ptr);
+
+  auto bool_extended_type = llvm::Type::getInt8Ty(cgen_state_->context_);
+  auto array_null_extended =
+      cgen_state_->ir_builder_.CreateZExt(array_null, bool_extended_type);
+  auto array_is_null_ptr =
+      cgen_state_->ir_builder_.CreateStructGEP(array_abstraction, alloc_mem, 2);
+  cgen_state_->ir_builder_.CreateStore(array_null_extended, array_is_null_ptr);
+  output_args.push_back(alloc_mem);
+}
+
 // Generate CAST operations for arguments in `orig_arg_lvs` to the types required by
 // `ext_func_sig`.
 std::vector<llvm::Value*> CodeGenerator::codegenFunctionOperCastArgs(
@@ -365,10 +422,29 @@ std::vector<llvm::Value*> CodeGenerator::codegenFunctionOperCastArgs(
                             {orig_arg_lvs[k],
                              posArg(arg),
                              cgen_state_->llInt(log2_bytes(elem_ti.get_logical_size()))});
-      args.push_back(castArrayPointer(ptr_lv, elem_ti));
-      args.push_back(cgen_state_->ir_builder_.CreateZExt(
-          len_lv, get_int_type(64, cgen_state_->context_)));
-      j++;
+
+      if (!is_ext_arg_type_array(ext_func_args[i])) {
+        args.push_back(castArrayPointer(ptr_lv, elem_ti));
+        args.push_back(cgen_state_->ir_builder_.CreateZExt(
+            len_lv, get_int_type(64, cgen_state_->context_)));
+        j++;
+      } else {
+        auto array_buf_arg = castArrayPointer(ptr_lv, elem_ti);
+        auto builder = cgen_state_->ir_builder_;
+        auto array_size_arg =
+            builder.CreateZExt(len_lv, get_int_type(64, cgen_state_->context_));
+        auto array_null_arg =
+            cgen_state_->emitExternalCall("array_is_null",
+                                          get_int_type(1, cgen_state_->context_),
+                                          {orig_arg_lvs[k], posArg(arg)});
+        codegenArrayArgs(ext_func_sig->getName(),
+                         k,
+                         array_buf_arg,
+                         array_size_arg,
+                         array_null_arg,
+                         args);
+      }
+
     } else if (arg_ti.is_geometry()) {
       // Coords
       bool const_arr = (const_arr_size.count(orig_arg_lvs[k]) > 0);

QueryEngine/NativeCodegen.cpp

@@ -1635,7 +1635,14 @@ Executor::compileWorkUnit(const std::vector<InputTableInfo>& query_infos,
   } else {
     rt_module_copy->setDataLayout(get_gpu_data_layout());
     rt_module_copy->setTargetTriple(get_gpu_target_triple_string());
+
     if (is_udf_module_present()) {
+      llvm::Triple gpu_triple(udf_gpu_module->getTargetTriple());
+
+      if (!gpu_triple.isNVPTX()) {
+        throw QueryMustRunOnCpu();
+      }
+
       link_udf_module(udf_gpu_module, *rt_module_copy, cgen_state_.get());
     }
     if (is_rt_udf_module_present()) {

Tests/Udf/udf_sample.cpp

@@ -1,4 +1,7 @@
 #include <cstdint>
+#include <limits>
+#include <type_traits>
+
 #if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__)
 #define DEVICE __device__
 #else
@@ -21,6 +24,91 @@
 
 #define EXTENSION_INLINE extern "C" ALWAYS_INLINE DEVICE
 
+// use std::size_t;
+
+template <typename T>
+struct Array {
+  T* ptr;
+  std::size_t sz;
+  bool is_null;
+
+  DEVICE T operator()(const std::size_t index) {
+    if (index < sz)
+      return ptr[index];
+    else
+      return 0;  // see array_at
+  }
+
+  DEVICE std::size_t getSize() const { return sz; }
+
+  DEVICE bool isNull() const { return is_null; }
+
+  DEVICE constexpr inline T null_value() {
+    return std::is_signed<T>::value ? std::numeric_limits<T>::min()
+                                    : std::numeric_limits<T>::max();
+  }
+};
+
+EXTENSION_NOINLINE
+bool array_is_null_double(Array<double> arr) {
+  return arr.isNull();
+}
+
+EXTENSION_NOINLINE
+int32_t array_sz_double(Array<double> arr) {
+  return arr.getSize();
+}
+
+EXTENSION_NOINLINE
+double array_at_double(Array<double> arr, std::size_t idx) {
+  return arr(idx);
+}
+
+EXTENSION_NOINLINE
+bool array_is_null_int32(Array<int32_t> arr) {
+  return arr.isNull();
+}
+
+EXTENSION_NOINLINE
+int32_t array_sz_int32(Array<int32_t> arr) {
+  return (int32_t)arr.getSize();
+}
+
+EXTENSION_NOINLINE
+int32_t array_at_int32(Array<int32_t> arr, std::size_t idx) {
+  return arr(idx);
+}
+
+EXTENSION_NOINLINE
+int8_t array_at_int32_is_null(Array<int32_t> arr, std::size_t idx) {
+  return (int8_t)(array_at_int32(arr, idx) == arr.null_value());
+}
+
+EXTENSION_NOINLINE
+bool array_is_null_int64(Array<int64_t> arr) {
+  return arr.isNull();
+}
+
+EXTENSION_NOINLINE
+int64_t array_sz_int64(Array<int64_t> arr) {
+  return arr.getSize();
+}
+
+EXTENSION_NOINLINE
+int64_t array_at_int64(Array<int64_t> arr, std::size_t idx) {
+  return arr(idx);
+}
+
+EXTENSION_NOINLINE
+int8_t array_at_int64_is_null(Array<int64_t> arr, std::size_t idx) {
+  return (int8_t)(array_at_int64(arr, idx) == arr.null_value());
+}
+
+EXTENSION_NOINLINE
+int32_t udf_diff(const int32_t x, const int32_t y) {
+  return x - y;
+}
+
 EXTENSION_NOINLINE
 double udf_range(const double high_price, const double low_price) {
   return high_price - low_price;