Skip to content

Commit

Permalink
Merge pull request #892 from devavret/enh-ext-jitify-binops
Browse files Browse the repository at this point in the history
[REVIEW] Jitify versions of binaryops for non-homogeneous types
  • Loading branch information
jrhemstad authored Mar 9, 2019
2 parents 8ccaf8b + 2bfa21e commit 0676d5f
Show file tree
Hide file tree
Showing 35 changed files with 2,549 additions and 111 deletions.
4 changes: 4 additions & 0 deletions .gitmodules
Original file line number Diff line number Diff line change
Expand Up @@ -10,3 +10,7 @@
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
branch = cudf
2 changes: 1 addition & 1 deletion CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -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_
Expand All @@ -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`
Expand Down
38 changes: 36 additions & 2 deletions cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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}")
Expand Down Expand Up @@ -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
Expand All @@ -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 ---------------------------------------------------------------------------------

Expand All @@ -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 --------------------------------------------------------------------------
Expand Down Expand Up @@ -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
Expand Down
54 changes: 54 additions & 0 deletions cpp/include/cudf/functions.h
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
48 changes: 48 additions & 0 deletions cpp/include/cudf/types.h
Original file line number Diff line number Diff line change
Expand Up @@ -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.
*
Expand Down Expand Up @@ -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.
Expand Down
37 changes: 37 additions & 0 deletions cpp/src/binary/jit/code/code.h
Original file line number Diff line number Diff line change
@@ -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
69 changes: 69 additions & 0 deletions cpp/src/binary/jit/code/kernel.cpp
Original file line number Diff line number Diff line change
@@ -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
Loading

0 comments on commit 0676d5f

Please sign in to comment.