diff --git a/.gitignore b/.gitignore index 770c680d..e3f21998 100644 --- a/.gitignore +++ b/.gitignore @@ -18,3 +18,4 @@ python/libgdf_cffi/libgdf_cffi.py ## eclipse .project +build-optimized diff --git a/.gitmodules b/.gitmodules index f9570491..be9c5403 100644 --- a/.gitmodules +++ b/.gitmodules @@ -4,3 +4,6 @@ [submodule "thirdparty/moderngpu"] path = thirdparty/moderngpu url = https://github.com/moderngpu/moderngpu.git +[submodule "thirdparty/jitify"] + path = thirdparty/jitify + url = https://github.com/NVIDIA/jitify.git diff --git a/CMakeLists.txt b/CMakeLists.txt index 970d29b2..a4f36c96 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -31,6 +31,7 @@ set(CMAKE_CXX_STANDARD 11) message(STATUS "Using C++ standard: c++${CMAKE_CXX_STANDARD}") set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/Modules/" ${CMAKE_MODULE_PATH}) +set(CMAKE_MODULE_PATH "${CMAKE_MODULE_PATH}" "${CMAKE_CURRENT_SOURCE_DIR}/cmake/Configuration/") message(STATUS "CMAKE_MODULE_PATH:" "${CMAKE_MODULE_PATH}") IF(CMAKE_COMPILER_IS_GNUCXX) @@ -66,6 +67,9 @@ if(CUDA_FOUND) set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-Werror cross-execution-space-call;-Xcompiler -Wall,-Werror) message(STATUS "CUDA_NVCC_FLAGS: ${CUDA_NVCC_FLAGS}") + + include(ConfigCuda RESULT_VARIABLE result) + message(STATUS "ConfigCuda: ${result}") else() message(FATAL_ERROR "CUDA not found, please check your settings.") endif() @@ -84,24 +88,18 @@ else() endif() include_directories( + "${CMAKE_CURRENT_SOURCE_DIR}/src" "${CMAKE_CURRENT_SOURCE_DIR}/include" "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/cub" + "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/jitify" "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/moderngpu/src" "${CUDA_INCLUDE_DIRS}" "${ARROW_INCLUDEDIR}" ) -IF(CUDA_VERSION_MAJOR GREATER 7) - set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode;arch=compute_60,code=sm_60) -ENDIF(CUDA_VERSION_MAJOR GREATER 7) - -IF(CUDA_VERSION_MAJOR GREATER 8) - set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode;arch=compute_70,code=sm_70) -ENDIF(CUDA_VERSION_MAJOR GREATER 8) +link_directories(${CUDA_LIBRARY_DIR} ${CUDA_LIBRARY_STUBS_DIR}) -IF(CMAKE_BUILD_TYPE MATCHES DEBUG) - set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-g;-G;) -ENDIF(CMAKE_BUILD_TYPE MATCHES DEBUG) +set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode;arch=compute_61,code=sm_61) # switch between hash-join and sort-join option(HASH_JOIN "hash-based join" ON) @@ -119,7 +117,18 @@ if(HT_LEGACY_ALLOCATOR) endif() -cuda_add_library(gdf SHARED +## Benchmark +include(ConfigBenchmark RESULT_VARIABLE result) +message(STATUS "ConfigBenchmark: ${result}") + + +## Binary Operations +add_subdirectory(src/binary-operation) +get_directory_property(binary_operation_source_files DIRECTORY src/binary-operation DEFINITION source_files) + + +## Variable 'libgdf_source_files' with all source files. +set(libgdf_source_files src/binaryops.cu src/bitmaskops.cu src/column.cpp @@ -142,11 +151,34 @@ cuda_add_library(gdf SHARED #src/windowedops.cu src/quantiles.cu src/io/csv/csv-reader.cu - src/io/convert/gdf-to-csr.cu + src/io/convert/gdf-to-csr.cu src/validops.cu src/nvtx_utils.cpp + ${binary_operation_source_files} ) + +cuda_add_library(gdf SHARED ${libgdf_source_files}) + + +## Position Independent Code for shared library +set_target_properties(gdf PROPERTIES POSITION_INDEPENDENT_CODE ON) +set_target_properties(gdf PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + + +## Build static library +## By default is disabled. +## Add -DLIBGDF_STATIC_LIB=ON as cmake parameter to create the static library. +option(LIBGDF_STATIC_LIB "Build static library for libgdf" OFF) +if(LIBGDF_STATIC_LIB) + message(STATUS "libgdf static library enabled") + cuda_add_library(gdf_static STATIC ${libgdf_source_files}) + set_target_properties(gdf_static PROPERTIES OUTPUT_NAME gdf) +else() + message(STATUS "libgdf static library disabled") +endif() + + # Switch to enable NVTX ranges for profiling option(USE_NVTX "Build with NVTX" ON) if(USE_NVTX) @@ -154,13 +186,26 @@ if(USE_NVTX) find_library(NVTX_LIBRARY nvToolsExt HINTS /usr/local/cuda/lib64) target_link_libraries(gdf ${NVTX_LIBRARY}) set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS}; -DUSE_NVTX) + + if(LIBGDF_STATIC_LIB) + target_link_libraries(gdf_static ${NVTX_LIBRARY}) + endif() endif(USE_NVTX) + if (ARROW_NUMERIC_VERSION EQUAL 900) # arrow-0.9 requires boost_regex - target_link_libraries(gdf arrow boost_regex) + target_link_libraries(gdf arrow boost_regex ${CUDA_CUDA_LIB} ${CUDA_NVRTC_LIB}) + + if(LIBGDF_STATIC_LIB) + target_link_libraries(gdf_static arrow boost_regex ${CUDA_CUDA_LIB} ${CUDA_NVRTC_LIB}) + endif() else() - target_link_libraries(gdf arrow) + target_link_libraries(gdf arrow ${CUDA_CUDA_LIB} ${CUDA_NVRTC_LIB}) + + if(LIBGDF_STATIC_LIB) + target_link_libraries(gdf_static arrow ${CUDA_CUDA_LIB} ${CUDA_NVRTC_LIB}) + endif() endif() # Command to symlink files into the build directory @@ -193,6 +238,9 @@ add_custom_command(TARGET pytest POST_BUILD # The install target install(TARGETS gdf LIBRARY DESTINATION lib) install(DIRECTORY include/gdf DESTINATION include) +if(LIBGDF_STATIC_LIB) + install(TARGETS gdf_static ARCHIVE DESTINATION lib) +endif() # Configure the C++ tests find_package(GTest QUIET) diff --git a/README.md b/README.md index 283415f0..0e809e56 100644 --- a/README.md +++ b/README.md @@ -66,7 +66,7 @@ not necessary. You can install Boost C++ 1.58 from sources as well: https://www.boost.org/doc/libs/1_58_0/more/getting_started/unix-variants.html -To run the python tests it is recommended to setup a conda environment for +To run the python tests it is recommended to setup a conda environment for the dependencies. ```bash @@ -114,6 +114,12 @@ $ cmake -DCMAKE_INSTALL_PREFIX=$CONDA_PREFIX .. To build the C/C++ code, run `make`. This should produce a shared library named `libgdf.so` or `libgdf.dylib`. +To build the static library for libgdf just run cmake with: +```bash +$ cmake -DLIBGDF_STATIC_LIB=ON -DCMAKE_INSTALL_PREFIX=$CONDA_PREFIX .. +``` +Note: By default the option LIBGDF_STATIC_LIB is disabled (default value OFF). + If you run into compile errors about missing header files: ```bash @@ -138,3 +144,36 @@ conda environment activated), run below to exceute test: ```bash $ make pytest # this auto trigger target "copy_python" ``` + +### Benchmark + +It is used [Google Benchmark](https://github.com/google/benchmark) for this purpose. +Refer to that repository in order to get more information related to its functionality. + +In order to enable the benchmark functionality, which by default is disabled, +it is required to add the option BENCHMARK in the command line when 'cmake' is executed. +It is also recommended to use release build type. + +```bash +$ cmake -DCMAKE_BUILD_TYPE=Release -DBENCHMARK=ON path/to/libgdf +``` + +When the build process is completed, the benchmark binary files will be located in the 'bench' folder. +In order to get the different input arguments for a benchmark binary file, +add the argument '--help' in the command line when the file is executed. + +The benchmark is applied to the following features: +* Binary Operations (NVRTC - Jitify). + +###### Binary Operations + +Recommended input arguments when this benchmark is executed. +```bash +$ BinaryOperationsBench --benchmark_counters_tabular=true +``` + +### Benchmark - Troubleshoot + +When a message related to CPU scaling (`***WARNING*** CPU scaling is enabled`) appeared, +refer to the section [Disable CPU frequency scaling](https://github.com/google/benchmark#disable-cpu-frequency-scaling) +in the Google Benchmark documentation. diff --git a/cmake/Configuration/ConfigBenchmark.cmake b/cmake/Configuration/ConfigBenchmark.cmake new file mode 100644 index 00000000..b7124207 --- /dev/null +++ b/cmake/Configuration/ConfigBenchmark.cmake @@ -0,0 +1,28 @@ +#============================================================================= +# 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. +#============================================================================= + +## Build Benchmark +## By default is disabled. +## Add -DBENCHMARK:BOOL=ON as cmake parameter to create the benchmark. + +option(BENCHMARK "Benchmark" OFF) + +if(BENCHMARK) + message(STATUS "Benchmark is enabled") + include(GoogleBenchmark) + add_subdirectory(src/bench) +endif() diff --git a/cmake/Configuration/ConfigCuda.cmake b/cmake/Configuration/ConfigCuda.cmake new file mode 100644 index 00000000..5a740152 --- /dev/null +++ b/cmake/Configuration/ConfigCuda.cmake @@ -0,0 +1,38 @@ +#============================================================================= +# 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. +#============================================================================= + +set(CUDA_LIBRARY_DIR ${CUDA_TOOLKIT_ROOT_DIR}) +set(CUDA_LIBRARY_STUBS_DIR ${CUDA_TOOLKIT_ROOT_DIR}) + +## Based on NVRTC (Runtime Compilation) - CUDA Toolkit Documentation - v9.2.148 +## 2.2. Installation +if(CMAKE_HOST_APPLE) + set(CUDA_LIBRARY_DIR ${CUDA_TOOLKIT_ROOT_DIR}/lib) + set(CUDA_LIBRARY_STUBS_DIR ${CUDA_TOOLKIT_ROOT_DIR}/lib/stubs) +elseif(CMAKE_HOST_UNIX) + set(CUDA_LIBRARY_DIR ${CUDA_TOOLKIT_ROOT_DIR}/lib64) + set(CUDA_LIBRARY_STUBS_DIR ${CUDA_TOOLKIT_ROOT_DIR}/lib64/stubs) +elseif(CMAKE_HOST_WIN32) + set(CUDA_LIBRARY_DIR ${CUDA_TOOLKIT_ROOT_DIR}\lib\x64) + set(CUDA_LIBRARY_STUBS_DIR ${CUDA_TOOLKIT_ROOT_DIR}\lib\x64\stubs) +endif() + +set(CUDA_CUDA_LIB cuda) +set(CUDA_NVRTC_LIB nvrtc) + +message(STATUS "CUDA_LIBRARY_DIR: ${CUDA_LIBRARY_DIR}") +message(STATUS "CUDA_LIBRARY_STUBS_DIR: ${CUDA_LIBRARY_STUBS_DIR}") diff --git a/cmake/Modules/GoogleBenchmark.cmake b/cmake/Modules/GoogleBenchmark.cmake new file mode 100644 index 00000000..4eb12766 --- /dev/null +++ b/cmake/Modules/GoogleBenchmark.cmake @@ -0,0 +1,48 @@ +#============================================================================= +# 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. +#============================================================================= + +# Download and unpack google-benchmark at configure time +configure_file(${CMAKE_SOURCE_DIR}/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake + ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/download/CMakeLists.txt) + +execute_process( + COMMAND ${CMAKE_COMMAND} -G "${CMAKE_GENERATOR}" . + RESULT_VARIABLE result + WORKING_DIRECTORY ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/download +) + +if(result) + message(FATAL_ERROR "CMake step for google benchmark failed: ${result}") +endif() + + +execute_process( + COMMAND ${CMAKE_COMMAND} --build . + RESULT_VARIABLE result + WORKING_DIRECTORY ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/download +) + +if(result) + message(FATAL_ERROR "Build step for google benchmark failed: ${result}") +endif() + + +set(GOOGLE_BENCHMARK_LIB "benchmark") +set(GOOGLE_BENCHMARK_MAIN_LIB "benchmark_main") +set(GOOGLE_BENCHMARK_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/install") +set(GOOGLE_BENCHMARK_INCLUDE_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/install/include") +set(GOOGLE_BENCHMARK_LIBRARY_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/install/lib") diff --git a/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake b/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake new file mode 100644 index 00000000..8e175867 --- /dev/null +++ b/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake @@ -0,0 +1,36 @@ +#============================================================================= +# 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. +#============================================================================= + +cmake_minimum_required(VERSION 3.0) + +project(google-benchmark) + +include(ExternalProject) + +ExternalProject_Add( + google-benchmark + GIT_REPOSITORY https://github.com/google/benchmark.git + GIT_TAG master + SOURCE_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/source" + BINARY_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/build" + INSTALL_DIR "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/install" + CMAKE_ARGS -DCMAKE_BUILD_TYPE=Release + -DBENCHMARK_ENABLE_TESTING=OFF + -DBENCHMARK_ENABLE_GTEST_TESTS=OFF + -DBENCHMARK_DOWNLOAD_DEPENDENCIES=OFF + -DCMAKE_INSTALL_PREFIX:PATH=${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/thirdparty/google-benchmark/install +) \ No newline at end of file diff --git a/conda_environments/dev_py35.yml b/conda_environments/dev_py35.yml index c879875e..0e8b0c25 100644 --- a/conda_environments/dev_py35.yml +++ b/conda_environments/dev_py35.yml @@ -10,7 +10,7 @@ dependencies: - numpy=1.12.1=py35_0 - openssl=1.0.2k=2 - pip=9.0.1=py35_1 -- py=1.4.33=py35_0 +- py=1.5.2 - pycparser=2.17=py35_0 - pytest=3.0.7=py35_0 - python=3.5.3=1 diff --git a/include/gdf/cffi/functions.h b/include/gdf/cffi/functions.h index 610e756d..590def44 100644 --- a/include/gdf/cffi/functions.h +++ b/include/gdf/cffi/functions.h @@ -478,8 +478,152 @@ gdf_error gdf_extract_datetime_minute(gdf_column *input, gdf_column *output); gdf_error gdf_extract_datetime_second(gdf_column *input, gdf_column *output); -/* binary operators */ +/** + * @brief Binary operation function between gdf_scalar and gdf_column structs. + * + * The function performs the binary operation of a gdf_scalar operand and a + * gdf_column operand. + * + * The valid field in the gdf_column output will be 1 (by bit) when the two + * operands (vax and vay) are not null. Otherwise, it will be 0 (by bit). + * + * It is required to set in an appropriate manner the fields in the gdf_scalar and + * gdf_column structs due to that the binary operation will not be performed. + * + * @param out (gdf_column) Output of the operation. + * @param vax (gdf_scalar) First operand of the operation. + * @param vay (gdf_column) Second operand of the operation. + * @param ope (enum) The binary operator that is going to be used in the operation. + * @return GDF_SUCCESS if the operation was successful, otherwise an appropriate + * error code + */ +gdf_error gdf_binary_operation_v_s_v(gdf_column* out, gdf_scalar* vax, gdf_column* vay, gdf_binary_operator ope); + +/** + * @brief Binary operation function between gdf_column and gdf_scalar structs. + * + * The function performs the binary operation of a gdf_column operand and a + * gdf_scalar operand. + * + * The valid field in the gdf_column output will be 1 (by bit) when the two + * operands (vax and vay) are not null. Otherwise, it will be 0 (by bit). + * + * It is required to set in an appropriate manner the fields in the gdf_scalar and + * gdf_column structs due to that the binary operation will not be performed. + * + * @param out (gdf_column) Output of the operation. + * @param vax (gdf_column) First operand of the operation. + * @param vay (gdf_scalar) Second operand of the operation. + * @param ope (enum) The binary operator that is going to be used in the operation. + * @return GDF_SUCCESS if the operation was successful, otherwise an appropriate + * error code + */ +gdf_error gdf_binary_operation_v_v_s(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_binary_operator ope); +/** + * @brief Binary operation function between two gdf_column structs. + * + * The function performs the binary operation of two gdf_column operands. + * + * The valid field in the gdf_column output will be 1 (by bit) when the two + * operands (vax and vay) are not null. Otherwise, it will be 0 (by bit). + * + * It is required to set in an appropriate manner the fields in the gdf_column + * struct due to that the binary operation will not be performed. + * + * @param out (gdf_column) Output of the operation. + * @param vax (gdf_column) First operand of the operation. + * @param vay (gdf_column) Second operand of the operation. + * @param ope (enum) The binary operator that is going to be used in the operation. + * @return GDF_SUCCESS if the operation was successful, otherwise an appropriate + * error code + */ +gdf_error gdf_binary_operation_v_v_v(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_binary_operator ope); + +/** + * @brief Binary operation function between gdf_scalar and gdf_column structs. + * + * The function performs the binary operation of a gdf_scalar operand and a + * gdf_column operand. A default scalar operand is used to replace an operand + * in the binary operation when such operand is null. + * + * Whether any operand (vax or vay) is null, then it will be replaced with the + * default scalar operand value (def). In case both operands (vax and vay) are + * null, each of them will be replaced with the default scalar operand value. + * + * The valid field in the gdf_column output will be 1 (by bit) when two or three + * operands (vax, vay and def) are not null. Otherwise, it will be 0 (by bit). + * + * It is required to set in an appropriate manner the fields in the gdf_scalar and + * gdf_column structs due to that the binary operation will not be performed. + * + * @param out (gdf_column) Output of the operation. + * @param vax (gdf_column) First operand of the operation. + * @param vay (gdf_scalar) Second operand of the operation - gdf_scalar. + * @param def (gdf_scalar) Default operand used to replace a null operand. + * @param ope (enum) The binary operator that is going to be used in the operation. + * @return GDF_SUCCESS if the operation was successful, otherwise an appropriate + * error code + */ +gdf_error gdf_binary_operation_v_s_v_d(gdf_column* out, gdf_scalar* vax, gdf_column* vay, gdf_scalar* def, gdf_binary_operator ope); + +/** + * @brief Binary operation function between gdf_column and gdf_scalar structs. + * + * The function performs the binary operation of a gdf_column operand and a + * gdf_scalar operand. A default scalar operand is used to replace an operand + * in the binary operation when such operand is null. + * + * Whether any operand (vax or vay) is null, then it will be replaced with the + * default scalar operand value (def). In case both operands (vax and vay) are + * null, each of them will be replaced with the default scalar operand value. + * + * The valid field in the gdf_column output will be 1 (by bit) when two or three + * operands (vax, vay and def) are not null. Otherwise, it will be 0 (by bit). + * + * It is required to set in an appropriate manner the fields in the gdf_scalar and + * gdf_column structs due to that the binary operation will not be performed. + * + * @param out (gdf_column) Output of the operation. + * @param vax (gdf_column) First operand of the operation. + * @param vay (gdf_scalar) Second operand of the operation - gdf_scalar. + * @param def (gdf_scalar) Default operand used to replace a null operand. + * @param ope (enum) The binary operator that is going to be used in the operation. + * @return GDF_SUCCESS if the operation was successful, otherwise an appropriate + * error code + */ +gdf_error gdf_binary_operation_v_v_s_d(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_scalar* def, gdf_binary_operator ope); + +/** + * @brief Binary operation function between two gdf_column structs. + * + * The function performs the binary operation of two gdf_column operands. A default + * scalar operand is used to replace an operand in the binary operation when such + * operand is null. + * + * Whether any gdf_column (vax or vay) is null, then it will be replaced with the + * default scalar operand value (def). In case both operands (vax and vay) are null, + * each of them will be replaced with the default scalar operand value. + * + * The valid field in the gdf_column output will be 1 (by bit) when two or three + * operands (vax, vay and def) are not null. Otherwise, it will be 0 (by bit). + * + * It is required to set in an appropriate manner the fields in the gdf_scalar and + * gdf_column structs due to that the binary operation will not be performed. + * + * @param out (gdf_column) Output of the operation. + * @param vax (gdf_column) First operand of the operation. + * @param vay (gdf_column) Second operand of the operation. + * @param def (gdf_scalar) Default operand used to replace a null operand. + * @param ope (enum) The binary operator that is going to be used in the operation. + * @return GDF_SUCCESS if the operation was successful, otherwise an appropriate + * error code. + */ +gdf_error gdf_binary_operation_v_v_v_d(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_scalar* def, gdf_binary_operator ope); + + + +/* binary operators */ /* arith */ gdf_error gdf_add_generic(gdf_column *lhs, gdf_column *rhs, gdf_column *output); @@ -740,3 +884,11 @@ gdf_error gdf_quantile_aprrox( gdf_column* col_in, //input column; double q, //requested quantile in [0,1] void* t_erased_res, //type-erased result of same type as column; gdf_context* ctxt); //context info + + + +gdf_error gdf_order_by_asc_desc( + gdf_column * input_columns, //pointers to pointers of input columns + size_t num_inputs, //number of pointeres in the first parameter (e.g. number of columsn to sort by + gdf_column * output_indices, //a gdf_column that is pre allocated for storing sorted indices + gdf_valid_type * asc_desc_bitmask); //asc_desc bitmask e.g. 101 would mean sort the first and last oclumns ascending and the second one descending diff --git a/include/gdf/cffi/types.h b/include/gdf/cffi/types.h index 0b251a73..e8134c32 100644 --- a/include/gdf/cffi/types.h +++ b/include/gdf/cffi/types.h @@ -25,10 +25,37 @@ typedef enum { GDF_TIMESTAMP, /**< Exact timestamp encoded with int64 since UNIX epoch (Default unit millisecond) */ GDF_CATEGORY, GDF_STRING, + GDF_UINT8, + GDF_UINT16, + GDF_UINT32, + GDF_UINT64, N_GDF_TYPES, /* additional types should go BEFORE N_GDF_TYPES */ } gdf_dtype; +/** + * @union gdf_data + * @brief Union used for scalar type. + * It stores a unique value for scalar type. + * It has a direct relationship with the gdf_dtype. + */ +typedef union { + int8_t si08; /**< GDF_INT8 */ + int16_t si16; /**< GDF_INT16 */ + int32_t si32; /**< GDF_INT32 */ + int64_t si64; /**< GDF_INT64 */ + uint8_t ui08; /**< GDF_UINT8 */ + uint16_t ui16; /**< GDF_UINT16 */ + uint32_t ui32; /**< GDF_UINT32 */ + uint64_t ui64; /**< GDF_UINT64 */ + float fp32; /**< GDF_FLOAT32 */ + double fp64; /**< GDF_FLOAT64 */ + int32_t dt32; /**< GDF_DATE32 */ + int64_t dt64; /**< GDF_DATE64 */ + int64_t tmst; /**< GDF_TIMESTAMP */ +} gdf_data; + + /* --------------------------------------------------------------------------*/ /** * @Synopsis These are all possible gdf error codes that can be returned from @@ -57,8 +84,8 @@ typedef enum { GDF_UNSUPPORTED_JOIN_TYPE, /**< The type of join requested is unsupported */ GDF_UNDEFINED_NVTX_COLOR, /**< The requested color used to define an NVTX range is not defined */ GDF_NULL_NVTX_NAME, /**< The requested name for an NVTX range cannot be nullptr */ - GDF_C_ERROR, /**< C error not related to CUDA */ - GDF_FILE_ERROR, /**< error processing sepcified file */ + GDF_C_ERROR, /**< C error not related to CUDA */ + GDF_FILE_ERROR, /**< error processing sepcified file */ } gdf_error; typedef enum { @@ -79,6 +106,24 @@ typedef struct { // here we can also hold info for decimal datatype or any other datatype that requires additional information } gdf_dtype_extra_info; + +/** + * @struct gdf_scalar + * @brief literal or variable + * + * The struct is used as a literal or a variable in the libgdf library. + * + * @var data A union that represents the value. + * @var dtype An enum that represents the type of the value. + * @var is_valid A boolean that represents whether the scalar is null. + */ +typedef struct { + gdf_data data; + gdf_dtype dtype; + bool is_valid; +} gdf_scalar; + + typedef struct gdf_column_{ void *data; /**< Pointer to the columns data */ gdf_valid_type *valid; /**< Pointer to the columns validity bit mask where the 'i'th bit indicates if the 'i'th row is NULL */ @@ -130,7 +175,7 @@ typedef enum { /* --------------------------------------------------------------------------*/ -/** +/** * @Synopsis Colors for use with NVTX ranges. * * These enumerations are the available pre-defined colors for use with @@ -138,7 +183,7 @@ typedef enum { */ /* ----------------------------------------------------------------------------*/ typedef enum { - GDF_GREEN = 0, + GDF_GREEN = 0, GDF_BLUE, GDF_YELLOW, GDF_PURPLE, @@ -150,6 +195,41 @@ typedef enum { GDF_NUM_COLORS, /** Add new colors above this line */ } gdf_color; + +/** + * @enum gdf_binary_operator + * It contains the different operations that can be performed in the binary operations. + * The enumeration is used in the following functions: + * - gdf_binary_operation_v_s_v + * - gdf_binary_operation_v_v_s + * - gdf_binary_operation_v_v_v + * - gdf_binary_operation_v_s_v_d + * - gdf_binary_operation_v_v_s_d + * - gdf_binary_operation_v_v_v_d + */ +typedef enum { + 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, + //GDF_COMBINE, + //GDF_COMBINE_FIRST, + //GDF_ROUND, + //GDF_PRODUCT, + //GDF_DOT +} gdf_binary_operator; + + /* --------------------------------------------------------------------------*/ /** * @Synopsis This struct holds various information about how an operation should be diff --git a/include/sqls_rtti_comp.hpp b/include/sqls_rtti_comp.hpp index 8b641276..40275b85 100644 --- a/include/sqls_rtti_comp.hpp +++ b/include/sqls_rtti_comp.hpp @@ -24,6 +24,8 @@ #include #include +#include "gdf/utils.h" + //for int_t: // #include @@ -49,7 +51,8 @@ struct LesserRTTI columns_(cols), rtti_(types), sz_(sz), - vals_(nullptr) + vals_(nullptr), + asc_desc_bitmask_(nullptr) { } @@ -61,11 +64,65 @@ struct LesserRTTI columns_(cols), rtti_(types), sz_(sz), - vals_(vals) + vals_(vals), + asc_desc_bitmask_(nullptr) { } + __host__ __device__ + LesserRTTI(void* const* cols, + int* const types, + size_t sz, + gdf_valid_type* const asc_desc_bitmask): + columns_(cols), + rtti_(types), + sz_(sz), + vals_(nullptr), + asc_desc_bitmask_(asc_desc_bitmask) + { + } +/** + * Should be used when you want to sort multiple columns using asc / desc flags for each column + * + * + */ + __host__ __device__ + bool asc_desc_comparison(IndexT row1, IndexT row2) const + { + for(size_t col_index = 0; col_index < sz_; ++col_index) + { + gdf_dtype col_type = static_cast(rtti_[col_index]); + bool asc; + + if(asc_desc_bitmask_ == nullptr){ + asc = true; + }else{ + asc = gdf_is_valid(asc_desc_bitmask_, col_index); + } + //if flag == true + + State state; + if(asc){ + OpLess less(row1, row2); + state =type_dispatcher(less, col_type, col_index); + }else{ + OpGreater greater(row1, row2); + state =type_dispatcher(greater, col_type, col_index); + } + + switch( state ) + { + case State::False: + return false; + case State::True: + return true; + case State::Undecided: + break; + } + } + return false; + } __host__ __device__ bool equal(IndexT row1, IndexT row2) const @@ -187,6 +244,36 @@ struct LesserRTTI IndexT row2_; }; + struct OpGreater + { + __host__ __device__ + OpGreater(IndexT row1, IndexT row2): + row1_(row1), + row2_(row2) + { + } + + template + __host__ __device__ + State operator() (int col_index, + const void* const * columns, + ColType ) + { + ColType res1 = LesserRTTI::at(col_index, row1_, columns); + ColType res2 = LesserRTTI::at(col_index, row2_, columns); + + if( res1 > res2 ) + return State::True; + else if( res1 == res2 ) + return State::Undecided; + else + return State::False; + } + private: + IndexT row1_; + IndexT row2_; + }; + struct OpEqual { __host__ __device__ @@ -336,6 +423,7 @@ struct LesserRTTI const int* const rtti_; size_t sz_; const void* const * vals_; //for filtering + const gdf_valid_type* asc_desc_bitmask_; //a bitmask that allows us to know whether or not a column should be sorted ascending or descending }; //########################################################################### @@ -379,6 +467,8 @@ void multi_col_order_by(size_t nrows, }); } + + //########################################################################### //# Multi-column Filter: # //########################################################################### @@ -709,4 +799,28 @@ size_t multi_col_group_by_avg_sort(size_t nrows, return new_sz; } +#include +template +void multi_col_order_by_asc_desc( + void* const * d_col_data, + int* d_col_types, + size_t num_inputs, + gdf_valid_type * asc_desc_bitmask, + IndexT* d_indx, + size_t nrows, + cudaStream_t stream = NULL){ + + LesserRTTI f(d_col_data, d_col_types, num_inputs,asc_desc_bitmask); + + thrust::sequence(thrust::cuda::par.on(stream), d_indx, d_indx+nrows, 0); + + thrust::sort(thrust::cuda::par.on(stream), + d_indx, d_indx+nrows, + [f] __host__ __device__ (IndexT i1, IndexT i2){ + + return f.asc_desc_comparison(i1, i2); + }); + + +} diff --git a/src/bench/CMakeLists.txt b/src/bench/CMakeLists.txt new file mode 100644 index 00000000..957aeeae --- /dev/null +++ b/src/bench/CMakeLists.txt @@ -0,0 +1,18 @@ +#============================================================================= +# 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. +#============================================================================= + +add_subdirectory(binary-operation) diff --git a/src/bench/binary-operation/CMakeLists.txt b/src/bench/binary-operation/CMakeLists.txt new file mode 100644 index 00000000..c81f0257 --- /dev/null +++ b/src/bench/binary-operation/CMakeLists.txt @@ -0,0 +1,66 @@ +#============================================================================= +# 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. +#============================================================================= + +## Test name +set(GDF_APP_NAME BinaryOperationsBench) +set(GDF_BENCH_DIR bench) + + +## Set sources +set(GDF_SOURCE_FILES + "${PROJECT_SOURCE_DIR}/src/tests/binary-operation/util/types.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/binary-operation-benchmark.cpp") + + +## Add link directory +link_directories(${CUDA_LIBRARY_DIR}) +link_directories(${CUDA_LIBRARY_STUBS_DIR}) +link_directories(${GOOGLE_BENCHMARK_LIBRARY_DIR}) + + +## Add target +add_executable(${GDF_APP_NAME} ${GDF_SOURCE_FILES}) + + +## Clear include header property +set_target_properties(${GDF_APP_NAME} + PROPERTIES + INCLUDE_DIRECTORIES "") +set_target_properties(${GDF_APP_NAME} + PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/${GDF_BENCH_DIR}") + + +## Include headers +target_include_directories(${GDF_APP_NAME} PRIVATE "${CUDA_INCLUDE_DIRS}") +target_include_directories(${GDF_APP_NAME} PRIVATE "${PROJECT_SOURCE_DIR}/include") +target_include_directories(${GDF_APP_NAME} PRIVATE "${PROJECT_SOURCE_DIR}/src") +target_include_directories(${GDF_APP_NAME} PRIVATE "${GOOGLE_BENCHMARK_INCLUDE_DIR}") + + +## Link libraries +target_link_libraries(${GDF_APP_NAME} ${CMAKE_THREAD_LIBS_INIT}) +target_link_libraries(${GDF_APP_NAME} ${GOOGLE_BENCHMARK_LIB}) +target_link_libraries(${GDF_APP_NAME} ${GOOGLE_BENCHMARK_MAIN_LIB}) +target_link_libraries(${GDF_APP_NAME} gdf) +target_link_libraries(${GDF_APP_NAME} ${CUDA_CUDA_LIB}) +target_link_libraries(${GDF_APP_NAME} ${CUDA_NVRTC_LIB}) + + +message(STATUS "benchmark for binary operations") +message(STATUS "Execute help: ${GDF_BENCH_DIR}/${GDF_APP_NAME} --help") +message(STATUS "Execute benchmark: ${GDF_BENCH_DIR}/${GDF_APP_NAME} --benchmark_counters_tabular=true") diff --git a/src/bench/binary-operation/binary-operation-benchmark.cpp b/src/bench/binary-operation/binary-operation-benchmark.cpp new file mode 100644 index 00000000..dbaf57d7 --- /dev/null +++ b/src/bench/binary-operation/binary-operation-benchmark.cpp @@ -0,0 +1,203 @@ +/* + * 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 +#include +#include +#include "benchmark/benchmark.h" +#include "tests/binary-operation/util/scalar.h" +#include "tests/binary-operation/util/vector.h" + +namespace bench { + static constexpr int NUMBER_COUNTERS = 6; + + template + struct BinaryOperation : public benchmark::Fixture { + public: + virtual void SetUp(benchmark::State& state) { + if (id != state.range(0)) { + id = state.range(0); + sampler.counter = 0; + sampler.limit = state.range(1); + sampler.data.clear(); + } + + vax.rangeData((Type)state.range(2), (Type)state.range(3), (Type)state.range(4)); + vax.rangeValid(true, (int)state.range(2), 3); + + vay.rangeData((Type)state.range(2), (Type)state.range(3), (Type)state.range(4)); + vay.rangeValid(true, (int)state.range(2), 4); + + out.emplaceVector(vax.dataSize()); + def.setValue((Type)state.range(3)); + } + + virtual void TearDown(benchmark::State& state) { + vax.clearGpu(); + vay.clearGpu(); + out.clearGpu(); + } + + void startTime() { + startPointClock = std::chrono::high_resolution_clock::now(); + } + + void stopTime() { + stopPointClock = std::chrono::high_resolution_clock::now(); + } + + void saveTime(benchmark::State& state) { + auto elapsedTimeSeconds = std::chrono::duration_cast>(stopPointClock - startPointClock); + state.SetIterationTime(elapsedTimeSeconds.count()); + } + + void makeSamplerData(benchmark::State& state) { + if (sampler.counter < sampler.limit) { + auto elapsedTimeSeconds = std::chrono::duration_cast>(stopPointClock - startPointClock); + sampler.data.push_back(elapsedTimeSeconds.count()); + sampler.counter++; + } + } + + void saveSamplerData(benchmark::State& state) { + for (int k = 0; k < sampler.data.size(); ++k) { + state.counters.insert({sampler.name + std::to_string(k), {sampler.data[k]}}); + } + } + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + std::chrono::time_point startPointClock; + std::chrono::time_point stopPointClock; + + struct Sampler { + int limit = 0; + int counter = 0; + std::vector data; + const std::string name { "Iter" }; + }; + + int id = 0; + Sampler sampler; + }; + + BENCHMARK_TEMPLATE_DEFINE_F(BinaryOperation, KernelInteger, uint64_t)(benchmark::State& state) { + for (auto _ : state) { + startTime(); + gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_ADD); + stopTime(); + saveTime(state); + makeSamplerData(state); + } + saveSamplerData(state); + } + + BENCHMARK_TEMPLATE_DEFINE_F(BinaryOperation, KernelFloat, float)(benchmark::State& state) { + for (auto _ : state) { + startTime(); + gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_ADD); + stopTime(); + saveTime(state); + makeSamplerData(state); + } + saveSamplerData(state); + } + + BENCHMARK_TEMPLATE_DEFINE_F(BinaryOperation, KernelDouble, double)(benchmark::State& state) { + for (auto _ : state) { + startTime(); + gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_ADD); + stopTime(); + saveTime(state); + makeSamplerData(state); + } + saveSamplerData(state); + } + + BENCHMARK_TEMPLATE_DEFINE_F(BinaryOperation, KernelIntegerDefault, uint64_t)(benchmark::State& state) { + for (auto _ : state) { + startTime(); + gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + stopTime(); + saveTime(state); + makeSamplerData(state); + } + saveSamplerData(state); + } + + BENCHMARK_TEMPLATE_DEFINE_F(BinaryOperation, KernelFloatDefault, float)(benchmark::State& state) { + for (auto _ : state) { + startTime(); + gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + stopTime(); + saveTime(state); + makeSamplerData(state); + } + saveSamplerData(state); + } + + BENCHMARK_TEMPLATE_DEFINE_F(BinaryOperation, KernelDoubleDefault, double)(benchmark::State& state) { + for (auto _ : state) { + startTime(); + gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + stopTime(); + saveTime(state); + makeSamplerData(state); + } + saveSamplerData(state); + } + + + BENCHMARK_REGISTER_F(BinaryOperation, KernelInteger) + ->UseManualTime() + ->Args({1, NUMBER_COUNTERS, 0, 100, 1}) + ->Args({2, NUMBER_COUNTERS, 0, 10000, 1}) + ->Args({3, NUMBER_COUNTERS, 0, 1000000, 1}); + + BENCHMARK_REGISTER_F(BinaryOperation, KernelFloat) + ->UseManualTime() + ->Args({1, NUMBER_COUNTERS, 0, 100, 1}) + ->Args({2, NUMBER_COUNTERS, 0, 10000, 1}) + ->Args({3, NUMBER_COUNTERS, 0, 1000000, 1}); + + BENCHMARK_REGISTER_F(BinaryOperation, KernelDouble) + ->UseManualTime() + ->Args({1, NUMBER_COUNTERS, 0, 100, 1}) + ->Args({2, NUMBER_COUNTERS, 0, 10000, 1}) + ->Args({3, NUMBER_COUNTERS, 0, 1000000, 1}); + + BENCHMARK_REGISTER_F(BinaryOperation, KernelIntegerDefault) + ->UseManualTime() + ->Args({1, NUMBER_COUNTERS, 0, 100, 1}) + ->Args({2, NUMBER_COUNTERS, 0, 10000, 1}) + ->Args({3, NUMBER_COUNTERS, 0, 1000000, 1}); + + BENCHMARK_REGISTER_F(BinaryOperation, KernelFloatDefault) + ->UseManualTime() + ->Args({1, NUMBER_COUNTERS, 0, 100, 1}) + ->Args({2, NUMBER_COUNTERS, 0, 10000, 1}) + ->Args({3, NUMBER_COUNTERS, 0, 1000000, 1}); + + BENCHMARK_REGISTER_F(BinaryOperation, KernelDoubleDefault) + ->UseManualTime() + ->Args({1, NUMBER_COUNTERS, 0, 100, 1}) + ->Args({2, NUMBER_COUNTERS, 0, 10000, 1}) + ->Args({3, NUMBER_COUNTERS, 0, 1000000, 1}); +} diff --git a/src/binary-operation/CMakeLists.txt b/src/binary-operation/CMakeLists.txt new file mode 100644 index 00000000..c5e1df3a --- /dev/null +++ b/src/binary-operation/CMakeLists.txt @@ -0,0 +1,27 @@ +#============================================================================= +# 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. +#============================================================================= + +set(source_files + src/binary-operation/jit/code/gdf-data.cpp + src/binary-operation/jit/code/kernel.cpp + src/binary-operation/jit/code/operation.cpp + src/binary-operation/jit/code/traits.cpp + src/binary-operation/jit/core/binop.cpp + src/binary-operation/jit/core/launcher.cpp + src/binary-operation/jit/util/operator.cpp + src/binary-operation/jit/util/type.cpp +) diff --git a/src/binary-operation/jit/code/code.h b/src/binary-operation/jit/code/code.h new file mode 100644 index 00000000..e714bae0 --- /dev/null +++ b/src/binary-operation/jit/code/code.h @@ -0,0 +1,36 @@ +/* + * 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 gdf { +namespace binops { +namespace jit { +namespace code { + + extern const char* kernel; + extern const char* traits; + extern const char* operation; + extern const char* gdf_data; + +} +} +} +} + +#endif diff --git a/src/binary-operation/jit/code/gdf-data.cpp b/src/binary-operation/jit/code/gdf-data.cpp new file mode 100644 index 00000000..bdc65a22 --- /dev/null +++ b/src/binary-operation/jit/code/gdf-data.cpp @@ -0,0 +1,85 @@ +/* + * 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 gdf { +namespace binops { +namespace jit { +namespace code { + +const char* gdf_data = +R"***( +#pragma once + + union gdf_data { + int8_t si08; + int16_t si16; + int32_t si32; + int64_t si64; + uint8_t ui08; + uint16_t ui16; + uint32_t ui32; + uint64_t ui64; + float fp32; + double fp64; + + operator int8_t() const { + return si08; + } + + operator int16_t() const { + return si16; + } + + operator int32_t() const { + return si32; + } + + operator int64_t() const { + return si64; + } + + operator uint8_t() const { + return ui08; + } + + operator uint16_t() const { + return ui16; + } + + operator uint32_t() const { + return ui32; + } + + operator uint64_t() const { + return ui64; + } + + operator float() const { + return fp32; + } + + operator double() const { + return fp64; + } + }; + +)***"; + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/code/kernel.cpp b/src/binary-operation/jit/code/kernel.cpp new file mode 100644 index 00000000..72c44b8a --- /dev/null +++ b/src/binary-operation/jit/code/kernel.cpp @@ -0,0 +1,153 @@ +/* + * 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 gdf { +namespace binops { +namespace jit { +namespace code { + +const char* kernel = +R"***( + #include + #include "traits.h" + #include "operation.h" + #include "gdf_data.h" + + template + __global__ + void kernel_v_s(int size, + TypeOut* out_data, TypeVax* vax_data, gdf_data vay_data, + uint32_t* out_valid, uint32_t* vax_valid, uint32_t vay_valid) { + 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 (int i=start; i(vax_data[i], (TypeVay)vay_data); + + if ((i % warpSize) == 0) { + int index = i / warpSize; + out_valid[index] = vax_valid[index] & vay_valid; + } + } + } + + template + __global__ + void kernel_v_v(int size, + TypeOut* out_data, TypeVax* vax_data, TypeVay* vay_data, + uint32_t* out_valid, uint32_t* vax_valid, uint32_t* vay_valid) { + 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 (int i=start; i(vax_data[i], vay_data[i]); + + if ((i % warpSize) == 0) { + int index = i / warpSize; + out_valid[index] = vax_valid[index] & vay_valid[index]; + } + } + } + + template + __global__ + void kernel_v_s_d(int size, + TypeOut* out_data, TypeVax* vax_data, gdf_data vay_data, gdf_data def_data, + uint32_t* out_valid, uint32_t* vax_valid, uint32_t vay_valid, uint32_t def_valid) { + 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 (int i=start; i> position) & 1; + TypeVax vax_data_aux = ((TypeVax)sel_vax * vax_data[i]) + + ((TypeVax)(sel_vax ^ 1) * (TypeVax)((TypeVal)def_data)); + + TypeVay vay_data_aux = ((TypeVay)(vay_valid & 1) * (TypeVay)vay_data) + + ((TypeVay)(vay_valid + 1) * (TypeVay)((TypeVal)def_data)); + + out_data[i] = TypeOpe::template operate(vax_data_aux, vay_data_aux); + + if ((i % warpSize) == 0) { + out_valid[index] = (vax_valid[index] & vay_valid) | + (vax_valid[index] & def_valid) | + (vay_valid & def_valid); + } + } + } + + template + __global__ + void kernel_v_v_d(int size, + TypeOut* out_data, TypeVax* vax_data, TypeVay* vay_data, gdf_data def_data, + uint32_t* out_valid, uint32_t* vax_valid, uint32_t* vay_valid, uint32_t def_valid) { + 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 (int i=start; i> position) & 1; + TypeVax vax_data_aux = ((TypeVax)sel_vax * vax_data[i]) + + ((TypeVax)(sel_vax ^ 1) * (TypeVax)((TypeVal)def_data)); + + uint32_t sel_vay = (is_vay_valid >> position) & 1; + TypeVay vay_data_aux = ((TypeVay)sel_vay * vay_data[i]) + + ((TypeVay)(sel_vay ^ 1) * (TypeVay)((TypeVal)def_data)); + + out_data[i] = TypeOpe::template operate(vax_data_aux, vay_data_aux); + + if ((i % warpSize) == 0) { + out_valid[index] = (vax_valid[index] & vay_valid[index]) | + (vax_valid[index] & def_valid) | + (vay_valid[index] & def_valid); + } + } + } +)***"; + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/code/operation.cpp b/src/binary-operation/jit/code/operation.cpp new file mode 100644 index 00000000..78122639 --- /dev/null +++ b/src/binary-operation/jit/code/operation.cpp @@ -0,0 +1,306 @@ +/* + * 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 gdf { +namespace binops { +namespace jit { +namespace code { + +const char* operation = +R"***( +#pragma once + + struct Add { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x + (Common)y); + } + }; + + using RAdd = Add; + + struct Sub { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x - (Common)y); + } + }; + + struct RSub { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)y - (Common)x); + } + }; + + struct Mul { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x * (Common)y); + } + }; + + using RMul = Mul; + + struct Div { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x / (Common)y); + } + }; + + struct RDiv { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)y / (Common)x); + } + }; + + struct TrueDiv { + template + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)((double)x / (double)y); + } + }; + + struct RTrueDiv { + template + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)((double)y / (double)x); + } + }; + + struct FloorDiv { + template + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)floor((double)x / (double)y); + } + }; + + struct RFloorDiv { + template + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)floor((double)y / (double)x); + } + }; + + struct Mod { + template , + enableIf<(isIntegral)>* = nullptr> + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)((Common)x % (Common)y); + } + + template , + enableIf<(isFloat)>* = nullptr> + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)fmodf((Common)x, (Common)y); + } + + template , + enableIf<(isDouble)>* = nullptr> + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)fmod((Common)x, (Common)y); + } + }; + + struct RMod { + template , + enableIf<(isIntegral)>* = nullptr> + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)((Common)y % (Common)x); + } + + template , + enableIf<(isFloat)>* = nullptr> + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)fmodf((Common)y, (Common)x); + } + + template , + enableIf<(isDouble)>* = nullptr> + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)fmod((Common)y, (Common)x); + } + }; + + struct Pow { + template + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)pow((double)x, (double)y); + } + }; + + struct RPow { + template + static TypeOut operate(TypeVax x, TypeVay y) { + return (TypeOut)pow((double)y, (double)x); + } + }; + + struct Equal { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x == (Common)y); + } + }; + + using REqual = Equal; + + struct NotEqual { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x != (Common)y); + } + }; + + using RNotEqual = NotEqual; + + struct Less { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x < (Common)y); + } + }; + + struct RLess { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)y < (Common)x); + } + }; + + struct Greater { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x > (Common)y); + } + }; + + struct RGreater { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)y > (Common)x); + } + }; + + struct LessEqual { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x <= (Common)y); + } + }; + + struct RLessEqual { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)y <= (Common)x); + } + }; + + struct GreaterEqual { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)x >= (Common)y); + } + }; + + struct RGreaterEqual { + template + static TypeOut operate(TypeVax x, TypeVay y) { + using Common = CommonNumber; + return (TypeOut)((Common)y >= (Common)x); + } + }; +)***"; + +/* + * The following code could be used to detect overflow or underflow + * using 'Bit Hacks' in the operations, that's why the operation is + * divided into signed, unsigned and double functions. It's required + * to create a new field on gdf_column for this feature. + * + * struct Add { + * template , + * enableIf<(isIntegralSigned)>* = nullptr> + * __device__ + * TypeOut operate(TypeVax x, TypeVay y) { + * return (TypeOut)((Common)x + (Common)y); + * } + * + * template , + * enableIf<(isIntegralUnsigned)>* = nullptr> + * __device__ + * TypeOut operate(TypeVax x, TypeVay y) { + * return (TypeOut)((Common)x + (Common)y); + * } + * + * template , + * enableIf<(isFloatingPoint)>* = nullptr> + * __device__ + * TypeOut operate(TypeVax x, TypeVay y) { + * return (TypeOut)((Common)x + (Common)y); + * } + * }; + */ + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/code/traits.cpp b/src/binary-operation/jit/code/traits.cpp new file mode 100644 index 00000000..9e01d40f --- /dev/null +++ b/src/binary-operation/jit/code/traits.cpp @@ -0,0 +1,237 @@ +/* + * 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 gdf { +namespace binops { +namespace jit { +namespace code { + +const char* traits = +R"***( +#pragma once + + struct IntegralSigned {}; + + struct IntegralUnsigned {}; + + + template + constexpr bool isIntegral = false; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + template <> + constexpr bool isIntegral = true; + + + + template + constexpr bool isFloatingPoint = false; + + template <> + constexpr bool isFloatingPoint = true; + + template <> + constexpr bool isFloatingPoint = true; + + + + template + constexpr bool isIntegralSigned = false; + + template <> + constexpr bool isIntegralSigned = true; + + template <> + constexpr bool isIntegralSigned = true; + + template <> + constexpr bool isIntegralSigned = true; + + template <> + constexpr bool isIntegralSigned = true; + + + + template + constexpr bool isIntegralUnsigned = false; + + template <> + constexpr bool isIntegralUnsigned = true; + + template <> + constexpr bool isIntegralUnsigned = true; + + template <> + constexpr bool isIntegralUnsigned = true; + + template <> + constexpr bool isIntegralUnsigned = true; + + + template + constexpr bool isFloat = false; + + template <> + constexpr bool isFloat = true; + + + template + constexpr bool isDouble = false; + + template <> + constexpr bool isDouble = true; + + + template + constexpr int MaxSize = ((sizeof(X) < sizeof(Y)) ? sizeof(Y) : sizeof(X)); + + template + struct HelperIntegralMap; + + template <> + struct HelperIntegralMap<1, IntegralSigned> { + using Type = int8_t; + }; + + template <> + struct HelperIntegralMap<2, IntegralSigned> { + using Type = int16_t; + }; + + template <> + struct HelperIntegralMap<4, IntegralSigned> { + using Type = int32_t; + }; + + template <> + struct HelperIntegralMap<8, IntegralSigned> { + using Type = int64_t; + }; + + template <> + struct HelperIntegralMap<1, IntegralUnsigned> { + using Type = uint8_t; + }; + + template <> + struct HelperIntegralMap<2, IntegralUnsigned> { + using Type = uint16_t; + }; + + template <> + struct HelperIntegralMap<4, IntegralUnsigned> { + using Type = uint32_t; + }; + + template <> + struct HelperIntegralMap<8, IntegralUnsigned> { + using Type = uint64_t; + }; + + template + using IntegralMap = typename HelperIntegralMap::Type; + + + + template + struct helperIf; + + template + struct helperIf { + using type = T; + }; + + template + struct helperIf { + using type = F; + }; + + template + using If = typename helperIf::type; + + + + template + struct helperEnableIf + {}; + + template + struct helperEnableIf { + using type = T; + }; + + template + using enableIf = typename helperEnableIf::type; + + + + template + struct HelperCommonNumber {}; + + template + struct HelperCommonNumber || isFloatingPoint)>> { + using Type = If<(sizeof(Vax) == 8 || sizeof(Vay) == 8), double, float>; + }; + + template + struct HelperCommonNumber && isIntegralSigned)>> { + using Type = IntegralMap<(MaxSize), IntegralSigned>; + }; + + template + struct HelperCommonNumber && isIntegralUnsigned)>> { + using Type = IntegralMap<(MaxSize), IntegralSigned>; + }; + + template + struct HelperCommonNumber && isIntegralSigned)>> { + using Type = IntegralMap<(MaxSize), IntegralSigned>; + }; + + template + struct HelperCommonNumber && isIntegralUnsigned)>> { + using Type = IntegralMap<(MaxSize), IntegralUnsigned>; + }; + + template + using CommonNumber = typename HelperCommonNumber::Type; +)***"; + +} // namespace code +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/core/binop.cpp b/src/binary-operation/jit/core/binop.cpp new file mode 100644 index 00000000..dedcf605 --- /dev/null +++ b/src/binary-operation/jit/core/binop.cpp @@ -0,0 +1,237 @@ +/* + * 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 "gdf/gdf.h" +#include "binary-operation/jit/core/launcher.h" +#include "binary-operation/jit/util/operator.h" + +namespace gdf { +namespace binops { +namespace jit { + + struct Option { + Option(bool state, gdf_error value) + : is_correct{state}, gdf_error_value{value} + { } + + operator bool() { + return is_correct; + } + + gdf_error get_gdf_error() { + return gdf_error_value; + } + + private: + bool is_correct; + gdf_error gdf_error_value; + }; + + Option verify_scalar(gdf_scalar* scalar) { + if (scalar == nullptr) { + return Option(false, GDF_DATASET_EMPTY); + } + if ((scalar->dtype <= GDF_invalid) || (N_GDF_TYPES <= scalar->dtype)) { + return Option(false, GDF_UNSUPPORTED_DTYPE); + } + return Option(true, GDF_SUCCESS); + } + + Option verify_column(gdf_column* vector) { + if (vector == nullptr) { + return Option(false, GDF_DATASET_EMPTY); + } + if (vector->size == 0) { + return Option(false, GDF_SUCCESS); + } + if (vector->data == nullptr) { + return Option(false, GDF_DATASET_EMPTY); + } + if ((vector->dtype <= GDF_invalid) || (N_GDF_TYPES <= vector->dtype)) { + return Option(false, GDF_UNSUPPORTED_DTYPE); + } + return Option(true, GDF_SUCCESS); + } + + Option verify_column(gdf_column* out, gdf_column* vax) { + auto result = verify_column(out); + if (!result) { + return result; + } + result = verify_column(vax); + if (!result) { + return result; + } + if (out->size < vax->size) { + return Option(false, GDF_COLUMN_SIZE_MISMATCH); + } + return Option(true, GDF_SUCCESS); + } + + Option verify_column(gdf_column* out, gdf_column* vax, gdf_column* vay) { + auto result = verify_column(out); + if (!result) { + return result; + } + result = verify_column(vax); + if (!result) { + return result; + } + result = verify_column(vay); + if (!result) { + return result; + } + if ((out->size < vax->size) || (out->size < vay->size) || (vay->size != vax->size)) { + return Option(false, GDF_COLUMN_SIZE_MISMATCH); + } + return Option(true, GDF_SUCCESS); + } + + gdf_error binary_operation(gdf_column* out, gdf_scalar* vax, gdf_column* vay, gdf_binary_operator ope) { + auto option_scalar = verify_scalar(vax); + if (!option_scalar) { + return option_scalar.get_gdf_error(); + } + auto option_column = verify_column(out, vay); + if (!option_column) { + return option_column.get_gdf_error(); + } + + Launcher::launch().kernel("kernel_v_s") + .instantiate(ope, Operator::Type::Reverse, out, vay, vax) + .launch(out, vay, vax); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_binary_operator ope) { + auto option_scalar = verify_scalar(vay); + if (!option_scalar) { + return option_scalar.get_gdf_error(); + } + auto option_column = verify_column(out, vax); + if (!option_column) { + return option_column.get_gdf_error(); + } + + Launcher::launch().kernel("kernel_v_s") + .instantiate(ope, Operator::Type::Direct, out, vax, vay) + .launch(out, vax, vay); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_binary_operator ope) { + auto option_column = verify_column(out, vax, vay); + if (!option_column) { + return option_column.get_gdf_error(); + } + + Launcher::launch().kernel("kernel_v_v") + .instantiate(ope, Operator::Type::Direct, out, vax, vay) + .launch(out, vax, vay); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_scalar* vax, gdf_column* vay, gdf_scalar* def, gdf_binary_operator ope) { + auto option_column = verify_column(out, vay); + if (!option_column) { + return option_column.get_gdf_error(); + } + auto option_scalar_vax = verify_scalar(vax); + if (!option_scalar_vax) { + return option_scalar_vax.get_gdf_error(); + } + auto option_scalar_def = verify_scalar(def); + if (!option_scalar_def) { + return option_scalar_def.get_gdf_error(); + } + + Launcher::launch().kernel("kernel_v_s_d") + .instantiate(ope, Operator::Type::Reverse, out, vay, vax, def) + .launch(out, vay, vax, def); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_scalar* def, gdf_binary_operator ope) { + auto option_column = verify_column(out, vax); + if (!option_column) { + return option_column.get_gdf_error(); + } + auto option_scalar_vax = verify_scalar(vay); + if (!option_scalar_vax) { + return option_scalar_vax.get_gdf_error(); + } + auto option_scalar_def = verify_scalar(def); + if (!option_scalar_def) { + return option_scalar_def.get_gdf_error(); + } + + Launcher::launch().kernel("kernel_v_s_d") + .instantiate(ope, Operator::Type::Direct, out, vax, vay, def) + .launch(out, vax, vay, def); + + return GDF_SUCCESS; + } + + gdf_error binary_operation(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_scalar* def, gdf_binary_operator ope) { + auto option_column = verify_column(out, vax, vay); + if (!option_column) { + return option_column.get_gdf_error(); + } + auto option_scalar = verify_scalar(def); + if (!option_scalar) { + return option_scalar.get_gdf_error(); + } + + Launcher::launch().kernel("kernel_v_v_d") + .instantiate(ope, Operator::Type::Direct, out, vax, vay, def) + .launch(out, vax, vay, def); + + return GDF_SUCCESS; + } + +} // namespace jit +} // namespace binops +} // namespace gdf + + +gdf_error gdf_binary_operation_v_s_v(gdf_column* out, gdf_scalar* vax, gdf_column* vay, gdf_binary_operator ope) { + return gdf::binops::jit::binary_operation(out, vax, vay, ope); +} + +gdf_error gdf_binary_operation_v_v_s(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_binary_operator ope) { + return gdf::binops::jit::binary_operation(out, vax, vay, ope); +} + +gdf_error gdf_binary_operation_v_v_v(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_binary_operator ope) { + return gdf::binops::jit::binary_operation(out, vax, vay, ope); +} + +gdf_error gdf_binary_operation_v_s_v_d(gdf_column* out, gdf_scalar* vax, gdf_column* vay, gdf_scalar* def, gdf_binary_operator ope) { + return gdf::binops::jit::binary_operation(out, vax, vay, def, ope); +} + +gdf_error gdf_binary_operation_v_v_s_d(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_scalar* def, gdf_binary_operator ope) { + return gdf::binops::jit::binary_operation(out, vax, vay, def, ope); +} + +gdf_error gdf_binary_operation_v_v_v_d(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_scalar* def, gdf_binary_operator ope) { + return gdf::binops::jit::binary_operation(out, vax, vay, def, ope); +} diff --git a/src/binary-operation/jit/core/launcher.cpp b/src/binary-operation/jit/core/launcher.cpp new file mode 100644 index 00000000..ad7d7d73 --- /dev/null +++ b/src/binary-operation/jit/core/launcher.cpp @@ -0,0 +1,112 @@ +/* + * 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 +#include "binary-operation/jit/core/launcher.h" +#include "binary-operation/jit/code/code.h" + +namespace gdf { +namespace binops { +namespace jit { + + static thread_local jitify::JitCache JitCache; + + 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; + } + if (filename == "gdf_data.h") { + stream << code::gdf_data; + 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* vax, gdf_scalar* vay) { + uint32_t vay_valid = (vay->is_valid ? UINT32_MAX : 0); + + program.kernel(kernelName.c_str()) + .instantiate(arguments) + .configure_1d_max_occupancy() + .launch(out->size, + out->data, vax->data, vay->data, + out->valid, vax->valid, vay_valid); + + return GDF_SUCCESS; + } + + gdf_error Launcher::launch(gdf_column* out, gdf_column* vax, gdf_column* vay) { + program.kernel(kernelName.c_str()) + .instantiate(arguments) + .configure_1d_max_occupancy() + .launch(out->size, + out->data, vax->data, vay->data, + out->valid, vax->valid, vay->valid); + + return GDF_SUCCESS; + } + + gdf_error Launcher::launch(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_scalar* def) + { + uint32_t vay_valid = (vay->is_valid ? UINT32_MAX : 0); + uint32_t def_valid = (def->is_valid ? UINT32_MAX : 0); + + program.kernel(kernelName) + .instantiate(arguments) + .configure_1d_max_occupancy() + .launch(out->size, + out->data, vax->data, vay->data, def->data, + out->valid, vax->valid, vay_valid, def_valid); + + return GDF_SUCCESS; + } + + gdf_error Launcher::launch(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_scalar* def) + { + uint32_t def_valid = (def->is_valid ? UINT32_MAX : 0); + + program.kernel(kernelName) + .instantiate(arguments) + .configure_1d_max_occupancy() + .launch(out->size, + out->data, vax->data, vay->data, def->data, + out->valid, vax->valid, vay->valid, def_valid); + + return GDF_SUCCESS; + } + +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/core/launcher.h b/src/binary-operation/jit/core/launcher.h new file mode 100644 index 00000000..e3d9fee6 --- /dev/null +++ b/src/binary-operation/jit/core/launcher.h @@ -0,0 +1,83 @@ +/* + * 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 +#include "binary-operation/jit/util/type.h" +#include "binary-operation/jit/util/operator.h" + +namespace gdf { +namespace binops { +namespace jit { + + std::istream* headersCode(std::string filename, std::iostream& stream); + + 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: + Launcher& kernel(std::string&& value); + + 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; + } + + gdf_error launch(gdf_column* out, gdf_column* vax, gdf_scalar* vay); + + gdf_error launch(gdf_column* out, gdf_column* vax, gdf_column* vay); + + gdf_error launch(gdf_column* out, gdf_column* vax, gdf_scalar* vay, gdf_scalar* def); + + gdf_error launch(gdf_column* out, gdf_column* vax, gdf_column* vay, gdf_scalar* def); + + private: + std::vector compilerFlags { "-std=c++14" }; + std::vector headersName { "operation.h" , "traits.h" , "gdf_data.h" }; + + private: + jitify::Program program; + + private: + std::string kernelName; + std::vector arguments; + }; + +} // namespace jit +} // namespace binops +} // namespace gdf + +#endif diff --git a/src/binary-operation/jit/util/operator.cpp b/src/binary-operation/jit/util/operator.cpp new file mode 100644 index 00000000..181c40a5 --- /dev/null +++ b/src/binary-operation/jit/util/operator.cpp @@ -0,0 +1,43 @@ +/* + * 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 +#include "binary-operation/jit/util/operator.h" +#include "binary-operation/jit/util/type.h" + +namespace gdf { +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)); + return buffer; + } + +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/util/operator.h b/src/binary-operation/jit/util/operator.h new file mode 100644 index 00000000..24958928 --- /dev/null +++ b/src/binary-operation/jit/util/operator.h @@ -0,0 +1,48 @@ +/* + * 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 "gdf/gdf.h" + +namespace gdf { +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 gdf + +#endif diff --git a/src/binary-operation/jit/util/type.cpp b/src/binary-operation/jit/util/type.cpp new file mode 100644 index 00000000..b6fc7cec --- /dev/null +++ b/src/binary-operation/jit/util/type.cpp @@ -0,0 +1,96 @@ +/* + * 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-operation/jit/util/type.h" + +namespace gdf { +namespace binops { +namespace jit { + + const char* getTypeName(gdf_dtype type) { + switch (type) { + case GDF_INT8: + return "int8_t"; + case GDF_INT16: + return "int16_t"; + case GDF_INT32: + case GDF_DATE32: + return "int32_t"; + case GDF_INT64: + case GDF_DATE64: + case GDF_TIMESTAMP: + return "int64_t"; + case GDF_UINT8: + return "uint8_t"; + case GDF_UINT16: + return "uint16_t"; + case GDF_UINT32: + return "uint32_t"; + case GDF_UINT64: + return "uint64_t"; + case GDF_FLOAT32: + return "float"; + case GDF_FLOAT64: + return "double"; + default: + return "double"; + } + } + + const char* 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_COMBINE: + //case GDF_COMBINE_FIRST: + //case GDF_ROUND: + 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"; + //GDF_PRODUCT, + //GDF_DOT + default: + return "None"; + } + } + +} // namespace jit +} // namespace binops +} // namespace gdf diff --git a/src/binary-operation/jit/util/type.h b/src/binary-operation/jit/util/type.h new file mode 100644 index 00000000..61e31f1c --- /dev/null +++ b/src/binary-operation/jit/util/type.h @@ -0,0 +1,35 @@ +/* + * 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 "gdf/gdf.h" + +namespace gdf { +namespace binops { +namespace jit { + + const char* getTypeName(gdf_dtype type); + + const char* getOperatorName(gdf_binary_operator ope); + +} +} +} + +#endif diff --git a/src/column.cpp b/src/column.cpp index 44ec1d26..8fe09109 100644 --- a/src/column.cpp +++ b/src/column.cpp @@ -72,6 +72,18 @@ gdf_error get_column_byte_width(gdf_column * col, int * width){ case GDF_TIMESTAMP : *width = 8; break; + case GDF_UINT8 : + *width = 1; + break; + case GDF_UINT16: + *width = 2; + break; + case GDF_UINT32: + *width = 4; + break; + case GDF_UINT64: + *width = 8; + break; default : *width = -1; return GDF_UNSUPPORTED_DTYPE; diff --git a/src/sqls_ops.cu b/src/sqls_ops.cu index 2d5307fc..0e84ef69 100644 --- a/src/sqls_ops.cu +++ b/src/sqls_ops.cu @@ -38,6 +38,23 @@ namespace{ //annonymus cudaMemcpy(d_types, h_types, ncols*sizeof(int), cudaMemcpyHostToDevice);//TODO: add streams } + //does not assume the columns were contiguous in memory but arer rather pointers to columns + void soa_col_info_arr(gdf_column** cols, size_t ncols, void** d_cols, int* d_types) + { + std::vector v_cols(ncols,nullptr); + std::vector v_types(ncols, 0); + for(size_t i=0;idata; + v_types[i] = cols[i]->dtype; + } + + void** h_cols = &v_cols[0]; + int* h_types = &v_types[0]; + cudaMemcpy(d_cols, h_cols, ncols*sizeof(void*), cudaMemcpyHostToDevice);//TODO: add streams + cudaMemcpy(d_types, h_types, ncols*sizeof(int), cudaMemcpyHostToDevice);//TODO: add streams + } + template using Vector = thrust::device_vector; @@ -1375,3 +1392,41 @@ gdf_error gdf_group_by_count(int ncols, // # columns } +gdf_error gdf_order_by_asc_desc( + gdf_column * input_columns, + size_t num_inputs, + gdf_column * output_indices, + gdf_valid_type * asc_desc_bitmask){ + + + //TODO: don't assume type of output is size_t + typedef size_t IndexT; + //TODO: make these allocations happen with the new memory manager when we can + //also we are kind of assuming they will just work, yeesh! + thrust::device_vector test(2); + + + thrust::device_vector d_cols(num_inputs); + thrust::device_vector d_types(num_inputs, 0); + + void** d_col_data = d_cols.data().get(); + int* d_col_types = d_types.data().get(); + + + + soa_col_info(input_columns, num_inputs, d_col_data, d_col_types); + + + multi_col_order_by_asc_desc( + d_col_data, + d_col_types, + num_inputs, + asc_desc_bitmask, + (size_t *) output_indices->data, + input_columns[0].size); + + return GDF_SUCCESS; + +} + + diff --git a/src/streamcompactionops.cu b/src/streamcompactionops.cu index bdbc0356..b96f2677 100644 --- a/src/streamcompactionops.cu +++ b/src/streamcompactionops.cu @@ -34,6 +34,7 @@ #include #include +#include //std lib #include @@ -74,28 +75,17 @@ private: typedef repeat_iterator > > gdf_valid_iterator; gdf_size_type get_number_of_bytes_for_valid (gdf_size_type column_size) { - return sizeof(gdf_valid_type) * (column_size + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE; + return sizeof(gdf_valid_type) * (column_size + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE; } // note: functor inherits from unary_function struct modulus_bit_width : public thrust::unary_function { - gdf_size_type n_bytes; - gdf_size_type column_size; - - modulus_bit_width (gdf_size_type b_nytes, gdf_size_type column_size) { - this->n_bytes = n_bytes; - this->column_size = column_size; - } __host__ __device__ gdf_size_type operator()(gdf_size_type x) const { - gdf_size_type col_position = x / 8; - gdf_size_type length_col = n_bytes != col_position+1 ? GDF_VALID_BITSIZE : column_size - GDF_VALID_BITSIZE * (n_bytes - 1); - //return x % GDF_VALID_BITSIZE; - return (length_col - 1) - (x % 8); - // x << + return x % GDF_VALID_BITSIZE; } }; @@ -108,11 +98,11 @@ struct shift_left: public thrust::unary_function } - __host__ __device__ - gdf_valid_type operator()(gdf_valid_type x) const - { - return x << num_bits; - } + __host__ __device__ + gdf_valid_type operator()(gdf_valid_type x) const + { + return x << num_bits; + } }; struct shift_right: public thrust::unary_function @@ -121,25 +111,25 @@ struct shift_right: public thrust::unary_function gdf_valid_type num_bits; bool not_too_many; shift_right(gdf_valid_type num_bits, bool not_too_many) - : num_bits(num_bits), not_too_many(not_too_many){ + : num_bits(num_bits), not_too_many(not_too_many){ } - __host__ __device__ - gdf_valid_type operator()(gdf_valid_type x) const - { - //if you want to force the shift to be fill bits with 0 you need to use an unsigned type - /*if (not_too_many) { // is the last - return x; + __host__ __device__ + gdf_valid_type operator()(gdf_valid_type x) const + { + //if you want to force the shift to be fill bits with 0 you need to use an unsigned type + /*if (not_too_many) { // is the last + return x; }*/ - return *((unsigned char *) &x) >> num_bits; + return *((unsigned char *) &x) >> num_bits; - } + } }; - + struct bit_or: public thrust::unary_function,gdf_valid_type> { - + __host__ __device__ gdf_valid_type operator()(thrust::tuple x) const @@ -147,7 +137,7 @@ struct bit_or: public thrust::unary_function(x) | thrust::get<1>(x); } }; - + typedef thrust::transform_iterator > bit_position_iterator; @@ -173,23 +163,35 @@ struct is_bit_set return ((thrust::get<0>(value) >> position) & 1); } -}; +}; struct bit_mask_pack_op : public thrust::unary_function { __host__ __device__ - gdf_valid_type operator()(const int64_t expanded) - { - gdf_valid_type result = 0; - for(unsigned int i = 0; i < GDF_VALID_BITSIZE; i++){ - // 0, 8, 16, ....,48, 56 - unsigned char byte = (expanded >> ( (GDF_VALID_BITSIZE - 1 - i ) * 8)); - result |= (byte & 1) << i; - } - return (result); + gdf_valid_type operator()(const int64_t expanded) + { + gdf_valid_type result = 0; + for(size_t i = 0; i < GDF_VALID_BITSIZE; i++){ + unsigned char byte = (expanded >> (i * 8)); + result |= (byte & 1) << i; } + return result; + } }; +struct is_valid_functor: public thrust::unary_function { + + + int8_t * data; + is_valid_functor(void * _data){ + data = (int8_t *) _data; + } + __device__ + int8_t operator()(const size_t index){ + return data[index]; + //return data[index / 8] & (1 << (index % 8)); + } +}; std::map column_type_width = {{GDF_INT8, sizeof(int8_t)}, {GDF_INT16, sizeof(int16_t)},{GDF_INT32, sizeof(int32_t)}, {GDF_INT64, sizeof(int64_t)}, {GDF_FLOAT32, sizeof(float)}, {GDF_FLOAT64, sizeof(double)} }; @@ -198,6 +200,8 @@ std::map column_type_width = {{GDF_INT8, sizeof(int8_t)}, {G //storing a map from gdf_type to width //TODO: add a way for the space where we store temp bitmaps for compaction be allocated //on the outside + +#include gdf_error gpu_apply_stencil(gdf_column *lhs, gdf_column * stencil, gdf_column * output){ //OK: add a rquire here that output and lhs are the same size GDF_REQUIRE(output->size == lhs->size, GDF_COLUMN_SIZE_MISMATCH); @@ -207,117 +211,158 @@ gdf_error gpu_apply_stencil(gdf_column *lhs, gdf_column * stencil, gdf_column * auto searched_item = column_type_width.find(lhs->dtype); int16_t width = searched_item->second; //width in bytes - searched_item = column_type_width.find(stencil->dtype); - int16_t stencil_width= searched_item->second; //width in bytes + //searched_item = column_type_width.find(stencil->dtype); + //int16_t stencil_width= searched_item->second; //width in bytes cudaStream_t stream; cudaStreamCreate(&stream); size_t n_bytes = get_number_of_bytes_for_valid(stencil->size); - bit_position_iterator bit_position_iter(thrust::make_counting_iterator(0), modulus_bit_width(n_bytes, stencil->size)); - gdf_valid_iterator valid_iterator(thrust::detail::make_normal_iterator(thrust::device_pointer_cast(stencil->valid)),GDF_VALID_BITSIZE); - //TODO: can probably make this happen with some kind of iterator so it can work on any width size - - //zip the stencil and the valid iterator together - typedef thrust::tuple >,gdf_valid_iterator, bit_position_iterator > zipped_stencil_tuple; - typedef thrust::zip_iterator zipped_stencil_iterator; - - //what kind of shit is that you might wonder? - //well basically we are zipping up an iterator to the stencil, one to the bit masks, and one which lets us get the bit position based on our index - zipped_stencil_iterator zipped_stencil_iter( - thrust::make_tuple( - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int8_t * )stencil->data)), - valid_iterator, - thrust::make_transform_iterator >( - thrust::make_counting_iterator(0), - modulus_bit_width(n_bytes, stencil->size)) - )); //NOTE!!!! the output column is getting set to a specific size but we are NOT compacting the allocation, //whoever calls that should handle that + + //temp storage for cub call + int *d_num_selected_out; + cudaMalloc(&d_num_selected_out,sizeof(int)); + void *d_temp_storage = NULL; + size_t temp_storage_bytes = 0; + int out_size; if(width == 1){ - thrust::detail::normal_iterator > input_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int8_t *) lhs->data)); - thrust::detail::normal_iterator > output_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int8_t *) output->data)); - thrust::detail::normal_iterator > output_end = - thrust::copy_if(thrust::cuda::par.on(stream),input_start,input_start + lhs->size,zipped_stencil_iter,output_start,is_stencil_true >::value_type >()); - output->size = output_end - output_start; + typedef int8_t data_type; + cub::DeviceSelect::Flagged(d_temp_storage, temp_storage_bytes, + (data_type *) lhs->data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + // Allocate temporary storage + cudaError_t err = cudaMalloc(&d_temp_storage, temp_storage_bytes); + if(err != cudaSuccess){ + std::cout<<"couldnt allocate temp space"<data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + + cudaMemcpyAsync(&out_size,d_num_selected_out,sizeof(int),cudaMemcpyDeviceToHost,stream); + output->size = out_size; + cudaStreamSynchronize(stream); + }else if(width == 2){ - thrust::detail::normal_iterator > input_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int16_t *) lhs->data)); - thrust::detail::normal_iterator > output_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int16_t *) output->data)); - thrust::detail::normal_iterator > output_end = - thrust::copy_if(thrust::cuda::par.on(stream),input_start,input_start + lhs->size,zipped_stencil_iter,output_start,is_stencil_true >::value_type >()); - output->size = output_end - output_start; + typedef int16_t data_type; + cub::DeviceSelect::Flagged(d_temp_storage, temp_storage_bytes, + (data_type *) lhs->data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + // Allocate temporary storage + cudaError_t err = cudaMalloc(&d_temp_storage, temp_storage_bytes); + if(err != cudaSuccess){ + std::cout<<"couldnt allocate temp space"<data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + + cudaMemcpyAsync(&out_size,d_num_selected_out,sizeof(int),cudaMemcpyDeviceToHost,stream); + output->size = out_size; + cudaStreamSynchronize(stream); }else if(width == 4){ - thrust::detail::normal_iterator > input_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int32_t *) lhs->data)); - thrust::detail::normal_iterator > output_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int32_t *) output->data)); - thrust::detail::normal_iterator > output_end = - thrust::copy_if(thrust::cuda::par.on(stream),input_start,input_start + lhs->size,zipped_stencil_iter,output_start,is_stencil_true >::value_type >()); - output->size = output_end - output_start; + typedef int32_t data_type; + cub::DeviceSelect::Flagged(d_temp_storage, temp_storage_bytes, + (data_type *) lhs->data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + // Allocate temporary storage + cudaError_t err = cudaMalloc(&d_temp_storage, temp_storage_bytes); + if(err != cudaSuccess){ + std::cout<<"couldnt allocate temp space"<data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + + cudaMemcpyAsync(&out_size,d_num_selected_out,sizeof(int),cudaMemcpyDeviceToHost,stream); + output->size = out_size; + cudaStreamSynchronize(stream); }else if(width == 8){ - thrust::detail::normal_iterator > input_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int64_t *) lhs->data)); - thrust::detail::normal_iterator > output_start = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int64_t *) output->data)); - thrust::detail::normal_iterator > output_end = - thrust::copy_if(thrust::cuda::par.on(stream),input_start,input_start + lhs->size,zipped_stencil_iter,output_start,is_stencil_true >::value_type >()); - output->size = output_end - output_start; - } - - gdf_size_type num_values = lhs->size; - //TODO:BRING OVER THE BITMASK!!! - //need to store a prefix sum - //align to size 8 - thrust::device_vector valid_bit_mask; //we are expanding the bit mask to an int8 because I can't envision an algorithm that operates on the bitmask that - if(num_values % GDF_VALID_BITSIZE != 0){ - valid_bit_mask.resize(num_values + (GDF_VALID_BITSIZE - (num_values % GDF_VALID_BITSIZE))); //align this allocation on GDF_VALID_BITSIZE so we don't have to bounds check - }else{ - valid_bit_mask.resize(num_values); + typedef int64_t data_type; + cub::DeviceSelect::Flagged(d_temp_storage, temp_storage_bytes, + (data_type *) lhs->data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + // Allocate temporary storage + cudaError_t err = cudaMalloc(&d_temp_storage, temp_storage_bytes); + if(err != cudaSuccess){ + std::cout<<"couldnt allocate temp space"<data, + (int8_t *) stencil->data, + (data_type *) output->data, d_num_selected_out, lhs->size,stream); + + cudaMemcpyAsync(&out_size,d_num_selected_out,sizeof(int),cudaMemcpyDeviceToHost,stream); + output->size = out_size; + cudaStreamSynchronize(stream); } + /* + if(false){ + gdf_size_type num_values = lhs->size; + //TODO:BRING OVER THE BITMASK!!! + //need to store a prefix sum + //align to size 8 + thrust::device_vector valid_bit_mask; //we are expanding the bit mask to an int8 because I can't envision an algorithm that operates on the bitmask that + if(num_values % GDF_VALID_BITSIZE != 0){ + valid_bit_mask.resize(num_values + (GDF_VALID_BITSIZE - (num_values % GDF_VALID_BITSIZE))); //align this allocation on GDF_VALID_BITSIZE so we don't have to bounds check + }else{ + valid_bit_mask.resize(num_values); + } - // doesn't require the use for a prefix sum which will have size 8 * num rows which is much larger than this + // doesn't require the use for a prefix sum which will have size 8 * num rows which is much larger than this - typedef thrust::tuple mask_tuple; - typedef thrust::zip_iterator zipped_mask; + typedef thrust::tuple mask_tuple; + typedef thrust::zip_iterator zipped_mask; - zipped_mask zipped_mask_iter( - thrust::make_tuple( - valid_iterator, - thrust::make_transform_iterator >( - thrust::make_counting_iterator(0), - modulus_bit_width(n_bytes, stencil->size)) - ) - ); + zipped_mask zipped_mask_iter( + thrust::make_tuple( + valid_iterator, + thrust::make_transform_iterator >( + thrust::make_counting_iterator(0), + modulus_bit_width()) + ) + ); - typedef thrust::transform_iterator bit_set_iterator; - bit_set_iterator bit_set_iter = thrust::make_transform_iterator( - zipped_mask_iter, - is_bit_set() - ); + typedef thrust::transform_iterator bit_set_iterator; + bit_set_iterator bit_set_iter = thrust::make_transform_iterator( + zipped_mask_iter, + is_bit_set() + ); - //copy the bitmask to device_vector of int8 - thrust::copy(thrust::cuda::par.on(stream), bit_set_iter, bit_set_iter + num_values, valid_bit_mask.begin()); + //copy the bitmask to device_vector of int8 + thrust::copy(thrust::cuda::par.on(stream), bit_set_iter, bit_set_iter + num_values, valid_bit_mask.begin()); - //remove the values that don't pass the stencil - thrust::remove_if(thrust::cuda::par.on(stream),valid_bit_mask.begin(), valid_bit_mask.begin() + num_values,zipped_stencil_iter, is_stencil_true >::value_type >()); + //remove the values that don't pass the stencil + thrust::remove_if(thrust::cuda::par.on(stream),valid_bit_mask.begin(), valid_bit_mask.begin() + num_values,zipped_stencil_iter, is_stencil_true >::value_type >()); - //recompact the values and store them in the output bitmask - //we can group them into pieces of 8 because we aligned this earlier on when we made the device_vector - thrust::detail::normal_iterator > valid_bit_mask_group_8_iter = - thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int64_t *) valid_bit_mask.data().get())); + //recompact the values and store them in the output bitmask + //we can group them into pieces of 8 because we aligned this earlier on when we made the device_vector + thrust::detail::normal_iterator > valid_bit_mask_group_8_iter = + thrust::detail::make_normal_iterator(thrust::device_pointer_cast((int64_t *) valid_bit_mask.data().get())); - //you may notice that we can write out more bytes than our valid_num_bytes, this only happens when we are not aligned to GDF_VALID_BITSIZE bytes, becasue the - //arrow standard requires 64 byte alignment, this is a safe assumption to make - thrust::transform(thrust::cuda::par.on(stream), valid_bit_mask_group_8_iter, valid_bit_mask_group_8_iter + ((num_values + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE), - thrust::detail::make_normal_iterator(thrust::device_pointer_cast(output->valid)),bit_mask_pack_op()); + //you may notice that we can write out more bytes than our valid_num_bytes, this only happens when we are not aligned to GDF_VALID_BITSIZE bytes, becasue the + //arrow standard requires 64 byte alignment, this is a safe assumption to make + thrust::transform(thrust::cuda::par.on(stream), valid_bit_mask_group_8_iter, valid_bit_mask_group_8_iter + ((num_values + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE), + thrust::detail::make_normal_iterator(thrust::device_pointer_cast(output->valid)),bit_mask_pack_op()); + } + */ cudaStreamSynchronize(stream); @@ -325,54 +370,54 @@ gdf_error gpu_apply_stencil(gdf_column *lhs, gdf_column * stencil, gdf_column * return GDF_SUCCESS; -} +} size_t get_last_byte_length(size_t column_size) { - size_t n_bytes = get_number_of_bytes_for_valid(column_size); - size_t length = column_size - GDF_VALID_BITSIZE * (n_bytes - 1); - if (n_bytes == 1 ) { - length = column_size; - } - return length; + size_t n_bytes = get_number_of_bytes_for_valid(column_size); + size_t length = column_size - GDF_VALID_BITSIZE * (n_bytes - 1); + if (n_bytes == 1 ) { + length = column_size; + } + return length; } size_t get_right_byte_length(size_t column_size, size_t iter, size_t left_length) { - size_t n_bytes = get_number_of_bytes_for_valid(column_size); - size_t length = column_size - GDF_VALID_BITSIZE * (n_bytes - 1); - if (iter == n_bytes - 1) { // the last one - if (left_length + length > GDF_VALID_BITSIZE) { - length = GDF_VALID_BITSIZE - left_length; - } - } - else { - length = GDF_VALID_BITSIZE - left_length; - } - return length; + size_t n_bytes = get_number_of_bytes_for_valid(column_size); + size_t length = column_size - GDF_VALID_BITSIZE * (n_bytes - 1); + if (iter == n_bytes - 1) { // the last one + if (left_length + length > GDF_VALID_BITSIZE) { + length = GDF_VALID_BITSIZE - left_length; + } + } + else { + length = GDF_VALID_BITSIZE - left_length; + } + return length; } - - - bool last_with_too_many_bits(size_t column_size, size_t iter, size_t left_length) { - size_t n_bytes = get_number_of_bytes_for_valid(column_size); - size_t length = column_size - GDF_VALID_BITSIZE * (n_bytes - 1); - if (iter == n_bytes) { // the last one - // the last one has to many bits - if (left_length + length > GDF_VALID_BITSIZE) { - return true; - } - } - return false; + + +bool last_with_too_many_bits(size_t column_size, size_t iter, size_t left_length) { + size_t n_bytes = get_number_of_bytes_for_valid(column_size); + size_t length = column_size - GDF_VALID_BITSIZE * (n_bytes - 1); + if (iter == n_bytes) { // the last one + // the last one has to many bits + if (left_length + length > GDF_VALID_BITSIZE) { + return true; + } + } + return false; } - gdf_valid_type concat_bins (gdf_valid_type A, gdf_valid_type B, int len_a, int len_b, bool has_next, size_t right_length){ - A = A << len_b; - if (!has_next) { - B = B << len_a; - B = B >> len_a; - } else { - B = B >> right_length - len_b; - } - return (A | B); +gdf_valid_type concat_bins (gdf_valid_type A, gdf_valid_type B, int len_a, int len_b, bool has_next, size_t right_length){ + A = A << len_b; + if (!has_next) { + B = B << len_a; + B = B >> len_a; + } else { + B = B >> right_length - len_b; + } + return (A | B); } gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) @@ -386,17 +431,17 @@ gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) cudaMemcpyAsync(output->data, lhs->data, type_width * lhs->size, cudaMemcpyDeviceToDevice, stream); cudaMemcpyAsync( (void *)( (int8_t*) (output->data) + type_width * lhs->size), rhs->data, type_width * rhs->size, cudaMemcpyDeviceToDevice, stream); - + int left_num_chars = get_number_of_bytes_for_valid(lhs->size); int right_num_chars = get_number_of_bytes_for_valid(rhs->size); - int output_num_chars = get_number_of_bytes_for_valid(output->size); - + int output_num_chars = get_number_of_bytes_for_valid(output->size); + thrust::device_ptr left_device_bits = thrust::device_pointer_cast((gdf_valid_type *)lhs->valid); thrust::device_ptr right_device_bits = thrust::device_pointer_cast((gdf_valid_type *)rhs->valid); thrust::device_ptr output_device_bits = thrust::device_pointer_cast((gdf_valid_type *)output->valid); thrust::copy(left_device_bits, left_device_bits + left_num_chars, output_device_bits); - + gdf_valid_type shift_bits = (GDF_VALID_BITSIZE - (lhs->size % GDF_VALID_BITSIZE)); if(shift_bits == 8){ shift_bits = 0; @@ -405,14 +450,14 @@ gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) size_t prev_len = get_last_byte_length(lhs->size); // copy all the rnbytes bytes from right column - if (shift_bits == 0) { + if (shift_bits == 0) { thrust::copy(right_device_bits, right_device_bits + right_num_chars, output_device_bits + left_num_chars); } - else { + else { thrust::host_vector last_byte (2); thrust::copy (left_device_bits + left_num_chars - 1, left_device_bits + left_num_chars, last_byte.begin()); thrust::copy (right_device_bits, right_device_bits + 1, last_byte.begin() + 1); - + size_t curr_len = get_right_byte_length(rhs->size, 0, prev_len); if (1 != right_num_chars) { @@ -423,7 +468,7 @@ gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) last_byte[0] = concat_bins(last_byte[0], last_byte[1], prev_len, curr_len, flag, last_right_byte_length); thrust::copy( last_byte.begin(), last_byte.begin() + 1, output_device_bits + left_num_chars - 1); - + if(right_num_chars > 1) { using first_iterator_type = thrust::transform_iterator::iterator>; using second_iterator_type = thrust::transform_iterator::iterator>; @@ -433,22 +478,22 @@ gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) auto too_many_bits = last_with_too_many_bits(rhs->size, right_num_chars, prev_len); size_t last_byte_length = get_last_byte_length(rhs->size); - if (last_byte_length >= (GDF_VALID_BITSIZE - shift_bits)) { // + if (last_byte_length >= (GDF_VALID_BITSIZE - shift_bits)) { // thrust::host_vector last_byte (right_device_bits + right_num_chars - 1, right_device_bits + right_num_chars); last_byte[0] = last_byte[0] << GDF_VALID_BITSIZE - last_byte_length; thrust::copy( last_byte.begin(), last_byte.begin() + 1, right_device_bits + right_num_chars - 1); } - + zipped_offset zipped_offset_iter( thrust::make_tuple( thrust::make_transform_iterator::iterator >( right_device_bits, shift_left(shift_bits)), - - thrust::make_transform_iterator::iterator >( - right_device_bits + 1, - shift_right(GDF_VALID_BITSIZE - shift_bits, !too_many_bits)) - ) + + thrust::make_transform_iterator::iterator >( + right_device_bits + 1, + shift_right(GDF_VALID_BITSIZE - shift_bits, !too_many_bits)) + ) ); //so what this does is give you an iterator which gives you a tuple where you have your char, and the char after you, so you can get the last bits! using transformed_or = thrust::transform_iterator; @@ -471,7 +516,7 @@ gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) thrust::copy (right_device_bits + right_num_chars - 1, right_device_bits + right_num_chars, last_byte.begin() + 1); last_byte[0] = last_byte[0] << last_byte_length | last_byte[1]; thrust::copy( last_byte.begin(), last_byte.begin() + 1, output_device_bits + output_num_chars - 1); - } + } } } if( last_with_too_many_bits(rhs->size, right_num_chars, prev_len)){ @@ -488,5 +533,3 @@ gdf_error gpu_concat(gdf_column *lhs, gdf_column *rhs, gdf_column *output) cudaStreamDestroy(stream); return GDF_SUCCESS; } - - diff --git a/src/tests/CMakeLists.txt b/src/tests/CMakeLists.txt index 4b4a5cf8..793be8bb 100644 --- a/src/tests/CMakeLists.txt +++ b/src/tests/CMakeLists.txt @@ -21,8 +21,15 @@ function(configure_test TEST_NAME Tests_SRCS) message(STATUS "${TEST_NAME} will link against: gdf arrow") + link_directories(${CUDA_LIBRARY_DIR}) + link_directories(${CUDA_LIBRARY_STUBS_DIR}) + cuda_add_executable(${TEST_NAME} ${Tests_SRCS}) + target_link_libraries(${TEST_NAME} gmock_main gmock GTest::GTest gdf arrow) + target_link_libraries(${TEST_NAME} ${CUDA_CUDA_LIB}) + target_link_libraries(${TEST_NAME} ${CUDA_NVRTC_LIB}) + set_target_properties(${TEST_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/gtests/") @@ -38,6 +45,7 @@ enable_testing() # ------------------- message(STATUS "******** Configuring tests ********") +add_subdirectory(binary-operation) add_subdirectory(foo-sample) add_subdirectory(datetime) add_subdirectory(hashing) diff --git a/src/tests/binary-operation/CMakeLists.txt b/src/tests/binary-operation/CMakeLists.txt new file mode 100644 index 00000000..b72829fa --- /dev/null +++ b/src/tests/binary-operation/CMakeLists.txt @@ -0,0 +1,19 @@ +#============================================================================= +# 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. +#============================================================================= + +add_subdirectory(integration) +add_subdirectory(unit) diff --git a/src/tests/binary-operation/integration/CMakeLists.txt b/src/tests/binary-operation/integration/CMakeLists.txt new file mode 100644 index 00000000..864c25a0 --- /dev/null +++ b/src/tests/binary-operation/integration/CMakeLists.txt @@ -0,0 +1,58 @@ +#============================================================================= +# 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. +#============================================================================= + +function(ConfigureTest test_name test_source) + ## Add link directory + link_directories(${CUDA_LIBRARY_DIR}) + link_directories(${CUDA_LIBRARY_STUBS_DIR}) + + ## Add test + add_executable(${test_name} ${test_source}) + add_test(NAME ${test_name} COMMAND ${test_name}) + + ## Clear include header property + set_target_properties(${test_name} PROPERTIES INCLUDE_DIRECTORIES "") + set_target_properties(${test_name} PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/gtests/") + + ## Include headers + target_include_directories(${test_name} PRIVATE "${GTEST_INCLUDE_DIRS}") + target_include_directories(${test_name} PRIVATE "${CUDA_INCLUDE_DIRS}") + target_include_directories(${test_name} PRIVATE "${PROJECT_SOURCE_DIR}/include") + target_include_directories(${test_name} PRIVATE "${PROJECT_SOURCE_DIR}/src") + + ## Link libraries + target_link_libraries(${test_name} ${GTEST_BOTH_LIBRARIES}) + target_link_libraries(${test_name} ${GTEST_LIBRARIES}) + target_link_libraries(${test_name} gdf) + target_link_libraries(${test_name} ${CUDA_CUDA_LIB}) + target_link_libraries(${test_name} ${CUDA_NVRTC_LIB}) +endfunction() + + +## BinaryOperationIntegrationTest +set(binop_source + "${PROJECT_SOURCE_DIR}/src/tests/binary-operation/util/types.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/binary-operation-integration-test.cpp") + +ConfigureTest("BinaryOperationIntegrationTest" "${binop_source}") + +## BinaryOperationOperandsNullTest +set(binop_source + "${PROJECT_SOURCE_DIR}/src/tests/binary-operation/util/types.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/binary-operation-operands-null-test.cpp") + +ConfigureTest("BinaryOperationOperandsNullTest" "${binop_source}") diff --git a/src/tests/binary-operation/integration/assert-binops.h b/src/tests/binary-operation/integration/assert-binops.h new file mode 100644 index 00000000..10b5e5aa --- /dev/null +++ b/src/tests/binary-operation/integration/assert-binops.h @@ -0,0 +1,247 @@ +/* + * 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 "gtest/gtest.h" +#include "tests/binary-operation/util/scalar.h" +#include "tests/binary-operation/util/vector.h" +#include "tests/binary-operation/util/operation.h" + +namespace gdf { +namespace test { +namespace binop { + +template +void ASSERT_BINOP(gdf::library::Vector& out, + gdf::library::Scalar& vax, + gdf::library::Vector& vay, + TypeOpe&& ope) { + ASSERT_TRUE(out.dataSize() == vay.dataSize()); + for (int index = 0; index < out.dataSize(); ++index) { + ASSERT_TRUE(out.data[index] == (TypeOut)(ope((TypeVay) vax.getValue(), vay.data[index]))); + } + + uint32_t vax_valid = (vax.isValid() ? UINT32_MAX : 0); + ASSERT_TRUE(out.validSize() == vay.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + ASSERT_TRUE(out.valid[index] == (vax_valid & vay.valid[index])); + } +} + +template +void ASSERT_BINOP(gdf::library::Vector& out, + gdf::library::Vector& vax, + gdf::library::Scalar& vay, + TypeOpe&& ope) { + ASSERT_TRUE(out.dataSize() == vax.dataSize()); + for (int index = 0; index < out.dataSize(); ++index) { + ASSERT_TRUE(out.data[index] == (TypeOut)(ope(vax.data[index], (TypeVay) vay.getValue()))); + } + + uint32_t vay_valid = (vay.isValid() ? UINT32_MAX : 0); + ASSERT_TRUE(out.validSize() == vax.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + ASSERT_TRUE(out.valid[index] == (vax.valid[index] & vay_valid)); + } +} + +template +void ASSERT_BINOP(gdf::library::Vector& out, + gdf::library::Vector& vax, + gdf::library::Vector& vay, + TypeOpe&& ope) { + ASSERT_TRUE(out.dataSize() == vax.dataSize()); + ASSERT_TRUE(out.dataSize() == vay.dataSize()); + for (int index = 0; index < out.dataSize(); ++index) { + ASSERT_TRUE(out.data[index] == (TypeOut)(ope(vax.data[index], vay.data[index]))); + } + + ASSERT_TRUE(out.validSize() == vax.validSize()); + ASSERT_TRUE(out.validSize() == vay.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + ASSERT_TRUE(out.valid[index] == vax.valid[index] | vay.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(gdf::library::Vector& out, + gdf::library::Vector& vax, + gdf::library::Vector& vay, + gdf::library::operation::Pow&& ope) { + const int ULP = 2.0; + ASSERT_TRUE(out.dataSize() == vax.dataSize()); + ASSERT_TRUE(out.dataSize() == vay.dataSize()); + for (int index = 0; index < out.dataSize(); ++index) { + ASSERT_TRUE(abs(out.data[index] - (TypeOut)(ope(vax.data[index], vay.data[index]))) < ULP); + } + + ASSERT_TRUE(out.validSize() == vax.validSize()); + ASSERT_TRUE(out.validSize() == vay.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + ASSERT_TRUE(out.valid[index] == (vax.valid[index] & vay.valid[index])); + } +} + +template +void ASSERT_BINOP(gdf::library::Vector& out, + gdf::library::Scalar& vax, + gdf::library::Vector& vay, + gdf::library::Scalar& def, + TypeOpe&& ope) { + using ValidType = typename gdf::library::Vector::ValidType; + int ValidSize = gdf::library::Vector::ValidSize; + + ASSERT_TRUE(out.dataSize() == vay.dataSize()); + ASSERT_TRUE(out.validSize() == vay.validSize()); + + ValidType mask = 1; + int index_valid = 0; + for (int index = 0; index < out.dataSize(); ++index) { + if (!(index % ValidSize)) { + mask = 1; + index_valid = index / ValidSize; + } else { + mask <<= 1; + } + + TypeVay vax_aux = (TypeVay)vax; + if (!vax.isValid()) { + vax_aux = (TypeVay)((TypeVal) def.getValue()); + } + + TypeVax vay_aux = vay.data[index]; + if ((vay.valid[index_valid] & mask) == 0) { + vay_aux = (TypeVal) def.getValue(); + } + + ASSERT_TRUE(out.data[index] == (TypeOut)(ope(vax_aux, vay_aux))); + } + + uint32_t vax_valid = (vax.isValid() ? UINT32_MAX : 0); + uint32_t def_valid = (def.isValid() ? UINT32_MAX : 0); + ASSERT_TRUE(out.validSize() == vay.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + uint32_t output = (vay.valid[index] & vax_valid) | + (vay.valid[index] & def_valid) | + (vax_valid & def_valid); + ASSERT_TRUE(out.valid[index] == output); + } +} + +template +void ASSERT_BINOP(gdf::library::Vector& out, + gdf::library::Vector& vax, + gdf::library::Scalar& vay, + gdf::library::Scalar& def, + TypeOpe&& ope) { + using ValidType = typename gdf::library::Vector::ValidType; + int ValidSize = gdf::library::Vector::ValidSize; + + ASSERT_TRUE(out.dataSize() == vax.dataSize()); + ASSERT_TRUE(out.validSize() == vax.validSize()); + + ValidType mask = 1; + int index_valid = 0; + for (int index = 0; index < out.dataSize(); ++index) { + if (!(index % ValidSize)) { + mask = 1; + index_valid = index / ValidSize; + } else { + mask <<= 1; + } + + TypeVax vax_aux = vax.data[index]; + if ((vax.valid[index_valid] & mask) == 0) { + vax_aux = (TypeVal) def.getValue(); + } + + TypeVay vay_aux = (TypeVay)vay; + if (!vay.isValid()) { + vay_aux = (TypeVay)((TypeVal) def.getValue()); + } + + ASSERT_TRUE(out.data[index] == (TypeOut)(ope(vax_aux, vay_aux))); + } + + uint32_t vay_valid = (vay.isValid() ? UINT32_MAX : 0); + uint32_t def_valid = (def.isValid() ? UINT32_MAX : 0); + ASSERT_TRUE(out.validSize() == vax.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + uint32_t output = (vax.valid[index] & vay_valid) | + (vax.valid[index] & def_valid) | + (vay_valid & def_valid); + ASSERT_TRUE(out.valid[index] == output); + } +} + +template +void ASSERT_BINOP(gdf::library::Vector& out, + gdf::library::Vector& vax, + gdf::library::Vector& vay, + gdf::library::Scalar& def, + TypeOpe&& ope) { + using ValidType = typename gdf::library::Vector::ValidType; + int ValidSize = gdf::library::Vector::ValidSize; + + ASSERT_TRUE(out.dataSize() == vax.dataSize()); + ASSERT_TRUE(out.dataSize() == vay.dataSize()); + + ValidType mask = 1; + int index_valid = 0; + for (int index = 0; index < out.dataSize(); ++index) { + if (!(index % ValidSize)) { + mask = 1; + index_valid = index / ValidSize; + } else { + mask <<= 1; + } + + TypeVax vax_aux = vax.data[index]; + if ((vax.valid[index_valid] & mask) == 0) { + vax_aux = (TypeVax) def.getValue(); + } + + TypeVay vay_aux = vay.data[index]; + if ((vay.valid[index_valid] & mask) == 0) { + vay_aux = (TypeVay) def.getValue(); + } + + ASSERT_TRUE(out.data[index] == (TypeOut)(ope(vax_aux, vay_aux))); + } + + uint32_t def_valid = (def.isValid() ? UINT32_MAX : 0); + ASSERT_TRUE(out.validSize() == vax.validSize()); + ASSERT_TRUE(out.validSize() == vay.validSize()); + for (int index = 0; index < out.validSize(); ++index) { + ASSERT_TRUE(out.valid[index] == ((vax.valid[index] & vay.valid[index]) | + (vax.valid[index] & def_valid) | + (vay.valid[index] & def_valid))); + } +} + +} // namespace binop +} // namespace test +} // namespace gdf + +#endif diff --git a/src/tests/binary-operation/integration/binary-operation-integration-test.cpp b/src/tests/binary-operation/integration/binary-operation-integration-test.cpp new file mode 100644 index 00000000..a1da9947 --- /dev/null +++ b/src/tests/binary-operation/integration/binary-operation-integration-test.cpp @@ -0,0 +1,453 @@ +/* + * 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-operation/integration/assert-binops.h" + +namespace gdf { +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_UI32) { + using SI32 = gdf::library::GdfEnumType; + using FP32 = gdf::library::GdfEnumType; + using UI32 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Scalar vax; + gdf::library::Vector vay; + + vay.rangeData(0, 100000, 1) + .rangeValid(false, 0, 4); + vax.setValue(100); + out.emplaceVector(vay.dataSize()); + + auto result = gdf_binary_operation_v_s_v(out.column(), vax.scalar(), vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Sub_Scalar_Vector_SI32_FP32_UI32) { + using SI32 = gdf::library::GdfEnumType; + using FP32 = gdf::library::GdfEnumType; + using UI32 = gdf::library::GdfEnumType; + using SUB = gdf::library::operation::Sub; + + gdf::library::Vector out; + gdf::library::Scalar vax; + gdf::library::Vector vay; + + vay.rangeData(0, 100000, 1) + .rangeValid(false, 0, 4); + vax.setValue(10000); + out.emplaceVector(vay.dataSize()); + + auto result = gdf_binary_operation_v_s_v(out.column(), vax.scalar(), vay.column(), GDF_SUB); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, SUB()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Vector_Scalar_SI08_UI16_SI16) { + using SI08 = gdf::library::GdfEnumType; + using UI16 = gdf::library::GdfEnumType; + using SI16 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 6); + vay.setValue(100); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s(out.column(), vax.column(), vay.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Vector_Vector_UI32_FP64_SI08) { + using UI32 = gdf::library::GdfEnumType; + using FP64 = gdf::library::GdfEnumType; + using SI08 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(0.0, 200.0, 2.0) + .rangeValid(false, 0, 3); + vay.rangeData(0, 100, 1) + .rangeValid(false, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Scalar_Vector_Default_SI32_SI16_UI64_SI64) { + using SI32 = gdf::library::GdfEnumType; + using SI16 = gdf::library::GdfEnumType; + using UI64 = gdf::library::GdfEnumType; + using SI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Scalar vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.setValue(50); + vay.rangeData(0, 10000, 2) + .rangeValid(false, 0, 4); + def.setValue(1000); + out.emplaceVector(vay.dataSize()); + + auto result = gdf_binary_operation_v_s_v_d(out.column(), vax.scalar(), vay.column(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Sub_Scalar_Vector_Default_SI32_SI16_UI64_SI64) { + using SI32 = gdf::library::GdfEnumType; + using SI16 = gdf::library::GdfEnumType; + using UI64 = gdf::library::GdfEnumType; + using SI64 = gdf::library::GdfEnumType; + using SUB = gdf::library::operation::Sub; + + gdf::library::Vector out; + gdf::library::Scalar vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.setValue(500); + vay.rangeData(0, 10000, 2) + .rangeValid(false, 0, 4); + def.setValue(1000); + out.emplaceVector(vay.dataSize()); + + auto result = gdf_binary_operation_v_s_v_d(out.column(), vax.scalar(), vay.column(), def.scalar(), GDF_SUB); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, SUB()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Vector_Scalar_Default_FP32_SI16_UI08_UI32) { + using FP32 = gdf::library::GdfEnumType; + using SI16 = gdf::library::GdfEnumType; + using UI08 = gdf::library::GdfEnumType; + using UI32 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 30000, 3) + .rangeValid(false, 0, 4); + vay.setValue(50); + def.setValue(150); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Add_Vector_Vector_Default_FP64_SI32_UI32_UI16) { + using FP64 = gdf::library::GdfEnumType; + using SI32 = gdf::library::GdfEnumType; + using UI32 = gdf::library::GdfEnumType; + using UI16 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100000, 1) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200000, 2) + .rangeValid(false, 0, 4); + def.setValue(150); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Sub_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using SUB = gdf::library::operation::Sub; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(100000, 200000, 2) + .rangeValid(true, 0, 4); + vay.rangeData(50000, 100000, 1) + .rangeValid(false, 0, 3); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_SUB); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, SUB()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mul_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using MUL = gdf::library::operation::Mul; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(100000, 200000, 2) + .rangeValid(false, 0 , 3); + vay.rangeData(50000, 100000, 1) + .rangeValid(false, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_MUL); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, MUL()); +} + + +TEST_F(BinaryOperationIntegrationTest, Div_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using DIV = gdf::library::operation::Div; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(100000, 200000, 2) + .rangeValid(false, 0, 6); + vay.rangeData(50000, 100000, 1) + .rangeValid(false, 0, 8); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_DIV); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, DIV()); +} + + +TEST_F(BinaryOperationIntegrationTest, TrueDiv_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using TRUEDIV = gdf::library::operation::TrueDiv; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(100000, 200000, 2) + .rangeValid(true, 0, 3); + vay.rangeData(50000, 100000, 1) + .rangeValid(true, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_TRUE_DIV); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, TRUEDIV()); +} + + +TEST_F(BinaryOperationIntegrationTest, FloorDiv_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using FLOORDIV = gdf::library::operation::FloorDiv; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(100000, 200000, 2) + .rangeValid(false, 0, 6); + vay.rangeData(50000, 100000, 1) + .rangeValid(false, 0, 8); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_FLOOR_DIV); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, FLOORDIV()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mod_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using MOD = gdf::library::operation::Mod; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(120, 220, 2) + .rangeValid(false, 0, 3); + vay.rangeData(50, 100, 1) + .rangeValid(false, 0, 5); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_MOD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, MOD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mod_Vector_Vector_FP32) { + using FP32 = gdf::library::GdfEnumType; + using MOD = gdf::library::operation::Mod; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(120, 220, 2) + .rangeValid(false, 0, 4); + vay.rangeData(50, 100, 1) + .rangeValid(false, 0, 6); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_MOD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, MOD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Mod_Vector_Vector_FP64) { + using FP64 = gdf::library::GdfEnumType; + using MOD = gdf::library::operation::Mod; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(120, 220, 2) + .rangeValid(true, 0, 3); + vay.rangeData(50, 100, 1) + .rangeValid(false, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_MOD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, MOD()); +} + + +TEST_F(BinaryOperationIntegrationTest, Pow_Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using POW = gdf::library::operation::Pow; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(0, 500, 1) + .rangeValid(false, 0, 6); + vay.fillData(500, 2) + .rangeValid(false, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_POW); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, POW()); +} + +} // namespace binop +} // namespace test +} // namespace gdf diff --git a/src/tests/binary-operation/integration/binary-operation-operands-null-test.cpp b/src/tests/binary-operation/integration/binary-operation-operands-null-test.cpp new file mode 100644 index 00000000..22f27461 --- /dev/null +++ b/src/tests/binary-operation/integration/binary-operation-operands-null-test.cpp @@ -0,0 +1,514 @@ +/* + * 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-operation/integration/assert-binops.h" + +namespace gdf { +namespace test { +namespace binop { + +struct BinaryOperationOperandsNullTest : public ::testing::Test { + BinaryOperationOperandsNullTest() { + } + + virtual ~BinaryOperationOperandsNullTest() { + } + + virtual void SetUp() { + } + + virtual void TearDown() { + } +}; + +/* + * Kernels v_v_s, using UI64 + * Output:Vector, OperandX:Vector, OperandY:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_UI64_WithScalarOperandNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.setValue(500) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s(out.column(), vax.column(), vay.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_UI64_WithScalarOperandNotNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.setValue(500) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s(out.column(), vax.column(), vay.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + +/* + * Kernels v_v_s, using FP64 + * Output:Vector, OperandX:Vector, OperandY:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_FP64_WithScalarOperandNull) { + using FP64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.setValue(500.0) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s(out.column(), vax.column(), vay.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_FP64_WithScalarOperandNotNull) { + using FP64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.setValue(500.0) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s(out.column(), vax.column(), vay.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + +/* + * Kernels v_v_v, using UI64 + * Output:Vector, OperandX:Vector, OperandY:Vector + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_UI64) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200, 2) + .rangeValid(false, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + +/* + * Kernels v_v_v, using FP64 + * Output:Vector, OperandX:Vector, OperandY:Vector + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_FP64) { + using FP64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200.0, 2.0) + .rangeValid(false, 0, 4); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v(out.column(), vax.column(), vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, ADD()); +} + +/* + * Kernels v_v_s_d, using UI64 + * Output:Vector, OperandX:Vector, OperandY:Scalar, OperandDefault:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_UI64_WithAllOperandsNotNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.setValue(222) + .setValid(true); + def.setValue(555) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_UI64_WithScalarOperandNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.setValue(1000) + .setValid(false); + def.setValue(500) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_UI64_WithDefaultOperandNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.setValue(250) + .setValid(true); + def.setValue(750) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_UI64_WithAllOperandsNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.setValue(500) + .setValid(false); + def.setValue(1000) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + +/* + * Kernels v_v_s_d, using FP32 + * Output:Vector, OperandX:Vector, OperandY:Scalar, OperandDefault:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_FP32_WithAllOperandsNotNull) { + using FP32 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.setValue(222.0) + .setValid(true); + def.setValue(555.0) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_FP32_WithScalarOperandNull) { + using FP32 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.setValue(1000.0) + .setValid(false); + def.setValue(500.0) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_FP32_WithDefaultOperandNull) { + using FP32 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.setValue(250.0) + .setValid(true); + def.setValue(750.0) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_FP32_WithAllOperandsNull) { + using FP32 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Scalar vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.setValue(500.0) + .setValid(false); + def.setValue(1000.0) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_s_d(out.column(), vax.column(), vay.scalar(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + +/* + * Kernels v_v_v_d, using UI64 + * Output:Vector, OperandX:Vector, OperandY:Vector, OperandDefault:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_Default_UI64_WithDefaultOperandNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200, 2) + .rangeValid(false, 0, 4); + def.setValue(666) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_Default_UI64_WithDefaultOperandNotNull) { + using UI64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100, 1) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200, 2) + .rangeValid(false, 0, 4); + def.setValue(222) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + +/* + * Kernels v_v_v_d, using FP64 + * Output:Vector, OperandX:Vector, OperandY:Vector, OperandDefault:Scalar + */ +TEST_F(BinaryOperationOperandsNullTest, Vector_Scalar_Default_FP64_WithDefaultOperandNull) { + using FP64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200.0, 2.0) + .rangeValid(false, 0, 4); + def.setValue(555.0) + .setValid(false); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + + +TEST_F(BinaryOperationOperandsNullTest, Vector_Vector_Default_FP64_WithDefaultOperandNotNull) { + using FP64 = gdf::library::GdfEnumType; + using ADD = gdf::library::operation::Add; + + gdf::library::Vector out; + gdf::library::Vector vax; + gdf::library::Vector vay; + gdf::library::Scalar def; + + vax.rangeData(0, 100.0, 1.0) + .rangeValid(false, 0, 3); + vay.rangeData(0, 200.0, 2.0) + .rangeValid(false, 0, 4); + def.setValue(555.0) + .setValid(true); + out.emplaceVector(vax.dataSize()); + + auto result = gdf_binary_operation_v_v_v_d(out.column(), vax.column(), vay.column(), def.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); + + out.readVector(); + + ASSERT_BINOP(out, vax, vay, def, ADD()); +} + +} // namespace binop +} // namespace test +} // namespace gdf diff --git a/src/tests/binary-operation/unit/CMakeLists.txt b/src/tests/binary-operation/unit/CMakeLists.txt new file mode 100644 index 00000000..7a2375e4 --- /dev/null +++ b/src/tests/binary-operation/unit/CMakeLists.txt @@ -0,0 +1,54 @@ +#============================================================================= +# 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. +#============================================================================= + +## Test name +set(GDF_TEST_NAME BinaryOperationVerifyInputTest) + + +## Set sources +set(GDF_SOURCE_FILES + "${PROJECT_SOURCE_DIR}/src/tests/binary-operation/util/types.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/binop-verify-input-test.cpp") + + +## Add link directory +link_directories(${CUDA_LIBRARY_DIR}) +link_directories(${CUDA_LIBRARY_STUBS_DIR}) + + +## Add test +add_executable(${GDF_TEST_NAME} ${GDF_SOURCE_FILES}) +add_test(NAME ${GDF_TEST_NAME} COMMAND ${GDF_TEST_NAME}) + + +## Clear include header property +set_target_properties(${GDF_TEST_NAME} PROPERTIES INCLUDE_DIRECTORIES "") +set_target_properties(${GDF_TEST_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/gtests/") + + +## Include headers +target_include_directories(${GDF_TEST_NAME} PRIVATE "${GTEST_INCLUDE_DIRS}") +target_include_directories(${GDF_TEST_NAME} PRIVATE "${CUDA_INCLUDE_DIRS}") +target_include_directories(${GDF_TEST_NAME} PRIVATE "${PROJECT_SOURCE_DIR}/include") +target_include_directories(${GDF_TEST_NAME} PRIVATE "${PROJECT_SOURCE_DIR}/src") + + +## Link libraries +target_link_libraries(${GDF_TEST_NAME} ${GTEST_BOTH_LIBRARIES}) +target_link_libraries(${GDF_TEST_NAME} gdf) +target_link_libraries(${GDF_TEST_NAME} ${CUDA_CUDA_LIB}) +target_link_libraries(${GDF_TEST_NAME} ${CUDA_NVRTC_LIB}) diff --git a/src/tests/binary-operation/unit/binop-verify-input-test.cpp b/src/tests/binary-operation/unit/binop-verify-input-test.cpp new file mode 100644 index 00000000..706a934a --- /dev/null +++ b/src/tests/binary-operation/unit/binop-verify-input-test.cpp @@ -0,0 +1,571 @@ +/* + * 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 "gtest/gtest.h" +#include "tests/binary-operation/util/scalar.h" +#include "tests/binary-operation/util/vector.h" + +struct BinopVerifyInputTest : public ::testing::Test { + BinopVerifyInputTest() { + } + + virtual ~BinopVerifyInputTest() { + } + + virtual void SetUp() { + } + + virtual void TearDown() { + } +}; + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOutputVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.fillData(0, 0); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), vector_vax.column(), scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.fillData(0, 0); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), vector_vax.column(), scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOutputVectorNull) { + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + auto result = gdf_binary_operation_v_v_s(nullptr, vector_vax.column(), scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandVectorNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), nullptr, scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandScalarNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), vector_vax.column(), nullptr, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOutputVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + vector_out.column()->dtype = (gdf_dtype)100; + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), vector_vax.column(), scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + vector_vax.column()->dtype = (gdf_dtype)100; + + gdf::library::Scalar scalar; + scalar.setValue(100); + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), vector_vax.column(), scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_ErrorOperandScalarType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + scalar.scalar()->dtype = (gdf_dtype)100; + + auto result = gdf_binary_operation_v_v_s(vector_out.column(), vector_vax.column(), scalar.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorOutputVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.fillData(0, 0); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorFirstOperandVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.fillData(0, 0); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorSecondOperandVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.fillData(0, 0); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorOutputVectorNull) { + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(nullptr, vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorFirstOperandVectorNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), nullptr, vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorSecondOperandVectorNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), nullptr, GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorOutputVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + vector_out.column()->dtype = (gdf_dtype)100; + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorFirstOperandVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + vector_vax.column()->dtype = (gdf_dtype)100; + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_ErrorSecondOperandVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + vector_vay.column()->dtype = (gdf_dtype)100; + + auto result = gdf_binary_operation_v_v_v(vector_out.column(), vector_vax.column(), vector_vay.column(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOutputVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.fillData(0, 0); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), vector_vax.column(), scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOperandVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.fillData(0, 0); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), vector_vax.column(), scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOutputVectorNull) { + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(nullptr, vector_vax.column(), scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOperandVectorNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), nullptr, scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOperandScalarNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), vector_vax.column(), nullptr, defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOutputVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + vector_out.column()->dtype = (gdf_dtype)100; + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), vector_vax.column(), scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOperandVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + vector_vax.column()->dtype = (gdf_dtype)100; + + gdf::library::Scalar scalar; + scalar.setValue(100); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), vector_vax.column(), scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Scalar_Default_ErrorOperandScalarType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar scalar; + scalar.setValue(100); + scalar.scalar()->dtype = (gdf_dtype)100; + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_s_d(vector_out.column(), vector_vax.column(), scalar.scalar(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorOutputVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.fillData(0, 0); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorFirstOperandVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.fillData(0, 0); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorSecondOperandVectorZeroSize) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.fillData(0, 0); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_SUCCESS); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorOutputVectorNull) { + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(nullptr, vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorFirstOperandVectorNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), nullptr, vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorSecondOperandVectorNull) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), nullptr, defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_DATASET_EMPTY); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorOutputVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + vector_out.column()->dtype = (gdf_dtype)100; + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorFirstOperandVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + vector_vax.column()->dtype = (gdf_dtype)100; + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} + + +TEST_F(BinopVerifyInputTest, Vector_Vector_Default_ErrorSecondOperandVectorType) { + gdf::library::Vector vector_out; + vector_out.rangeData(1, 10, 1); + + gdf::library::Vector vector_vax; + vector_vax.rangeData(1, 10, 1); + + gdf::library::Vector vector_vay; + vector_vay.rangeData(1, 10, 1); + vector_vay.column()->dtype = (gdf_dtype)100; + + gdf::library::Scalar defvalue; + defvalue.setValue(100); + + auto result = gdf_binary_operation_v_v_v_d(vector_out.column(), vector_vax.column(), vector_vay.column(), defvalue.scalar(), GDF_ADD); + ASSERT_TRUE(result == GDF_UNSUPPORTED_DTYPE); +} diff --git a/src/tests/binary-operation/util/field.h b/src/tests/binary-operation/util/field.h new file mode 100644 index 00000000..47d8016d --- /dev/null +++ b/src/tests/binary-operation/util/field.h @@ -0,0 +1,113 @@ +/* + * 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_FIELD_H +#define GDF_TESTS_BINARY_OPERATION_UTIL_FIELD_H + +#include +#include +#include + +namespace gdf { +namespace library { + + template + class Field { + public: + ~Field() { + destroy(); + } + + public: + void clear() { + mCpuData.clear(); + destroy(); + } + + void resize(int size) { + int sizeBytes = size * sizeof(Type); + if (sizeBytes != mSizeAllocBytes) { + mCpuData.resize(size); + destroy(); + create(size); + } + } + + public: + auto getGpuData() -> Type* { + return mGpuData; + } + + public: + auto begin() -> typename std::vector::iterator { + return mCpuData.begin(); + } + + auto end() -> typename std::vector::iterator { + return mCpuData.end(); + } + + auto back() -> typename std::vector::reference { + return mCpuData.back(); + } + + public: + auto size() -> std::size_t { + return mCpuData.size(); + } + + auto operator[](int index) -> Type& { + assert(index < mCpuData.size()); + return mCpuData[index]; + } + + public: + void write() { + if (mSizeAllocBytes) { + cudaMemcpy(mGpuData, mCpuData.data(), mSizeAllocBytes, cudaMemcpyHostToDevice); + } + } + + void read() { + if (mSizeAllocBytes) { + cudaMemcpy(mCpuData.data(), mGpuData, mSizeAllocBytes, cudaMemcpyDeviceToHost); + } + } + + protected: + void create(int size) { + mSizeAllocBytes = size * sizeof(Type); + cudaMalloc((void**)&(mGpuData), mSizeAllocBytes); + } + + void destroy() { + if (mSizeAllocBytes) { + mSizeAllocBytes = 0; + cudaFree(mGpuData); + } + } + + private: + int mSizeAllocBytes {0}; + Type* mGpuData {nullptr}; + std::vector mCpuData; + }; + +} // namespace library +} // namespace gdf + +#endif diff --git a/src/tests/binary-operation/util/operation.h b/src/tests/binary-operation/util/operation.h new file mode 100644 index 00000000..4db30d04 --- /dev/null +++ b/src/tests/binary-operation/util/operation.h @@ -0,0 +1,111 @@ +/* + * 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 + +namespace gdf { +namespace library { +namespace operation { + + template + struct Add { + TypeOut operator()(TypeVax vax, TypeVay vay) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)vax + (TypeCommon)vay); + } + }; + + template + struct Sub { + TypeOut operator()(TypeVax vax, TypeVay vay) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)vax - (TypeCommon)vay); + } + }; + + template + struct Mul { + TypeOut operator()(TypeVax vax, TypeVay vay) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)vax * (TypeCommon)vay); + } + }; + + template + struct Div { + TypeOut operator()(TypeVax vax, TypeVay vay) { + using TypeCommon = typename std::common_type::type; + return (TypeOut)((TypeCommon)vax / (TypeCommon)vay); + } + }; + + template + struct TrueDiv { + TypeOut operator()(TypeVax vax, TypeVay vay) { + return (TypeOut)((double)vax / (double)vay); + } + }; + + template + struct FloorDiv { + TypeOut operator()(TypeVax vax, TypeVay vay) { + return (TypeOut)floor((double)vax / (double)vay); + } + }; + + template ::type> + struct Mod; + + template + struct Mod { + TypeOut operator()(TypeVax x, TypeVay y) { + return (TypeOut)((uint64_t)x % (uint64_t)y); + } + }; + + template + struct Mod { + TypeOut operator()(TypeVax x, TypeVay y) { + return (TypeOut)fmod((float)x, (float)y); + } + }; + + template + struct Mod { + TypeOut operator()(TypeVax x, TypeVay y) { + return (TypeOut)fmod((double)x, (double)y); + } + }; + + template + struct Pow { + TypeOut operator()(TypeVax vax, TypeVay vay) { + return (TypeOut)pow((double)vax, (double)vay); + } + }; + +} // namespace operation +} // namespace library +} // namespace gdf + +#endif diff --git a/src/tests/binary-operation/util/scalar.h b/src/tests/binary-operation/util/scalar.h new file mode 100644 index 00000000..defea611 --- /dev/null +++ b/src/tests/binary-operation/util/scalar.h @@ -0,0 +1,108 @@ +/* + * 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_SCALAR_H +#define GDF_TESTS_BINARY_OPERATION_UTIL_SCALAR_H + +#include "gdf/gdf.h" +#include "tests/binary-operation/util/types.h" + +namespace gdf { +namespace library { + + template + class Scalar { + public: + Scalar& setValue(Type value) { + mScalar.dtype = gdf::library::GdfDataType::Value; + gdf::library::setScalar(mScalar, value); + mScalar.is_valid = true; + return *this; + } + + Scalar& setValid(bool value) { + mScalar.is_valid = value; + return *this; + } + + public: + Type getValue() { + return (Type)*this; + } + + gdf_dtype getType() { + return mScalar.dtype; + } + + bool isValid() { + return mScalar.is_valid; + } + + public: + gdf_scalar* scalar() { + return &mScalar; + } + + public: + operator int8_t() const { + return mScalar.data.si08; + } + + operator int16_t() const { + return mScalar.data.si16; + } + + operator int32_t() const { + return mScalar.data.si32; + } + + operator int64_t() const { + return mScalar.data.si64; + } + + operator uint8_t() const { + return mScalar.data.ui08; + } + + operator uint16_t() const { + return mScalar.data.ui16; + } + + operator uint32_t() const { + return mScalar.data.ui32; + } + + operator uint64_t() const { + return mScalar.data.ui64; + } + + operator float() const { + return mScalar.data.fp32; + } + + operator double() const { + return mScalar.data.fp64; + } + + private: + gdf_scalar mScalar; + }; + +} // namespace library +} // namespace gdf + +#endif diff --git a/src/tests/binary-operation/util/types.cpp b/src/tests/binary-operation/util/types.cpp new file mode 100644 index 00000000..4245a758 --- /dev/null +++ b/src/tests/binary-operation/util/types.cpp @@ -0,0 +1,137 @@ +/* + * 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-operation/util/types.h" + +namespace gdf { +namespace library { + + const char* getTypeName(gdf_dtype type) { + switch (type) { + case GDF_INT8: + return "GDF_INT8"; + case GDF_INT16: + return "GDF_INT16"; + case GDF_INT32: + return "GDF_INT32"; + case GDF_INT64: + return "GDF_INT64"; + case GDF_UINT8: + return "GDF_UINT8"; + case GDF_UINT16: + return "GDF_UINT16"; + case GDF_UINT32: + return "GDF_UINT32"; + case GDF_UINT64: + return "GDF_UINT64"; + case GDF_FLOAT32: + return "GDF_FLOAT32"; + case GDF_FLOAT64: + return "GDF_FLOAT64"; + case GDF_DATE32: + return "GDF_DATE32"; + case GDF_DATE64: + return "GDF_DATE64"; + case GDF_TIMESTAMP: + return "GDF_TIMESTAMP"; + } + } + + namespace helper { + void setScalar(gdf_scalar& scalar, int8_t value) { + scalar.data.si08 = value; + } + + void setScalar(gdf_scalar& scalar, int16_t value) { + scalar.data.si16 = value; + } + + void setScalar(gdf_scalar& scalar, int32_t value) { + scalar.data.si32 = value; + } + + void setScalar(gdf_scalar& scalar, int64_t value) { + scalar.data.si64 = value; + } + + void setScalar(gdf_scalar& scalar, uint8_t value) { + scalar.data.ui08 = value; + } + + void setScalar(gdf_scalar& scalar, uint16_t value) { + scalar.data.ui16 = value; + } + + void setScalar(gdf_scalar& scalar, uint32_t value) { + scalar.data.ui32 = value; + } + + void setScalar(gdf_scalar& scalar, uint64_t value) { + scalar.data.ui64 = value; + } + + void setScalar(gdf_scalar& scalar, float value) { + scalar.data.fp32 = value; + } + + void setScalar(gdf_scalar& scalar, double value) { + scalar.data.fp64 = value; + } + } + + int8_t getScalar(int8_t, gdf_scalar* scalar) { + return scalar->data.si08; + } + + int16_t getScalar(int16_t, gdf_scalar* scalar) { + return scalar->data.si16; + } + + int32_t getScalar(int32_t, gdf_scalar* scalar) { + return scalar->data.si32; + } + + int64_t getScalar(int64_t, gdf_scalar* scalar) { + return scalar->data.si64; + } + + uint8_t getScalar(uint8_t, gdf_scalar* scalar) { + return scalar->data.ui08; + } + + uint16_t getScalar(uint16_t, gdf_scalar* scalar) { + return scalar->data.ui16; + } + + uint32_t getScalar(uint32_t, gdf_scalar* scalar) { + return scalar->data.ui32; + } + + uint64_t getScalar(uint64_t, gdf_scalar* scalar) { + return scalar->data.ui64; + } + + float getScalar(float, gdf_scalar* scalar) { + return scalar->data.fp32; + } + + double getScalar(double, gdf_scalar* scalar) { + return scalar->data.fp64; + } + +} // namespace library +} // namespace gdf diff --git a/src/tests/binary-operation/util/types.h b/src/tests/binary-operation/util/types.h new file mode 100644 index 00000000..6faa362b --- /dev/null +++ b/src/tests/binary-operation/util/types.h @@ -0,0 +1,199 @@ +/* + * 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_TYPES_H +#define GDF_TESTS_BINARY_OPERATION_UTIL_TYPES_H + +#include "gdf/gdf.h" + +namespace gdf { +namespace library { + + namespace helper { + void setScalar(gdf_scalar& scalar, int8_t value); + + void setScalar(gdf_scalar& scalar, int16_t value); + + void setScalar(gdf_scalar& scalar, int32_t value); + + void setScalar(gdf_scalar& scalar, int64_t value); + + void setScalar(gdf_scalar& scalar, uint8_t value); + + void setScalar(gdf_scalar& scalar, uint16_t value); + + void setScalar(gdf_scalar& scalar, uint32_t value); + + void setScalar(gdf_scalar& scalar, uint64_t value); + + void setScalar(gdf_scalar& scalar, float value); + + void setScalar(gdf_scalar& scalar, double value); + } + + namespace helper { + template + struct GdfEnumType; + + template <> + struct GdfEnumType { + using Type = void; + }; + + template <> + struct GdfEnumType { + using Type = int8_t; + }; + + template <> + struct GdfEnumType { + using Type = int16_t; + }; + + template <> + struct GdfEnumType { + using Type = int32_t; + }; + + template <> + struct GdfEnumType { + using Type = int64_t; + }; + + template <> + struct GdfEnumType { + using Type = uint8_t; + }; + + template <> + struct GdfEnumType { + using Type = uint16_t; + }; + + template <> + struct GdfEnumType { + using Type = uint32_t; + }; + + template <> + struct GdfEnumType { + using Type = uint64_t; + }; + + template <> + struct GdfEnumType { + using Type = float; + }; + + template <> + struct GdfEnumType { + using Type = double; + }; + } + + namespace helper { + template + struct GdfDataType; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_INT8; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_INT16; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_INT32; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_INT64; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_UINT8; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_UINT16; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_UINT32; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_UINT64; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_FLOAT32; + }; + + template <> + struct GdfDataType { + static constexpr gdf_dtype Value = GDF_FLOAT64; + }; + } + + template + using GdfEnumType = typename helper::GdfEnumType::Type; + + template + using GdfDataType = helper::GdfDataType; + + template + void setScalar(gdf_scalar& scalar, Type value) { + helper::setScalar(scalar, value); + } + + const char* getTypeName(gdf_dtype type); + + int8_t getScalar(int8_t, gdf_scalar* scalar); + + int16_t getScalar(int16_t, gdf_scalar* scalar); + + int32_t getScalar(int32_t, gdf_scalar* scalar); + + int64_t getScalar(int64_t, gdf_scalar* scalar); + + uint8_t getScalar(uint8_t, gdf_scalar* scalar); + + uint16_t getScalar(uint16_t, gdf_scalar* scalar); + + uint32_t getScalar(uint32_t, gdf_scalar* scalar); + + uint64_t getScalar(uint64_t, gdf_scalar* scalar); + + float getScalar(float, gdf_scalar* scalar); + + double getScalar(double, gdf_scalar* scalar); + +} // namespace library +} // namespace gdf + +#endif diff --git a/src/tests/binary-operation/util/vector.h b/src/tests/binary-operation/util/vector.h new file mode 100644 index 00000000..6214d88b --- /dev/null +++ b/src/tests/binary-operation/util/vector.h @@ -0,0 +1,185 @@ +/* + * 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_VECTOR_H +#define GDF_TESTS_BINARY_OPERATION_UTIL_VECTOR_H + +#include "gdf/gdf.h" +#include "tests/binary-operation/util/types.h" +#include "tests/binary-operation/util/field.h" + +namespace gdf { +namespace library { + + template + class Vector { + public: + using ValidType = int32_t; + static constexpr int ValidSize = 32; + + private: + template + class InnerWrapper { + public: + InnerWrapper(Field& container) + : mField (container) + { } + + T operator[](int index) { + return mField[index]; + } + + private: + Field& mField; + }; + + public: + ~Vector() { + eraseGpu(); + } + + Vector& clearGpu() { + eraseGpu(); + return *this; + } + + Vector& rangeData(Type init, Type final, Type step) { + assert((Type)0 < step); + assert(init < final); + + int size = (final - init) / step; + mData.resize(size); + for (int k = 0; k < size; ++k) { + mData[k] = init; + init += step; + } + mData.write(); + updateData(); + return *this; + } + + Vector& fillData(int size, Type value) { + mData.resize(size); + std::fill(mData.begin(), mData.end(), value); + mData.write(); + updateData(); + return *this; + } + + Vector& rangeValid(bool value) { + int size = (mData.size() / ValidSize) + ((mData.size() % ValidSize) ? 1 : 0); + mValid.resize(size); + + std::generate(mValid.begin(), mValid.end(), [value] { return -(ValidType)value; }); + clearPaddingBits(); + + mValid.write(); + updateValid(); + return *this; + } + + Vector& rangeValid(bool value, int init, int step) { + int final = mData.size(); + int size = (final / ValidSize) + ((final % ValidSize) ? 1 : 0); + mValid.resize(size); + + for (int index = 0; index < size; ++index) { + ValidType val = 0; + while (((init / ValidSize) == index) && (init < final)) { + val |= (1 << (init % ValidSize)); + init += step; + } + if (value) { + mValid[index] = val; + } else { + mValid[index] = ~val; + } + } + clearPaddingBits(); + + mValid.write(); + updateValid(); + return *this; + } + + void emplaceVector(int size) { + int validSize = (size / ValidSize) + ((size % ValidSize) ? 1 : 0); + mData.resize(size); + mValid.resize(validSize); + updateData(); + updateValid(); + } + + void readVector() { + mData.read(); + mValid.read(); + } + + public: + int dataSize() { + return mData.size(); + } + + int validSize() { + return mValid.size(); + } + + gdf_column* column() { + return &mColumn; + } + + public: + InnerWrapper data{mData}; + + InnerWrapper valid{mValid}; + + private: + void eraseGpu() { + mData.clear(); + mValid.clear(); + } + + void updateData() { + mColumn.size = mData.size(); + mColumn.dtype = GdfDataType::Value; + mColumn.data = (void*)mData.getGpuData(); + } + + void updateValid() { + mColumn.valid = (gdf_valid_type*)mValid.getGpuData(); + } + + void clearPaddingBits() { + int padding = mData.size() % ValidSize; + if (padding) { + padding = (1 << padding) - 1; + mValid.back() &= padding; + } + } + + private: + gdf_column mColumn; + + private: + Field mData; + Field mValid; + }; + +} // namespace library +} // namespace gdf + +#endif diff --git a/src/tests/filterops_numeric/CMakeLists.txt b/src/tests/filterops_numeric/CMakeLists.txt index f535393e..42eaa3ed 100644 --- a/src/tests/filterops_numeric/CMakeLists.txt +++ b/src/tests/filterops_numeric/CMakeLists.txt @@ -1,10 +1,10 @@ set(filterops_test_SRCS helper/utils.cuh helper/utils.cu - + test_example.cu test_filterops.cu - test_gpu_concat.cu + # test_gpu_concat.cu ) configure_test(filterops_test "${filterops_test_SRCS}") diff --git a/src/tests/filterops_numeric/helper/utils.cu b/src/tests/filterops_numeric/helper/utils.cu index d16db563..b61064c9 100644 --- a/src/tests/filterops_numeric/helper/utils.cu +++ b/src/tests/filterops_numeric/helper/utils.cu @@ -6,29 +6,17 @@ #include #include #include "utils.cuh" +#include "../../../util/bit_util.cuh" gdf_valid_type * get_gdf_valid_from_device(gdf_column* column) { gdf_valid_type * host_valid_out; - size_t n_bytes = get_number_of_bytes_for_valid(column->size); + size_t n_bytes = gdf_get_num_chars_bitmask(column->size); host_valid_out = new gdf_valid_type[n_bytes]; cudaMemcpy(host_valid_out,column->valid, n_bytes, cudaMemcpyDeviceToHost); return host_valid_out; } -std::string gdf_valid_to_str(gdf_valid_type *valid, size_t column_size) -{ - size_t n_bytes = get_number_of_bytes_for_valid(column_size); - std::string response; - for (size_t i = 0; i < n_bytes; i++) - { - size_t length = (n_bytes != i + 1) ? GDF_VALID_BITSIZE : (column_size - GDF_VALID_BITSIZE * (n_bytes - 1)); - auto result = chartobin(valid[i], length); - response += std::string(result); - } - return response; -} - gdf_valid_type* gen_gdf_valid(size_t column_size, size_t init_value) { gdf_valid_type *valid = nullptr; @@ -38,7 +26,7 @@ gdf_valid_type* gen_gdf_valid(size_t column_size, size_t init_value) } else { - size_t n_bytes = get_number_of_bytes_for_valid (column_size); + size_t n_bytes = gdf_get_num_chars_bitmask (column_size); valid = new gdf_valid_type[n_bytes]; size_t i; for (i = 0; i < n_bytes - 1; ++i) @@ -58,7 +46,7 @@ void delete_gdf_column(gdf_column * column){ } gdf_size_type count_zero_bits(gdf_valid_type *valid, size_t column_size) -{ +{ size_t numbits = 0; auto bin = gdf_valid_to_str(valid, column_size); @@ -69,20 +57,6 @@ gdf_size_type count_zero_bits(gdf_valid_type *valid, size_t column_size) return numbits; } -std::string chartobin(gdf_valid_type c, int size/* = 8*/) -{ - std::string bin; - bin.resize(size); - bin[0] = 0; - int i; - for (i = size - 1; i >= 0; i--) - { - bin[i] = (c % 2) + '0'; - c /= 2; - } - return bin; -} - auto print_binary(gdf_valid_type n, int size) -> void { std::cout << chartobin(n) << "\t sz: " << size << "\tbinary: " << chartobin(n, size) << std::endl; } diff --git a/src/tests/filterops_numeric/helper/utils.cuh b/src/tests/filterops_numeric/helper/utils.cuh index 7348f4da..72e5ee8f 100644 --- a/src/tests/filterops_numeric/helper/utils.cuh +++ b/src/tests/filterops_numeric/helper/utils.cuh @@ -14,6 +14,11 @@ #include #include #include "gdf/gdf.h" +#include "gdf/utils.h" +#include "../../../util/bit_util.cuh" +#include "../../test_utils/gdf_test_utils.cuh" + +using namespace gdf::util; template inline gdf_dtype gdf_enum_type_for() @@ -57,23 +62,16 @@ inline gdf_dtype gdf_enum_type_for() return GDF_FLOAT64; } -inline auto get_number_of_bytes_for_valid (size_t column_size) -> size_t { - return sizeof(gdf_valid_type) * (column_size + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE; -} - auto print_binary(gdf_valid_type n, int size = 8) -> void ; -auto chartobin(gdf_valid_type n, int size = 8) -> std::string; - gdf_size_type count_zero_bits(gdf_valid_type *valid, size_t column_size); -auto delete_gdf_column(gdf_column * column) -> void; - +auto delete_gdf_column(gdf_column * column) -> void; + auto gen_gdf_valid(size_t column_size, size_t init_value) -> gdf_valid_type *; gdf_valid_type * get_gdf_valid_from_device(gdf_column* column) ; -std::string gdf_valid_to_str(gdf_valid_type *valid, size_t column_size); template auto init_device_vector(gdf_size_type num_elements) -> std::tuple> @@ -112,12 +110,12 @@ gdf_column convert_to_device_gdf_column (gdf_column *column) { char *raw_pointer; thrust::device_ptr device_pointer; std::tie(raw_pointer, device_pointer) = init_device_vector(column_size); - + void* host_out = column->data; cudaMemcpy(raw_pointer, host_out, sizeof(ValueType) * column->size, cudaMemcpyHostToDevice); gdf_valid_type *host_valid = column->valid; - size_t n_bytes = get_number_of_bytes_for_valid(column_size); + size_t n_bytes = gdf_get_num_chars_bitmask(column_size); gdf_valid_type *valid_value_pointer; cudaMalloc((void **)&valid_value_pointer, n_bytes); @@ -132,61 +130,35 @@ template gdf_column convert_to_host_gdf_column (gdf_column *column) { auto host_out = get_gdf_data_from_device(column); auto host_valid_out = get_gdf_valid_from_device(column); - + auto output = *column; output.data = host_out; output.valid = host_valid_out; return output; } - -template -auto print_column(gdf_column * column) -> void { - auto host_out = get_gdf_data_from_device(column); - auto host_valid_out = get_gdf_valid_from_device(column); - std::cout<<"Printing Column\t null_count:" << column->null_count << "\t type " << column->dtype << std::endl; - int n_bytes = sizeof(int8_t) * (column->size + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE; - // for(int i = 0; i < column->size; i++) { - // int col_position = i / 8; - // int length_col = n_bytes != col_position+1 ? GDF_VALID_BITSIZE : column->size - GDF_VALID_BITSIZE * (n_bytes - 1); - // int bit_offset = (length_col - 1) - (i % 8); - // if (sizeof(ValueType) == 1) { - // std::cout << "host_out[" << i << "] = " << ((int)host_out[i])<<" valid="<<((host_valid_out[col_position] >> bit_offset ) & 1)<(column_size); - // std::cout << "1. gen_gdb_column\n"; - + // std::cout << "1. gen_gdb_column\n"; + using thrust::detail::make_normal_iterator; thrust::fill(make_normal_iterator(device_pointer), make_normal_iterator(device_pointer + column_size), init_value); - //std::cout << "2. gen_gdb_column\n"; - + //std::cout << "2. gen_gdb_column\n"; + gdf_valid_type *host_valid = gen_gdf_valid(column_size, init_value); - size_t n_bytes = get_number_of_bytes_for_valid(column_size); + size_t n_bytes = gdf_get_num_chars_bitmask(column_size); gdf_valid_type *valid_value_pointer; cudaMalloc((void **)&valid_value_pointer, n_bytes); cudaMemcpy(valid_value_pointer, host_valid, n_bytes, cudaMemcpyHostToDevice); - // std::cout << "3. gen_gdb_column\n"; - + // std::cout << "3. gen_gdb_column\n"; + gdf_column output; auto zero_bits = output.null_count = count_zero_bits(host_valid, column_size); @@ -195,12 +167,20 @@ gdf_column gen_gdb_column(size_t column_size, ValueType init_value) column_size, gdf_enum_type_value, zero_bits); - //std::cout << "4. gen_gdb_column\n"; - + //std::cout << "4. gen_gdb_column\n"; + delete []host_valid; return output; } + + +static inline void print_column(const gdf_column * col ) +{ + print_gdf_column(col); +} + + template void check_column_for_stencil_operation(gdf_column *column, gdf_column *stencil, gdf_column *output_op) { gdf_column host_column = convert_to_host_gdf_column(column); @@ -210,7 +190,7 @@ void check_column_for_stencil_operation(gdf_column *column, gdf_column *stencil, EXPECT_EQ(host_column.size, host_stencil.size); //EXPECT_EQ(host_column.dtype == host_output_op.dtype); // it must have the same type - + int n_bytes = sizeof(int8_t) * (column->size + GDF_VALID_BITSIZE - 1) / GDF_VALID_BITSIZE; std::vector indexes; for(size_t i = 0; i < host_stencil.size; i++) { @@ -222,14 +202,14 @@ void check_column_for_stencil_operation(gdf_column *column, gdf_column *stencil, indexes.push_back(i); } } - - for(size_t i = 0; i < indexes.size(); i++) + + for(size_t i = 0; i < indexes.size(); i++) { int index = indexes[i]; LeftValueType value = ((LeftValueType *)(host_column.data))[index]; std::cout << "filtered values: " << index << "** " << "\t value: " << (int)value << std::endl; EXPECT_EQ( ((RightValueType*)host_output_op.data)[i], value); - + int col_position = i / 8; int length_col = n_bytes != col_position+1 ? GDF_VALID_BITSIZE : output_op->size - GDF_VALID_BITSIZE * (n_bytes - 1); int bit_offset = (length_col - 1) - (i % 8); @@ -245,23 +225,23 @@ void check_column_for_comparison_operation(gdf_column *lhs, gdf_column *rhs, gdf auto lhs_valid = get_gdf_valid_from_device(lhs); auto rhs_valid = get_gdf_valid_from_device(rhs); auto output_valid = get_gdf_valid_from_device(output); - - size_t n_bytes = get_number_of_bytes_for_valid(output->size); - EXPECT_EQ(lhs->size, rhs->size); - + size_t n_bytes = gdf_get_num_chars_bitmask(output->size); + + EXPECT_EQ(lhs->size, rhs->size); + for(size_t i = 0; i < output->size; i++) { size_t col_position = i / 8; size_t length_col = n_bytes != col_position+1 ? GDF_VALID_BITSIZE : output->size - GDF_VALID_BITSIZE * (n_bytes - 1); size_t bit_offset = (length_col - 1) - (i % 8); - + EXPECT_EQ( ((lhs_valid[col_position] >> bit_offset ) & 1) & ((rhs_valid[col_position] >> bit_offset ) & 1), ((output_valid[col_position] >> bit_offset ) & 1) ); } - + delete[] lhs_valid; delete[] rhs_valid; - delete[] output_valid; + delete[] output_valid; } { @@ -269,17 +249,17 @@ void check_column_for_comparison_operation(gdf_column *lhs, gdf_column *rhs, gdf auto rhs_data = get_gdf_data_from_device(rhs); auto output_data = get_gdf_data_from_device(output); - EXPECT_EQ(lhs->size, rhs->size); + EXPECT_EQ(lhs->size, rhs->size); for(size_t i = 0; i < lhs->size; i++) { - EXPECT_EQ(lhs_data[i] == rhs_data[i] ? 1 : 0, output_data[i]); + EXPECT_EQ(lhs_data[i] == rhs_data[i] ? 1 : 0, output_data[i]); } - + delete[] lhs_data; delete[] rhs_data; - delete[] output_data; + delete[] output_data; } - + } template @@ -289,16 +269,29 @@ void check_column_for_concat_operation(gdf_column *lhs, gdf_column *rhs, gdf_col auto lhs_valid = get_gdf_valid_from_device(lhs); auto rhs_valid = get_gdf_valid_from_device(rhs); auto output_valid = get_gdf_valid_from_device(output); - - auto computed = gdf_valid_to_str(output_valid, output->size); + + + // auto gdf_valid_to_str = [](gdf_valid_type* valid, size_t column_size) + // { + // size_t last_byte = gdf::util::last_byte_index(column_size); + // std::string response; + // for (size_t i = 0; i < last_byte; i++) { + // size_t n_bits = last_byte != i + 1 ? 8 : column_size - 8 * (last_byte - 1); + // auto result = chartobin(valid[i], n_bits); + // response += std::string(result); + // } + // return response; + // } auto expected = gdf_valid_to_str(lhs_valid, lhs->size) + gdf_valid_to_str(rhs_valid, rhs->size); + auto computed = gdf_valid_to_str(output_valid, output->size); + //std::cout << "computed: " << computed << std::endl; //std::cout << "expected: " << expected << std::endl; delete[] lhs_valid; delete[] rhs_valid; - delete[] output_valid; + delete[] output_valid; EXPECT_EQ(computed, expected); } @@ -311,10 +304,10 @@ void check_column_for_concat_operation(gdf_column *lhs, gdf_column *rhs, gdf_col auto expected = gdf_data_to_str(lhs_data, lhs->size) + gdf_data_to_str(rhs_data, rhs->size); delete[] lhs_data; delete[] rhs_data; - delete[] output_data; - EXPECT_EQ(computed, expected); + delete[] output_data; + EXPECT_EQ(computed, expected); } - + } diff --git a/src/tests/filterops_numeric/test_filterops.cu b/src/tests/filterops_numeric/test_filterops.cu index 9a1c1260..fef0e43e 100644 --- a/src/tests/filterops_numeric/test_filterops.cu +++ b/src/tests/filterops_numeric/test_filterops.cu @@ -43,7 +43,7 @@ TEST(FilterOperationsTest, usage_example) { gdf_comparison_operator gdf_operator = GDF_EQUALS; gdf_column lhs = gen_gdb_column(column_size, init_value); // 4, 2, 0 - + gdf_column rhs = gen_gdb_column(column_size, 0.01 + max_size - init_value); // 0, 2, 4 gdf_column output = gen_gdb_column(column_size, 0); @@ -52,13 +52,13 @@ TEST(FilterOperationsTest, usage_example) { EXPECT_TRUE(error == GDF_SUCCESS); std::cout << "Left" << std::endl; - print_column(&lhs); + print_column(&lhs); std::cout << "Right" << std::endl; - print_column(&rhs); + print_column(&rhs); std::cout << "Output" << std::endl; - print_column(&output); + print_column(&output); check_column_for_comparison_operation(&lhs, &rhs, &output, gdf_operator); @@ -73,9 +73,11 @@ TEST(FilterOperationsTest, usage_example) { } -template +template void test_filterops_using_templates(gdf_comparison_operator gdf_operator = GDF_EQUALS) { + + typedef LeftValueType RightValueType; //0, ..., 100, //100, 10000, 10000, 100000 for (int column_size = 0; column_size < 10; column_size += 1) @@ -110,66 +112,34 @@ void test_filterops_using_templates(gdf_comparison_operator gdf_operator = GDF_E TEST(FilterOperationsTest, WithInt8AndOthers) { - test_filterops_using_templates(); - test_filterops_using_templates(); - - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); + test_filterops_using_templates(); } TEST(FilterOperationsTest, WithInt16AndOthers) { - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - + test_filterops_using_templates(); + } TEST(FilterOperationsTest, WithInt32AndOthers) { - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - + test_filterops_using_templates(); + } TEST(FilterOperationsTest, WithInt64AndOthers) { - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - + test_filterops_using_templates(); } TEST(FilterOperationsTest, WithFloat32AndOthers) { - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - + test_filterops_using_templates(); + } TEST(FilterOperationsTest, WithFloat64AndOthers) { - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - test_filterops_using_templates(); - + test_filterops_using_templates(); + } diff --git a/src/tests/filterops_numeric/test_gpu_concat.cu b/src/tests/filterops_numeric/test_gpu_concat.cu index e50aacb4..e774d47d 100644 --- a/src/tests/filterops_numeric/test_gpu_concat.cu +++ b/src/tests/filterops_numeric/test_gpu_concat.cu @@ -37,7 +37,7 @@ using ValueType = int16_t; */ TEST(GdfConcat, usage_example) { const size_t lhs_size = 10; - const size_t rhs_size = 20; + const size_t rhs_size = 5; gdf_column lhs = gen_gdb_column(lhs_size, 2); gdf_column rhs = gen_gdb_column(rhs_size, 3); std::cout << "*****left**************\n"; @@ -45,7 +45,7 @@ TEST(GdfConcat, usage_example) { std::cout << "*****right**************\n"; print_column(&rhs); std::cout << "*******************\n"; - + // reserve space for gdf_column output gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); @@ -53,10 +53,10 @@ TEST(GdfConcat, usage_example) { gpu_concat(&lhs, &rhs, &output); std::cout << "*****output**************\n"; print_column(&output); - + // check results check_column_for_concat_operation(&lhs, &rhs, &output); - + delete_gdf_column(&lhs); delete_gdf_column(&rhs); delete_gdf_column(&output); @@ -150,188 +150,188 @@ TEST(GdfConcat, CaseWithOutputOfTwoBytes) delete_gdf_column(&output); } -TEST(GdfConcat, CaseWithInput_2_2_Output3) -{ - // 8, 3 + 8, 1 - // 8, 3|5, 3|1 - - const size_t lhs_size = 8 + 3; - const size_t rhs_size = 8 + 1; - - gdf_column lhs = gen_gdb_column(lhs_size, 2); - gdf_column rhs = gen_gdb_column(rhs_size, 3); - gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); - - std::cout << "*****left**************\n"; - print_column(&lhs); - std::cout << "*****right**************\n"; - print_column(&rhs); - std::cout << "*******************\n"; - - gpu_concat(&lhs, &rhs, &output); - check_column_for_concat_operation(&lhs, &rhs, &output); - delete_gdf_column(&lhs); - delete_gdf_column(&rhs); - delete_gdf_column(&output); -} - -TEST(GdfConcat, CaseWithInput_2_5_Output5) -{ - // 8, 2 + 8, 8, 8, 8, 5 - // 8, 2|6, 2|6, 2|6, 2|5 - - const size_t lhs_size = 8 + 2; - const size_t rhs_size = 8 + 8 + 8 + 8 + 5; - - gdf_column lhs = gen_gdb_column(lhs_size, 2); - gdf_column rhs = gen_gdb_column(rhs_size, 3); - gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); - - gpu_concat(&lhs, &rhs, &output); - check_column_for_concat_operation(&lhs, &rhs, &output); - delete_gdf_column(&lhs); - delete_gdf_column(&rhs); - delete_gdf_column(&output); -} - -TEST(GdfConcat, CaseWithInput_1_4_Output5) -{ - // 3 + 8, 8, 8, 7 // caso especial - // 3|5, 3|5, 3|5, 3|5, 2 - - // 100 - // 10101111 10101111 10101111 10000 00 - // 100 10101 - // - // 11110101 - // 11110101 - // 11110000 - // 00 - - const size_t lhs_size = 3; - const size_t rhs_size = 8 + 8 + 8 + 7; - - gdf_column lhs = gen_gdb_column(lhs_size, 2); - gdf_column rhs = gen_gdb_column(rhs_size, 3); - gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); - std::cout << "*****left**************\n"; - print_column(&lhs); - std::cout << "*****right**************\n"; - print_column(&rhs); - std::cout << "*******************\n"; - - gpu_concat(&lhs, &rhs, &output); - - - std::cout << "*****output**************\n"; - print_column(&output); - std::cout << "*******************\n"; - - check_column_for_concat_operation(&lhs, &rhs, &output); - - delete_gdf_column(&lhs); - delete_gdf_column(&rhs); - delete_gdf_column(&output); -} - - -TEST(GdfConcat, CaseWithInput_0_9_Output2) -{ - // 3 + 8, 8, 8, 7 // caso especial - // 3|5, 3|5, 3|5, 3|5, 2 - - // 100 - // 10101111 10101111 10101111 10000 00 - // 100 10101 - // - // 11110101 - // 11110101 - // 11110000 - // 00 - - const size_t lhs_size = 0; - const size_t rhs_size = 8 + 1; - - gdf_column lhs = gen_gdb_column(lhs_size, 2); - gdf_column rhs = gen_gdb_column(rhs_size, 3); - gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); - std::cout << "*****left**************\n"; - print_column(&lhs); - std::cout << "*****right**************\n"; - print_column(&rhs); - std::cout << "*******************\n"; - - gpu_concat(&lhs, &rhs, &output); - check_column_for_concat_operation(&lhs, &rhs, &output); - - delete_gdf_column(&lhs); - delete_gdf_column(&rhs); - delete_gdf_column(&output); -} - -TEST(GdfConcat, CaseWithInput_5_11_Output22) -{ - // 3 + 8, 8, 8, 7 // caso especial - // 3|5, 3|5, 3|5, 3|5, 2 - - // 100 - // 10101111 10101111 10101111 10000 00 - // 100 10101 - // - // 11110101 - // 11110101 - // 11110000 - // 00 - - const size_t lhs_size = 5; - const size_t rhs_size = 8 + 5; - - gdf_column lhs = gen_gdb_column(lhs_size, 2); - gdf_column rhs = gen_gdb_column(rhs_size, 3); - gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); - std::cout << "*****left**************\n"; - print_column(&lhs); - std::cout << "*****right**************\n"; - print_column(&rhs); - std::cout << "*******************\n"; - - gpu_concat(&lhs, &rhs, &output); - std::cout << "*****output**************\n"; - print_column(&output); - std::cout << "*******************\n"; - - check_column_for_concat_operation(&lhs, &rhs, &output); - - delete_gdf_column(&lhs); - delete_gdf_column(&rhs); - delete_gdf_column(&output); -} - - -TEST(GpuConcatTest, WithDifferentColumnSizes) -{ - using ValueType = int16_t; - //0, ..., 100, - //100, 10000, 10000, 100000 - for (int lhs_size = 0; lhs_size < 100; lhs_size += 1) - { - for (int rhs_size = 0; rhs_size < 100; rhs_size += 1) - { - gdf_column lhs = gen_gdb_column(lhs_size, 2); - - gdf_column rhs = gen_gdb_column(rhs_size, 3); - - gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); - gdf_error error = gpu_concat(&lhs, &rhs, &output); - - // std::cout << "Output" << std::endl; - // print_column(&output); - - check_column_for_concat_operation(&lhs, &rhs, &output); - - delete_gdf_column(&lhs); - delete_gdf_column(&rhs); - delete_gdf_column(&output); - } - } -} \ No newline at end of file +// TEST(GdfConcat, CaseWithInput_2_2_Output3) +// { +// // 8, 3 + 8, 1 +// // 8, 3|5, 3|1 + +// const size_t lhs_size = 8 + 3; +// const size_t rhs_size = 8 + 1; + +// gdf_column lhs = gen_gdb_column(lhs_size, 2); +// gdf_column rhs = gen_gdb_column(rhs_size, 3); +// gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); + +// std::cout << "*****left**************\n"; +// print_column(&lhs); +// std::cout << "*****right**************\n"; +// print_column(&rhs); +// std::cout << "*******************\n"; + +// gpu_concat(&lhs, &rhs, &output); +// check_column_for_concat_operation(&lhs, &rhs, &output); +// delete_gdf_column(&lhs); +// delete_gdf_column(&rhs); +// delete_gdf_column(&output); +// } + +// TEST(GdfConcat, CaseWithInput_2_5_Output5) +// { +// // 8, 2 + 8, 8, 8, 8, 5 +// // 8, 2|6, 2|6, 2|6, 2|5 + +// const size_t lhs_size = 8 + 2; +// const size_t rhs_size = 8 + 8 + 8 + 8 + 5; + +// gdf_column lhs = gen_gdb_column(lhs_size, 2); +// gdf_column rhs = gen_gdb_column(rhs_size, 3); +// gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); + +// gpu_concat(&lhs, &rhs, &output); +// check_column_for_concat_operation(&lhs, &rhs, &output); +// delete_gdf_column(&lhs); +// delete_gdf_column(&rhs); +// delete_gdf_column(&output); +// } + +// TEST(GdfConcat, CaseWithInput_1_4_Output5) +// { +// // 3 + 8, 8, 8, 7 // caso especial +// // 3|5, 3|5, 3|5, 3|5, 2 + +// // 100 +// // 10101111 10101111 10101111 10000 00 +// // 100 10101 +// // +// // 11110101 +// // 11110101 +// // 11110000 +// // 00 + +// const size_t lhs_size = 3; +// const size_t rhs_size = 8 + 8 + 8 + 7; + +// gdf_column lhs = gen_gdb_column(lhs_size, 2); +// gdf_column rhs = gen_gdb_column(rhs_size, 3); +// gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); +// std::cout << "*****left**************\n"; +// print_column(&lhs); +// std::cout << "*****right**************\n"; +// print_column(&rhs); +// std::cout << "*******************\n"; + +// gpu_concat(&lhs, &rhs, &output); + + +// std::cout << "*****output**************\n"; +// print_column(&output); +// std::cout << "*******************\n"; + +// check_column_for_concat_operation(&lhs, &rhs, &output); + +// delete_gdf_column(&lhs); +// delete_gdf_column(&rhs); +// delete_gdf_column(&output); +// } + + +// TEST(GdfConcat, CaseWithInput_0_9_Output2) +// { +// // 3 + 8, 8, 8, 7 // caso especial +// // 3|5, 3|5, 3|5, 3|5, 2 + +// // 100 +// // 10101111 10101111 10101111 10000 00 +// // 100 10101 +// // +// // 11110101 +// // 11110101 +// // 11110000 +// // 00 + +// const size_t lhs_size = 0; +// const size_t rhs_size = 8 + 1; + +// gdf_column lhs = gen_gdb_column(lhs_size, 2); +// gdf_column rhs = gen_gdb_column(rhs_size, 3); +// gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); +// std::cout << "*****left**************\n"; +// print_column(&lhs); +// std::cout << "*****right**************\n"; +// print_column(&rhs); +// std::cout << "*******************\n"; + +// gpu_concat(&lhs, &rhs, &output); +// check_column_for_concat_operation(&lhs, &rhs, &output); + +// delete_gdf_column(&lhs); +// delete_gdf_column(&rhs); +// delete_gdf_column(&output); +// } + +// TEST(GdfConcat, CaseWithInput_5_11_Output22) +// { +// // 3 + 8, 8, 8, 7 // caso especial +// // 3|5, 3|5, 3|5, 3|5, 2 + +// // 100 +// // 10101111 10101111 10101111 10000 00 +// // 100 10101 +// // +// // 11110101 +// // 11110101 +// // 11110000 +// // 00 + +// const size_t lhs_size = 5; +// const size_t rhs_size = 8 + 5; + +// gdf_column lhs = gen_gdb_column(lhs_size, 2); +// gdf_column rhs = gen_gdb_column(rhs_size, 3); +// gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); +// std::cout << "*****left**************\n"; +// print_column(&lhs); +// std::cout << "*****right**************\n"; +// print_column(&rhs); +// std::cout << "*******************\n"; + +// gpu_concat(&lhs, &rhs, &output); +// std::cout << "*****output**************\n"; +// print_column(&output); +// std::cout << "*******************\n"; + +// check_column_for_concat_operation(&lhs, &rhs, &output); + +// delete_gdf_column(&lhs); +// delete_gdf_column(&rhs); +// delete_gdf_column(&output); +// } + + +// TEST(GpuConcatTest, WithDifferentColumnSizes) +// { +// using ValueType = int16_t; +// //0, ..., 100, +// //100, 10000, 10000, 100000 +// for (int lhs_size = 0; lhs_size < 100; lhs_size += 1) +// { +// for (int rhs_size = 0; rhs_size < 100; rhs_size += 1) +// { +// gdf_column lhs = gen_gdb_column(lhs_size, 2); + +// gdf_column rhs = gen_gdb_column(rhs_size, 3); + +// gdf_column output = gen_gdb_column(lhs_size + rhs_size, 0); +// gdf_error error = gpu_concat(&lhs, &rhs, &output); + +// // std::cout << "Output" << std::endl; +// // print_column(&output); + +// check_column_for_concat_operation(&lhs, &rhs, &output); + +// delete_gdf_column(&lhs); +// delete_gdf_column(&rhs); +// delete_gdf_column(&output); +// } +// } +// } \ No newline at end of file diff --git a/src/tests/test_utils/gdf_test_utils.cuh b/src/tests/test_utils/gdf_test_utils.cuh index e038fd81..b3da8a5f 100644 --- a/src/tests/test_utils/gdf_test_utils.cuh +++ b/src/tests/test_utils/gdf_test_utils.cuh @@ -23,15 +23,35 @@ #include #include #include "../../util/bit_util.cuh" +using namespace gdf::util; // Type for a unique_ptr to a gdf_column with a custom deleter // Custom deleter is defined at construction -using gdf_col_pointer = typename std::unique_ptr>; + + +// Prints a vector and valids +template +void print_vector_and_valid(T * v, + gdf_valid_type * valid, + const size_t num_rows) +{ + auto functor = [&valid, &v](int index) -> std::string { + if (gdf_is_valid(valid, index)) + return std::to_string((int)v[index]); + return std::string("@"); + }; + std::vector indexes(num_rows); + std::iota(std::begin(indexes), std::end(indexes), 0); + std::transform(indexes.begin(), indexes.end(), std::ostream_iterator(std::cout, ", "), functor); + std::cout << std::endl; +} + template -void print_typed_column(col_type * col_data, - gdf_valid_type * validity_mask, +void print_typed_column(col_type * col_data, + gdf_valid_type * validity_mask, const size_t num_rows) { @@ -40,33 +60,36 @@ void print_typed_column(col_type * col_data, const size_t num_masks = gdf_get_num_chars_bitmask(num_rows); - std::vector h_mask(num_masks); - if(nullptr != validity_mask) - { - cudaMemcpy(h_mask.data(), validity_mask, num_masks * sizeof(gdf_valid_type), cudaMemcpyDeviceToHost); - } + std::cout << "column :\n"; + std::vector h_mask(num_masks); + cudaMemcpy(h_mask.data(), validity_mask, num_masks * sizeof(gdf_valid_type), cudaMemcpyDeviceToHost); + print_vector_and_valid(h_data.data(), h_mask.data(), num_rows); - for(size_t i = 0; i < num_rows; ++i) - { - // If the element is valid, print it's value - if(true == gdf_is_valid(h_mask.data(), i)) - { - std::cout << h_data[i] << " "; - } - // Otherwise, print an @ to represent a null value - else + std::cout << "\n"; + if (validity_mask != nullptr) { + auto gdf_valid_to_str_lambda = [](gdf_valid_type *valid, size_t column_size) { - std::cout << "@" << " "; - } + size_t n_bytes = gdf_get_num_chars_bitmask(column_size); + std::string response; + for (size_t i = 0; i < n_bytes; i++) + { + size_t length = n_bytes != i + 1 ? GDF_VALID_BITSIZE : column_size - GDF_VALID_BITSIZE * (n_bytes - 1); + auto result = chartobin(valid[i], length); + response += std::string(result) + '|'; + } + return response; + }; + auto res = gdf_valid_to_str_lambda(h_mask.data(), num_rows); + std::cout << "mask : " << std::endl; + std::cout << res << std::endl; } std::cout << std::endl; } -void print_gdf_column(gdf_column const * the_column) +static inline void print_gdf_column(gdf_column const * the_column) { const size_t num_rows = the_column->size; - const gdf_dtype gdf_col_type = the_column->dtype; switch(gdf_col_type) { @@ -148,8 +171,8 @@ gdf_col_pointer create_gdf_column(std::vector const & host_vector, // Create a new instance of a gdf_column with a custom deleter that will free // the associated device memory when it eventually goes out of scope auto deleter = [](gdf_column* col){ - col->size = 0; - if(nullptr != col->data){cudaFree(col->data);} + col->size = 0; + if(nullptr != col->data){cudaFree(col->data);} if(nullptr != col->valid){cudaFree(col->valid);} }; gdf_col_pointer the_column{new gdf_column, deleter}; @@ -159,7 +182,7 @@ gdf_col_pointer create_gdf_column(std::vector const & host_vector, cudaMemcpy(the_column->data, host_vector.data(), host_vector.size() * sizeof(col_type), cudaMemcpyHostToDevice); - // If a validity bitmask vector was passed in, allocate device storage + // If a validity bitmask vector was passed in, allocate device storage // and copy its contents from the host vector if(valid_vector.size() > 0) { @@ -185,7 +208,7 @@ gdf_col_pointer create_gdf_column(std::vector const & host_vector, // a gdf_column and append it to a vector of gdf_columns template inline typename std::enable_if::type -convert_tuple_to_gdf_columns(std::vector &gdf_columns,std::tuple...>& t, +convert_tuple_to_gdf_columns(std::vector &gdf_columns,std::tuple...>& t, valid_initializer_t bit_initializer) { //bottom of compile-time recursion @@ -221,7 +244,7 @@ convert_tuple_to_gdf_columns(std::vector &gdf_columns,std::tupl // Converts a tuple of host vectors into a vector of gdf_columns template -std::vector initialize_gdf_columns(std::tuple...> & host_columns, +std::vector initialize_gdf_columns(std::tuple...> & host_columns, valid_initializer_t bit_initializer) { std::vector gdf_columns; @@ -230,12 +253,12 @@ std::vector initialize_gdf_columns(std::tuple.. } -// Overload for default initialization of validity bitmasks which +// Overload for default initialization of validity bitmasks which // sets every element to valid template std::vector initialize_gdf_columns(std::tuple...> & host_columns ) { - return initialize_gdf_columns(host_columns, + return initialize_gdf_columns(host_columns, [](const size_t row, const size_t col){return true;}); } diff --git a/thirdparty/jitify b/thirdparty/jitify new file mode 160000 index 00000000..e25d6f3b --- /dev/null +++ b/thirdparty/jitify @@ -0,0 +1 @@ +Subproject commit e25d6f3b77f61aef4f6e4fe7c3822ab600c02b9e diff --git a/travisci/install-cuda-trusty.sh b/travisci/install-cuda-trusty.sh index 1994b07e..ee1c79ea 100755 --- a/travisci/install-cuda-trusty.sh +++ b/travisci/install-cuda-trusty.sh @@ -45,5 +45,15 @@ else exit 1 fi export CUDA_HOME=/usr/local/cuda-${CUDA:0:3} -export LD_LIBRARY_PATH=${CUDA_HOME}/lib64:${LD_LIBRARY_PATH} export PATH=${CUDA_HOME}/bin:${PATH} + +LIBRARY_PATH_VAR=${CUDA_HOME}/lib64:${CUDA_HOME}/lib64/stubs +if [ -z ${LD_LIBRARY_PATH} ]; then + export LD_LIBRARY_PATH=${LIBRARY_PATH_VAR} +else + export LD_LIBRARY_PATH=${LIBRARY_PATH_VAR}:${LD_LIBRARY_PATH} +fi + +if [ -f ${CUDA_HOME}/lib64/stubs/libcuda.so ] && [ ! -f ${CUDA_HOME}/lib64/stubs/libcuda.so.1 ]; then + sudo ln -s ${CUDA_HOME}/lib64/stubs/libcuda.so ${CUDA_HOME}/lib64/stubs/libcuda.so.1 +fi