Skip to content

Commit f7f790f

Browse files
Yifei-yang7Yifei-yang7
committed
Wrap gpu column type as a struct to allow for more type info, such as width/scale for decimal types.
1 parent 5eef478 commit f7f790f

28 files changed

Lines changed: 363 additions & 353 deletions

‎src/cuda/cudf/cudf_aggregate.cu‎

Lines changed: 17 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ void cudf_aggregate(vector<shared_ptr<GPUColumn>>& column, uint64_t num_aggregat
3333
if (agg_mode[agg_idx] == AggregationType::COUNT_STAR || agg_mode[agg_idx] == AggregationType::COUNT) {
3434
uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
3535
cudaMemset(temp, 0, sizeof(uint64_t));
36-
column[agg_idx] = make_shared_ptr<GPUColumn>(1, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
36+
column[agg_idx] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp));
3737
} else {
3838
column[agg_idx] = make_shared_ptr<GPUColumn>(0, column[agg_idx]->data_wrapper.type, column[agg_idx]->data_wrapper.data);
3939
}
@@ -58,17 +58,17 @@ void cudf_aggregate(vector<shared_ptr<GPUColumn>>& column, uint64_t num_aggregat
5858
if (column[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::COUNT && column[agg]->column_length == 0) {
5959
uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
6060
cudaMemset(temp, 0, sizeof(uint64_t));
61-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
61+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp));
6262
} else if (column[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::SUM && column[agg]->column_length == 0) {
6363
uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
6464
cudaMemset(temp, 0, sizeof(uint64_t));
65-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
65+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp));
6666
} else if (column[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::COUNT_STAR && column[agg]->column_length != 0) {
6767
uint64_t* res = gpuBufferManager->customCudaHostAlloc<uint64_t>(1);
6868
res[0] = size;
6969
uint64_t* result_temp = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
7070
cudaMemcpy(result_temp, res, sizeof(uint64_t), cudaMemcpyHostToDevice);
71-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::INT64, reinterpret_cast<uint8_t*>(result_temp));
71+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(result_temp));
7272
} else if (agg_mode[agg] == AggregationType::SUM) {
7373
auto aggregate = make_reduce_aggregation<cudf::reduce_aggregation::SUM>();
7474
auto cudf_column = column[agg]->convertToCudfColumn();
@@ -94,29 +94,29 @@ void cudf_aggregate(vector<shared_ptr<GPUColumn>>& column, uint64_t num_aggregat
9494
res[0] = size;
9595
uint64_t* result_temp = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
9696
cudaMemcpy(result_temp, res, sizeof(uint64_t), cudaMemcpyHostToDevice);
97-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::INT64, reinterpret_cast<uint8_t*>(result_temp));
97+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(result_temp));
9898
} else if (agg_mode[agg] == AggregationType::FIRST) {
99-
if (column[agg]->data_wrapper.type == ColumnType::INT64) {
99+
if (column[agg]->data_wrapper.type.id() == GPUColumnTypeId::INT64) {
100100
uint64_t* result_temp = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
101101
cudaMemcpy(result_temp, reinterpret_cast<uint64_t*>(column[agg]->data_wrapper.data), sizeof(uint64_t), cudaMemcpyDeviceToDevice);
102-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::INT64, reinterpret_cast<uint8_t*>(result_temp));
103-
} else if (column[agg]->data_wrapper.type == ColumnType::INT32) {
102+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(result_temp));
103+
} else if (column[agg]->data_wrapper.type.id() == GPUColumnTypeId::INT32) {
104104
int32_t* result_temp = gpuBufferManager->customCudaMalloc<int32_t>(1, 0, 0);
105105
cudaMemcpy(result_temp, reinterpret_cast<int32_t*>(column[agg]->data_wrapper.data), sizeof(int32_t), cudaMemcpyDeviceToDevice);
106-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::INT32, reinterpret_cast<uint8_t*>(result_temp));
107-
} else if (column[agg]->data_wrapper.type == ColumnType::FLOAT32) {
106+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::INT32), reinterpret_cast<uint8_t*>(result_temp));
107+
} else if (column[agg]->data_wrapper.type.id() == GPUColumnTypeId::FLOAT32) {
108108
float* result_temp = gpuBufferManager->customCudaMalloc<float>(1, 0, 0);
109109
cudaMemcpy(result_temp, reinterpret_cast<float*>(column[agg]->data_wrapper.data), sizeof(float), cudaMemcpyDeviceToDevice);
110-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::FLOAT32, reinterpret_cast<uint8_t*>(result_temp));
111-
} else if (column[agg]->data_wrapper.type == ColumnType::FLOAT64) {
110+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::FLOAT32), reinterpret_cast<uint8_t*>(result_temp));
111+
} else if (column[agg]->data_wrapper.type.id() == GPUColumnTypeId::FLOAT64) {
112112
double* result_temp = gpuBufferManager->customCudaMalloc<double>(1, 0, 0);
113113
cudaMemcpy(result_temp, reinterpret_cast<double*>(column[agg]->data_wrapper.data), sizeof(double), cudaMemcpyDeviceToDevice);
114-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::FLOAT64, reinterpret_cast<uint8_t*>(result_temp));
115-
} else if (column[agg]->data_wrapper.type == ColumnType::BOOLEAN) {
114+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::FLOAT64), reinterpret_cast<uint8_t*>(result_temp));
115+
} else if (column[agg]->data_wrapper.type.id() == GPUColumnTypeId::BOOLEAN) {
116116
uint8_t* result_temp = gpuBufferManager->customCudaMalloc<uint8_t>(1, 0, 0);
117117
cudaMemcpy(result_temp, reinterpret_cast<uint8_t*>(column[agg]->data_wrapper.data), sizeof(uint8_t), cudaMemcpyDeviceToDevice);
118-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::BOOLEAN, reinterpret_cast<uint8_t*>(result_temp));
119-
} else if (column[agg]->data_wrapper.type == ColumnType::VARCHAR) {
118+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::BOOLEAN), reinterpret_cast<uint8_t*>(result_temp));
119+
} else if (column[agg]->data_wrapper.type.id() == GPUColumnTypeId::VARCHAR) {
120120
uint64_t* length = gpuBufferManager->customCudaHostAlloc<uint64_t>(1);
121121
cudaMemcpy(length, column[agg]->data_wrapper.offset + 1, sizeof(uint64_t), cudaMemcpyDeviceToHost);
122122

@@ -126,7 +126,7 @@ void cudf_aggregate(vector<shared_ptr<GPUColumn>>& column, uint64_t num_aggregat
126126
uint64_t* new_offset = gpuBufferManager->customCudaMalloc<uint64_t>(2, 0, 0);
127127
cudaMemcpy(new_offset, column[agg]->data_wrapper.offset, 2 * sizeof(uint64_t), cudaMemcpyDeviceToDevice);
128128

129-
column[agg] = make_shared_ptr<GPUColumn>(1, ColumnType::VARCHAR, reinterpret_cast<uint8_t*>(result_temp), new_offset, length[0], true);
129+
column[agg] = make_shared_ptr<GPUColumn>(1, GPUColumnType(GPUColumnTypeId::VARCHAR), reinterpret_cast<uint8_t*>(result_temp), new_offset, length[0], true);
130130
}
131131
}
132132
else {

‎src/cuda/cudf/cudf_groupby.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@ void cudf_groupby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
1212
SIRIUS_LOG_DEBUG("Input size is 0");
1313
for (idx_t group = 0; group < num_keys; group++) {
1414
bool old_unique = keys[group]->is_unique;
15-
if (keys[group]->data_wrapper.type == ColumnType::VARCHAR) {
15+
if (keys[group]->data_wrapper.type.id() == GPUColumnTypeId::VARCHAR) {
1616
keys[group] = make_shared_ptr<GPUColumn>(0, keys[group]->data_wrapper.type, keys[group]->data_wrapper.data, keys[group]->data_wrapper.offset, 0, true);
1717
} else {
1818
keys[group] = make_shared_ptr<GPUColumn>(0, keys[group]->data_wrapper.type, keys[group]->data_wrapper.data);
@@ -22,7 +22,7 @@ void cudf_groupby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
2222

2323
for (int agg_idx = 0; agg_idx < num_aggregates; agg_idx++) {
2424
if (agg_mode[agg_idx] == AggregationType::COUNT_STAR || agg_mode[agg_idx] == AggregationType::COUNT) {
25-
aggregate_keys[agg_idx] = make_shared_ptr<GPUColumn>(0, ColumnType::INT64, aggregate_keys[agg_idx]->data_wrapper.data);
25+
aggregate_keys[agg_idx] = make_shared_ptr<GPUColumn>(0, GPUColumnType(GPUColumnTypeId::INT64), aggregate_keys[agg_idx]->data_wrapper.data);
2626
} else {
2727
aggregate_keys[agg_idx] = make_shared_ptr<GPUColumn>(0, aggregate_keys[agg_idx]->data_wrapper.type, aggregate_keys[agg_idx]->data_wrapper.data);
2828
}
@@ -64,21 +64,21 @@ void cudf_groupby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
6464
requests[agg].aggregations.push_back(std::move(aggregate));
6565
uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(size, 0, 0);
6666
cudaMemset(temp, 0, size * sizeof(uint64_t));
67-
shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
67+
shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp));
6868
requests[agg].values = temp_column->convertToCudfColumn();
6969
} else if (aggregate_keys[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::SUM && aggregate_keys[agg]->column_length == 0) {
7070
auto aggregate = cudf::make_sum_aggregation<cudf::groupby_aggregation>();
7171
requests[agg].aggregations.push_back(std::move(aggregate));
7272
uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(size, 0, 0);
7373
cudaMemset(temp, 0, size * sizeof(uint64_t));
74-
shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
74+
shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp));
7575
requests[agg].values = temp_column->convertToCudfColumn();
7676
} else if (aggregate_keys[agg]->data_wrapper.data == nullptr && agg_mode[agg] == AggregationType::COUNT_STAR && aggregate_keys[agg]->column_length != 0) {
7777
auto aggregate = cudf::make_count_aggregation<cudf::groupby_aggregation>(cudf::null_policy::EXCLUDE);
7878
requests[agg].aggregations.push_back(std::move(aggregate));
7979
uint64_t* temp = gpuBufferManager->customCudaMalloc<uint64_t>(size, 0, 0);
8080
cudaMemset(temp, 0, size * sizeof(uint64_t));
81-
shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, ColumnType::INT64, reinterpret_cast<uint8_t*>(temp));
81+
shared_ptr<GPUColumn> temp_column = make_shared_ptr<GPUColumn>(size, GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp));
8282
requests[agg].values = temp_column->convertToCudfColumn();
8383
} else if (agg_mode[agg] == AggregationType::SUM) {
8484
auto aggregate = cudf::make_sum_aggregation<cudf::groupby_aggregation>();
@@ -119,7 +119,7 @@ void cudf_groupby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
119119
if (agg_mode[agg] == AggregationType::COUNT || agg_mode[agg] == AggregationType::COUNT_STAR) {
120120
auto agg_val_view = agg_val->view();
121121
auto temp_data = convertInt32ToUInt64(const_cast<int32_t*>(agg_val_view.data<int32_t>()), agg_val_view.size());
122-
aggregate_keys[agg] = make_shared_ptr<GPUColumn>(agg_val_view.size(), ColumnType::INT64, reinterpret_cast<uint8_t*>(temp_data));
122+
aggregate_keys[agg] = make_shared_ptr<GPUColumn>(agg_val_view.size(), GPUColumnType(GPUColumnTypeId::INT64), reinterpret_cast<uint8_t*>(temp_data));
123123
} else {
124124
aggregate_keys[agg]->setFromCudfColumn(*agg_val, false, nullptr, 0, gpuBufferManager);
125125
// aggregate_keys[agg] = gpuBufferManager->copyDataFromcuDFColumn(agg_val_view, 0);

‎src/cuda/cudf/cudf_orderby.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@ void cudf_orderby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
1212
SIRIUS_LOG_DEBUG("Input size is 0");
1313
for (idx_t col = 0; col < num_projections; col++) {
1414
bool old_unique = projection[col]->is_unique;
15-
if (projection[col]->data_wrapper.type == ColumnType::VARCHAR) {
15+
if (projection[col]->data_wrapper.type.id() == GPUColumnTypeId::VARCHAR) {
1616
projection[col] = make_shared_ptr<GPUColumn>(0, projection[col]->data_wrapper.type, projection[col]->data_wrapper.data, projection[col]->data_wrapper.offset, 0, true);
1717
} else {
1818
projection[col] = make_shared_ptr<GPUColumn>(0, projection[col]->data_wrapper.type, projection[col]->data_wrapper.data);
@@ -75,7 +75,7 @@ void cudf_orderby(vector<shared_ptr<GPUColumn>>& keys, vector<shared_ptr<GPUColu
7575
//copy the projection columns to a new array
7676
// GPUColumn** projection_columns = new GPUColumn*[num_projections];
7777
// for (int projection_idx = 0; projection_idx < num_projections; projection_idx++) {
78-
// if (projection[projection_idx]->data_wrapper.type == ColumnType::VARCHAR) {
78+
// if (projection[projection_idx]->data_wrapper.type.id() == GPUColumnTypeId::VARCHAR) {
7979
// uint64_t* temp_offset = gpuBufferManager->customCudaMalloc<uint64_t>(projection[projection_idx]->column_length, 0, false);
8080
// uint8_t* temp_column = gpuBufferManager->customCudaMalloc<uint8_t>(projection[projection_idx]->data_wrapper.num_bytes, 0, false);
8181
// callCudaMemcpyDeviceToDevice<uint64_t>(temp_offset, projection[projection_idx]->data_wrapper.offset, projection[projection_idx]->column_length, 0);

‎src/cuda/cudf/test_sorting.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -108,7 +108,7 @@ int main() {
108108
// //copy the projection columns to a new array
109109
// GPUColumn** projection_columns = new GPUColumn*[num_projections];
110110
// for (int projection_idx = 0; projection_idx < num_projections; projection_idx++) {
111-
// if (projection[projection_idx]->data_wrapper.type == ColumnType::VARCHAR) {
111+
// if (projection[projection_idx]->data_wrapper.type.id() == GPUColumnTypeId::VARCHAR) {
112112
// uint64_t* temp_offset = gpuBufferManager->customCudaMalloc<uint64_t>(projection[projection_idx]->column_length, 0, false);
113113
// uint8_t* temp_column = gpuBufferManager->customCudaMalloc<uint8_t>(projection[projection_idx]->data_wrapper.num_bytes, 0, false);
114114
// callCudaMemcpyDeviceToDevice<uint64_t>(temp_offset, projection[projection_idx]->data_wrapper.offset, projection[projection_idx]->column_length, 0);

‎src/cuda/expression_executor/gpu_dispatch_materialize.cu‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -173,47 +173,47 @@ std::unique_ptr<cudf::column> GpuDispatcher::DispatchMaterialize(const GPUColumn
173173
const auto* input_data = input->data_wrapper.data;
174174
const auto* input_offsets = input->data_wrapper.offset; // Maybe unused
175175

176-
switch (input->data_wrapper.type)
176+
switch (input->data_wrapper.type.id())
177177
{
178-
case ColumnType::INT32:
178+
case GPUColumnTypeId::INT32:
179179
return MaterializeNumeric<int32_t>::Do(reinterpret_cast<const int32_t*>(input_data),
180180
input->row_ids,
181181
input->row_id_count,
182182
mr);
183-
case ColumnType::INT64:
183+
case GPUColumnTypeId::INT64:
184184
return MaterializeNumeric<uint64_t>::Do(reinterpret_cast<const uint64_t*>(input_data),
185185
input->row_ids,
186186
input->row_id_count,
187187
mr);
188-
case ColumnType::FLOAT32:
188+
case GPUColumnTypeId::FLOAT32:
189189
return MaterializeNumeric<float_t>::Do(reinterpret_cast<const float_t*>(input_data),
190190
input->row_ids,
191191
input->row_id_count,
192192
mr);
193-
case ColumnType::FLOAT64:
193+
case GPUColumnTypeId::FLOAT64:
194194
return MaterializeNumeric<double_t>::Do(reinterpret_cast<const double_t*>(input_data),
195195
input->row_ids,
196196
input->row_id_count,
197197
mr);
198-
case ColumnType::BOOLEAN:
198+
case GPUColumnTypeId::BOOLEAN:
199199
return MaterializeNumeric<bool>::Do(reinterpret_cast<const bool*>(input_data),
200200
input->row_ids,
201201
input->row_id_count,
202202
mr);
203-
case ColumnType::DATE:
203+
case GPUColumnTypeId::DATE:
204204
return MaterializeNumeric<cudf::timestamp_D>::Do(reinterpret_cast<const cudf::timestamp_D*>(input_data),
205205
input->row_ids,
206206
input->row_id_count,
207207
mr);
208-
case ColumnType::VARCHAR:
208+
case GPUColumnTypeId::VARCHAR:
209209
return MaterializeString::Do(input_data,
210210
input_offsets,
211211
input->row_ids,
212212
input->row_id_count,
213213
mr);
214214
default:
215215
throw InternalException("Unsupported sirius column type in `Dispatch[Materialize]`: %d",
216-
static_cast<int>(input->data_wrapper.type));
216+
static_cast<int>(input->data_wrapper.type.id()));
217217
}
218218
}
219219

‎src/expression_executor/specializations/gpu_execute_conjunction.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ std::unique_ptr<cudf::column> GpuExpressionExecutor::Execute(const BoundConjunct
3030
std::unique_ptr<cudf::column> output_column;
3131
for (idx_t i = 0; i < expr.children.size(); i++)
3232
{
33-
D_ASSERT(state->intermediate_columns[i].data_wrapper.type == ColumnType::BOOLEAN);
33+
D_ASSERT(state->intermediate_columns[i].data_wrapper.type.id() == GPUColumnTypeId::BOOLEAN);
3434

3535
auto current_result = Execute(*expr.children[i], state->child_states[i].get());
3636

‎src/expression_executor/specializations/gpu_execute_reference.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ std::unique_ptr<cudf::column> GpuExpressionExecutor::Execute(const BoundReferenc
2929
// DispatchMaterialize will handle byte overflow from the materialized offsets
3030
return GpuDispatcher::DispatchMaterialize(input_column.get(), resource_ref);
3131
}
32-
if (input_column->data_wrapper.type == ColumnType::VARCHAR &&
32+
if (input_column->data_wrapper.type.id() == GPUColumnTypeId::VARCHAR &&
3333
input_column->data_wrapper.num_bytes > INT32_MAX)
3434
{
3535
throw NotImplementedException(

0 commit comments

Comments
 (0)