Skip to content

Commit f2eb2cb

Browse files
Yifei-yang7Yifei-yang7
committed
Support reading column of decimal type
1 parent f7f790f commit f2eb2cb

6 files changed

Lines changed: 101 additions & 1 deletion

File tree

‎src/cuda/operator/materialize.cu‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,8 @@ template
6060
__global__ void materialize_expression<double, BLOCK_THREADS, ITEMS_PER_THREAD>(const double *a, double* result, uint64_t *row_ids, uint64_t N);
6161
template
6262
__global__ void materialize_expression<uint8_t, BLOCK_THREADS, ITEMS_PER_THREAD>(const uint8_t *a, uint8_t* result, uint64_t *row_ids, uint64_t N);
63+
template
64+
__global__ void materialize_expression<int64_t, BLOCK_THREADS, ITEMS_PER_THREAD>(const int64_t *a, int64_t* result, uint64_t *row_ids, uint64_t N);
6365

6466
template <typename T>
6567
void materializeExpression(T *a, T*& result, uint64_t *row_ids, uint64_t result_len, uint64_t input_len) {
@@ -189,5 +191,7 @@ template
189191
void materializeExpression<double>(double *a, double*& result, uint64_t *row_ids, uint64_t result_len, uint64_t input_size);
190192
template
191193
void materializeExpression<uint8_t>(uint8_t *a, uint8_t*& result, uint64_t *row_ids, uint64_t result_len, uint64_t input_size);
194+
template
195+
void materializeExpression<int64_t>(int64_t *a, int64_t*& result, uint64_t *row_ids, uint64_t result_len, uint64_t input_size);
192196

193197
} // namespace duckdb

‎src/gpu_buffer_manager.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,9 @@ GPUBufferManager::customCudaMalloc<int>(size_t size, int gpu, bool caching);
1818
template uint64_t*
1919
GPUBufferManager::customCudaMalloc<uint64_t>(size_t size, int gpu, bool caching);
2020

21+
template int64_t*
22+
GPUBufferManager::customCudaMalloc<int64_t>(size_t size, int gpu, bool caching);
23+
2124
template uint8_t*
2225
GPUBufferManager::customCudaMalloc<uint8_t>(size_t size, int gpu, bool caching);
2326

‎src/gpu_columns.cpp‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,25 @@
11
#include "gpu_columns.hpp"
22
#include "gpu_buffer_manager.hpp"
33
#include "log/logging.hpp"
4+
#include "duckdb/common/types/decimal.hpp"
45

56
namespace duckdb {
67

8+
size_t GPUDecimalTypeInfo::GetDecimalTypeSize() const {
9+
if (width_ <= Decimal::MAX_WIDTH_INT16) {
10+
return sizeof(int16_t);
11+
} else if (width_ <= Decimal::MAX_WIDTH_INT32) {
12+
return sizeof(int32_t);
13+
} else if (width_ <= Decimal::MAX_WIDTH_INT64) {
14+
return sizeof(int64_t);
15+
} else if (width_ <= Decimal::MAX_WIDTH_INT128) {
16+
return sizeof(__int128_t);
17+
} else {
18+
throw InternalException("Decimal has a width of %d which is bigger than the maximum supported width of %d",
19+
width_, DecimalType::MaxWidth());
20+
}
21+
}
22+
723
DataWrapper::DataWrapper(GPUColumnType _type, uint8_t* _data, size_t _size) : data(_data), size(_size) {
824
type = _type;
925
num_bytes = size * getColumnTypeSize();
@@ -31,6 +47,13 @@ DataWrapper::getColumnTypeSize() {
3147
return sizeof(uint8_t);
3248
case GPUColumnTypeId::VARCHAR:
3349
return 128;
50+
case GPUColumnTypeId::DECIMAL: {
51+
GPUDecimalTypeInfo* decimal_type_info = type.GetDecimalTypeInfo();
52+
if (decimal_type_info == nullptr) {
53+
throw InternalException("`decimal_type_info` not set for DECIMAL type in `getColumnTypeSize`");
54+
}
55+
return decimal_type_info->GetDecimalTypeSize();
56+
}
3457
default:
3558
throw duckdb::InternalException("Unsupported sirius column type in `getColumnTypeSize()`: %d",
3659
static_cast<int>(type.id()));

‎src/include/gpu_columns.hpp‎

Lines changed: 28 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,20 @@ enum class GPUColumnTypeId {
2424
BOOLEAN,
2525
DATE,
2626
VARCHAR,
27-
INT128
27+
INT128,
28+
DECIMAL
29+
};
30+
31+
struct GPUDecimalTypeInfo {
32+
GPUDecimalTypeInfo(uint8_t width, uint8_t scale)
33+
: width_(width), scale_(scale) {
34+
D_ASSERT(width >= scale);
35+
}
36+
37+
size_t GetDecimalTypeSize() const;
38+
39+
uint8_t width_;
40+
uint8_t scale_;
2841
};
2942

3043
struct GPUColumnType {
@@ -37,8 +50,17 @@ struct GPUColumnType {
3750
return id_;
3851
}
3952

53+
inline GPUDecimalTypeInfo* GetDecimalTypeInfo() const {
54+
return decimal_type_info_.get();
55+
}
56+
57+
void SetDecimalTypeInfo(uint8_t width, uint8_t scale) {
58+
decimal_type_info_ = make_shared_ptr<GPUDecimalTypeInfo>(width, scale);
59+
}
60+
4061
private:
4162
GPUColumnTypeId id_;
63+
shared_ptr<GPUDecimalTypeInfo> decimal_type_info_;
4264
};
4365

4466
inline GPUColumnType convertLogicalTypeToColumnType(LogicalType type) {
@@ -57,6 +79,11 @@ inline GPUColumnType convertLogicalTypeToColumnType(LogicalType type) {
5779
return GPUColumnType(GPUColumnTypeId::DATE);
5880
case LogicalTypeId::VARCHAR:
5981
return GPUColumnType(GPUColumnTypeId::VARCHAR);
82+
case LogicalTypeId::DECIMAL: {
83+
GPUColumnType column_type(GPUColumnTypeId::DECIMAL);
84+
column_type.SetDecimalTypeInfo(DecimalType::GetWidth(type), DecimalType::GetScale(type));
85+
return column_type;
86+
}
6087
default:
6188
throw InvalidInputException("Unsupported duckdb column type in `convertLogicalTypeToColumnType`: %d",
6289
static_cast<int>(type.id()));

‎src/operator/gpu_materialize.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,17 @@ HandleMaterializeExpression(shared_ptr<GPUColumn> column, BoundReferenceExpressi
6969
return ResolveTypeMaterializeExpression<uint8_t>(column, bound_ref, gpuBufferManager);
7070
case GPUColumnTypeId::VARCHAR:
7171
return ResolveTypeMaterializeString(column, bound_ref, gpuBufferManager);
72+
case GPUColumnTypeId::DECIMAL: {
73+
switch (column->data_wrapper.getColumnTypeSize()) {
74+
case sizeof(int32_t):
75+
return ResolveTypeMaterializeExpression<int32_t>(column, bound_ref, gpuBufferManager);
76+
case sizeof(int64_t):
77+
return ResolveTypeMaterializeExpression<int64_t>(column, bound_ref, gpuBufferManager);
78+
default:
79+
throw NotImplementedException("Unsupported sirius DECIMAL column type size in `HandleMaterializeExpression`: %zu",
80+
column->data_wrapper.getColumnTypeSize());
81+
}
82+
}
7283
default:
7384
throw NotImplementedException("Unsupported sirius column type in `HandleMaterializeExpression`: %d",
7485
static_cast<int>(column->data_wrapper.type.id()));

‎src/operator/gpu_physical_result_collector.cpp‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -154,6 +154,23 @@ GPUPhysicalMaterializedCollector::FinalMaterialize(GPUIntermediateRelation input
154154
case GPUColumnTypeId::VARCHAR:
155155
FinalMaterializeString(input_relation, output_relation, col);
156156
break;
157+
case GPUColumnTypeId::DECIMAL: {
158+
switch (input_relation.columns[col]->data_wrapper.getColumnTypeSize()) {
159+
case sizeof(int32_t): {
160+
FinalMaterializeInternal<int32_t>(input_relation, output_relation, col);
161+
size_bytes = output_relation.columns[col]->column_length * sizeof(int32_t);
162+
break;
163+
}
164+
case sizeof(int64_t): {
165+
FinalMaterializeInternal<int64_t>(input_relation, output_relation, col);
166+
size_bytes = output_relation.columns[col]->column_length * sizeof(int64_t);
167+
break;
168+
}
169+
throw NotImplementedException("Unsupported sirius DECIMAL column type size in `FinalMaterialize`: %zu",
170+
input_relation.columns[col]->data_wrapper.getColumnTypeSize());
171+
}
172+
break;
173+
}
157174
default:
158175
throw NotImplementedException("Unsupported sirius column type in `FinalMaterialize`: %d",
159176
static_cast<int>(input_relation.columns[col]->data_wrapper.type.id()));
@@ -181,6 +198,13 @@ LogicalType ColumnTypeToLogicalType(const GPUColumnType& type) {
181198
return LogicalType::VARCHAR;
182199
case GPUColumnTypeId::INT128:
183200
return LogicalType::HUGEINT;
201+
case GPUColumnTypeId::DECIMAL: {
202+
GPUDecimalTypeInfo* decimal_type_info = type.GetDecimalTypeInfo();
203+
if (decimal_type_info == nullptr) {
204+
throw InternalException("`decimal_type_info` not set for DECIMAL type in `ColumnTypeToLogicalType`");
205+
}
206+
return LogicalType::DECIMAL(decimal_type_info->width_, decimal_type_info->scale_);
207+
}
184208
default:
185209
throw NotImplementedException("Unsupported sirius column type in `ColumnTypeToLogicalType`: %d",
186210
static_cast<int>(type.id()));
@@ -203,6 +227,14 @@ Vector rawDataToVector(uint8_t* host_data, size_t vector_offset, const GPUColumn
203227
sizeof_type = sizeof(uint8_t); break;
204228
case GPUColumnTypeId::INT128:
205229
sizeof_type = 2 * sizeof(uint64_t); break;
230+
case GPUColumnTypeId::DECIMAL: {
231+
GPUDecimalTypeInfo* decimal_type_info = type.GetDecimalTypeInfo();
232+
if (decimal_type_info == nullptr) {
233+
throw InternalException("`decimal_type_info` not set for DECIMAL type in `rawDataToVector`");
234+
}
235+
sizeof_type = decimal_type_info->GetDecimalTypeSize();
236+
break;
237+
}
206238
default:
207239
throw NotImplementedException("Unsupported sirius column type in `rawDataToVector`: %d",
208240
static_cast<int>(type.id()));

0 commit comments

Comments
 (0)