[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/rapidsai/jitify.git

CHANGELOG.md

@@ -59,6 +59,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_
@@ -75,7 +76,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
@@ -215,6 +223,31 @@ add_library(cudf SHARED
 #Override RPATH for cudf
 SET_TARGET_PROPERTIES(cudf PROPERTIES BUILD_RPATH "\$ORIGIN")
 
+###################################################################################################
+# - jitify ----------------------------------------------------------------------------------------
+
+# Creates executable stringify and uses it to convert types.h to c-str for use in JIT code
+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 "Run stringify on header types.h to convert it to c-str for use in JIT compiled code"
+                   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)
+
+option(JITIFY_PROCESS_CACHE "Use a process level (instead of thread level) cache for JIT compiled kernels" ON)
+if(JITIFY_PROCESS_CACHE)
+    message(STATUS "Using process level cache for JIT compiled kernels")
+    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --define-macro JITIFY_THREAD_SAFE")
+endif(JITIFY_PROCESS_CACHE)
+
 ###################################################################################################
 # - build options ---------------------------------------------------------------------------------
 
@@ -236,7 +269,8 @@ endif(HT_LEGACY_ALLOCATOR)
 ###################################################################################################
 # - 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 --------------------------------------------------------------------------
@@ -268,7 +302,7 @@ add_custom_command(OUTPUT INSTALL_PYTHON_CFFI
 
 add_custom_target(install_python DEPENDS cudf rmm_install_python PYTHON_CFFI INSTALL_PYTHON_CFFI)
 
-	###################################################################################################
+###################################################################################################
 # - make documentation ----------------------------------------------------------------------------
 
 add_custom_command(OUTPUT CUDF_DOXYGEN

cpp/include/cudf/functions.h

@@ -1969,6 +1969,60 @@ gdf_error gdf_extract_datetime_second(gdf_column *input, gdf_column *output);
 
 /* binary operators */
 
+/**
+ * @brief Performs a binary operation between a gdf_scalar and a gdf_column.
+ *
+ * The desired output type must 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 mask of rhs gdf_column and is_valid
+ * bool of lhs gdf_scalar
+ *
+ * @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 to use
+ * @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 Performs a binary operation between a gdf_column and a gdf_scalar.
+ *
+ * The desired output type must 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 mask of lhs gdf_column and is_valid
+ * bool of rhs gdf_scalar
+ *
+ * @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 to use
+ * @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 Performs a binary operation between two gdf_columns.
+ *
+ * The desired output type must 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_columns
+ *
+ * @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 to use
+ * @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

@@ -93,6 +93,33 @@ typedef struct {
 } gdf_dtype_extra_info;
 
 
+
+/**
+ * @brief Union used to store single value for scalar type
+ */
+// TODO: #1119 Use traits to set `gdf_data` elements
+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   */
+  gdf_date32    dt32;  /**< GDF_DATE32    */
+  gdf_date64    dt64;  /**< GDF_DATE64    */
+  gdf_timestamp tmst;  /**< GDF_TIMESTAMP */
+} gdf_data;
+
+/**
+ * @brief A struct to hold a scalar (single) value and its type information
+ */
+typedef struct {
+  gdf_data  data;      /**< A union that represents the value */
+  gdf_dtype dtype;     /**< The datatype of the scalar's data */
+  bool      is_valid;  /**< False if the value is null */
+} gdf_scalar;
+
+
 /**
  * @brief The C representation of a column in CUDF. This is the main unit of operation.
  *
@@ -167,6 +194,27 @@ typedef enum {
 } 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 
  * performed as well as additional information about the input data.

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