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

simoneves · 2026-06-11T17:48:53Z

Decimal expression code refactoring and tidying moved from Part 3.

@bdice's comments moved from #16751.

netlify · 2026-06-11T17:49:46Z

✅ Deploy Preview for meta-velox canceled.

Name	Link
🔨 Latest commit	`d9ec8ac`
🔍 Latest deploy log	https://app.netlify.com/projects/meta-velox/deploys/6a2c7ee01c2afa0008cf0bbd

simoneves · 2026-06-11T17:50:29Z

+  }
+  int sign = 1;
+  if (numerator < 0) {
+    numerator = -numerator;


Per @bdice:
This is UB if numerator == INT128_MIN.

simoneves · 2026-06-11T17:51:29Z

+  __int128_t scaled = numerator * scale;
+  __int128_t quotient = scaled / denom;
+  __int128_t remainder = scaled % denom;
+  if (remainder * 2 >= denom) {


Per @bdice:
This may help guard against overflow.
Suggested change

if (remainder * 2 >= denom) {

if (remainder >= denom - remainder) {

Also per @bdice:
Just verifying: Is -1.5 supposed to round "half up" to -1 or "half away from zero" to -2? This does the latter. Please add a comment indicating the convention being followed.

simoneves · 2026-06-11T17:52:42Z

+
+template <typename OutT>
+__device__ OutT
+decimalDivideImpl(__int128_t numerator, __int128_t denom, __int128_t scale) {


Per @bdice:
The parameter named scale is actually something like 10 ^ scale. We should name it more clearly, since scale is typically the exponent and not the value ten-to-the-exponent.

simoneves · 2026-06-11T17:53:13Z

+    denom = -denom;
+    sign = -sign;
+  }
+  __int128_t scaled = numerator * scale;


Per @bdice:
This function should explicitly document that the rescaling has the potential for UB (overflows).

simoneves · 2026-06-11T17:53:50Z

+  OutT* out;
+  __int128_t scale;
+
+  __device__ void operator()(int32_t idx) const {


Per @bdice:
As before, int32_t indicates a row index? Use cudf::size_type here and in the counting iterators if so.

simoneves · 2026-06-11T17:54:17Z

+    int32_t aRescale,
+    rmm::cuda_stream_view stream) {
+  if (inType == cudf::type_id::DECIMAL64) {
+    if (outType == cudf::type_id::DECIMAL64) {


Per @bdice:
We could use cudf::double_type_dispatcher here (potentially with a custom type map that only accepts numeric::decimal64 / numeric::decimal128), then use cudf::device_storage_type_t to recover the underlying types. It sounds like a lot of logic but it's reusable across the repeating dispatches for column/scalar variants.

Done, but not using a custom type map, as it's not clear to me how to do that with a double_type_dispatcher and not sure it would make the code any cleaner, as we'd still need to further validate the type combinations even if the custom map only allowed certain types in at all. Adding CUDF_UNREACHABLE in the other-type-combinations stub instead.

simoneves · 2026-06-11T17:54:52Z


 target_compile_options(velox_cudf_expression PRIVATE -Wno-missing-field-initializers)
-
-set_target_properties(velox_cudf_expression PROPERTIES CUDA_STANDARD 20 CUDA_STANDARD_REQUIRED ON)


Moved to root CMakeLists.txt (possible contentious)

simoneves · 2026-06-11T17:56:26Z

  enable_language(CUDA)
+  # Use same C++ standard throughout
+  set(CMAKE_CUDA_STANDARD ${CMAKE_CXX_STANDARD})
+  set(CMAKE_CUDA_STANDARD_REQUIRED ${CMAKE_CXX_STANDARD_REQUIRED})


Not sure why this was deemed necessary. It might have been a leftover from Part 2, although Part 2 obviously works just fine without it. Should it still be moved here (from the CMakeLists.txt in cudf/expression)? Will it break Wave, or at least offend those devs?

github-actions · 2026-06-11T17:59:11Z

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

This reverts commit e1a8cbf.

(cherry picked from commit f5ca598)

(cherry picked from commit 224bd70)

(cherry picked from commit 8f976af)

(cherry picked from commit a34c540)

github-actions · 2026-06-12T21:52:23Z

CI Failure Analysis

Auto-generated by the CI Failure Analysis workflow. This comment is updated in place each time CI fails on a new commit, so it always reflects the latest run — re-pushing or re-running CI will refresh the analysis below. Last updated 2026-06-12 23:06:08 UTC from workflow run 27444049307.

❌ Linux release with adapters — BUILD Failure View logs

Build errors:

Linker error when building velox_hive_connector_test — missing ORC reader factory symbols:

FAILED: velox/connectors/hive/tests/velox_hive_connector_test

ld: FileConnectorUtilTest.cpp.o: in function `FileConnectorUtilTest::TearDown()':
  undefined reference to `facebook::velox::orc::unregisterOrcReaderFactory()'

ld: FileConnectorUtilTest.cpp.o: in function `FileConnectorUtilTest::SetUp()':
  undefined reference to `facebook::velox::orc::registerOrcReaderFactory()'

ld: HiveConnectorUtilTest.cpp.o: in function `HiveConnectorUtilTest::TearDown()':
  undefined reference to `facebook::velox::orc::unregisterOrcReaderFactory()'

ld: HiveConnectorUtilTest.cpp.o: in function `HiveConnectorUtilTest::SetUp()':
  undefined reference to `facebook::velox::orc::registerOrcReaderFactory()'

collect2: error: ld returned 1 exit status

The velox_hive_connector_test binary links against test files (FileConnectorUtilTest.cpp, HiveConnectorUtilTest.cpp) that call registerOrcReaderFactory() / unregisterOrcReaderFactory(), but the ORC reader library (velox_dwio_orc_reader) is not included in the link dependencies.

Correlation with PR changes:

This failure is not caused by this PR. PR #17806 only modifies files under velox/experimental/cudf/expression/ (CUDA decimal kernels) and the top-level CMakeLists.txt (CUDA standard settings). It does not touch any Hive connector code, ORC reader code, or the Hive connector test CMakeLists.txt.

Known issues:

🔗 Tracked by #17816 — "Link error for velox_hive_connector_test in Linux adapter CI"
This is a pre-existing failure on the main branch. The most recent main CI run (27445031257) also fails on the same job with the same error.
A fix was merged to main in commit 16d2e3b ("fix: Link ORC reader into Hive connector tests (fix: Link ORC reader into Hive connector tests #17817)"). Rebasing this PR onto the latest main should resolve this failure.

Recommended fix:

Rebase this PR branch onto the latest main to pick up the ORC reader link fix from commit 16d2e3b:

git fetch upstream && git rebase upstream/main

simoneves requested review from bdice, devavret, jinchengchenghh, karthikeyann, majetideepak, mhaseeb123 and shrshi as code owners June 11, 2026 17:48

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 Jun 11, 2026

simoneves commented Jun 11, 2026

View reviewed changes

simoneves mentioned this pull request Jun 11, 2026

feat(cudf): GPU Decimal (Part 3) #16751

Open

simoneves force-pushed the simoneves/decimal_pr4 branch 3 times, most recently from 65a3c66 to 95b71f3 Compare June 11, 2026 18:26

GregoryKimball added this to libcudf Jun 11, 2026

Expression and CMake changes moved from Part 3

80fc52b

This reverts commit e1a8cbf.

simoneves force-pushed the simoneves/decimal_pr4 branch from 95b71f3 to 80fc52b Compare June 11, 2026 21:24

simoneves mentioned this pull request Jun 12, 2026

Proposed changes to Decimal Part 4 simoneves/velox#10

Open

simoneves added 4 commits June 12, 2026 13:54

Refactor to use cudf::type_dispatcher

cc9fac3

(cherry picked from commit f5ca598)

UB and overflow avoidance (part 1)

558a968

(cherry picked from commit 224bd70)

Match Velox CPU overflow behavior

aef1a92

(cherry picked from commit 8f976af)

Use cuda::size_type and cuda::counting_iterator

d9ec8ac

(cherry picked from commit a34c540)

github-actions Bot added the cudf cudf related - GPU acceleration label Jun 12, 2026

simoneves added 2 commits June 12, 2026 13:58

Merge branch 'main' into simoneves/decimal_pr4

a9ca6a4

Add CUDF_FAILs

6ab05db


		target_compile_options(velox_cudf_expression PRIVATE -Wno-missing-field-initializers)

		set_target_properties(velox_cudf_expression PROPERTIES CUDA_STANDARD 20 CUDA_STANDARD_REQUIRED ON)

Conversation

simoneves commented Jun 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

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

Uh oh!

✅ Deploy Preview for meta-velox canceled.

Uh oh!

simoneves Jun 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

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

Uh oh!

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

simoneves Jun 11, 2026

Choose a reason for hiding this comment

Uh oh!

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

Uh oh!

Selective Build Plan

Uh oh!

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

Uh oh!

CI Failure Analysis

❌ Linux release with adapters — BUILD Failure View logs

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

simoneves commented Jun 11, 2026 •

edited

Loading

netlify Bot commented Jun 11, 2026 •

edited

Loading

simoneves Jun 11, 2026 •

edited

Loading

simoneves Jun 11, 2026 •

edited

Loading

simoneves Jun 12, 2026 •

edited

Loading

github-actions Bot commented Jun 11, 2026 •

edited

Loading

github-actions Bot commented Jun 12, 2026 •

edited

Loading