Merge remote-tracking branch 'origin' into type_date

Author: Yifei-yang7

CMakeLists.txt

@@ -59,6 +59,7 @@ set(EXTENSION_SOURCES
   src/gpu_buffer_manager.cpp
   src/gpu_columns.cpp
   src/gpu_expression_executor.cpp
+  src/gpu_query_result.cpp
 )
 add_subdirectory(src/operator)
 add_subdirectory(src/plan)

docs/README.md

@@ -43,12 +43,12 @@ conda install -c rapidsai -c conda-forge -c nvidia rapidsai::libcudf
 Set the environment variables `USE_CUDF` to 1 and `LIBCUDF_ENV_PREFIX` to the conda environment's path. For example, if we installed miniconda in `~/miniconda3` and installed libcudf in the conda environment `libcudf-env`, then we would set the `LIBCUDF_ENV_PREFIX` to `~/miniconda3/envs/libcudf-env`.
 ```
 export USE_CUDF=1
-export LIBCUDF_ENV_PREFIX = {PATH to libcudf-env}
+export LIBCUDF_ENV_PREFIX={PATH to libcudf-env}
 ```
 
 ### Clone the Sirius repository
 ```
-git clone --recurse-submodules https://github.com/bwyogatama/sirius.git
+git clone --recurse-submodules https://github.com/sirius-db/sirius.git
 cd sirius
 export SIRIUS_HOME_PATH=`pwd`
 cd duckdb
@@ -146,4 +146,4 @@ Sirius is still under major development and we are working on adding more featur
 
 Sirius always welcomes new contributors! If you are interested, check our [website](https://www.sirius-db.com/), reach out to our [email](siriusdb@cs.wisc.edu), or join our [slack channel](https://join.slack.com/t/sirius-db/shared_invite/zt-33tuwt1sk-aa2dk0EU_dNjklSjIGW3vg).
 
-**Let's kickstart the GPU eras for Data Analytics!**
+**Let's kickstart the GPU eras for Data Analytics!**

queries/devesh_test_queries/q10_group_by_test.sql

@@ -47,7 +47,7 @@ void cudf_aggregate(vector<shared_ptr<GPUColumn>>& column, uint64_t num_aggregat
 
     uint64_t size = 0;
     for (int agg = 0; agg < num_aggregates; agg++) {
-        if (column[agg]->data_wrapper.data != nullptr) {
+        if (column[agg]->data_wrapper.data != nullptr || (column[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::COUNT_STAR && column[agg]->column_length > 0)) {
             size = column[agg]->column_length;
             break;
         }

queries/devesh_test_queries/string_group_and_order_by.txt

@@ -60,29 +60,22 @@ void cudf_groupby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
   for (int agg = 0; agg < num_aggregates; agg++) {
     requests.emplace_back(cudf::groupby::aggregation_request());
     if (aggregate_keys[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::COUNT && aggregate_keys[agg]->column_length == 0) {
-      SIRIUS_LOG_DEBUG("Count aggregation");
       auto aggregate = cudf::make_sum_aggregation<cudf::groupby_aggregation>();
       requests[agg].aggregations.push_back(std::move(aggregate));
-      // auto const_scalar = cudf::make_fixed_width_scalar<uint64_t>(static_cast<uint64_t>(0));
-      // requests[agg].values = cudf::make_column_from_scalar(*const_scalar, size)->view();
       uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(size, 0, 0);
       cudaMemset(temp, 0, size * sizeof(uint64_t));
       shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
       requests[agg].values = temp_column->convertToCudfColumn();
     } else if (aggregate_keys[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::SUM && aggregate_keys[agg]->column_length == 0) {
       auto aggregate = cudf::make_sum_aggregation<cudf::groupby_aggregation>();
       requests[agg].aggregations.push_back(std::move(aggregate));
-      // auto const_scalar = cudf::make_fixed_width_scalar<uint64_t>(static_cast<uint64_t>(0));
-      // requests[agg].values = cudf::make_column_from_scalar(*const_scalar, size)->view();
       uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(size, 0, 0);
       cudaMemset(temp, 0, size * sizeof(uint64_t));
       shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
       requests[agg].values = temp_column->convertToCudfColumn();
     } else if (aggregate_keys[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::COUNT_STAR && aggregate_keys[agg]->column_length != 0) {
       auto aggregate = cudf::make_count_aggregation<cudf::groupby_aggregation>(cudf::null_policy::EXCLUDE);
       requests[agg].aggregations.push_back(std::move(aggregate));
-      // auto const_scalar = cudf::make_fixed_width_scalar<uint64_t>(static_cast<uint64_t>(1));
-      // requests[agg].values = cudf::make_column_from_scalar(*const_scalar, size)->view();
       uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(size, 0, 0);
       cudaMemset(temp, 0, size * sizeof(uint64_t));
       shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));

src/cuda/cudf/cudf_aggregate.cu

@@ -199,8 +199,8 @@ __global__ void probe_multikey<BLOCK_THREADS, ITEMS_PER_THREAD>(uint64_t **keys,
 
 void buildHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uint64_t N, int* condition_mode, int num_keys, bool is_right) {
     CHECK_ERROR();
-    if (N == 0) {
-        SIRIUS_LOG_DEBUG("Input size is 0");
+    if (N == 0 || ht_len == 0) {
+        SIRIUS_LOG_DEBUG("Input size is 0 or hash table is empty");
         return;
     }
     SIRIUS_LOG_DEBUG("Launching Build Kernel");
@@ -239,11 +239,11 @@ void buildHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uin
 void probeHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uint64_t* &row_ids_left, uint64_t* &row_ids_right, uint64_t* &count, uint64_t N, int* condition_mode, int num_keys, bool is_right) {
     CHECK_ERROR();
     GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
-    if (N == 0) {
+    if (N == 0 || ht_len == 0) {
         uint64_t* h_count = gpuBufferManager->customCudaHostAlloc<uint64_t>(1);
         h_count[0] = 0;
         count = h_count;
-        SIRIUS_LOG_DEBUG("Input size is 0");
+        SIRIUS_LOG_DEBUG("Input size is 0 or hash table is empty");
         return;
     }
     SIRIUS_LOG_DEBUG("Launching Probe Kernel");

src/cuda/cudf/cudf_groupby.cu

@@ -181,8 +181,8 @@ void scanHashTableRight(unsigned long long* ht, uint64_t ht_len, uint64_t* &row_
 
 void probeHashTableRightSemiAnti(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uint64_t N, int* condition_mode, int num_keys) {
     CHECK_ERROR();
-    if (N == 0) {
-        SIRIUS_LOG_DEBUG("Input size is 0");
+    if (N == 0 || ht_len == 0) {
+        SIRIUS_LOG_DEBUG("Input size is 0 or hash table is empty");
         return;
     }
     SIRIUS_LOG_DEBUG("Launching Probe Kernel");

src/cuda/operator/hash_join_inner.cu

@@ -224,11 +224,11 @@ void probeHashTableSingleMatch(uint8_t **keys, unsigned long long* ht, uint64_t
             uint64_t* &count, uint64_t N, int* condition_mode, int num_keys, int join_mode) {
     CHECK_ERROR();
     GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
-    if (N == 0) {
+    if (N == 0 || ht_len == 0) {
         uint64_t* h_count = gpuBufferManager->customCudaHostAlloc<uint64_t>(1);
         h_count[0] = 0;
         count = h_count;
-        SIRIUS_LOG_DEBUG("Input size is 0");
+        SIRIUS_LOG_DEBUG("Input size is 0 or hash table is empty");
         return;
     }
     SIRIUS_LOG_DEBUG("Launching Probe Kernel Unique Join");
@@ -292,8 +292,8 @@ void probeHashTableSingleMatch(uint8_t **keys, unsigned long long* ht, uint64_t
 
 void probeHashTableRightSemiAntiSingleMatch(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uint64_t N, int* condition_mode, int num_keys) {
     CHECK_ERROR();
-    if (N == 0) {
-        SIRIUS_LOG_DEBUG("Input size is 0");
+    if (N == 0 || ht_len == 0) {
+        SIRIUS_LOG_DEBUG("Input size is 0 or hash table is empty");
         return;
     }
     SIRIUS_LOG_DEBUG("Launching Probe Kernel Unique Join");
@@ -332,14 +332,16 @@ void probeHashTableRightSemiAntiSingleMatch(uint8_t **keys, unsigned long long*
 
 void probeHashTableMark(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uint8_t* &output, uint64_t N, int* condition_mode, int num_keys) {
     CHECK_ERROR();
-    if (N == 0) {
-        SIRIUS_LOG_DEBUG("Input size is 0");
+    GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
+    if (N == 0 || ht_len == 0) {
+        output = gpuBufferManager->customCudaMalloc<uint8_t>(N, 0, 0);
+        cudaMemset(output, 0, N * sizeof(uint8_t));
+        SIRIUS_LOG_DEBUG("Input size is 0 or hash table is empty");
         return;
     }
     SIRIUS_LOG_DEBUG("Launching Probe Kernel Mark");
     SETUP_TIMING();
     START_TIMER();
-    GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
 
     //reinterpret cast the keys to uint64_t
     uint64_t** keys_data = gpuBufferManager->customCudaHostAlloc<uint64_t*>(num_keys);

src/cuda/operator/hash_join_right.cu

@@ -153,7 +153,7 @@ __global__ void create_cpu_strings(duckdb_string_type* gpu_strings, char* cpu_ch
     }
 }
 
-void materializeStringColumnToDuckdbFormat(GPUColumn* column, char* column_char_write_buffer, string_t* column_string_write_buffer) {
+void materializeStringColumnToDuckdbFormat(shared_ptr<GPUColumn> column, char* column_char_write_buffer, string_t* column_string_write_buffer) {
     // First copy the characters from the GPU to the CPU
     SIRIUS_LOG_DEBUG("Materialize String Column to Duckdb format");
     SETUP_TIMING();