diff --git a/.gitmodules b/.gitmodules index 7a64546f4d78..c0f28a452af9 100644 --- a/.gitmodules +++ b/.gitmodules @@ -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 \ No newline at end of file diff --git a/CHANGELOG.md b/CHANGELOG.md index 3868feff4d0f..690acac91fe6 100644 --- a/CHANGELOG.md +++ b/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` diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index ff20545fbaa1..99c42b12b8a0 100644 --- a/cpp/CMakeLists.txt +++ b/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 diff --git a/cpp/include/cudf/functions.h b/cpp/include/cudf/functions.h index 70a853b63373..6acf05b00a31 100644 --- a/cpp/include/cudf/functions.h +++ b/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); diff --git a/cpp/include/cudf/types.h b/cpp/include/cudf/types.h index 6d805156c214..967ee6c6d351 100644 --- a/cpp/include/cudf/types.h +++ b/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. diff --git a/cpp/src/binary/jit/code/code.h b/cpp/src/binary/jit/code/code.h new file mode 100644 index 000000000000..7fdb63ab928c --- /dev/null +++ b/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 + * + * 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 diff --git a/cpp/src/binary/jit/code/kernel.cpp b/cpp/src/binary/jit/code/kernel.cpp new file mode 100644 index 000000000000..bb52a8b76bfd --- /dev/null +++ b/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 + * Copyright 2018 Rommel Quintanilla + * + * 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 + __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(lhs_data[i], *reinterpret_cast(&rhs_data)); + } + } + + template + __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(lhs_data[i], rhs_data[i]); + } + } +)***"; + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace cudf diff --git a/cpp/src/binary/jit/code/operation.cpp b/cpp/src/binary/jit/code/operation.cpp new file mode 100644 index 000000000000..4591a89c51ce --- /dev/null +++ b/cpp/src/binary/jit/code/operation.cpp @@ -0,0 +1,248 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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* operation = +R"***( +#pragma once + #include "traits.h" + + struct Add { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(x) + static_cast(y)); + } + }; + + using RAdd = Add; + + struct Sub { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(x) - static_cast(y)); + } + }; + + struct RSub { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(y) - static_cast(x)); + } + }; + + struct Mul { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(x) * static_cast(y)); + } + }; + + using RMul = Mul; + + struct Div { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(x) / static_cast(y)); + } + }; + + struct RDiv { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(y) / static_cast(x)); + } + }; + + struct TrueDiv { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(x) / static_cast(y)); + } + }; + + struct RTrueDiv { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(y) / static_cast(x)); + } + }; + + struct FloorDiv { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return floor(static_cast(x) / static_cast(y)); + } + }; + + struct RFloorDiv { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return floor(static_cast(y) / static_cast(x)); + } + }; + + struct Mod { + template )>* = nullptr> + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(x) % static_cast(y)); + } + + template )>* = nullptr> + static TypeOut operate(TypeLhs x, TypeRhs y) { + return fmodf(static_cast(x), static_cast(y)); + } + + template )>* = nullptr> + static TypeOut operate(TypeLhs x, TypeRhs y) { + return fmod(static_cast(x), static_cast(y)); + } + }; + + struct RMod { + template )>* = nullptr> + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (static_cast(y) % static_cast(x)); + } + + template )>* = nullptr> + static TypeOut operate(TypeLhs x, TypeRhs y) { + return fmodf(static_cast(y), static_cast(x)); + } + + template )>* = nullptr> + static TypeOut operate(TypeLhs x, TypeRhs y) { + return fmod(static_cast(y), static_cast(x)); + } + }; + + struct Pow { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return pow(static_cast(x), static_cast(y)); + } + }; + + struct RPow { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return pow(static_cast(y), static_cast(x)); + } + }; + + struct Equal { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x == y); + } + }; + + using REqual = Equal; + + struct NotEqual { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x != y); + } + }; + + using RNotEqual = NotEqual; + + struct Less { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x < y); + } + }; + + struct RLess { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (y < x); + } + }; + + struct Greater { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x > y); + } + }; + + struct RGreater { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (y > x); + } + }; + + struct LessEqual { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x <= y); + } + }; + + struct RLessEqual { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (y <= x); + } + }; + + struct GreaterEqual { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (x >= y); + } + }; + + struct RGreaterEqual { + template + static TypeOut operate(TypeLhs x, TypeRhs y) { + return (y >= x); + } + }; +)***"; + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace cudf diff --git a/cpp/src/binary/jit/code/traits.cpp b/cpp/src/binary/jit/code/traits.cpp new file mode 100644 index 000000000000..c434a18281aa --- /dev/null +++ b/cpp/src/binary/jit/code/traits.cpp @@ -0,0 +1,55 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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* traits = +R"***( +#pragma once + #include + #include + + // ------------------------------------------------------------------------- + // Simplifying std::is_integral + template + constexpr bool is_integral_v = std::is_integral::value; + + // ------------------------------------------------------------------------- + // type_traits cannot tell the difference between float and double + template + constexpr bool isFloat = false; + + template <> + constexpr bool isFloat = true; + + + template + constexpr bool isDouble = false; + + template <> + constexpr bool isDouble = true; +)***"; + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace cudf diff --git a/cpp/src/binary/jit/core/binop.cpp b/cpp/src/binary/jit/core/binop.cpp new file mode 100644 index 000000000000..dddf9ea407ad --- /dev/null +++ b/cpp/src/binary/jit/core/binop.cpp @@ -0,0 +1,229 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "binary/jit/core/launcher.h" +#include "binary/jit/util/operator.h" +#include "bitmask/bitmask_ops.h" +#include "utilities/error_utils.hpp" +#include "utilities/cudf_utils.h" +#include "cudf.h" + +namespace cudf { +namespace binops { + + /**---------------------------------------------------------------------------* + * @brief Computes bitwise AND of two input valid masks + * + * This is just a wrapper on apply_bitmask_to_bitmask that can also handle + * cases when one or both of the input masks are nullptr, in which case, it + * copies the mask from the non nullptr input or sets all the output mask to + * valid respectively + * + * @param out_null_coun[out] number of nulls in output + * @param valid_out preallocated output mask + * @param valid_left input mask 1 + * @param valid_right input mask 2 + * @param num_values number of values in each input mask valid_left and valid_right + * @return gdf_error + *---------------------------------------------------------------------------**/ + gdf_error binary_valid_mask_and(gdf_size_type & out_null_count, + gdf_valid_type * valid_out, + gdf_valid_type * valid_left, + gdf_valid_type * valid_right, + gdf_size_type num_values) { + if (num_values == 0) { + out_null_count = 0; + return GDF_SUCCESS; + } + + if (valid_out == nullptr && valid_left == nullptr && valid_right == nullptr) { + // if both in cols have no mask, then out col is allowed to have no mask + out_null_count = 0; + return GDF_SUCCESS; + } + + GDF_REQUIRE((valid_out != nullptr), GDF_DATASET_EMPTY) + + if ( valid_left != nullptr && valid_right != nullptr ) { + cudaStream_t stream; + CUDA_TRY( cudaStreamCreate(&stream) ); + auto error = apply_bitmask_to_bitmask(out_null_count, valid_out, valid_left, valid_right, stream, num_values); + CUDA_TRY(cudaStreamSynchronize(stream)); + CUDA_TRY(cudaStreamDestroy(stream)); + return error; + } + + gdf_size_type num_chars_bitmask = gdf_get_num_chars_bitmask( num_values ); + + if ( valid_left == nullptr && valid_right != nullptr ) { + CUDA_TRY( cudaMemcpy(valid_out, valid_right, num_chars_bitmask, cudaMemcpyDeviceToDevice) ); + } + else if ( valid_left != nullptr && valid_right == nullptr ) { + CUDA_TRY( cudaMemcpy(valid_out, valid_left, num_chars_bitmask, cudaMemcpyDeviceToDevice) ); + } + else if ( valid_left == nullptr && valid_right == nullptr ) { + CUDA_TRY( cudaMemset(valid_out, 0xff, num_chars_bitmask) ); + } + + gdf_size_type non_nulls; + auto error = gdf_count_nonzero_mask(valid_out, num_values, &non_nulls); + out_null_count = num_values - non_nulls; + return error; + } + + /**---------------------------------------------------------------------------* + * @brief Computes output valid mask for op between a column and a scalar + * + * @param out_null_coun[out] number of nulls in output + * @param valid_out preallocated output mask + * @param valid_col input mask of column + * @param valid_scalar bool indicating if scalar is valid + * @param num_values number of values in input mask valid_col + * @return gdf_error + *---------------------------------------------------------------------------**/ + gdf_error scalar_col_valid_mask_and(gdf_size_type & out_null_count, + gdf_valid_type * valid_out, + gdf_valid_type * valid_col, + bool valid_scalar, + gdf_size_type num_values) + { + if (num_values == 0) { + out_null_count = 0; + return GDF_SUCCESS; + } + + if (valid_out == nullptr && valid_col == nullptr && valid_scalar == true) { + // if in col has no mask and scalar is valid, then out col is allowed to have no mask + out_null_count = 0; + return GDF_SUCCESS; + } + + GDF_REQUIRE((valid_out != nullptr), GDF_DATASET_EMPTY) + + gdf_size_type num_chars_bitmask = gdf_get_num_chars_bitmask( num_values ); + + if ( valid_scalar == false ) { + CUDA_TRY( cudaMemset(valid_out, 0x00, num_chars_bitmask) ); + } + else if ( valid_scalar == true && valid_col != nullptr ) { + CUDA_TRY( cudaMemcpy(valid_out, valid_col, num_chars_bitmask, cudaMemcpyDeviceToDevice) ); + } + else if ( valid_scalar == true && valid_col == nullptr ) { + CUDA_TRY( cudaMemset(valid_out, 0xff, num_chars_bitmask) ); + } + + gdf_size_type non_nulls; + auto error = gdf_count_nonzero_mask(valid_out, num_values, &non_nulls); + out_null_count = num_values - non_nulls; + return error; + } + +namespace jit { + + gdf_error binary_operation(gdf_column* out, gdf_scalar* lhs, gdf_column* rhs, gdf_binary_operator ope) { + + // Check for null pointers in input + GDF_REQUIRE((out != nullptr) && (lhs != nullptr) && (rhs != nullptr), GDF_DATASET_EMPTY) + + // Check for 0 sized data + GDF_REQUIRE((out->size != 0) && (rhs->size != 0), GDF_SUCCESS) + GDF_REQUIRE((out->size == rhs->size), GDF_COLUMN_SIZE_MISMATCH) + + // Check for null data pointer + GDF_REQUIRE((out->data != nullptr) && (rhs->data != nullptr), GDF_DATASET_EMPTY) + + // Check for datatype + GDF_REQUIRE((out->dtype > GDF_invalid) && (lhs->dtype > GDF_invalid) && (rhs->dtype > GDF_invalid), GDF_UNSUPPORTED_DTYPE) + GDF_REQUIRE((out->dtype < N_GDF_TYPES) && (lhs->dtype < N_GDF_TYPES) && (rhs->dtype < N_GDF_TYPES), GDF_UNSUPPORTED_DTYPE) + + scalar_col_valid_mask_and(out->null_count, out->valid, rhs->valid, lhs->is_valid, rhs->size); + + Launcher::launch().kernel("kernel_v_s") + .instantiate(ope, Operator::Type::Reverse, out, rhs, lhs) + .launch(out, rhs, lhs); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs, gdf_binary_operator ope) { + + // Check for null pointers in input + GDF_REQUIRE((out != nullptr) && (lhs != nullptr) && (rhs != nullptr), GDF_DATASET_EMPTY) + + // Check for 0 sized data + GDF_REQUIRE((out->size != 0) && (lhs->size != 0), GDF_SUCCESS) + GDF_REQUIRE((out->size == lhs->size), GDF_COLUMN_SIZE_MISMATCH) + + // Check for null data pointer + GDF_REQUIRE((out->data != nullptr) && (lhs->data != nullptr), GDF_DATASET_EMPTY) + + // Check for datatype + GDF_REQUIRE((out->dtype > GDF_invalid) && (lhs->dtype > GDF_invalid) && (rhs->dtype > GDF_invalid), GDF_UNSUPPORTED_DTYPE) + GDF_REQUIRE((out->dtype < N_GDF_TYPES) && (lhs->dtype < N_GDF_TYPES) && (rhs->dtype < N_GDF_TYPES), GDF_UNSUPPORTED_DTYPE) + + scalar_col_valid_mask_and(out->null_count, out->valid, lhs->valid, rhs->is_valid, lhs->size); + + Launcher::launch().kernel("kernel_v_s") + .instantiate(ope, Operator::Type::Direct, out, lhs, rhs) + .launch(out, lhs, rhs); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_column* lhs, gdf_column* rhs, gdf_binary_operator ope) { + + // Check for null pointers in input + GDF_REQUIRE((out != nullptr) && (lhs != nullptr) && (rhs != nullptr), GDF_DATASET_EMPTY) + + // Check for 0 sized data + GDF_REQUIRE((out->size != 0) && (lhs->size != 0) && (rhs->size != 0), GDF_SUCCESS) + GDF_REQUIRE((out->size == lhs->size) && (lhs->size == rhs->size), GDF_COLUMN_SIZE_MISMATCH) + + // Check for null data pointer + GDF_REQUIRE((out->data != nullptr) && (lhs->data != nullptr) && (rhs->data != nullptr), GDF_DATASET_EMPTY) + + // Check for datatype + GDF_REQUIRE((out->dtype > GDF_invalid) && (lhs->dtype > GDF_invalid) && (rhs->dtype > GDF_invalid), GDF_UNSUPPORTED_DTYPE) + GDF_REQUIRE((out->dtype < N_GDF_TYPES) && (lhs->dtype < N_GDF_TYPES) && (rhs->dtype < N_GDF_TYPES), GDF_UNSUPPORTED_DTYPE) + + binary_valid_mask_and(out->null_count, out->valid, lhs->valid, rhs->valid, rhs->size); + + Launcher::launch().kernel("kernel_v_v") + .instantiate(ope, Operator::Type::Direct, out, lhs, rhs) + .launch(out, lhs, rhs); + + return GDF_SUCCESS; + } + +} // namespace jit +} // namespace binops +} // namespace cudf + + +gdf_error gdf_binary_operation_s_v(gdf_column* out, gdf_scalar* lhs, gdf_column* rhs, gdf_binary_operator ope) { + return cudf::binops::jit::binary_operation(out, lhs, rhs, ope); +} + +gdf_error gdf_binary_operation_v_s(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs, gdf_binary_operator ope) { + return cudf::binops::jit::binary_operation(out, lhs, rhs, ope); +} + +gdf_error gdf_binary_operation_v_v(gdf_column* out, gdf_column* lhs, gdf_column* rhs, gdf_binary_operator ope) { + return cudf::binops::jit::binary_operation(out, lhs, rhs, ope); +} diff --git a/cpp/src/binary/jit/core/launcher.cpp b/cpp/src/binary/jit/core/launcher.cpp new file mode 100644 index 000000000000..619642f26556 --- /dev/null +++ b/cpp/src/binary/jit/core/launcher.cpp @@ -0,0 +1,100 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "binary/jit/core/launcher.h" +#include "binary/jit/code/code.h" +#include "types.h.jit" +#include + +namespace cudf { +namespace binops { +namespace jit { + +/**---------------------------------------------------------------------------* + * @brief Cache to hold previously compiled code. If JITIFY_THREAD_SAFE is + * defined, then the cache is held globally in the process, otherwise it is + * held locally in each thread + * + *---------------------------------------------------------------------------**/ +#ifdef JITIFY_THREAD_SAFE + static jitify::JitCache JitCache; +#else + static thread_local jitify::JitCache JitCache; +#endif + + const std::vector Launcher::compilerFlags { "-std=c++14" }; + const std::vector Launcher::headersName + { "operation.h" , "traits.h" , cudf_types_h }; + + /**---------------------------------------------------------------------------* + * @brief Used to provide Jitify with strings that should be used as headers + * during JIT compilation. + * + * @param filename file which was requested to include in source + * @param stream stream to pass string of the requested header to + * @return std::istream* + *---------------------------------------------------------------------------**/ + std::istream* headersCode(std::string filename, std::iostream& stream) { + if (filename == "operation.h") { + stream << code::operation; + return &stream; + } + if (filename == "traits.h") { + stream << code::traits; + return &stream; + } + return nullptr; + } + + Launcher::Launcher() + : program {JitCache.program(code::kernel, headersName, compilerFlags, headersCode)} + { } + + Launcher::Launcher(Launcher&& launcher) + : program {std::move(launcher.program)} + { } + + Launcher& Launcher::kernel(std::string&& value) { + kernelName = value; + return *this; + } + + gdf_error Launcher::launch(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs) { + program.kernel(kernelName.c_str()) + .instantiate(arguments) + .configure_1d_max_occupancy() + .launch(out->size, + out->data, lhs->data, rhs->data); + + return GDF_SUCCESS; + } + + gdf_error Launcher::launch(gdf_column* out, gdf_column* lhs, gdf_column* rhs) { + program.kernel(kernelName.c_str()) + .instantiate(arguments) + .configure_1d_max_occupancy() + .launch(out->size, + out->data, lhs->data, rhs->data); + + return GDF_SUCCESS; + } + +} // namespace jit +} // namespace binops +} // namespace cudf diff --git a/cpp/src/binary/jit/core/launcher.h b/cpp/src/binary/jit/core/launcher.h new file mode 100644 index 000000000000..4d548ac735df --- /dev/null +++ b/cpp/src/binary/jit/core/launcher.h @@ -0,0 +1,117 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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_CORE_LAUNCHER_H +#define GDF_BINARY_OPERATION_JIT_CORE_LAUNCHER_H + +#include "binary/jit/util/type.h" +#include "binary/jit/util/operator.h" +#include + +namespace cudf { +namespace binops { +namespace jit { + + std::istream* headersCode(std::string filename, std::iostream& stream); + + /**---------------------------------------------------------------------------* + * @brief Class used to handle compilation and execution of JIT kernels + * + *---------------------------------------------------------------------------**/ + class Launcher { + public: + static Launcher launch() { + return Launcher(); + } + + public: + Launcher(); + + Launcher(Launcher&&); + + public: + Launcher(const Launcher&) = delete; + + Launcher& operator=(Launcher&&) = delete; + + Launcher& operator=(const Launcher&) = delete; + + public: + /**---------------------------------------------------------------------------* + * @brief Set the kernel name that this launcher will compile and launch + * + * @param value kernel name + *---------------------------------------------------------------------------**/ + Launcher& kernel(std::string&& value); + + /**---------------------------------------------------------------------------* + * @brief Method to generate vector containing all template types for a JIT + * kernel. This vector is used to instantiate the kernel code for one set of types + * + * @tparam Args Output dtype, LHS dtype, RHS dtype + * @param type Operator type (direct (lhs op rhs) or reverse (rhs op lhs)) + * @param args gdf_column* output, lhs, rhs + * @return Launcher& ref to this launcehr object + *---------------------------------------------------------------------------**/ + template + Launcher& instantiate(gdf_binary_operator ope, Operator::Type type, Args&& ... args) { + Operator operatorSelector; + arguments.assign({getTypeName(args->dtype)..., operatorSelector.getOperatorName(ope, type)}); + return *this; + } + + /**---------------------------------------------------------------------------* + * @brief Handle the Jitify API to instantiate and launch using information + * contained in the members of `this` + * + * @param out[out] Output column + * @param lhs[in] LHS column + * @param rhs[in] RHS scalar (single value) + * @return gdf_error + *---------------------------------------------------------------------------**/ + gdf_error launch(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs); + + /**---------------------------------------------------------------------------* + * @brief Handle the Jitify API to instantiate and launch using information + * contained in the members of `this` + * + * @param out[out] Output column + * @param lhs[in] LHS column + * @param rhs[in] RHS column + * @return gdf_error + *---------------------------------------------------------------------------**/ + gdf_error launch(gdf_column* out, gdf_column* lhs, gdf_column* rhs); + + private: + static const std::vector compilerFlags; + static const std::vector headersName; + + private: + jitify::Program program; + + private: + std::string kernelName; + std::vector arguments; + }; + +} // namespace jit +} // namespace binops +} // namespace cudf + +#endif diff --git a/cpp/src/binary/jit/util/operator.cpp b/cpp/src/binary/jit/util/operator.cpp new file mode 100644 index 000000000000..fe986c730d1a --- /dev/null +++ b/cpp/src/binary/jit/util/operator.cpp @@ -0,0 +1,45 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "binary/jit/util/operator.h" +#include "binary/jit/util/type.h" +#include + +namespace cudf { +namespace binops { +namespace jit { + + Operator::Operator() + : buffer{'\0'} + { } + + char* Operator::getOperatorName(gdf_binary_operator ope, Operator::Type type) { + if (type == Operator::Type::Direct) { + buffer[0] = '\0'; + } else { + buffer[0] = 'R'; + buffer[1] = '\0'; + } + strcat(buffer, jit::getOperatorName(ope).data()); + return buffer; + } + +} // namespace jit +} // namespace binops +} // namespace cudf diff --git a/cpp/src/binary/jit/util/operator.h b/cpp/src/binary/jit/util/operator.h new file mode 100644 index 000000000000..9e5e576b4380 --- /dev/null +++ b/cpp/src/binary/jit/util/operator.h @@ -0,0 +1,50 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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_UTIL_OPERATOR_H +#define GDF_BINARY_OPERATION_JIT_UTIL_OPERATOR_H + +#include "cudf.h" + +namespace cudf { +namespace binops { +namespace jit { + + class Operator { + public: + enum class Type { + Direct, + Reverse + }; + + public: + Operator(); + + public: + char* getOperatorName(gdf_binary_operator ope, Operator::Type type); + + private: + char buffer[16]; + }; + +} // namespace jit +} // namespace binops +} // namespace cudf + +#endif diff --git a/cpp/src/binary/jit/util/type.cpp b/cpp/src/binary/jit/util/type.cpp new file mode 100644 index 000000000000..8ec7cdc70b8a --- /dev/null +++ b/cpp/src/binary/jit/util/type.cpp @@ -0,0 +1,119 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "binary/jit/util/type.h" +#include "utilities/type_dispatcher.hpp" + +namespace cudf { +namespace binops { +namespace jit { + + /**---------------------------------------------------------------------------* + * @brief Functor to get type name in string + * + * This functor uses the unofficial compiler macro __PRETTY_FUNCTION__ + * to obtain the function signature which contains the template type name. + * The type name is searched and extracted from the string. + * + * Example (clang): __PRETTY_FUNCTION__ = + * std::string type_name::operator()() [T = short] + * returns std::string("short") + * + * Example (gcc): __PRETTY_FUNCTION__ = + * std::__cxx11::string type_name::operator()() [with T = short int; std::__cxx11::string = std::__cxx11::basic_string] + * returns std::string("short int") + * + * In case the type is wrapped using `wrapper`, the extra string "wrapper<" is + * also skipped to get the underlying wrapped type. + * + *---------------------------------------------------------------------------**/ + struct type_name { + template + CUDA_HOST_DEVICE_CALLABLE + std::string operator()() { +#if defined(__clang__) || defined(__GNUC__) + std::string p = __PRETTY_FUNCTION__; + std::string search_str = "T = "; + size_t start_pos = p.find(search_str) + search_str.size(); + std::string wrapper_str = "wrapper<"; + size_t wrapper_pos = p.find(wrapper_str, start_pos); + if (wrapper_pos != std::string::npos) + start_pos = wrapper_pos + wrapper_str.size(); + size_t end_pos = p.find_first_of(",;]", start_pos); + return p.substr(start_pos, end_pos - start_pos); +#else +# error Only clang and gcc supported +#endif + } + }; + + /**---------------------------------------------------------------------------* + * @brief Get the Type Name + * + * @param type The data type + * @return std::string Name of the data type in string + *---------------------------------------------------------------------------**/ + std::string getTypeName(gdf_dtype type) { + return type_dispatcher(type, type_name()); + } + + /**---------------------------------------------------------------------------* + * @brief Get the Operator Name + * + * @param ope (enum) The binary operator as enum of type gdf_binary_operator + * @return std::string The name of the operator as string + *---------------------------------------------------------------------------**/ + std::string getOperatorName(gdf_binary_operator ope) { + switch (ope) { + case GDF_ADD: + return "Add"; + case GDF_SUB: + return "Sub"; + case GDF_MUL: + return "Mul"; + case GDF_DIV: + return "Div"; + case GDF_TRUE_DIV: + return "TrueDiv"; + case GDF_FLOOR_DIV: + return "FloorDiv"; + case GDF_MOD: + return "Mod"; + case GDF_POW: + return "Pow"; + case GDF_EQUAL: + return "Equal"; + case GDF_NOT_EQUAL: + return "NotEqual"; + case GDF_LESS: + return "Less"; + case GDF_GREATER: + return "Greater"; + case GDF_LESS_EQUAL: + return "LessEqual"; + case GDF_GREATER_EQUAL: + return "GreaterEqual"; + default: + return "None"; + } + } + +} // namespace jit +} // namespace binops +} // namespace cudf diff --git a/cpp/src/binary/jit/util/type.h b/cpp/src/binary/jit/util/type.h new file mode 100644 index 000000000000..769d9bd267f8 --- /dev/null +++ b/cpp/src/binary/jit/util/type.h @@ -0,0 +1,38 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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_UTIL_TYPE_H +#define GDF_BINARY_OPERATION_JIT_UTIL_TYPE_H + +#include "cudf.h" +#include + +namespace cudf { +namespace binops { +namespace jit { + + std::string getTypeName(gdf_dtype type); + + std::string getOperatorName(gdf_binary_operator ope); + +} +} +} + +#endif diff --git a/cpp/src/bitmask/bitmask_ops.cu b/cpp/src/bitmask/bitmask_ops.cu index 45ac3e697b9c..08a1a145b34b 100644 --- a/cpp/src/bitmask/bitmask_ops.cu +++ b/cpp/src/bitmask/bitmask_ops.cu @@ -12,48 +12,15 @@ #include #include -/* - * bit_mask_null_counts Generated using the following code - -#include - - -int main() -{ - for (int i = 0 ; i != 256 ; i++) { - int count = 0; - for (int p = 0 ; p != 8 ; p++) { - if (i & (1 << p)) { - count++; - } - } - std::cout<<(8-count)<<", "; - } - std::cout< bit_mask_null_counts = { 8, 7, 7, 6, 7, 6, 6, 5, 7, 6, 6, 5, 6, 5, 5, 4, 7, 6, 6, 5, 6, 5, 5, 4, 6, 5, 5, 4, 5, 4, 4, 3, 7, 6, 6, 5, 6, 5, 5, 4, 6, 5, 5, 4, 5, 4, 4, 3, 6, 5, 5, 4, 5, 4, 4, 3, 5, 4, 4, 3, 4, 3, 3, 2, 7, 6, 6, 5, 6, 5, 5, 4, 6, 5, 5, 4, 5, 4, 4, 3, 6, 5, 5, 4, 5, 4, 4, 3, 5, 4, 4, 3, 4, 3, 3, 2, 6, 5, 5, 4, 5, 4, 4, 3, 5, 4, 4, 3, 4, 3, 3, 2, 5, 4, 4, 3, 4, 3, 3, 2, 4, 3, 3, 2, 3, 2, 2, 1, 7, 6, 6, 5, 6, 5, 5, 4, 6, 5, 5, 4, 5, 4, 4, 3, 6, 5, 5, 4, 5, 4, 4, 3, 5, 4, 4, 3, 4, 3, 3, 2, 6, 5, 5, 4, 5, 4, 4, 3, 5, 4, 4, 3, 4, 3, 3, 2, 5, 4, 4, 3, 4, 3, 3, 2, 4, 3, 3, 2, 3, 2, 2, 1, 6, 5, 5, 4, 5, 4, 4, 3, 5, 4, 4, 3, 4, 3, 3, 2, 5, 4, 4, 3, 4, 3, 3, 2, 4, 3, 3, 2, 3, 2, 2, 1, 5, 4, 4, 3, 4, 3, 3, 2, 4, 3, 3, 2, 3, 2, 2, 1, 4, 3, 3, 2, 3, 2, 2, 1, 3, 2, 2, 1, 2, 1, 1, 0 }; - -unsigned char gdf_num_bits_zero_after_pos(unsigned char number, int pos){ - //if pos == 0 then its aligned - if(pos == 0){ - return 0; - } - unsigned char count = 0; - for (int p = pos ; p != 8 ; p++) { - if (number & (number << p)) { - count++; - } - } - return (8 - pos) - count; -} gdf_error all_bitmask_on(gdf_valid_type * valid_out, gdf_size_type & out_null_count, gdf_size_type num_values, cudaStream_t stream){ - gdf_size_type num_chars_bitmask = ( ( num_values +( GDF_VALID_BITSIZE - 1)) / GDF_VALID_BITSIZE ); + gdf_size_type num_chars_bitmask = gdf_get_num_chars_bitmask( num_values ); - thrust::device_ptr valid_out_ptr = thrust::device_pointer_cast(valid_out); gdf_valid_type max_char = 255; - thrust::fill(rmm::exec_policy(stream)->on(stream),thrust::detail::make_normal_iterator(valid_out_ptr),thrust::detail::make_normal_iterator(valid_out_ptr + num_chars_bitmask),max_char); + thrust::fill(rmm::exec_policy(stream)->on(stream), + valid_out, + valid_out + num_chars_bitmask, + max_char); //we have no nulls so set all the bits in gdf_valid_type to 1 out_null_count = 0; return GDF_SUCCESS; @@ -62,34 +29,17 @@ gdf_error all_bitmask_on(gdf_valid_type * valid_out, gdf_size_type & out_null_co gdf_error apply_bitmask_to_bitmask(gdf_size_type & out_null_count, gdf_valid_type * valid_out, gdf_valid_type * valid_left, gdf_valid_type * valid_right, cudaStream_t stream, gdf_size_type num_values){ - gdf_size_type num_chars_bitmask = ( ( num_values +( GDF_VALID_BITSIZE - 1)) / GDF_VALID_BITSIZE ); - thrust::device_ptr valid_out_ptr = thrust::device_pointer_cast(valid_out); - thrust::device_ptr valid_left_ptr = thrust::device_pointer_cast(valid_left); - //here we are basically figuring out what is the last pointed to unsigned char that can contain part of the bitmask - thrust::device_ptr valid_left_end_ptr = thrust::device_pointer_cast(valid_left + num_chars_bitmask ); - thrust::device_ptr valid_right_ptr = thrust::device_pointer_cast(valid_right); - - - thrust::transform(rmm::exec_policy(stream)->on(stream), thrust::detail::make_normal_iterator(valid_left_ptr), - thrust::detail::make_normal_iterator(valid_left_end_ptr), thrust::detail::make_normal_iterator(valid_right_ptr), - thrust::detail::make_normal_iterator(valid_out_ptr), thrust::bit_and()); + gdf_size_type num_chars_bitmask = gdf_get_num_chars_bitmask( num_values ); + thrust::transform(rmm::exec_policy(stream)->on(stream), + valid_left, + valid_left + num_chars_bitmask, + valid_right, + valid_out, + thrust::bit_and()); - char * last_char = new char[1]; - cudaError_t error = cudaMemcpyAsync(last_char,valid_out + ( num_chars_bitmask-1),sizeof(gdf_valid_type),cudaMemcpyDeviceToHost,stream); - - - rmm::device_vector bit_mask_null_counts_device(bit_mask_null_counts); - - //this permutation iterator makes it so that each char basically gets replaced with its number of null counts - //so if you sum up this perm iterator you add up all of the counts for null values per unsigned char - thrust::permutation_iterator::iterator,thrust::detail::normal_iterator > > - null_counts_iter( bit_mask_null_counts_device.begin(),thrust::detail::make_normal_iterator(valid_out_ptr)); - - //you will notice that we subtract the number of zeros we found in the last character - out_null_count = thrust::reduce(rmm::exec_policy(stream)->on(stream),null_counts_iter, null_counts_iter + num_chars_bitmask) - gdf_num_bits_zero_after_pos(*last_char,num_values % GDF_VALID_BITSIZE ); - - delete[] last_char; - return GDF_SUCCESS; + gdf_size_type non_nulls; + auto error = gdf_count_nonzero_mask(valid_out, num_values, &non_nulls); + out_null_count = num_values - non_nulls; + return error; } - diff --git a/cpp/src/bitmask/bitmask_ops.h b/cpp/src/bitmask/bitmask_ops.h index dedf1fc6f506..b74faf3b552b 100644 --- a/cpp/src/bitmask/bitmask_ops.h +++ b/cpp/src/bitmask/bitmask_ops.h @@ -3,11 +3,28 @@ #include - +/**---------------------------------------------------------------------------* + * @brief Sets all bits in input valid mask to 1 + * + * @param valid_out preallocated output valid mask + * @param out_null_count number of nulls (0 bits) in valid mask. Always set to 0 + * @param num_values number of values in column associated with output mask + * @param stream cuda stream to run in + * @return gdf_error + *---------------------------------------------------------------------------**/ gdf_error all_bitmask_on(gdf_valid_type * valid_out, gdf_size_type & out_null_count, gdf_size_type num_values, cudaStream_t stream); - - +/**---------------------------------------------------------------------------* + * @brief Computes bitwise AND on two valid masks and sets it in output + * + * @param out_null_count number of nulls (0 bits) in output valid mask + * @param valid_out preallocated mask to set the result values in + * @param valid_left input valid mask 1 + * @param valid_right input valid mask 2 + * @param stream cuda stream to run in + * @param num_values number of values in each input mask valid_left and valid_right + * @return gdf_error + *---------------------------------------------------------------------------**/ gdf_error apply_bitmask_to_bitmask(gdf_size_type & out_null_count, gdf_valid_type * valid_out, gdf_valid_type * valid_left, gdf_valid_type * valid_right, cudaStream_t stream, gdf_size_type num_values); #endif diff --git a/cpp/src/bitmask/valid_ops.cu b/cpp/src/bitmask/valid_ops.cu index 376d41d02f7e..97d4b4314819 100644 --- a/cpp/src/bitmask/valid_ops.cu +++ b/cpp/src/bitmask/valid_ops.cu @@ -41,7 +41,6 @@ constexpr int BITS_PER_MASK32 = GDF_VALID_BITSIZE * RATIO; constexpr int block_size = 256; - /* --------------------------------------------------------------------------*/ /** * @brief Kernel to count the number of set bits in a column's validity buffer @@ -128,7 +127,7 @@ void count_valid_bits(valid32_t const * const masks32, gdf_error gdf_count_nonzero_mask(gdf_valid_type const *masks, gdf_size_type num_rows, gdf_size_type *count) { - + // TODO: add a default parameter cudaStream_t stream = 0 when we move API to C++ if((nullptr == count)){return GDF_DATASET_EMPTY;} diff --git a/cpp/src/utilities/cudf_utils.h b/cpp/src/utilities/cudf_utils.h index a7391b8ed88d..c80ec4164ceb 100644 --- a/cpp/src/utilities/cudf_utils.h +++ b/cpp/src/utilities/cudf_utils.h @@ -11,8 +11,8 @@ #define CUDA_DEVICE_CALLABLE __device__ __forceinline__ #define CUDA_LAUNCHABLE __global__ #else -#define CUDA_HOST_DEVICE_CALLABLE -#define CUDA_DEVICE_CALLABLE +#define CUDA_HOST_DEVICE_CALLABLE inline +#define CUDA_DEVICE_CALLABLE inline #define CUDA_LAUNCHABLE #endif diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 1d90cda80ea3..194daf932811 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -103,7 +103,7 @@ set(BITMASK_TEST_SRC ConfigureTest(BITMASK_TEST "${BITMASK_TEST_SRC}") ################################################################################################### -# - bit_mask tests --------------------------------------------------------------------------------- +# - bit_mask tests -------------------------------------------------------------------------------- set(BIT_MASK_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/bitmask/bit_mask_test.cu") @@ -165,6 +165,16 @@ set(UNARY_TEST_SRC ConfigureTest(UNARY_TEST "${UNARY_TEST_SRC}") +################################################################################################### +# - binary tests ---------------------------------------------------------------------------------- + +set(BINARY_TEST_SRC + "${CMAKE_CURRENT_SOURCE_DIR}/binary/unit/binop-verify-input-test.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/binary/integration/binary-operation-operands-null-test.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/binary/integration/binary-operation-integration-test.cpp") + +ConfigureTest(BINARY_TEST "${BINARY_TEST_SRC}") + ################################################################################################### # - csv tests ------------------------------------------------------------------------------------- @@ -174,14 +184,14 @@ set(CSV_TEST_SRC ConfigureTest(CSV_TEST "${CSV_TEST_SRC}") ################################################################################################### -# - sort tests ------------------------------------------------------------------------------------- +# - sort tests ------------------------------------------------------------------------------------ set(SORT_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/sort/digitize_test.cu") ConfigureTest(SORT_TEST "${SORT_TEST_SRC}") ################################################################################################### -# - types tests ------------------------------------------------------------------------------------- +# - types tests ----------------------------------------------------------------------------------- set(TYPES_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/types/types_test.cpp" @@ -190,7 +200,7 @@ set(TYPES_TEST_SRC ConfigureTest(TYPES_TEST "${TYPES_TEST_SRC}") ################################################################################################### -# - copying tests ------------------------------------------------------------------------------------- +# - copying tests --------------------------------------------------------------------------------- set(COPYING_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/copying/gather_tests.cu" @@ -199,7 +209,7 @@ set(COPYING_TEST_SRC ConfigureTest(COPYING_TEST "${COPYING_TEST_SRC}") ################################################################################################### -# - utilities tests ------------------------------------------------------------------------------------- +# - utilities tests ------------------------------------------------------------------------------- set(UTILITIES_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/utilities_tests/column_wrapper_tests.cu") @@ -214,7 +224,8 @@ set(TRANSPOSE_TEST_SRC ConfigureTest(TRANSPOSE_TEST "${TRANSPOSE_TEST_SRC}") -# - table tests ------------------------------------------------------------------------------- +################################################################################################### +# - table tests ----------------------------------------------------------------------------------- set(TABLE_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/table/table_tests.cu") diff --git a/cpp/tests/binary/integration/assert-binops.h b/cpp/tests/binary/integration/assert-binops.h new file mode 100644 index 000000000000..eda36ff32609 --- /dev/null +++ b/cpp/tests/binary/integration/assert-binops.h @@ -0,0 +1,152 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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_TESTS_BINARY_OPERATION_INTEGRATION_ASSERT_BINOPS_H +#define GDF_TESTS_BINARY_OPERATION_INTEGRATION_ASSERT_BINOPS_H + +#include "tests/binary/util/operation.h" +#include "tests/utilities/column_wrapper.cuh" +#include "tests/utilities/scalar_wrapper.cuh" +#include "gtest/gtest.h" + +namespace cudf { +namespace test { +namespace binop { + +template +void ASSERT_BINOP(cudf::test::column_wrapper& out, + cudf::test::scalar_wrapper& lhs, + cudf::test::column_wrapper& rhs, + TypeOpe&& ope) { + auto lhs_h = lhs.value(); + auto rhs_h = rhs.to_host(); + auto rhs_data = std::get<0>(rhs_h); + auto out_h = out.to_host(); + auto out_data = std::get<0>(out_h); + + ASSERT_TRUE(out_data.size() == rhs_data.size()); + for (int index = 0; index < out_data.size(); ++index) { + ASSERT_TRUE(out_data[index] == (TypeOut)(ope(lhs_h, rhs_data[index]))); + } + + auto rhs_valid = std::get<1>(rhs_h); + auto out_valid = std::get<1>(out_h); + + uint32_t lhs_valid = (lhs.is_valid() ? UINT32_MAX : 0); + ASSERT_TRUE(out_valid.size() == rhs_valid.size()); + for (int index = 0; index < out_valid.size(); ++index) { + ASSERT_TRUE(out_valid[index] == (lhs_valid & rhs_valid[index])); + } +} + +template +void ASSERT_BINOP(cudf::test::column_wrapper& out, + cudf::test::column_wrapper& lhs, + cudf::test::scalar_wrapper& rhs, + TypeOpe&& ope) { + auto rhs_h = rhs.value(); + auto lhs_h = lhs.to_host(); + auto lhs_data = std::get<0>(lhs_h); + auto out_h = out.to_host(); + auto out_data = std::get<0>(out_h); + + ASSERT_TRUE(out_data.size() == lhs_data.size()); + for (int index = 0; index < out_data.size(); ++index) { + ASSERT_TRUE(out_data[index] == (TypeOut)(ope(lhs_data[index], rhs_h))); + } + + auto lhs_valid = std::get<1>(lhs_h); + auto out_valid = std::get<1>(out_h); + + uint32_t rhs_valid = (rhs.is_valid() ? UINT32_MAX : 0); + ASSERT_TRUE(out_valid.size() == lhs_valid.size()); + for (int index = 0; index < out_valid.size(); ++index) { + ASSERT_TRUE(out_valid[index] == (lhs_valid[index] & rhs_valid)); + } +} + +template +void ASSERT_BINOP(cudf::test::column_wrapper& out, + cudf::test::column_wrapper& lhs, + cudf::test::column_wrapper& rhs, + TypeOpe&& ope) { + auto lhs_h = lhs.to_host(); + auto lhs_data = std::get<0>(lhs_h); + auto rhs_h = rhs.to_host(); + auto rhs_data = std::get<0>(rhs_h); + auto out_h = out.to_host(); + auto out_data = std::get<0>(out_h); + + ASSERT_TRUE(out_data.size() == lhs_data.size()); + ASSERT_TRUE(out_data.size() == rhs_data.size()); + for (int index = 0; index < out_data.size(); ++index) { + ASSERT_TRUE(out_data[index] == (TypeOut)(ope(lhs_data[index], rhs_data[index]))); + } + + auto lhs_valid = std::get<1>(lhs_h); + auto rhs_valid = std::get<1>(rhs_h); + auto out_valid = std::get<1>(out_h); + + ASSERT_TRUE(out_valid.size() == lhs_valid.size()); + ASSERT_TRUE(out_valid.size() == rhs_valid.size()); + for (int index = 0; index < out_valid.size(); ++index) { + ASSERT_TRUE(out_valid[index] == lhs_valid[index] | rhs_valid[index]); + } +} + +/** + * According to CUDA Programming Guide, 'E.1. Standard Functions', 'Table 7 - Double-Precision + * Mathematical Standard Library Functions with Maximum ULP Error' + * The pow function has 2 (full range) maximum ulp error. + */ +template +void ASSERT_BINOP(cudf::test::column_wrapper& out, + cudf::test::column_wrapper& lhs, + cudf::test::column_wrapper& rhs, + cudf::library::operation::Pow&& ope) { + auto lhs_h = lhs.to_host(); + auto lhs_data = std::get<0>(lhs_h); + auto rhs_h = rhs.to_host(); + auto rhs_data = std::get<0>(rhs_h); + auto out_h = out.to_host(); + auto out_data = std::get<0>(out_h); + + const int ULP = 2.0; + ASSERT_TRUE(out_data.size() == lhs_data.size()); + ASSERT_TRUE(out_data.size() == rhs_data.size()); + for (int index = 0; index < out_data.size(); ++index) { + ASSERT_TRUE(abs(out_data[index] - (TypeOut)(ope(lhs_data[index], rhs_data[index]))) < ULP); + } + + auto lhs_valid = std::get<1>(lhs_h); + auto rhs_valid = std::get<1>(rhs_h); + auto out_valid = std::get<1>(out_h); + + ASSERT_TRUE(out_valid.size() == lhs_valid.size()); + ASSERT_TRUE(out_valid.size() == rhs_valid.size()); + for (int index = 0; index < out_valid.size(); ++index) { + ASSERT_TRUE(out_valid[index] == (lhs_valid[index] & rhs_valid[index])); + } +} + +} // namespace binop +} // namespace test +} // namespace cudf + +#endif diff --git a/cpp/tests/binary/integration/binary-operation-integration-test.cpp b/cpp/tests/binary/integration/binary-operation-integration-test.cpp new file mode 100644 index 000000000000..3c926f36194d --- /dev/null +++ b/cpp/tests/binary/integration/binary-operation-integration-test.cpp @@ -0,0 +1,270 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "tests/binary/integration/assert-binops.h" + +namespace cudf { +namespace test { +namespace binop { + +struct BinaryOperationIntegrationTest : public ::testing::Test { + BinaryOperationIntegrationTest() { + } + + virtual ~BinaryOperationIntegrationTest() { + } + + virtual void SetUp() { + } + + virtual void TearDown() { + } +}; + + +TEST_F(BinaryOperationIntegrationTest, Add_Scalar_Vector_SI32_FP32_SI64) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::scalar_wrapper{100}; + auto rhs = cudf::test::column_wrapper{100000, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{rhs.get()->size, true}; + + auto result = gdf_binary_operation_s_v(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Sub_Scalar_Vector_SI32_FP32_SI64) { + using SUB = cudf::library::operation::Sub; + + auto lhs = cudf::test::scalar_wrapper{10000}; + auto rhs = cudf::test::column_wrapper{100000, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{rhs.get()->size, true}; + + auto result = gdf_binary_operation_s_v(out.get(), lhs.get(), rhs.get(), GDF_SUB); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, SUB()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Vector_Scalar_SI08_SI16_SI32) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto rhs = cudf::test::scalar_wrapper{100}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_s(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Vector_Vector_SI32_FP64_SI08) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row * 2.0;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row * 1;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Sub_Vector_Vector_SI64) { + using SUB = cudf::library::operation::Sub; + + auto lhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 100000 + row * 2;}, + [](gdf_size_type row) {return (row % 4 == 0);}}; + auto rhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 50000 + row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_SUB); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, SUB()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mul_Vector_Vector_SI64) { + using MUL = cudf::library::operation::Mul; + + auto lhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 100000 + row * 2;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 50000 + row;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_MUL); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, MUL()); +} + + +TEST_F(BinaryOperationIntegrationTest, Div_Vector_Vector_SI64) { + using DIV = cudf::library::operation::Div; + + auto lhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 100000 + row * 2;}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto rhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 50000 + row;}, + [](gdf_size_type row) {return (row % 8 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_DIV); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, DIV()); +} + + +TEST_F(BinaryOperationIntegrationTest, TrueDiv_Vector_Vector_SI64) { + using TRUEDIV = cudf::library::operation::TrueDiv; + + auto lhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 100000 + row * 2;}, + [](gdf_size_type row) {return (row % 3 == 0);}}; + auto rhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 50000 + row;}, + [](gdf_size_type row) {return (row % 4 == 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_TRUE_DIV); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, TRUEDIV()); +} + + +TEST_F(BinaryOperationIntegrationTest, FloorDiv_Vector_Vector_SI64) { + using FLOORDIV = cudf::library::operation::FloorDiv; + + auto lhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 100000 + row * 2;}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto rhs = cudf::test::column_wrapper{50000, + [](gdf_size_type row) {return 50000 + row;}, + [](gdf_size_type row) {return (row % 8 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_FLOOR_DIV); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, FLOORDIV()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mod_Vector_Vector_SI64) { + using MOD = cudf::library::operation::Mod; + + auto lhs = cudf::test::column_wrapper{50, + [](gdf_size_type row) {return 120 + row * 2;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::column_wrapper{50, + [](gdf_size_type row) {return 50 + row;}, + [](gdf_size_type row) {return (row % 5 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_MOD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, MOD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mod_Vector_Vector_FP32) { + using MOD = cudf::library::operation::Mod; + + auto lhs = cudf::test::column_wrapper{50, + [](gdf_size_type row) {return 120 + row * 2;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto rhs = cudf::test::column_wrapper{50, + [](gdf_size_type row) {return 50 + row;}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_MOD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, MOD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mod_Vector_Vector_FP64) { + using MOD = cudf::library::operation::Mod; + + auto lhs = cudf::test::column_wrapper{50, + [](gdf_size_type row) {return 120 + row * 2;}, + [](gdf_size_type row) {return (row % 3 == 0);}}; + auto rhs = cudf::test::column_wrapper{50, + [](gdf_size_type row) {return 50 + row;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_MOD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, MOD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Pow_Vector_Vector_SI64) { + using POW = cudf::library::operation::Pow; + + auto lhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 6 > 0);}}; + auto rhs = cudf::test::column_wrapper{500, + [](gdf_size_type row) {return 2;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_POW); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, POW()); +} + +} // namespace binop +} // namespace test +} // namespace cudf diff --git a/cpp/tests/binary/integration/binary-operation-operands-null-test.cpp b/cpp/tests/binary/integration/binary-operation-operands-null-test.cpp new file mode 100644 index 000000000000..d8bf5945d91b --- /dev/null +++ b/cpp/tests/binary/integration/binary-operation-operands-null-test.cpp @@ -0,0 +1,154 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "tests/binary/integration/assert-binops.h" + +namespace cudf { +namespace test { +namespace binop { + +struct BinaryOperationOperandsNullTest : public ::testing::Test { + BinaryOperationOperandsNullTest() { + } + + virtual ~BinaryOperationOperandsNullTest() { + } + + virtual void SetUp() { + } + + virtual void TearDown() { + } +}; + +/* + * Kernels v_v_s, using int64_t + * Output:Vector, OperandX:Vector, OperandY:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_SI64_WithScalarOperandNull) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::scalar_wrapper{500, false}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_s(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_SI64_WithScalarOperandNotNull) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::scalar_wrapper{500, true}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_s(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + +/* + * Kernels v_v_s, using double + * Output:Vector, OperandX:Vector, OperandY:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_FP64_WithScalarOperandNull) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::scalar_wrapper{500, false}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_s(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_FP64_WithScalarOperandNotNull) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::scalar_wrapper{500, true}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_s(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + +/* + * Kernels v_v_v, using int64_t + * Output:Vector, OperandX:Vector, OperandY:Vector + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_int64_t) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row * 2;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + +/* + * Kernels v_v_v, using double + * Output:Vector, OperandX:Vector, OperandY:Vector + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_FP64) { + using ADD = cudf::library::operation::Add; + + auto lhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return (row % 3 > 0);}}; + auto rhs = cudf::test::column_wrapper{100, + [](gdf_size_type row) {return row * 2;}, + [](gdf_size_type row) {return (row % 4 > 0);}}; + auto out = cudf::test::column_wrapper{lhs.get()->size, true}; + + auto result = gdf_binary_operation_v_v(out.get(), lhs.get(), rhs.get(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + ASSERT_BINOP(out, lhs, rhs, ADD()); +} + +} // namespace binop +} // namespace test +} // namespace gdf diff --git a/cpp/tests/binary/unit/binop-verify-input-test.cpp b/cpp/tests/binary/unit/binop-verify-input-test.cpp new file mode 100644 index 000000000000..fa4a5a7392a3 --- /dev/null +++ b/cpp/tests/binary/unit/binop-verify-input-test.cpp @@ -0,0 +1,302 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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. + */ + +#include "tests/utilities/column_wrapper.cuh" +#include "tests/utilities/scalar_wrapper.cuh" +#include "gtest/gtest.h" +#include +#include + +struct BinopVerifyInputTest : public ::testing::Test { + BinopVerifyInputTest() { + } + + virtual ~BinopVerifyInputTest() { + } + + virtual void SetUp() { + } + + virtual void TearDown() { + } +}; + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOutputVectorZeroSize) { + auto vector_out = cudf::test::column_wrapper{0}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto scalar = cudf::test::scalar_wrapper{100}; + + auto result = gdf_binary_operation_v_s(vector_out, vector_lhs, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandVectorZeroSize) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{0}; + + auto scalar = cudf::test::scalar_wrapper{100}; + + auto result = gdf_binary_operation_v_s(vector_out, vector_lhs, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOutputVectorNull) { + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto scalar = cudf::test::scalar_wrapper{100}; + + auto result = gdf_binary_operation_v_s(nullptr, vector_lhs, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandVectorNull) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto scalar = cudf::test::scalar_wrapper{100}; + + auto result = gdf_binary_operation_v_s(vector_out, nullptr, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandScalarNull) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_s(vector_out, vector_lhs, nullptr, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOutputVectorType) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + vector_out.get()->dtype = (gdf_dtype)100; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto scalar = cudf::test::scalar_wrapper{100}; + + auto result = gdf_binary_operation_v_s(vector_out, vector_lhs, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandVectorType) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + vector_lhs.get()->dtype = (gdf_dtype)100; + + auto scalar = cudf::test::scalar_wrapper{100}; + + auto result = gdf_binary_operation_v_s(vector_out, vector_lhs, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandScalarType) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto scalar = cudf::test::scalar_wrapper{100}; + scalar.get()->dtype = (gdf_dtype)100; + + auto result = gdf_binary_operation_v_s(vector_out, vector_lhs, scalar, GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorOutputVectorZeroSize) { + auto vector_out = cudf::test::column_wrapper{0}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorFirstOperandVectorZeroSize) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{0}; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorSecondOperandVectorZeroSize) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_rhs = cudf::test::column_wrapper{0}; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorOutputVectorNull) { + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(nullptr, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorFirstOperandVectorNull) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(vector_out, nullptr, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorSecondOperandVectorNull) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, nullptr, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorOutputVectorType) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + vector_out.get()->dtype = (gdf_dtype)100; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorFirstOperandVectorType) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + vector_lhs.get()->dtype = (gdf_dtype)100; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorSecondOperandVectorType) { + auto vector_out = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_lhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + + auto vector_rhs = cudf::test::column_wrapper{10, + [](gdf_size_type row) {return row;}, + [](gdf_size_type row) {return true;}}; + vector_rhs.get()->dtype = (gdf_dtype)100; + + auto result = gdf_binary_operation_v_v(vector_out, vector_lhs, vector_rhs, GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} diff --git a/cpp/tests/binary/util/operation.h b/cpp/tests/binary/util/operation.h new file mode 100644 index 000000000000..299f7347f9ed --- /dev/null +++ b/cpp/tests/binary/util/operation.h @@ -0,0 +1,115 @@ +/* + * Copyright (c) 2019, NVIDIA CORPORATION. + * + * Copyright 2018-2019 BlazingDB, Inc. + * Copyright 2018 Christian Noboa Mardini + * + * 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_TESTS_BINARY_OPERATION_UTIL_OPERATION_H +#define GDF_TESTS_BINARY_OPERATION_UTIL_OPERATION_H + +#include +#include +#include + +namespace cudf { +namespace library { +namespace operation { + + template + struct Add { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)lhs + (TypeCommon)rhs); + } + }; + + template + struct Sub { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)lhs - (TypeCommon)rhs); + } + }; + + template + struct Mul { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)lhs * (TypeCommon)rhs); + } + }; + + template + struct Div { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)lhs / (TypeCommon)rhs); + } + }; + + template + struct TrueDiv { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + return (TypeOut)((double)lhs / (double)rhs); + } + }; + + template + struct FloorDiv { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + return (TypeOut)floor((double)lhs / (double)rhs); + } + }; + + template ::type> + struct Mod; + + template + struct Mod { + TypeOut operator()(TypeLhs x, TypeRhs y) { + return (TypeOut)((int64_t)x % (int64_t)y); + } + }; + + template + struct Mod { + TypeOut operator()(TypeLhs x, TypeRhs y) { + return (TypeOut)fmod((float)x, (float)y); + } + }; + + template + struct Mod { + TypeOut operator()(TypeLhs x, TypeRhs y) { + return (TypeOut)fmod((double)x, (double)y); + } + }; + + template + struct Pow { + TypeOut operator()(TypeLhs lhs, TypeRhs rhs) { + return (TypeOut)pow((double)lhs, (double)rhs); + } + }; + +} // namespace operation +} // namespace library +} // namespace cudf + +#endif diff --git a/cpp/tests/utilities/scalar_wrapper.cuh b/cpp/tests/utilities/scalar_wrapper.cuh new file mode 100644 index 000000000000..166c4d21b726 --- /dev/null +++ b/cpp/tests/utilities/scalar_wrapper.cuh @@ -0,0 +1,122 @@ +/* + * Copyright (c) 2018, NVIDIA CORPORATION. + * + * 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 SCALAR_WRAPPER_H +#define SCALAR_WRAPPER_H + +#include "cudf.h" +#include "cudf_test_utils.cuh" +#include "rmm/rmm.h" +#include "utilities/bit_util.cuh" +#include "utilities/type_dispatcher.hpp" + +#include +#include +#include + +namespace cudf { +namespace test { + +/**---------------------------------------------------------------------------* + * @brief Wrapper for a gdf_scalar used for unit testing. + * + * An abstraction on top of a gdf_scalar that provides functionality for + * allocating, intiailizing, and otherwise managing gdf_scalars for passing to + * libcudf APIs in unit testing. + * + * @tparam ColumnType The underlying data type of the scalar + *---------------------------------------------------------------------------**/ +template +struct scalar_wrapper { + /**---------------------------------------------------------------------------* + * @brief Implicit conversion operator to a gdf_scalar pointer. + * + * In this way, a column_wrapper can be passed directly into a libcudf API + * and will be implicitly converted to a pointer to its underlying gdf_scalar + * without the need to use the `get()` member. + * + * @return gdf_scalar* Pointer to the underlying gdf_scalar + *---------------------------------------------------------------------------**/ + operator gdf_scalar*(){return &the_scalar;}; + + /**---------------------------------------------------------------------------* + * @brief Construct a new scalar wrapper object + * + * Constructs a scalar_wrapper using a value of type ScalarType. + * + * @param value The value to use for the scalar + * @param is_valid is the scalar valid + *---------------------------------------------------------------------------**/ + scalar_wrapper(ScalarType value, bool is_valid = true) { + auto dataptr = reinterpret_cast(&(the_scalar.data)); + *dataptr = value; + the_scalar.is_valid = is_valid; + the_scalar.dtype = gdf_dtype_of(); + } + + /**---------------------------------------------------------------------------* + * @brief Returns a pointer to the underlying gdf_scalar. + *---------------------------------------------------------------------------**/ + gdf_scalar* get() { return &the_scalar; } + gdf_scalar const* get() const { return &the_scalar; } + + /**---------------------------------------------------------------------------* + * @brief returns the value of the scalar + * + * @return ScalarType + *---------------------------------------------------------------------------**/ + ScalarType value() { + return *reinterpret_cast(&(the_scalar.data)); + } + + /**---------------------------------------------------------------------------* + * @brief returns the validity of the scalar + * + * @return bool + *---------------------------------------------------------------------------**/ + bool is_valid() { return the_scalar.is_valid; } + + /**---------------------------------------------------------------------------* + * @brief Prints the value of the underlying gdf_scalar. + *---------------------------------------------------------------------------**/ + void print() const { + ScalarType value = *reinterpret_cast(&(the_scalar.data)); + if (the_scalar.is_valid) + std::cout << value << std::endl; + else + std::cout << "null" << std::endl; + } + + /**---------------------------------------------------------------------------* + * @brief CompCompares this wrapper with another scalar_wrapper for equality. + * + * @param rhs The other scalar_wrapper to check for equality + * @return true The two scalars are equal + * @return false The two scalars are not equal + *---------------------------------------------------------------------------**/ + bool operator==(scalar_wrapper const& rhs) const { + return (the_scalar.data == rhs.data && + the_scalar.dtype == rhs.dtype && + the_scalar.is_valid == rhs.is_valid); + } + + private: + gdf_scalar the_scalar; +}; + +} // namespace test +} // namespace cudf +#endif diff --git a/cpp/tests/utilities/tuple_vectors.h b/cpp/tests/utilities/tuple_vectors.h index 8b01b067b434..af38d1a780df 100644 --- a/cpp/tests/utilities/tuple_vectors.h +++ b/cpp/tests/utilities/tuple_vectors.h @@ -20,6 +20,7 @@ #include #include #include +#include template using VTuple = std::tuple...>; diff --git a/python/cudf/bindings/binops.pyx b/python/cudf/bindings/binops.pyx new file mode 100644 index 000000000000..59a3b95208a9 --- /dev/null +++ b/python/cudf/bindings/binops.pyx @@ -0,0 +1,65 @@ +# Copyright (c) 2018, NVIDIA CORPORATION. + +# cython: profile=False +# distutils: language = c++ +# cython: embedsignature = True +# cython: language_level = 3 + +# Copyright (c) 2018, NVIDIA CORPORATION. + +from .cudf_cpp cimport * +from .cudf_cpp import * + +from librmm_cffi import librmm as rmm + +from libc.stdlib cimport free + + +_BINARY_OP = {} +_BINARY_OP['add'] = GDF_ADD +_BINARY_OP['sub'] = GDF_SUB +_BINARY_OP['mul'] = GDF_MUL +_BINARY_OP['div'] = GDF_DIV +_BINARY_OP['truediv'] = GDF_TRUE_DIV +_BINARY_OP['floordiv'] = GDF_FLOOR_DIV +_BINARY_OP['mod'] = GDF_MOD +_BINARY_OP['pow'] = GDF_POW +_BINARY_OP['eq'] = GDF_EQUAL +_BINARY_OP['ne'] = GDF_NOT_EQUAL +_BINARY_OP['lt'] = GDF_LESS +_BINARY_OP['gt'] = GDF_GREATER +_BINARY_OP['le'] = GDF_LESS_EQUAL +_BINARY_OP['ge'] = GDF_GREATER_EQUAL + + +def apply_op(lhs, rhs, out, op): + """ + Call JITified gdf binary ops. + """ + + check_gdf_compatibility(lhs) + check_gdf_compatibility(rhs) + check_gdf_compatibility(out) + + cdef gdf_column* c_lhs = column_view_from_column(lhs) + cdef gdf_column* c_rhs = column_view_from_column(rhs) + cdef gdf_column* c_out = column_view_from_column(out) + + cdef gdf_error result + cdef gdf_binary_operator c_op = _BINARY_OP[op] + with nogil: + result = gdf_binary_operation_v_v( + c_out, + c_lhs, + c_rhs, + c_op) + + cdef int nullct = c_out[0].null_count + + free(c_lhs) + free(c_rhs) + free(c_out) + + check_gdf_error(result) + + return nullct diff --git a/python/cudf/bindings/cudf_cpp.pxd b/python/cudf/bindings/cudf_cpp.pxd index eac5cd295edb..08b5ce3d494d 100644 --- a/python/cudf/bindings/cudf_cpp.pxd +++ b/python/cudf/bindings/cudf_cpp.pxd @@ -189,6 +189,28 @@ cdef extern from "cudf.h" nogil: GDF_WINDOW_VA + ctypedef enum gdf_binary_operator: + GDF_ADD, + GDF_SUB, + GDF_MUL, + GDF_DIV, + GDF_TRUE_DIV, + GDF_FLOOR_DIV, + GDF_MOD, + GDF_POW, + GDF_EQUAL, + GDF_NOT_EQUAL, + GDF_LESS, + GDF_GREATER, + GDF_LESS_EQUAL, + GDF_GREATER_EQUAL + + + ctypedef struct gdf_scalar: + void *data + gdf_dtype dtype + + cdef gdf_error gdf_nvtx_range_push(char * name, gdf_color color ) cdef gdf_error gdf_nvtx_range_push_hex(char * name, unsigned int color ) @@ -490,6 +512,10 @@ cdef extern from "cudf.h" nogil: cdef gdf_error gdf_extract_datetime_minute(gdf_column *input, gdf_column *output) cdef gdf_error gdf_extract_datetime_second(gdf_column *input, gdf_column *output) + cdef gdf_error gdf_binary_operation_s_v(gdf_column* out, gdf_scalar* lhs, gdf_column* rhs, gdf_binary_operator ope) + cdef gdf_error gdf_binary_operation_v_s(gdf_column* out, gdf_column* lhs, gdf_scalar* rhs, gdf_binary_operator ope) + cdef gdf_error gdf_binary_operation_v_v(gdf_column* out, gdf_column* lhs, gdf_column* rhs, gdf_binary_operator ope) + cdef gdf_error gdf_add_generic(gdf_column *lhs, gdf_column *rhs, gdf_column *output) cdef gdf_error gdf_add_i32(gdf_column *lhs, gdf_column *rhs, gdf_column *output) cdef gdf_error gdf_add_i64(gdf_column *lhs, gdf_column *rhs, gdf_column *output) diff --git a/python/cudf/dataframe/index.py b/python/cudf/dataframe/index.py index d2ea51c085bd..d295ffde02ca 100644 --- a/python/cudf/dataframe/index.py +++ b/python/cudf/dataframe/index.py @@ -145,6 +145,12 @@ def __mul__(self, other): def __rmul__(self, other): return self._apply_op('__rmul__', other) + def __mod__(self, other): + return self._apply_op('__mod__', other) + + def __rmod__(self, other): + return self._apply_op('__rmod__', other) + def __pow__(self, other): return self._apply_op('__pow__', other) diff --git a/python/cudf/dataframe/numerical.py b/python/cudf/dataframe/numerical.py index e80d7fb8eb8f..a1753fd9cf25 100644 --- a/python/cudf/dataframe/numerical.py +++ b/python/cudf/dataframe/numerical.py @@ -19,26 +19,21 @@ import cudf.bindings.reduce as cpp_reduce import cudf.bindings.replace as cpp_replace +import cudf.bindings.binops as cpp_binops import cudf.bindings.sort as cpp_sort # Operator mappings -# Unordered comparators -_unordered_impl = { +_binary_impl = { + # Unordered comparators 'eq': libgdf.gdf_eq_generic, 'ne': libgdf.gdf_ne_generic, -} - -# Ordered comparators -_ordered_impl = { + # Ordered comparators 'lt': libgdf.gdf_lt_generic, 'le': libgdf.gdf_le_generic, 'gt': libgdf.gdf_gt_generic, 'ge': libgdf.gdf_ge_generic, -} - -# Binary operators -_binary_impl = { + # Binary operators 'add': libgdf.gdf_add_generic, 'sub': libgdf.gdf_sub_generic, 'mul': libgdf.gdf_mul_generic, @@ -90,10 +85,9 @@ def deserialize(cls, deserialize, header, frames): def binary_operator(self, binop, rhs): if isinstance(rhs, NumericalColumn): - op = _binary_impl[binop] - lhs, rhs = numeric_normalize_types(self, rhs) - return numeric_column_binop(lhs=lhs, rhs=rhs, op=op, - out_dtype=lhs.dtype) + out_dtype = np.result_type(self.dtype, rhs.dtype) + return numeric_column_binop(lhs=self, rhs=rhs, op=binop, + out_dtype=out_dtype) else: msg = "{!r} operator not supported between {} and {}" raise TypeError(msg.format(binop, type(self), type(rhs))) @@ -103,12 +97,10 @@ def unary_operator(self, unaryop): out_dtype=self.dtype) def unordered_compare(self, cmpop, rhs): - lhs, rhs = numeric_normalize_types(self, rhs) - return numeric_column_compare(lhs, rhs, op=_unordered_impl[cmpop]) + return numeric_column_compare(self, rhs, op=cmpop) def ordered_compare(self, cmpop, rhs): - lhs, rhs = numeric_normalize_types(self, rhs) - return numeric_column_compare(lhs, rhs, op=_ordered_impl[cmpop]) + return numeric_column_compare(self, rhs, op=cmpop) def normalize_binop_value(self, other): other_dtype = np.min_scalar_type(other) @@ -402,15 +394,18 @@ def find_and_replace(self, to_replace, value): def numeric_column_binop(lhs, rhs, op, out_dtype): - if lhs.dtype != rhs.dtype: - raise TypeError('{} != {}'.format(lhs.dtype, rhs.dtype)) - nvtx_range_push("CUDF_BINARY_OP", "orange") # Allocate output masked = lhs.has_null_mask or rhs.has_null_mask out = columnops.column_empty_like(lhs, dtype=out_dtype, masked=masked) # Call and fix null_count - null_count = _gdf.apply_binaryop(op, lhs, rhs, out) + if lhs.dtype != rhs.dtype or op not in _binary_impl: + # Use JIT implementation + null_count = cpp_binops.apply_op(lhs=lhs, rhs=rhs, out=out, op=op) + else: + # Use compiled implementation + null_count = _gdf.apply_binaryop(_binary_impl[op], lhs, rhs, out) + out = out.replace(null_count=null_count) result = out.view(NumericalColumn, dtype=out_dtype) nvtx_range_pop() diff --git a/python/cudf/dataframe/series.py b/python/cudf/dataframe/series.py index ebd2135f0100..07e565edc515 100644 --- a/python/cudf/dataframe/series.py +++ b/python/cudf/dataframe/series.py @@ -402,11 +402,14 @@ def __mul__(self, other): def __rmul__(self, other): return self._rbinaryop(other, 'mul') + def __mod__(self, other): + return self._binaryop(other, 'mod') + + def __rmod__(self, other): + return self._rbinaryop(other, 'mod') + def __pow__(self, other): - if other == 2: - return self * self - else: - return NotImplemented + return self._binaryop(other, 'pow') def __floordiv__(self, other): return self._binaryop(other, 'floordiv') diff --git a/python/cudf/tests/test_binops.py b/python/cudf/tests/test_binops.py index f510754f6081..46752481e24c 100644 --- a/python/cudf/tests/test_binops.py +++ b/python/cudf/tests/test_binops.py @@ -21,6 +21,8 @@ operator.mul, operator.floordiv, operator.truediv, + operator.mod, + operator.pow, ] @@ -38,7 +40,7 @@ def test_series_binop(binop, obj_class): if obj_class == 'Index': result = Series(result) - np.testing.assert_equal(result.to_array(), binop(arr, arr)) + np.testing.assert_almost_equal(result.to_array(), binop(arr, arr)) @pytest.mark.parametrize('obj_class', ['Series', 'Index']) @@ -55,7 +57,7 @@ def test_series_binop_scalar(nelem, binop, obj_class): if obj_class == 'Index': result = Series(result) - np.testing.assert_equal(result.to_array(), binop(arr, rhs)) + np.testing.assert_almost_equal(result.to_array(), binop(arr, rhs)) _cmpops = [ @@ -70,7 +72,7 @@ def test_series_binop_scalar(nelem, binop, obj_class): @pytest.mark.parametrize('obj_class', ['Series', 'Index']) @pytest.mark.parametrize('cmpop', _cmpops) -@pytest.mark.parametrize('dtype', ['int8', 'int16', 'int32', 'int64', +@pytest.mark.parametrize('dtype', ['int8', 'int32', 'int64', 'float32', 'float64', 'datetime64[ms]']) def test_series_compare(cmpop, obj_class, dtype): arr1 = np.random.randint(0, 100, 100).astype(dtype) diff --git a/thirdparty/jitify b/thirdparty/jitify new file mode 160000 index 000000000000..c7c55b38333f --- /dev/null +++ b/thirdparty/jitify @@ -0,0 +1 @@ +Subproject commit c7c55b38333f5fa9ad5ec5e0804d5c9eedd14de2