Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
83 commits
Select commit Hold shift + click to select a range
b3043bc
feat(cudf): add decimal sum/avg kernels and aggregation support
simoneves Feb 25, 2026
256db01
Fix after rebase, and format
simoneves Apr 3, 2026
b3e3426
Avoid clang-tidy warnings about structs with virtual functions but no…
simoneves Apr 8, 2026
7512644
Merge branch 'main' into simoneves/decimal_pr3
simoneves Apr 24, 2026
e718ff7
Merge branch 'main' into simoneves/decimal_pr3 (broken)
simoneves Apr 24, 2026
f98a3d1
Reinstate agg registration changes
simoneves Apr 24, 2026
1e257e8
AI-assisted doReduce changes
simoneves Apr 24, 2026
6ebe3d6
Format
simoneves Apr 24, 2026
ea2aa51
Merge branch 'main' into simoneves/decimal_pr3
simoneves Apr 24, 2026
35c5b5f
Merge branch 'main' into simoneves/decimal_pr3
simoneves Apr 30, 2026
e432c95
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 2, 2026
7e68e64
Change ifs to CHECKs and tidy up flow, per @mattgara
simoneves May 2, 2026
ca006e4
Remove 'help' comment now that I understand
simoneves May 2, 2026
32365e8
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 11, 2026
bf5260f
Fix after update from main
simoneves May 11, 2026
7f222b4
Cursor refactor to split up aggregator class (struct)
simoneves May 4, 2026
890d792
Cursor refactor to split up Reduce aggegator in a similar way
simoneves May 4, 2026
837b463
Factor out common helper functions
simoneves May 4, 2026
30897a6
Refactor to remove state struct and associated unpacking at call sites
simoneves May 11, 2026
48575c1
Factor out remaining case body code to separate functions for clarity…
simoneves May 11, 2026
57c6593
Add comment about held state, per @mattgara, and format
simoneves May 11, 2026
37d434d
Rewrite agg kernels to use cub::DeviceFor, per @devavret
simoneves May 11, 2026
0f00a12
Format
simoneves May 11, 2026
62a5781
Refactor to keep CUDA/cub/Thrust code in .cu file and move everything…
simoneves May 12, 2026
de780a9
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 12, 2026
1338cf8
Refactor DecimalExpressionKernels in same manner as DecimalAggregatio…
simoneves May 12, 2026
b21a93e
Simplify some overcomplex and inconsistent generated code
simoneves May 12, 2026
999fde7
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 13, 2026
b1975cb
Fix after merge update
simoneves May 13, 2026
e49f979
Format. One day I'll remember do this before committing...
simoneves May 13, 2026
c29f7d0
Comment DecimalSumState fields, per @shrshi, and simplify
simoneves May 14, 2026
c2e1f90
Use VELOX_CHECK_EQ, per @karthikeyann
simoneves May 14, 2026
0a84963
outputType/resultType resolution, per @karthikeyann
simoneves May 14, 2026
486a0a6
Skip adding null mask if no nulls, add tests to validate, per @karthi…
simoneves May 14, 2026
67f3727
Add comments about scale in VARBINARY case, per @karthikeyann
simoneves May 14, 2026
25f06d0
Remove test that relies on unimplemented functionality, per @karthike…
simoneves May 14, 2026
1a421ec
Add failure values, per @majetideepak
simoneves May 14, 2026
c3eb088
Format
simoneves May 14, 2026
ad69d8e
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 14, 2026
cc45644
More error values, per @majetideepak
simoneves May 14, 2026
354c61d
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 18, 2026
30861b1
Alternative refactor, remove MR plumbing and pass get_output_mr() or …
simoneves May 20, 2026
2447a33
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 22, 2026
f51caff
Merge branch 'main' into simoneves/decimal_pr3
simoneves May 27, 2026
85c762e
Fix UB signed overflow issue
mattgara Jun 1, 2026
8252ef2
Ensure 64bit decimals are upcast to 128bit before aggregation to prevent
mattgara Jun 1, 2026
dbae047
Add comment clarifying `overflow` member
mattgara Jun 2, 2026
04a3541
Address resource accounting (i.e. use of `temp_mr` where appropriate.)
mattgara Jun 2, 2026
35a7ce2
Remove dead/redundant code
mattgara Jun 2, 2026
f76cb38
Groupby and reduce refactors and simplification
mattgara Jun 2, 2026
4e186bc
Address various PR comments on kernels
mattgara Jun 3, 2026
7618f6f
Address renaming comments on PR
mattgara Jun 3, 2026
160f46a
Pre-commit
mattgara Jun 3, 2026
82d29d2
Address additional comments
mattgara Jun 4, 2026
b3ad1f8
Address build config PR comments
mattgara Jun 4, 2026
f1d0621
Pre-commit
mattgara Jun 4, 2026
228ee4d
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 4, 2026
cc1fb10
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 4, 2026
8b17c49
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 4, 2026
c83a81f
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 9, 2026
f86c192
Remove off-style comments
simoneves Jun 9, 2026
bdafae0
Keep alphabetical order
simoneves Jun 9, 2026
e1e1823
Use cudf::size_type and cudf::type_to_name
simoneves Jun 9, 2026
052cfc2
Fix Doxygen function comments style
simoneves Jun 9, 2026
6aa5925
Move CUDA C++ standard setting (are we sure?)
simoneves Jun 9, 2026
5680f44
Use cuda::counting_iterator
simoneves Jun 10, 2026
e1a8cbf
Revert expression and CMake changes, moving to a Part 4 tidy-up
simoneves Jun 11, 2026
afbda57
Final get_output_mr() plumbing, per @bdice and @karthikeyann
simoneves Jun 11, 2026
baade9f
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 11, 2026
c43850f
Format
simoneves Jun 11, 2026
01685c0
Refactor to use cudf::type_dispatcher
simoneves Jun 10, 2026
3a0285f
Combine pairs of functions, per @bdice
simoneves Jun 12, 2026
f931ad6
Use C++20 requires/concepts instead of enable_if/traits, per @bdice
simoneves Jun 12, 2026
8ef386f
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 12, 2026
06011d3
Refactor the aggregation kernel/functor structure to match that of th…
simoneves Jun 12, 2026
9d1e8aa
Format
simoneves Jun 12, 2026
1fffd03
Use get_temp_mr() in castCountColumnToInt64 and change function comme…
simoneves Jun 16, 2026
35bd662
Use cudf::type_to_name, per @karthikeyann
simoneves Jun 16, 2026
80c1cce
Fix comparison, per @karthikeyann
simoneves Jun 16, 2026
db9c802
Add decimal sum state payload size check, per @karthikeyann
simoneves Jun 16, 2026
8aa0e2b
Combine conditional branches, per @karthikeyann
simoneves Jun 16, 2026
0c85fa1
Format
simoneves Jun 16, 2026
2e3e541
Merge branch 'main' into simoneves/decimal_pr3
simoneves Jun 16, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
109 changes: 106 additions & 3 deletions velox/experimental/cudf/exec/AggregationRegistry.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,37 @@ void registerCommonAggregationFunctions(
.argumentType("double")
.build()};

auto decimalSumSingle = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.integerVariable("a_precision")
.integerVariable("a_scale")
.returnType("decimal(38, a_scale)")
.argumentType("decimal(a_precision, a_scale)")
.build()};
auto decimalSumPartial = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.integerVariable("a_precision")
.integerVariable("a_scale")
.returnType("varbinary")
.argumentType("decimal(a_precision, a_scale)")
.build()};
auto decimalSumFinal = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.integerVariable("a_scale")
.returnType("decimal(38, a_scale)")
.argumentType("varbinary")
.build()};
auto decimalSumIntermediate =
std::vector<exec::FunctionSignaturePtr>{FunctionSignatureBuilder()
.returnType("varbinary")
.argumentType("varbinary")
.build()};

sumSingleSignatures.insert(
sumSingleSignatures.end(),
decimalSumSingle.begin(),
decimalSumSingle.end());

registerAggregationFunctionForStep(
registry,
prefix + "sum",
Expand Down Expand Up @@ -109,6 +140,12 @@ void registerCommonAggregationFunctions(
.returnType("double")
.argumentType("double")
.build()};

sumPartialSignatures.insert(
sumPartialSignatures.end(),
decimalSumPartial.begin(),
decimalSumPartial.end());

registerAggregationFunctionForStep(
registry,
prefix + "sum",
Expand All @@ -125,16 +162,27 @@ void registerCommonAggregationFunctions(
.argumentType("double")
.build()};

auto sumFinalSignatures = sumFinalIntermediateSignatures;
sumFinalSignatures.insert(
sumFinalSignatures.end(), decimalSumFinal.begin(), decimalSumFinal.end());

registerAggregationFunctionForStep(
registry,
prefix + "sum",
core::AggregationNode::Step::kFinal,
sumFinalIntermediateSignatures);
sumFinalSignatures);

auto sumIntermediateSignatures = sumFinalIntermediateSignatures;
sumIntermediateSignatures.insert(
sumIntermediateSignatures.end(),
decimalSumIntermediate.begin(),
decimalSumIntermediate.end());

registerAggregationFunctionForStep(
registry,
prefix + "sum",
core::AggregationNode::Step::kIntermediate,
sumFinalIntermediateSignatures);
sumIntermediateSignatures);

auto countSinglePartialSignatures = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
Expand Down Expand Up @@ -234,6 +282,12 @@ void registerCommonAggregationFunctions(
FunctionSignatureBuilder()
.returnType("varchar")
.argumentType("varchar")
.build(),
FunctionSignatureBuilder()
.integerVariable("p")
.integerVariable("s")
.returnType("decimal(p,s)")
.argumentType("decimal(p,s)")
.build()};

registerAggregationFunctionForStep(
Expand Down Expand Up @@ -296,6 +350,38 @@ void registerCommonAggregationFunctions(
.argumentType("double")
.build()};

auto decimalAvgSingle = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.integerVariable("a_precision")
.integerVariable("a_scale")
.returnType("decimal(a_precision, a_scale)")
.argumentType("decimal(a_precision, a_scale)")
.build()};
auto decimalAvgPartial = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.integerVariable("a_precision")
.integerVariable("a_scale")
.returnType("varbinary")
.argumentType("decimal(a_precision, a_scale)")
.build()};
auto decimalAvgFinal = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.integerVariable("a_precision")
.integerVariable("a_scale")
.returnType("decimal(a_precision, a_scale)")
.argumentType("varbinary")
.build()};
auto decimalAvgIntermediate =
std::vector<exec::FunctionSignaturePtr>{FunctionSignatureBuilder()
.returnType("varbinary")
.argumentType("varbinary")
.build()};

avgSingleSignatures.insert(
avgSingleSignatures.end(),
decimalAvgSingle.begin(),
decimalAvgSingle.end());

registerAggregationFunctionForStep(
registry,
prefix + "avg",
Expand Down Expand Up @@ -323,17 +409,28 @@ void registerCommonAggregationFunctions(
.returnType("row(double,bigint)")
.argumentType("double")
.build()};

avgPartialSignatures.insert(
avgPartialSignatures.end(),
decimalAvgPartial.begin(),
decimalAvgPartial.end());

registerAggregationFunctionForStep(
registry,
prefix + "avg",
core::AggregationNode::Step::kPartial,
avgPartialSignatures);

auto avgFinalSignatures = std::vector<exec::FunctionSignaturePtr>{
auto avgFinalIntermediateSignatures = std::vector<exec::FunctionSignaturePtr>{
FunctionSignatureBuilder()
.returnType("double")
.argumentType("row(double,bigint)")
.build()};

auto avgFinalSignatures = avgFinalIntermediateSignatures;
avgFinalSignatures.insert(
avgFinalSignatures.end(), decimalAvgFinal.begin(), decimalAvgFinal.end());

registerAggregationFunctionForStep(
registry,
prefix + "avg",
Expand All @@ -345,6 +442,12 @@ void registerCommonAggregationFunctions(
.returnType("row(double,bigint)")
.argumentType("row(double,bigint)")
.build()};

avgIntermediateSignatures.insert(
avgIntermediateSignatures.end(),
decimalAvgIntermediate.begin(),
decimalAvgIntermediate.end());

registerAggregationFunctionForStep(
registry,
prefix + "avg",
Expand Down
3 changes: 3 additions & 0 deletions velox/experimental/cudf/exec/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,9 @@ add_library(
CudfReduce.cpp
CudfTopN.cpp
DebugUtil.cpp
DecimalAggregationDevice.cu
DecimalAggregationHostOps.cpp
DecimalAggregationState.cpp
GpuResources.cpp
OperatorAdapters.cpp
PrestoAggregateFunctions.cpp
Expand Down
5 changes: 4 additions & 1 deletion velox/experimental/cudf/exec/CudfAggregation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -136,6 +136,8 @@ std::vector<ResolvedAggregateInfo> resolveAggregateInfos(
const auto resultType = exec::isPartialOutput(companionStep)
? exec::resolveIntermediateType(originalName, aggregate.rawInputTypes)
: outputType->childAt(numKeys + i);
const auto isDecimalAggregate = aggregate.rawInputTypes.size() == 1 &&
aggregate.rawInputTypes[0]->isDecimal();

params.emplace_back(
companionStep,
Expand All @@ -145,7 +147,8 @@ std::vector<ResolvedAggregateInfo> resolveAggregateInfos(
resultType,
isCountFunctionName(aggregate.call->name())
? std::make_optional(getCountInputKind(aggregate, constants[i]))
: std::nullopt);
: std::nullopt,
isDecimalAggregate);
}
return params;
}
Expand Down
4 changes: 4 additions & 0 deletions velox/experimental/cudf/exec/CudfAggregation.h
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,10 @@ struct ResolvedAggregateInfo {
VectorPtr constant;
TypePtr resultType;
std::optional<CountInputKind> countInputKind;
// True if the aggregate was declared on a decimal raw input in the plan.
// Routing keys off the function family, not the physical batch type (which is
// VARBINARY/STRING on intermediate and final steps).
bool isDecimalAggregate;
};

// Parse aggregate inputs from the aggregation node and resolve companion steps,
Expand Down
Loading
Loading