Merge branch 'main' into type_decimal

Author: Yifei-yang7

src/cuda/utils.cu

@@ -35,8 +35,8 @@ __global__ void convert_int64_to_int128(uint8_t *input, uint8_t *output, size_t
     // for (int i = 0; i < 8; ++i) {
       // output[idx * 16 + i] = input[idx * 8 + i];
       // output[idx * 16 + i + 8] = 0;
-      output_ptr[0] = input_ptr[0];
-      output_ptr[1] = 0; // Set the upper 64 bits to zero
+      // output_ptr[0] = input_ptr[0];
+      // output_ptr[1] = 0; // Set the upper 64 bits to zero
     // }
   }
 }
@@ -56,6 +56,16 @@ __global__ void convert_int32_to_int128(uint8_t *input, uint8_t *output, size_t
   }
 }
 
+__global__ void convert_int32_to_int64(uint8_t *input, uint8_t *output, size_t count) {
+  size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < count) {
+    // Store the converted value as 16 bytes in output
+    int32_t* input_ptr = reinterpret_cast<int32_t*>(input + idx * 4);
+    int64_t* output_ptr = reinterpret_cast<int64_t*>(output + idx * 8);
+    output_ptr[0] = input_ptr[0];
+  }
+}
+
 void warmup_gpu() {
   // Perform the warmup
   cudaFree(0);
@@ -85,4 +95,15 @@ void convertInt32ToInt128(uint8_t *input, uint8_t *output, size_t count) {
   cudaDeviceSynchronize();
 }
 
+void convertInt32ToInt64(uint8_t *input, uint8_t *output, size_t count) {
+  if (count == 0) return;
+
+  // Launch the kernel to convert the data
+  size_t threads_per_block = 256;
+  size_t blocks = (count + threads_per_block - 1) / threads_per_block;
+
+  convert_int32_to_int64<<<blocks, threads_per_block>>>(input, output, count);
+  cudaDeviceSynchronize();
+}
+
 } // namespace duckdb

src/include/utils.hpp

@@ -7,5 +7,6 @@ namespace duckdb {
 void warmup_gpu();
 void convertInt64ToInt128(uint8_t *input, uint8_t *output, size_t count);
 void convertInt32ToInt128(uint8_t *input, uint8_t *output, size_t count);
+void convertInt32ToInt64(uint8_t *input, uint8_t *output, size_t count);
 
 } // namespace duckdb

src/operator/gpu_physical_grouped_aggregate.cpp

@@ -4,6 +4,7 @@
 #include "gpu_buffer_manager.hpp"
 #include "gpu_materialize.hpp"
 #include "log/logging.hpp"
+#include "utils.hpp"
 
 namespace duckdb {
 
@@ -304,6 +305,14 @@ HandleGroupByAggregateCuDF(vector<shared_ptr<GPUColumn>> &group_by_keys, vector<
 			agg_mode[agg_idx] = AggregationType::SUM;
 		} else if (expr.function.name.compare("sum_no_overflow") == 0 && aggregate_keys[agg_idx]->data_wrapper.data != nullptr) {
 			agg_mode[agg_idx] = AggregationType::SUM;
+			if (aggregate_keys[agg_idx]->data_wrapper.type.id() == GPUColumnTypeId::INT32) {
+				SIRIUS_LOG_DEBUG("Converting INT32 to INT64 for sum_no_overflow");
+				uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(aggregate_keys[agg_idx]->column_length, 0, 0);
+				convertInt32ToInt64(aggregate_keys[agg_idx]->data_wrapper.data, reinterpret_cast<uint8_t*>(temp), aggregate_keys[agg_idx]->column_length);
+				aggregate_keys[agg_idx]->data_wrapper.data = reinterpret_cast<uint8_t*>(temp);
+				aggregate_keys[agg_idx]->data_wrapper.type = GPUColumnType(GPUColumnTypeId::INT64);
+				aggregate_keys[agg_idx]->data_wrapper.num_bytes = aggregate_keys[agg_idx]->data_wrapper.num_bytes * 2;
+			}
 		} else if (expr.function.name.compare("avg") == 0 && aggregate_keys[agg_idx]->data_wrapper.data != nullptr) {
 			agg_mode[agg_idx] = AggregationType::AVERAGE;
 		} else if (expr.function.name.compare("max") == 0 && aggregate_keys[agg_idx]->data_wrapper.data != nullptr) {

src/operator/gpu_physical_ungrouped_aggregate.cpp

@@ -119,6 +119,14 @@ HandleAggregateExpressionCuDF(vector<shared_ptr<GPUColumn>> &aggregate_keys, GPU
 			agg_mode[agg_idx] = AggregationType::SUM;
 		} else if (expr.function.name.compare("sum_no_overflow") == 0 && aggregate_keys[agg_idx]->data_wrapper.data != nullptr) {
 			agg_mode[agg_idx] = AggregationType::SUM;
+			if (aggregate_keys[agg_idx]->data_wrapper.type.id() == GPUColumnTypeId::INT32) {
+				SIRIUS_LOG_DEBUG("Converting INT32 to INT64 for sum_no_overflow");
+				uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(aggregate_keys[agg_idx]->column_length, 0, 0);
+				convertInt32ToInt64(aggregate_keys[agg_idx]->data_wrapper.data, reinterpret_cast<uint8_t*>(temp), aggregate_keys[agg_idx]->column_length);
+				aggregate_keys[agg_idx]->data_wrapper.data = reinterpret_cast<uint8_t*>(temp);
+				aggregate_keys[agg_idx]->data_wrapper.type = GPUColumnType(GPUColumnTypeId::INT64);
+				aggregate_keys[agg_idx]->data_wrapper.num_bytes = aggregate_keys[agg_idx]->data_wrapper.num_bytes * 2;
+			}
 		} else if (expr.function.name.compare("avg") == 0 && aggregate_keys[agg_idx]->data_wrapper.data != nullptr) {
 			agg_mode[agg_idx] = AggregationType::AVERAGE;
 		} else if (expr.function.name.compare("max") == 0 && aggregate_keys[agg_idx]->data_wrapper.data != nullptr) {