Skip to content

Commit e25ad41

Browse files
Yifei-yang7Yifei-yang7
committed
Support basic filtering and projections on decimal types
1 parent 34c0b9a commit e25ad41

6 files changed

Lines changed: 133 additions & 13 deletions

File tree

‎src/cuda/operator/arbitrary_expression.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -131,7 +131,7 @@ __global__ void table_scan_expression(uint8_t **col, uint64_t** offset, uint8_t
131131
for (int expr = 0; expr < num_expr; expr++) {
132132
if (threadIdx.x + ITEM * B < num_tile_items) {
133133

134-
if (data_type[expr] == INT32 || data_type[expr] == DATE) {
134+
if (data_type[expr] == INT32 || data_type[expr] == DATE || data_type[expr] == DECIMAL32) {
135135
uint64_t item_idx = tile_offset + threadIdx.x + ITEM * B;
136136
int item = (reinterpret_cast<int*>(col[expr]))[item_idx];
137137

@@ -141,15 +141,15 @@ __global__ void table_scan_expression(uint8_t **col, uint64_t** offset, uint8_t
141141

142142
selection_flags[ITEM] = device_comparison<int>(item, constant, compare_mode[expr]);
143143

144-
} else if (data_type[expr] == INT64) {
144+
} else if (data_type[expr] == INT64 || data_type[expr] == DECIMAL64) {
145145
uint64_t item_idx = tile_offset + threadIdx.x + ITEM * B;
146-
uint64_t item = (reinterpret_cast<uint64_t*>(col[expr]))[item_idx];
146+
int64_t item = (reinterpret_cast<int64_t*>(col[expr]))[item_idx];
147147

148148
uint64_t start_constant_offset = constant_offset[expr];
149-
uint64_t constant;
150-
memcpy(&constant, constant_compare + start_constant_offset, sizeof(uint64_t));
149+
int64_t constant;
150+
memcpy(&constant, constant_compare + start_constant_offset, sizeof(int64_t));
151151

152-
selection_flags[ITEM] = device_comparison<uint64_t>(item, constant, compare_mode[expr]);
152+
selection_flags[ITEM] = device_comparison<int64_t>(item, constant, compare_mode[expr]);
153153

154154
} else if (data_type[expr] == FLOAT32) {
155155
uint64_t item_idx = tile_offset + threadIdx.x + ITEM * B;

‎src/expression_executor/specializations/gpu_execute_function.cpp‎

Lines changed: 56 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -149,7 +149,7 @@ struct NumericBinaryFunctionDispatcher
149149
: executor(exec)
150150
{}
151151

152-
// Left scalar binary operator
152+
// Left scalar binary operator for numeric types
153153
template <typename T>
154154
std::unique_ptr<cudf::column> DoLeftScalarBinaryOp(const T& left_value,
155155
const cudf::column_view& right,
@@ -165,7 +165,24 @@ struct NumericBinaryFunctionDispatcher
165165
executor.resource_ref);
166166
}
167167

168-
// Right scalar binary operator
168+
// Left scalar binary operator for decimal types
169+
template <typename T>
170+
std::unique_ptr<cudf::column> DoLeftScalarBinaryOp(typename T::rep left_value,
171+
numeric::scale_type scale,
172+
const cudf::column_view& right,
173+
const cudf::data_type& return_type)
174+
{
175+
auto left_decimal_scalar = cudf::fixed_point_scalar<T>(
176+
left_value, scale, true, cudf::get_default_stream(), executor.resource_ref);
177+
return cudf::binary_operation(left_decimal_scalar,
178+
right,
179+
BinOp,
180+
return_type,
181+
cudf::get_default_stream(),
182+
executor.resource_ref);
183+
}
184+
185+
// Right scalar binary operator for numeric types
169186
template <typename T>
170187
std::unique_ptr<cudf::column> DoRightScalarBinaryOp(const cudf::column_view& left,
171188
const T& right_value,
@@ -181,6 +198,23 @@ struct NumericBinaryFunctionDispatcher
181198
executor.resource_ref);
182199
}
183200

201+
// Right scalar binary operator for decimal types
202+
template <typename T>
203+
std::unique_ptr<cudf::column> DoRightScalarBinaryOp(const cudf::column_view& left,
204+
typename T::rep right_value,
205+
numeric::scale_type scale,
206+
const cudf::data_type& return_type)
207+
{
208+
auto right_decimal_scalar = cudf::fixed_point_scalar<T>(
209+
right_value, scale, true, cudf::get_default_stream(), executor.resource_ref);
210+
return cudf::binary_operation(left,
211+
right_decimal_scalar,
212+
BinOp,
213+
return_type,
214+
cudf::get_default_stream(),
215+
executor.resource_ref);
216+
}
217+
184218
// Dispatch operator
185219
std::unique_ptr<cudf::column> operator()(const BoundFunctionExpression& expr,
186220
GpuExpressionState* state)
@@ -205,6 +239,16 @@ struct NumericBinaryFunctionDispatcher
205239
return DoLeftScalarBinaryOp(left_value.GetValue<float_t>(), right->view(), return_type);
206240
case cudf::type_id::FLOAT64:
207241
return DoLeftScalarBinaryOp(left_value.GetValue<double_t>(), right->view(), return_type);
242+
case cudf::type_id::DECIMAL32:
243+
return DoLeftScalarBinaryOp<numeric::decimal32>(
244+
left_value.GetValueUnsafe<int32_t>(),
245+
numeric::scale_type{duckdb::DecimalType::GetScale(left_value.type())},
246+
right->view(), return_type);
247+
case cudf::type_id::DECIMAL64:
248+
return DoLeftScalarBinaryOp<numeric::decimal64>(
249+
left_value.GetValueUnsafe<int64_t>(),
250+
numeric::scale_type{duckdb::DecimalType::GetScale(left_value.type())},
251+
right->view(), return_type);
208252
case cudf::type_id::BOOL8:
209253
throw NotImplementedException("Execute[Function]: Boolean types not supported for "
210254
"numeric binary operations!");
@@ -229,6 +273,16 @@ struct NumericBinaryFunctionDispatcher
229273
return DoRightScalarBinaryOp(left->view(), right_value.GetValue<float_t>(), return_type);
230274
case cudf::type_id::FLOAT64:
231275
return DoRightScalarBinaryOp(left->view(), right_value.GetValue<double_t>(), return_type);
276+
case cudf::type_id::DECIMAL32:
277+
return DoRightScalarBinaryOp<numeric::decimal32>(
278+
left->view(), right_value.GetValueUnsafe<int32_t>(),
279+
numeric::scale_type{duckdb::DecimalType::GetScale(right_value.type())},
280+
return_type);
281+
case cudf::type_id::DECIMAL64:
282+
return DoRightScalarBinaryOp<numeric::decimal64>(
283+
left->view(), right_value.GetValueUnsafe<int64_t>(),
284+
numeric::scale_type{duckdb::DecimalType::GetScale(right_value.type())},
285+
return_type);
232286
case cudf::type_id::BOOL8:
233287
throw NotImplementedException("Execute[Function]: Boolean types not supported for "
234288
"numeric binary operations!");

‎src/gpu_columns.cpp‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -137,6 +137,22 @@ GPUColumn::convertToCudfColumn() {
137137
std::move(children)
138138
);
139139
return str_col;
140+
} else if (data_wrapper.type.id() == GPUColumnTypeId::DECIMAL) {
141+
cudf::data_type cudf_type;
142+
switch (data_wrapper.getColumnTypeSize()) {
143+
case sizeof(int32_t): {
144+
cudf_type = cudf::data_type(cudf::type_id::DECIMAL32, data_wrapper.type.GetDecimalTypeInfo()->scale_);
145+
break;
146+
}
147+
case sizeof(int64_t): {
148+
cudf_type = cudf::data_type(cudf::type_id::DECIMAL64, data_wrapper.type.GetDecimalTypeInfo()->scale_);
149+
break;
150+
}
151+
default:
152+
throw duckdb::InternalException("Unsupported sirius DECIMAL column type size in `convertToCudfColumn()`: %zu",
153+
data_wrapper.getColumnTypeSize());
154+
}
155+
return cudf::column_view(cudf_type, size, reinterpret_cast<void*>(data_wrapper.data), nullptr, 0);
140156
}
141157
throw duckdb::InternalException("Unsupported sirius column type in `convertToCudfColumn()`: %d", data_wrapper.type.id());
142158
}

‎src/include/expression_executor/gpu_expression_executor_state.hpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -61,6 +61,17 @@ struct GpuExpressionState
6161
return cudf::data_type(cudf::type_id::TIMESTAMP_DAYS);
6262
case LogicalTypeId::VARCHAR:
6363
return cudf::data_type(cudf::type_id::STRING);
64+
case LogicalTypeId::DECIMAL: {
65+
switch (logical_type.InternalType()) {
66+
case PhysicalType::INT32:
67+
return cudf::data_type(cudf::type_id::DECIMAL32, DecimalType::GetScale(logical_type));
68+
case PhysicalType::INT64:
69+
return cudf::data_type(cudf::type_id::DECIMAL64, DecimalType::GetScale(logical_type));
70+
default:
71+
throw InvalidInputException("GetCudfType: Unsupported duckdb decimal physical type: %d",
72+
static_cast<int>(logical_type.InternalType()));
73+
}
74+
}
6475
default:
6576
throw InvalidInputException("GetCudfType: Unsupported duckdb type: %d", static_cast<int>(logical_type.id()));
6677
}

‎src/include/operator/gpu_physical_table_scan.hpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,9 @@ enum ScanDataType {
1717
FLOAT64,
1818
BOOLEAN,
1919
DATE,
20-
VARCHAR
20+
VARCHAR,
21+
DECIMAL32,
22+
DECIMAL64
2123
};
2224

2325
enum CompareType {

‎src/operator/gpu_physical_table_scan.cpp‎

Lines changed: 41 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -290,7 +290,7 @@ void HandleArbitraryConstantExpression(vector<shared_ptr<GPUColumn>> &column, ui
290290
data_type[expr] = INT32;
291291
break;
292292
} case GPUColumnTypeId::INT64: {
293-
total_bytes += sizeof(uint64_t);
293+
total_bytes += sizeof(int64_t);
294294
data_type[expr] = INT64;
295295
break;
296296
} case GPUColumnTypeId::FLOAT32: {
@@ -310,6 +310,22 @@ void HandleArbitraryConstantExpression(vector<shared_ptr<GPUColumn>> &column, ui
310310
total_bytes += lower_string.size();
311311
data_type[expr] = VARCHAR;
312312
break;
313+
} case GPUColumnTypeId::DECIMAL: {
314+
switch (column[expr]->data_wrapper.getColumnTypeSize()) {
315+
case sizeof(int32_t): {
316+
total_bytes += sizeof(int32_t);
317+
data_type[expr] = DECIMAL32;
318+
break;
319+
}
320+
case sizeof(int64_t): {
321+
total_bytes += sizeof(int64_t);
322+
data_type[expr] = DECIMAL64;
323+
break;
324+
}
325+
throw NotImplementedException("Unsupported sirius DECIMAL column type size in `HandleArbitraryConstantExpression`: %zu",
326+
column[expr]->data_wrapper.getColumnTypeSize());
327+
}
328+
break;
313329
} default: {
314330
throw NotImplementedException("Unsupported sirius column type in `HandleArbitraryConstantExpression`: %d",
315331
static_cast<int>(column[expr]->data_wrapper.type.id()));
@@ -333,10 +349,10 @@ void HandleArbitraryConstantExpression(vector<shared_ptr<GPUColumn>> &column, ui
333349
init_offset += sizeof(int);
334350
break;
335351
} case GPUColumnTypeId::INT64: {
336-
uint64_t temp = filter_constant[expr]->constant.GetValue<uint64_t>();
337-
memcpy(constant_compare + init_offset, &temp, sizeof(uint64_t));
352+
int64_t temp = filter_constant[expr]->constant.GetValue<int64_t>();
353+
memcpy(constant_compare + init_offset, &temp, sizeof(int64_t));
338354
constant_offset[expr] = init_offset;
339-
init_offset += sizeof(uint64_t);
355+
init_offset += sizeof(int64_t);
340356
break;
341357
} case GPUColumnTypeId::FLOAT32: {
342358
float temp = filter_constant[expr]->constant.GetValue<float>();
@@ -356,6 +372,27 @@ void HandleArbitraryConstantExpression(vector<shared_ptr<GPUColumn>> &column, ui
356372
constant_offset[expr] = init_offset;
357373
init_offset += lower_string.size();
358374
break;
375+
} case GPUColumnTypeId::DECIMAL: {
376+
switch (column[expr]->data_wrapper.getColumnTypeSize()) {
377+
case sizeof(int32_t): {
378+
// `GetValue()` cannot work since it will convert to double first, same for below
379+
int32_t temp = filter_constant[expr]->constant.GetValueUnsafe<int32_t>();
380+
memcpy(constant_compare + init_offset, &temp, sizeof(int32_t));
381+
constant_offset[expr] = init_offset;
382+
init_offset += sizeof(int32_t);
383+
break;
384+
}
385+
case sizeof(int64_t): {
386+
int64_t temp = filter_constant[expr]->constant.GetValueUnsafe<int64_t>();
387+
memcpy(constant_compare + init_offset, &temp, sizeof(int64_t));
388+
constant_offset[expr] = init_offset;
389+
init_offset += sizeof(int64_t);
390+
break;
391+
}
392+
throw NotImplementedException("Unsupported sirius DECIMAL column type size in `HandleArbitraryConstantExpression`: %zu",
393+
column[expr]->data_wrapper.getColumnTypeSize());
394+
}
395+
break;
359396
} default: {
360397
throw NotImplementedException("Unsupported sirius column type in `HandleArbitraryConstantExpression`: %d",
361398
column[expr]->data_wrapper.type.id());

0 commit comments

Comments
 (0)