feat(cudf): GPU Decimal (Part 3) by simoneves · Pull Request #16751 · facebookincubator/velox

simoneves · 2026-03-12T19:23:25Z

This is the third part of the GPU Decimal implementation. It adds support for SUM and AVG aggregates.

netlify · 2026-03-12T19:23:32Z

✅ Deploy Preview for meta-velox ready!

Name	Link
🔨 Latest commit	`2e3e541`
🔍 Latest deploy log	https://app.netlify.com/projects/meta-velox/deploys/6a31a91d9c1a410008e9ac22
😎 Deploy Preview	https://deploy-preview-16751--meta-velox.netlify.app
📱 Preview on mobile	Toggle QR Code... Use your smartphone camera to open QR code link.

To edit notification comments on pull requests, go to your Netlify project configuration.

simoneves · 2026-03-16T20:12:15Z

Needs rebasing on top of Part 2. Conflict resolution is WIP. Need to discuss with @devavret

Cherry-picked from f834458 (PR facebookincubator#16751 Part 3). Adds CUDA kernels for decimal SUM/AVG aggregation, decimal-aware type handling in hash aggregation with veloxToCudfDataType. Merged with ferd-dev's GpuGuard and pinned memory optimizations. Made-with: Cursor

simoneves · 2026-03-20T20:02:33Z

+  StateValidPredicate pred{*sumDeviceView, *countDeviceView};
+  auto begin = thrust::make_counting_iterator<cudf::size_type>(0);
+  auto end = begin + numRows;
+  return cudf::detail::valid_if(begin, end, pred, stream, mr);


Generally cudf::detail functions should not be used directly, as they may change at any time.

David Wendt's suggestion was (and I quote) "do a transform and build a device-vector of the predicate results and then call the public cudf::bools_to_mask"

This should be addressed, see @shrshi comment on the same thing below.

* feat(cudf): improve Arrow interop for decimals and varbinary Cherry-picked from 00889a2 (PR facebookincubator#16612 Part 1). Adds veloxToCudfDataType() preserving DECIMAL scale, extends Arrow bridge options for varbinary and decimal type width. Merged with ferd-dev's pinnedToArrowHost optimization. Made-with: Cursor * feat(cudf): add decimal expression support and tests Cherry-picked from 67e400d (PR facebookincubator#16750 Part 2). Adds CUDA kernels for decimal binary/unary/comparison ops, decimal type utilities, and extends expression evaluator with decimal-aware binary operators. Merged with ferd-dev's date_add and row_constructor functions. Made-with: Cursor * style: format decimal code (cherry-pick f5bc585) Made-with: Cursor * fix: fix after rebase for decimal tests (cherry-pick 564ae5c) Made-with: Cursor * feat(cudf): enable GPU ceil() (cherry-pick 2ac50f2) Made-with: Cursor * feat(cudf): add decimal sum/avg kernels and aggregation support Cherry-picked from f834458 (PR facebookincubator#16751 Part 3). Adds CUDA kernels for decimal SUM/AVG aggregation, decimal-aware type handling in hash aggregation with veloxToCudfDataType. Merged with ferd-dev's GpuGuard and pinned memory optimizations. Made-with: Cursor * feat(cudf): add TopNRowNumber GPU operator, xxhash64, and code quality fixes - Add CudfTopNRowNumber operator with GPU-based sort, groupby scan ranking, and TopNRowNumber kernel implementation - Add xxhash64_with_seed function for GPU hash computation - Unify memory resource usage to cudf::get_current_device_resource_ref() - Improve expression evaluator with priority-sorted candidate fallback and diagnostic logging when all evaluators fail - Fix indentation in CudfHashJoinProbe::initialize() - Remove redundant member variable writes in CudfTopNRowNumber constructor - Correct decimal type handling (veloxToCudfDataType) in NestedLoopJoin - Enhance AstExpressionUtils with VARCHAR-to-DATE constant folding and disambiguated field name resolution for join filters Made-with: Cursor * fix(cudf): fix xxhash64 tests, refactor InFunction, add varbinary-as-string export - Fix xxhash64_with_seed test expected values and add single-string test case (verified passing on GPU with cudf::hashing::xxhash_64) - Refactor IN expression from inline eval to proper InFunction CudfFunction class with type-cast support and null handling - Add exportVarbinaryAsString ArrowOption for cudf interop (cudf doesn't support Arrow binary type, export as utf-8 string instead) Made-with: Cursor * fix(cudf): handle field name mismatch in n{nodeId}_{colIdx} resolution (query31) When Gluten's Substrait plan uses different node IDs for the same column (e.g. expression references n11_1 but runtime schema has n10_1), the cuDF expression evaluator failed with "Field not found" and crashed the task. Two-part fix: 1. OperatorAdapters: add allFieldsResolvable() check in FilterProjectAdapter::canRunOnGPU() to detect unresolvable field references before creating CudfFilterProject. Falls back to CPU gracefully instead of crashing. 2. AstExpressionUtils: add colIdx-based fallback in pushExprToTree() and addPrecomputeInstruction() that matches n{X}_{Y} fields by their _{Y} suffix when exact name lookup fails. Only used when exactly one schema column shares the same colIdx suffix. Includes unit tests for both the successful fallback case and the no-match case. Made-with: Cursor * fix(cudf): add zero-size guards to prevent cudaErrorInvalidValue (query4/6/11/15/19/28/30/74/80) Add defensive checks for zero-row inputs that trigger invalid CUDA memory allocations through RMM's device_buffer: CudfHashJoinProbe: - Early return in getOutput() when probe table has 0 rows - Guard rightMatchedFlags_ init for empty build tables (n=0) - Guard leftSemiProjectJoin for leftNumRows=0 - Guard cudf::sequence calls in rightJoin/fullJoin for n=0 - Skip precomputeSubexpressions for empty right tables CudfFilterProject: - Skip project() when filter removes all rows (0-row result) Made-with: Cursor * fix(cudf): prevent Jitify error for STRING ops in AST evaluator (query83) cudf::compute_column / Jitify cannot handle variable-width types (STRING) in AST operations. Added two layers of protection: 1. isOpAndInputsSupported rejects non-fixed-width inputs, preventing the AST/JIT evaluator from claiming support for STRING operations at the top level. 2. pushExprToTree routes binary ops, IN, and BETWEEN with variable-width inputs to precompute instructions via FunctionExpression, handling the nested case (e.g., STRING ops inside AND expressions). Made-with: Cursor * fix(cudf): implement DATE→VARCHAR cast and might_contain bloom filter bypass - CastFunction: handle DATE to VARCHAR using cudf::strings::from_timestamps with "%Y-%m-%d" format instead of unsupported cudf::cast for this type pair - MightContainFunction: return all-true boolean column as safe placeholder for bloom filter probabilistic check (false positives are acceptable by design) - Update FunctionExpression::canEvaluate and create to route both expressions Fixes query83 (DATE cast), query11/15/74/4/80 (might_contain). Made-with: Cursor * fix(cudf): implement isnull/isnotnull expression for GPU evaluation (query44/76) IsNullFunction checks input column null mask using cudf::replace_nulls and cudf::unary_operation(NOT). Returns all-false/true constant when column has no nulls. Handles both isnull (negate=false) and isnotnull (negate=true) variants. Made-with: Cursor * fix(cudf): add field resolution check to CudfHashJoinBaseAdapter (query31 revisit) The previous fix (fe2bf9659) only added allFieldsResolvable check to FilterProjectAdapter. The actual crash was in CudfHashJoinProbe where precomputeSubexpressions evaluated a FunctionExpression with a per-side schema missing the n11_1 field. Add allFieldsResolvable validation to CudfHashJoinBaseAdapter::canRunOnGPU which checks the join filter's field references against the combined left+right schema. If any field is unresolvable, fall back to CPU. Made-with: Cursor * fix(cudf): add equalto/notequalto aliases for comparison operators Spark uses 'equalto' and 'notequalto' as expression names but cudf only registered 'equal'/'eq' and 'notequal'/'neq'. Add the missing aliases so these expressions are routed to GPU binary operators. Fixes query1, query71, query72 and other queries using equalto in filter pushdown expressions. Made-with: Cursor * fix(cudf): add kIntermediate step for mean aggregation in doReduce (query28) MeanAggregator::doReduce was missing the kIntermediate case, causing "Unsupported aggregation step for mean" when the query plan has a partial → intermediate → final aggregation pipeline. The intermediate step receives a struct(sum, count) from the previous stage, re-reduces (sum-of-sums, sum-of-counts), and outputs a new struct for the next stage. Made-with: Cursor * fix(cudf): skip STRING subfield filters from AST/Jitify in TableScan (query33/43/56/60/61) STRING/Bytes filter types (kBytesValues, kNegatedBytesValues, kBytesRange) cannot be evaluated by cudf AST or Jitify, causing "Jitify fatal error: Deserialization failed" / "Uninitialized" crashes in TableScan. Mark these filter kinds as VELOX_NYI in createAstFromSubfieldFilter so they are gracefully skipped. createAstFromSubfieldFilters now catches per-filter exceptions and continues with remaining (non-STRING) filters. If all filters are unsupported, CudfHiveDataSource catches the exception and leaves subfieldFilterExpr_ null, relying on the downstream FilterProject to apply the predicates on CPU. Made-with: Cursor * fix(cudf): prevent SIGSEGV in hash join probe for skewed partitions (query68) Convert GPU crashes (SIGSEGV/exit code 139) into clean VeloxRuntimeError exceptions that enable CPU fallback: - Wrap inner_join() call in try-catch to convert CUDA/RMM allocation failures into diagnostic Velox errors with probe/build row counts - Add VELOX_CHECK_LE on join output size vs cudf::size_type max to detect oversized join results before gather operations - Wrap gather/filter phase in try-catch for the same reason - Guard rightTables[0] access with empty-check in isBlocked() - Guard rightTables[0] access with empty-check in rightSemiFilterJoin() Root cause: partition 2 of stage 25 has extreme data skew, producing a massive join output that exhausts GPU memory or overflows size_type. In v11 this failed cleanly (cudaErrorInvalidValue); in v12 the zero-size guards let execution proceed deeper, converting the failure into an unrecoverable SIGSEGV. Made-with: Cursor (Agent D) * fix(cudf): lazy table_view construction to prevent column size mismatch (query18) ASTExpression::eval and JitExpression::eval eagerly constructed a cudf::table_view from input columns + precomputed columns, even when the root AST node was a column_reference that never used the table_view. Precomputed columns from nested evaluators may have different row counts (e.g., after filtering), triggering a spurious Column size mismatch error. Fix: defer table_view construction into the else branch where compute_column actually needs it. Also add output column size validation in CudfFilterProject::getOutput as a safety net. Made-with: Cursor * fix(cudf): graceful AST filter fallback + scalar precompute expansion CudfHashJoinProbe: wrap AST tree creation in try-catch so unsupported filter expressions (e.g. unresolvable fields) disable the AST filter instead of crashing the operator. AstExpressionUtils: expand 0/1-row precomputed scalar sub-expressions to match input row count, preventing "Column size mismatch" when constant expressions produce fewer rows than the input table. Made-with: Cursor * fix(cudf): validate all field references in precompute instructions (query4/11) When pushExprToTree creates a precompute instruction for a sub-expression, findExpressionSide may select a side whose schema does not contain all the expression's field references (e.g. n8_1 missing from probe side n7_*). This creates a FunctionExpression with an incomplete schema that crashes at eval time in precomputeSubexpressions (Field not found). Fix: In both precomputeVarWidthOp and the general precompute path, validate that ALL fields in the expression exist in the chosen side's schema. If not, return nullptr (varwidth) or throw VELOX_FAIL (general), causing createAstTree to fail at init time where the Leader's try-catch sets useAstFilter_=false and falls back to the combined-schema filterEvaluator_. Made-with: Cursor * fix(cudf): OOM-resilient probe splitting for hash join (query72/73/76) When a hash join OOM occurs (e.g. query72 inner_join needing 6.5GB for 723K×1.9M row cardinality explosion), the probe table is split in half and each half is retried independently. This iterates until the join fits in GPU memory or hits kMinSplitRows (1024). Changes: - innerJoin: re-throw std::bad_alloc so it propagates to the retry loop instead of being converted to VeloxRuntimeError - getOutput(): wrap the join dispatch in an iterative probe-splitting retry loop; on std::bad_alloc, split via cudf::split and re-enqueue - Only split for join types where it is semantics-preserving: Inner, Left, LeftSemiFilter, LeftSemiProject, Anti - Also applies Leader's precompute try-catch hardening to leftJoin (mirrors the existing innerJoin pattern) Made-with: Cursor * fix(cudf): add runtime Jitify error catch in CudfHiveDataSource compute_column Defense-in-depth for Issue 3 (Jitify Deserialization in TableScan). When cudf::compute_column throws a Jitify error (or any exception) during subfield filter evaluation in the experimental reader path, catch it and return the unfiltered data instead of crashing the task. The downstream FilterProject will handle the filtering on CPU. This complements b7ed2a6dc which prevents STRING filters from reaching Jitify at init time. Made-with: Cursor * fix(cudf): prevent cudf::cast on STRING columns in InFunction (query83) InFunction::shouldSkipCast incorrectly skipped DATE→STRING casts because cudf::is_supported_cast returns false for that pair, even though CastFunction handles it via cudf::strings::from_timestamps. This caused buildHaystackColumn to attempt cudf::cast(STRING, TIMESTAMP_DAYS) which fails at cast_ops.cu:411 ("Column type must be numeric or chrono or decimal32/64/128"). Fix: 1) Return false from shouldSkipCast for DATE→STRING since CastFunction supports it. 2) Add safety net in buildHaystackColumn: when haystack is STRING and target is a timestamp type, use cudf::strings::to_timestamps instead of cudf::cast. Made-with: Cursor * fix(cudf): add SIGSEGV signal handler and CUDA error synchronization at operator boundaries Agent D's try-catch fix (b7fcd02d4) didn't work because SIGSEGV is a signal, not a C++ exception. v14 still crashes on query18/29/45/68. Root cause analysis of v14 SIGSEGV crashes: - query68: single partition crash (~86s), data skew pattern unchanged - query29: v12 PASS(8s) → v14 FAIL, ALL tasks crash simultaneously - query45: v12 PASS(6s) → v14 FAIL, ALL executors crash within 3s - query18: ALL 28 partitions of stage 33 crash on ALL executors - All crashes show ~86s delay, consistent with asynchronous CUDA error propagation cascading into SIGSEGV Two-pronged fix: 1. Fatal signal handler (Utilities.cpp, ToCudf.cpp): Install SIGSEGV/SIGABRT/SIGBUS handler in registerCudf() that captures native backtrace via backtrace_symbols_fd() before re-raising to the JVM's handler. This gives us the exact crash location in executor stderr for ALL future SIGSEGV queries. 2. CUDA stream sync + error checking (operator boundaries): Add checkCudaOperationError() that synchronizes the CUDA stream and checks cudaPeekAtLastError() before each operator's main GPU work. This converts asynchronous CUDA errors (illegal address, OOM) into clean VeloxRuntimeError before they cascade into SIGSEGV. Added to: CudfFilterProject, CudfHashJoinBuild, CudfHashJoinProbe, CudfHashAggregation, CudfToVelox. Also fixes: pessimizing-move warning in CudfHiveDataSource.cpp. Made-with: Cursor * Revert "fix(cudf): add SIGSEGV signal handler and CUDA error synchronization at operator boundaries" This reverts commit d92407f679707b38746b068daa3445bfd29ac26c. * fix(cudf): skip subfield filter AST when any column is STRING/VARBINARY The normal parquet reader path (splitReader_->read_chunk()) applies the subfield filter AST internally via cudf::compute_column, which uses Jitify for string comparisons. When Jitify fails, the error is thrown inside the reader and cannot be caught externally. Fix: before building the subfield filter AST, check column types of all subfield filters. If any column is VARCHAR or VARBINARY, skip the entire GPU subfield filtering (set subfieldFilterExpr_ = nullptr). The downstream CPU FilterProject handles these predicates instead. This prevents Jitify errors in query8/33/43/56/60/61/91 by ensuring string-typed filters never reach the parquet reader's AST evaluator. Made-with: Cursor * fix(cudf): clear CUDA error state after OOM in hash join probe (Q72) After a CUDA OOM in CudfHashJoinProbe, the cuda_async_view_memory_resource pool enters an error state where ALL subsequent allocations fail with cudaErrorIllegalAddress — even 1-byte allocations. This cascade kills every concurrent task on the same GPU. Fix: call stream.synchronize_no_throw() + cudaGetLastError() in the std::bad_alloc catch block to clear the sticky CUDA error before the split-and-retry logic. Without this, probe splitting was never effective because retried allocations immediately failed on the corrupted pool. Also includes Agent B's pending improvements: pendingJoinOutputs_ streaming (avoids concatenation peak memory), empty table filtering. Made-with: Cursor * Revert "fix(cudf): clear CUDA error state after OOM in hash join probe (Q72)" This reverts commit a95d6dd. * fix(cudf): batched join output + CUDA error recovery for OOM (query72) Two root causes identified for Q72 CUDA OOM surviving probe splitting: 1. After std::bad_alloc from RMM, the CUDA async pool leaves a sticky error on the device. The subsequent cudf::split (which launches a null-count kernel) fails with cudaErrorIllegalAddress, escaping the catch handler entirely. Fix: call stream.synchronize_no_throw() and cudaGetLastError() to clear the sticky error before splitting. 2. Even when splitting succeeds, concatenateTables at the end allocates a contiguous buffer for ALL split results (~667MB for 112M rows), doubling peak GPU memory. Fix: adopt Spark RAPIDS JoinGatherer pattern — store multiple results in pendingJoinOutputs_ and return one per getOutput() call, avoiding the concatenation allocation. Also fix isFinished() to not report done while pendingJoinOutputs_ still has tables to emit. Made-with: Cursor * fix(cudf): reject AST expressions with non-matching operand types cuDF's AST expression parser requires all operands of a binary operation to have identical cudf::data_type (including decimal scale). When operand types differ (e.g., DECIMAL(15,2) vs DECIMAL(7,4)), cudf::compute_column throws "non-matching operand types" at runtime, crashing queries Q61/Q83. Three-layer fix: 1. isOpAndInputsSupported: reject binary ops with mismatched types early, preventing ASTExpression from being selected (falls to FunctionExpression) 2. pushExprToTree: for binary/between ops with type mismatch, try precompute via FunctionExpression; if that fails, throw to trigger full fallback 3. ASTExpression::eval: defensive catch around cudf::compute_column converts cuDF exceptions to VeloxException for clearer diagnostics Made-with: Cursor * fix(cudf): remove duplicate divide/greaterthan registration that drops decimal signatures (Q18) The `divide` and `greaterthan` functions were registered twice: first with only (double,double) signatures, then via registerBinaryOp/ registerComparisonOp with both double AND decimal signatures. Since overwrite=false, the second registration was silently dropped, meaning decimal divide and decimal greater-than never matched in canEvaluate. Also wrap canBeEvaluatedByCudf expression compilation in try-catch so that VeloxException from decimal type resolution (e.g. "Variable a_precision is not defined") gracefully falls back to CPU instead of potentially crashing the task. Made-with: Cursor * fix(cudf): disable JIT by default, fix JIT table view and SubfieldFilter isNull param - Disable jitExpressionEnabled by default (less mature than AST/Function) - Fix JitExpression to construct astInputTableView from inputColumnViews + precomputedColumns before passing to compute_column_jit - Add missing isNull parameter to makeScalarAndLiteral calls in SubfieldFiltersToAst.cpp Made-with: Cursor

github-actions · 2026-03-30T16:36:23Z

Selective Build Plan

Linux release with adapters is running a full build (changes touch velox/experimental/ or velox/external/). See the CI workflows README for what this means.

Selective build plan

…e expression code in Part 4

Changes discussed offline

devavret · 2026-06-16T17:33:22Z

+    std::unique_ptr<cudf::column> count,
+    rmm::cuda_stream_view stream,
+    rmm::device_async_resource_ref mr) {
+  count = castCountColumnToInt64(std::move(count), stream, mr);


I think you'd need to use temp mr here but it's not important. We'll fix all usages when we start to use the distinction.

devavret · 2026-06-16T17:36:33Z

+    rmm::cuda_stream_view stream,
+    rmm::device_async_resource_ref mr);
+
+// Finalizes AVG from intermediate SUM state: divides each sum by its count


nit: all function/class docs should have triple slash docs ///

karthikeyann · 2026-06-16T16:07:10Z

+          countIdx_,
+          decodedSum_,
+          decodedCount_);
+    } else if (step == core::AggregationNode::Step::kFinal) {


nitpick: kIntermediate and kFinal run identical code here (same call, same args) — would folding them into one branch read better?

karthikeyann · 2026-06-16T16:07:10Z

+  VELOX_CHECK(
+      countCol.type().id() == cudf::type_id::INT64,
+      "Decimal sum state requires INT64 count column (type is {})",
+      static_cast<int>(countCol.type().id()));


nitpick: the other type checks in this file print the type via cudf::type_to_name(...); this one uses static_cast<int>(...).

Suggested change

static_cast<int>(countCol.type().id()));

cudf::type_to_name(countCol.type()));

karthikeyann · 2026-06-16T16:07:10Z

+        return reduceIntermediateDecimalFromSerializedColumn(
+            inputCol, outputType, stream, mr);
+      case core::AggregationNode::Step::kFinal:
+        VELOX_CHECK(outputType == resultType, "outputType/resultType mismatch");


nitpick: outputType == resultType compares the shared_ptrs, not the types — it passes only because childAt returns the same stored pointer. *outputType == *resultType says what you mean.

Suggested change

VELOX_CHECK(outputType == resultType, "outputType/resultType mismatch");

VELOX_CHECK(*outputType == *resultType, "outputType/resultType mismatch");

karthikeyann · 2026-06-16T16:07:11Z

+  cudf::strings_column_view strings(stateCol);
+
+  auto offsetsView = strings.offsets();
+  auto charsPtr = reinterpret_cast<const uint8_t*>(strings.chars_begin(stream));


nitpick: would a VELOX_CHECK that the payload is exactly numRows * kDecimalSumStateSize bytes be worth adding before this? It'd catch a malformed/short state column before unpackDecimalSumState reads past the buffer.

@devavret

…nts to Doxygen style, per @devavret

karthikeyann

Approving the PR. Looks good.
Thanks @simoneves for proactively addressing all review comments. Thanks to all reviewers @devavret @bdice @shrshi @mattgara @majetideepak for continuous improvement of this major feature.

meta-codesync · 2026-06-16T23:05:43Z

@kagamiori has imported this pull request. If you are a Meta employee, you can view this in D108813872.

meta-codesync · 2026-06-19T04:37:19Z

@kagamiori merged this pull request in 5ded8d7.

simoneves requested review from bdice, devavret, karthikeyann and mhaseeb123 as code owners March 12, 2026 19:23

meta-cla Bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Mar 12, 2026

simoneves marked this pull request as draft March 12, 2026 19:23

simoneves mentioned this pull request Mar 18, 2026

feat(cudf): Support decimal types in CUDF GPU operators (REPLACED) #15825

Closed

simoneves force-pushed the simoneves/decimal_pr3 branch 2 times, most recently from f1befb7 to 5368b21 Compare March 20, 2026 19:26

simoneves commented Mar 20, 2026

View reviewed changes

karthikeyann added the cudf cudf related - GPU acceleration label Mar 23, 2026

GregoryKimball mentioned this pull request Mar 23, 2026

[cuDF] Expand GPU operator support for Presto TPC-DS #15772

Open

25 tasks

simoneves force-pushed the simoneves/decimal_pr3 branch 2 times, most recently from 9d28c71 to fdf00b6 Compare March 25, 2026 20:51

patdevinwilson mentioned this pull request Mar 27, 2026

feat(cudf): Add GPU NestedLoopJoin (cross join) implementation #16942

Closed

2 tasks

simoneves force-pushed the simoneves/decimal_pr3 branch 2 times, most recently from abbf158 to 21f210f Compare March 30, 2026 16:31

simoneves force-pushed the simoneves/decimal_pr3 branch 7 times, most recently from d81f9cf to 4f363fd Compare April 2, 2026 18:35

Merge branch 'main' into simoneves/decimal_pr3

baade9f

simoneves requested a review from bdice June 11, 2026 22:21

Format

c43850f

bdice changed the title ~~feat(cudf): GPU Decimal (Part 3 of 3)~~ feat(cudf): GPU Decimal (Part 3) Jun 12, 2026

simoneves added 6 commits June 12, 2026 13:57

Refactor to use cudf::type_dispatcher

01685c0

Combine pairs of functions, per @bdice

3a0285f

Use C++20 requires/concepts instead of enable_if/traits, per @bdice

f931ad6

Merge branch 'main' into simoneves/decimal_pr3

8ef386f

Refactor the aggregation kernel/functor structure to match that of th…

06011d3

…e expression code in Part 4

Format

9d1e8aa

mhaseeb123 removed their request for review June 13, 2026 01:35

devavret approved these changes Jun 16, 2026

View reviewed changes

karthikeyann reviewed Jun 16, 2026

View reviewed changes

simoneves added 7 commits June 16, 2026 11:50

Use get_temp_mr() in castCountColumnToInt64 and change function comme…

1fffd03

…nts to Doxygen style, per @devavret

Use cudf::type_to_name, per @karthikeyann

35bd662

Fix comparison, per @karthikeyann

80c1cce

Add decimal sum state payload size check, per @karthikeyann

db9c802

Combine conditional branches, per @karthikeyann

8aa0e2b

Format

0c85fa1

Merge branch 'main' into simoneves/decimal_pr3

2e3e541

simoneves requested a review from karthikeyann June 16, 2026 19:51

karthikeyann approved these changes Jun 16, 2026

View reviewed changes

karthikeyann added the ready-to-merge PR that have been reviewed and are ready for merging. PRs with this tag notify the Velox Meta oncall label Jun 16, 2026

meta-codesync Bot closed this in 5ded8d7 Jun 19, 2026

meta-codesync Bot added the Merged label Jun 19, 2026

GregoryKimball removed this from libcudf Jun 22, 2026

	static_cast<int>(countCol.type().id()));
	cudf::type_to_name(countCol.type()));

	VELOX_CHECK(outputType == resultType, "outputType/resultType mismatch");
	VELOX_CHECK(outputType == resultType, "outputType/resultType mismatch");

Uh oh!

Conversation

simoneves commented Mar 12, 2026 • edited by bdice Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

netlify Bot commented Mar 12, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ Deploy Preview for meta-velox ready!

Uh oh!

simoneves commented Mar 16, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Mar 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Selective Build Plan

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

karthikeyann left a comment

Choose a reason for hiding this comment

Uh oh!

meta-codesync Bot commented Jun 16, 2026

Uh oh!

meta-codesync Bot commented Jun 19, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

8 participants

simoneves commented Mar 12, 2026 •

edited by bdice

Loading

netlify Bot commented Mar 12, 2026 •

edited

Loading

github-actions Bot commented Mar 30, 2026 •

edited

Loading