[REVIEW] Jitify versions of binaryops for non-homogeneous types

.gitmodules

@@ -10,3 +10,6 @@
 	path = thirdparty/rmm
 	url = https://github.com/rapidsai/rmm.git
 	branch = branch-0.6
+[submodule "thirdparty/jitify"]
+	path = thirdparty/jitify
+	url = https://github.com/NVIDIA/jitify.git

CHANGELOG.md

@@ -41,6 +41,7 @@
 
 ## Improvements
 
+- PR #892 Add support for heterogeneous types in binary ops with JIT 
 - PR #730 Improve performance of `gdf_table` constructor
 - PR #561 Add Doxygen style comments to Join CUDA functions
 - PR #813 unified libcudf API functions by replacing gpu_ with gdf_
@@ -57,7 +58,6 @@
 - PR #909 CSV Reader: Avoid host->device->host copy for header row data
 - PR #916 Improved unit testing and error checking for `gdf_column_concat`
 - PR #941 Replace `numpy` call in `Series.hash_encode` with `numba`
-- PR #943 Updated `count_nonzero_mask` to return `num_rows` when the mask is null
 - PR #942 Added increment/decrement operators for wrapper types
 - PR #943 Updated `count_nonzero_mask` to return `num_rows` when the mask is null
 - PR #952 Added trait to map C++ type to `gdf_dtype`

cpp/CMakeLists.txt

@@ -141,6 +141,7 @@ include_directories("${ARROW_INCLUDE_DIR}"
                     "${CMAKE_SOURCE_DIR}/include"
                     "${CMAKE_SOURCE_DIR}/src"
                     "${CMAKE_SOURCE_DIR}/thirdparty/cub"
+                    "${CMAKE_SOURCE_DIR}/thirdparty/jitify"
                     "${CMAKE_SOURCE_DIR}/thirdparty/moderngpu/src"
                     "${CMAKE_SOURCE_DIR}/thirdparty/rmm/include"
                     "${ZLIB_INCLUDE_DIRS}")
@@ -189,6 +190,13 @@ add_library(cudf SHARED
             src/groupby/groupby.cu
             src/groupby/new_groupby.cu
             src/binary/binary_ops.cu
+            src/binary/jit/code/kernel.cpp
+            src/binary/jit/code/operation.cpp
+            src/binary/jit/code/traits.cpp
+            src/binary/jit/core/binop.cpp
+            src/binary/jit/core/launcher.cpp
+            src/binary/jit/util/operator.cpp
+            src/binary/jit/util/type.cpp
             src/bitmask/bitmask_ops.cu
             src/bitmask/valid_ops.cu
             src/compaction/stream_compaction_ops.cu
@@ -213,6 +221,24 @@ add_library(cudf SHARED
 #Override RPATH for cudf
 SET_TARGET_PROPERTIES(cudf PROPERTIES BUILD_RPATH "\$ORIGIN")
 
+###################################################################################################
+# - jitify ----------------------------------------------------------------------------------------
+
+add_executable(stringify "${CMAKE_SOURCE_DIR}/thirdparty/jitify/stringify.cpp")
+execute_process(WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    COMMAND ${CMAKE_COMMAND} -E make_directory ${CMAKE_BINARY_DIR}/include)
+
+add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/include/types.h.jit
+                   WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include
+                   COMMAND ${CMAKE_BINARY_DIR}/stringify cudf/types.h > ${CMAKE_BINARY_DIR}/include/types.h.jit
+                   COMMENT "stringify header types.h"
+                   DEPENDS stringify
+                   MAIN_DEPENDENCY ${CMAKE_CURRENT_SOURCE_DIR}/include/cudf/types.h)
+
+add_custom_target(stringify_run DEPENDS ${CMAKE_BINARY_DIR}/include/types.h.jit)
+
+add_dependencies(cudf stringify_run)
+
 ###################################################################################################
 # - build options ---------------------------------------------------------------------------------
 
@@ -230,11 +256,18 @@ if(HT_LEGACY_ALLOCATOR)
     set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --define-macro HT_LEGACY_ALLOCATOR")
 endif(HT_LEGACY_ALLOCATOR)
 
+option(JITIFY_THREAD_SAFE "Use a global cache for JIT compiled kernels" ON)
+if(JITIFY_THREAD_SAFE)
+    message(STATUS "Using global cache for JIT compiled kernels")
+    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --define-macro JITIFY_THREAD_SAFE")
+endif(JITIFY_THREAD_SAFE)
+
 
 ###################################################################################################
 # - link libraries --------------------------------------------------------------------------------
 
-target_link_libraries(cudf rmm "${ARROW_LIB}" ${ZLIB_LIBRARIES} NVStrings)
+# TODO: better nvrtc linking with optional variables
+target_link_libraries(cudf rmm "${ARROW_LIB}" ${ZLIB_LIBRARIES} NVStrings nvrtc)
 
 ###################################################################################################
 # - python cffi bindings --------------------------------------------------------------------------

cpp/include/cudf/functions.h

@@ -1960,6 +1960,66 @@ gdf_error gdf_extract_datetime_second(gdf_column *input, gdf_column *output);
 
 /* binary operators */
 
+/**
+ * @brief Binary operation function between gdf_scalar and gdf_column structs.
+ *
+ * The function performs the binary operation of a gdf_scalar operand and a
+ * gdf_column operand.
+ * 
+ * The desired output type needs to be specified in out->dtype
+ *
+ * If the valid field in the gdf_column output is not nullptr, then the valid
+ * mask from rhs gdf_column is copied into the data pointer to by out->valid
+ *
+ * @param out (gdf_column) Output of the operation.
+ * @param lhs (gdf_scalar) First operand of the operation.
+ * @param rhs (gdf_column) Second operand of the operation.
+ * @param ope (enum) The binary operator that is going to be used in the operation.
+ * @return    GDF_SUCCESS if the operation was successful, otherwise an appropriate
+ *            error code
+ */
+gdf_error gdf_binary_operation_s_v(gdf_column* out, gdf_scalar* lhs, gdf_column* rhs, gdf_binary_operator ope);
+
+/**
+ * @brief Binary operation function between gdf_column and gdf_scalar structs.
+ *
+ * The function performs the binary operation of a gdf_column operand and a
+ * gdf_scalar operand.
+ *
+ * The desired output type needs to be specified in out->dtype
+ *
+ * If the valid field in the gdf_column output is not nullptr, then the valid
+ * mask from lhs gdf_column is copied into the data pointer to by out->valid
+ *
+ * @param out (gdf_column) Output of the operation.
+ * @param lhs (gdf_column) First operand of the operation.
+ * @param rhs (gdf_scalar) Second operand of the operation.
+ * @param ope (enum) The binary operator that is going to be used in the operation.
+ * @return    GDF_SUCCESS if the operation was successful, otherwise an appropriate
+ *            error code
+ */
+gdf_error gdf_binary_operation_v_s(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs, gdf_binary_operator ope);
+
+/**
+ * @brief Binary operation function between two gdf_column structs.
+ *
+ * The function performs the binary operation of two gdf_column operands.
+ *
+ * The desired output type needs to be specified in out->dtype
+ *
+ * If the valid field in the gdf_column output is not nullptr, then it will be
+ * filled with the bitwise AND of the valid masks of lhs and rhs gdf_column's
+ *
+ * @param out (gdf_column) Output of the operation.
+ * @param lhs (gdf_column) First operand of the operation.
+ * @param rhs (gdf_column) Second operand of the operation.
+ * @param ope (enum) The binary operator that is going to be used in the operation.
+ * @return    GDF_SUCCESS if the operation was successful, otherwise an appropriate
+ *            error code
+ */
+gdf_error gdf_binary_operation_v_v(gdf_column* out, gdf_column* lhs, gdf_column* rhs, gdf_binary_operator ope);
+
+
 /* arith */
 
 gdf_error gdf_add_generic(gdf_column *lhs, gdf_column *rhs, gdf_column *output);

cpp/include/cudf/types.h

@@ -85,6 +85,35 @@ typedef struct {
 	// here we can also hold info for decimal datatype or any other datatype that requires additional information
 } gdf_dtype_extra_info;
 
+/**---------------------------------------------------------------------------*
+ * @union gdf_data
+ * @brief Union used for scalar type.
+ * It stores a unique value for scalar type.
+ * It has a direct relationship with the gdf_dtype.
+ *---------------------------------------------------------------------------**/
+typedef union {
+  int8_t   si08;  /**< GDF_INT8      */
+  int16_t  si16;  /**< GDF_INT16     */
+  int32_t  si32;  /**< GDF_INT32     */
+  int64_t  si64;  /**< GDF_INT64     */
+  float    fp32;  /**< GDF_FLOAT32   */
+  double   fp64;  /**< GDF_FLOAT64   */
+  int32_t  dt32;  /**< GDF_DATE32    */
+  int64_t  dt64;  /**< GDF_DATE64    */
+  int64_t  tmst;  /**< GDF_TIMESTAMP */
+} gdf_data;
+
+/**---------------------------------------------------------------------------*
+ * @brief A struct to hold a scalar (single) value and its type information
+ * 
+ *---------------------------------------------------------------------------**/
+typedef struct {
+  gdf_data  data;      /**< Pointer to the scalar data */
+  gdf_dtype dtype;     /**< The datatype of the scalar's data */
+  bool      is_valid;  /**< False if the value is null */
+} gdf_scalar;
+
+
 typedef struct gdf_column_{
     void *data;                       /**< Pointer to the columns data */ 
     gdf_valid_type *valid;            /**< Pointer to the columns validity bit mask where the 'i'th bit indicates if the 'i'th row is NULL */
@@ -156,6 +185,29 @@ typedef enum {
   GDF_NUM_COLORS, /** Add new colors above this line */
 } gdf_color;
 
+
+/**---------------------------------------------------------------------------*
+ * @brief Types of binary operations that can be performed on data.
+ * 
+ *---------------------------------------------------------------------------**/
+typedef enum {
+  GDF_ADD,            /**< operator + */
+  GDF_SUB,            /**< operator - */
+  GDF_MUL,            /**< operator * */
+  GDF_DIV,            /**< operator / using common type of lhs and rhs */
+  GDF_TRUE_DIV,       /**< operator / after promoting type to floating point*/
+  GDF_FLOOR_DIV,      /**< operator / after promoting to float and then flooring the result */
+  GDF_MOD,            /**< operator %  */
+  GDF_POW,            /**< lhs ^ rhs   */
+  GDF_EQUAL,          /**< operator == */
+  GDF_NOT_EQUAL,      /**< operator != */
+  GDF_LESS,           /**< operator <  */
+  GDF_GREATER,        /**< operator >  */
+  GDF_LESS_EQUAL,     /**< operator <= */
+  GDF_GREATER_EQUAL,  /**< operator >= */
+} gdf_binary_operator;
+
+
 /* --------------------------------------------------------------------------*/
 /** 
  * @brief  This struct holds various information about how an operation should be 

cpp/src/binary/jit/code/code.h

@@ -0,0 +1,37 @@
+/*
+ * Copyright (c) 2019, NVIDIA CORPORATION.
+ *
+ * Copyright 2018-2019 BlazingDB, Inc.
+ *     Copyright 2018 Christian Noboa Mardini <[email protected]>
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef GDF_BINARY_OPERATION_JIT_CODE_CODE_H
+#define GDF_BINARY_OPERATION_JIT_CODE_CODE_H
+
+namespace cudf {
+namespace binops {
+namespace jit {
+namespace code {
+
+    extern const char* kernel;
+    extern const char* traits;
+    extern const char* operation;
+
+}
+}
+}
+}
+
+#endif

cpp/src/binary/jit/code/kernel.cpp

@@ -0,0 +1,69 @@
+/*
+ * Copyright (c) 2019, NVIDIA CORPORATION.
+ *
+ * Copyright 2018-2019 BlazingDB, Inc.
+ *     Copyright 2018 Christian Noboa Mardini <[email protected]>
+ *     Copyright 2018 Rommel Quintanilla <[email protected]>
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+namespace cudf {
+namespace binops {
+namespace jit {
+namespace code {
+
+const char* kernel =
+R"***(
+    #include "operation.h"
+    #include "cudf/types.h"
+
+    template <typename TypeOut, typename TypeLhs, typename TypeRhs, typename TypeOpe>
+    __global__
+    void kernel_v_s(gdf_size_type size,
+                    TypeOut* out_data, TypeLhs* lhs_data, gdf_data rhs_data) {
+        int tid = threadIdx.x;
+        int blkid = blockIdx.x;
+        int blksz = blockDim.x;
+        int gridsz = gridDim.x;
+
+        int start = tid + blkid * blksz;
+        int step = blksz * gridsz;
+
+        for (gdf_size_type i=start; i<size; i+=step) {
+            out_data[i] = TypeOpe::template operate<TypeOut, TypeLhs, TypeRhs>(lhs_data[i], *reinterpret_cast<TypeRhs*>(&rhs_data));
+        }
+    }
+
+    template <typename TypeOut, typename TypeLhs, typename TypeRhs, typename TypeOpe>
+    __global__
+    void kernel_v_v(gdf_size_type size,
+                    TypeOut* out_data, TypeLhs* lhs_data, TypeRhs* rhs_data) {
+        int tid = threadIdx.x;
+        int blkid = blockIdx.x;
+        int blksz = blockDim.x;
+        int gridsz = gridDim.x;
+
+        int start = tid + blkid * blksz;
+        int step = blksz * gridsz;
+
+        for (gdf_size_type i=start; i<size; i+=step) {
+            out_data[i] = TypeOpe::template operate<TypeOut, TypeLhs, TypeRhs>(lhs_data[i], rhs_data[i]);
+        }
+    }
+)***";
+
+} // namespace code
+} // namespace jit
+} // namespace binops
+} // namespace cudf