[WIP][FEA] Support ANSI SQL operators in Transform by lamarrr · Pull Request #22224 · NVIDIA/cudf

lamarrr · 2026-04-20T23:40:12Z

Description

This is a POC PR to implement support for ANSI SQL semantics for CUDF JIT Transforms.

New Operators

Checklist

I am familiar with the Contributing Guidelines.
New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2026-04-20T23:40:15Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

Co-authored-by: Copilot <copilot@github.com>

- Mark `expression` struct with [[nodiscard]] to prevent unintended discards. - Update `operation` constructors and methods for improved clarity and consistency. - Introduce new JIT expressions: `nullify_if`, `coalesce`, `predicate`, and ANSI-compliant arithmetic operations (add, sub, mul, div) with error handling. - Implement tests for JIT expressions including `nullify_if`, `coalesce`, and ANSI arithmetic operations for both integer and decimal types. - Ensure proper handling of overflow scenarios in arithmetic operations. - Refactor null awareness and validity checks in the row IR to improve robustness. - Add additional casting and rescaling expressions with corresponding tests.

revans2

From a quick look through the public APIs it looks good to me, as something we could start using. It will not really be able to support decimal for us. There are a few issues here.

Spark knows about precision and scale, but this just knows about scale and bit-width. Precision and bit-width are related, but overflow checks in Spark require precision to work properly. Not width.
Spark also rounds the results to the desired output scale that it wants using half even rounding. When I look at rescaled it appears that it is not rounding at all. It is just adjusting the scale at a logical level. If that is true, then we don't have a way to represent the rounding to get to the scale that we need/want.
256-bit intermediate values. In order to support all decimal operations we need to be able to do the math at a wider width/precision that 128 bits supports. So with what we have we cannot do decimal operations for all 128 bit values and would have to fall back to our custom kernels for a lot of them.

None of these are blockers fro us to start using the code/testing it out. But it does mean that we will likely only have limited decimal support/testing.

lamarrr · 2026-05-05T14:28:48Z

Spark knows about precision and scale, but this just knows about scale and bit-width. Precision and bit-width are related, but overflow checks in Spark require precision to work properly. Not width.

The current functions are intended to preserve the existing CUDF operator semantic.
CUDF doesn't have a notion of precision but I added functions to help you cast to higher bit-width decimals and also perform a precision check ansi_precision_check and ansi_try_precision_check, which I had the impression would be sufficient.

Spark also rounds the results to the desired output scale that it wants using half even rounding. When I look at rescaled it appears that it is not rounding at all. It is just adjusting the scale at a logical level. If that is true, then we don't have a way to represent the rounding to get to the scale that we need/want.

rescaled uses the logical rescaling, I will add a half-even rounding method.

256-bit intermediate values. In order to support all decimal operations we need to be able to do the math at a wider width/precision that 128 bits supports. So with what we have we cannot do decimal operations for all 128 bit values and would have to fall back to our custom kernels for a lot of them.

I believe this is solely to detect overflows? We were able to perform overflow checks without using 256-bit intermediates. We currently don't support 256-bit decimals and I would defer implementing it unless absolutely necessary.

revans2 · 2026-05-05T15:38:52Z

I believe this is solely to detect overflows? We were able to perform overflow checks without using 256-bit intermediates. We currently don't support 256-bit decimals and I would defer implementing it unless absolutely necessary.

Yes and no. For add and subtract it is mostly for overflow detection, but not in all cases. DECIMAL(38, 5) + DECIMAL(38, 4). In spark we have a custom kernel to be able to do this because we have to normalize the two to add them, and to do that I would need DECIMAL(39, 5) + DECIMAL(39, 5) as the wider type. Then I would add them to a DECIMAL(39, 5) intermediate and finally cast it back to a DECIMAL(38, 5) with an overflow check.

For multiply and divide it is similar the intermediate value could grow up to a 256 bit value internally before it is cast back down to the desired scale and precision. I suppose with knowing the output scale and precision ahead of time you might be able to make the actual intermediate size smaller, but I don't know how you can do half even rounding at the end without at least having a precision of 39 so you can know that other digit to round with.

This is not a blocker like I said before. It just limits the range of expressions we can port over to this. We can keep our custom kernels for cases where we need them

igorpeshansky

Agree with @revans2 — most of these should be usable. Noticed a few more discrepancies besides the ones he pointed out.

Split from #22224 This Pull request: - refactors the ROW-IR codegen setup to prepare for adding ANSI SQL operator support. - begins splitting out the opcode from the AST operator - unifies the `node` type to use opcodes instead of dynamic dispatch - removes ambiguous AST `input` resolution logic - removes the redundant `join_column_accessor` and instead uses a table index attached to each `column_accessor` instead Authors: - Basit Ayantunde (https://github.com/lamarrr) Approvers: - Bradley Dice (https://github.com/bdice) URL: #22511

…22514) Split from #22224 Preceded by #22511 Story: #22598 This Pull request: - Implements error codes for row operators - Ports AST's operators to re-usable functions that can be used with JIT codegen - Adds new ANSI-compliant operators: - ANSI_ADD - ANSI_SUB - ANSI_MUL - ANSI_DIV - ANSI_MOD - ANSI_ABS - ANSI_NEG - Adds the `coalesce` operator - Introduces an operator `result` type to allow error returns from operators - Transitioned AST `operator_functors` to use the operator library Authors: - Basit Ayantunde (https://github.com/lamarrr) Approvers: - Lawrence Mitchell (https://github.com/wence-) - Yunsong Wang (https://github.com/PointKernel) URL: #22514

thirtiseven

Thank you for doing this! I'm working on a POC to integrate it into spark-rapids, and here are some findings from doing so (mostly from AI)

thirtiseven · 2026-06-11T10:09:07Z

-template <bool has_user_data, typename Args>
-__device__ void execute_transform_op(void* user_data, size_type element_idx, Args args)
+template <ops::error_mode mode, bool has_user_data, typename Args>
+__device__ void execute_transform_op(error_sink* __restrict__ error_sink,


In ANY_ROW mode, after a row reports an error, the transform kernel should not continue assigning that row's output. Could execute_transform_op return whether the UDF succeeded, and use that to skip output assignment and validity updates for failed rows?

that would not be okay; some value types still need to be assigned, as the output column is assumed to be uninitialized. For example, we can't leave a string_view or list_view column uninitialized.

thirtiseven · 2026-06-11T10:10:23Z


  auto finalized = finalize_outputs(is_null_aware, row_size, std::move(output_columns), stream, mr);
+
+  switch (error_handling_mode) {


Could we check the transform error sink immediately after jit_transform::run(...) and before finalize_outputs(...)? Otherwise failed transforms can still finalize outputs before throwing the reported error.

thirtiseven · 2026-06-11T10:10:58Z

+}
+
+template <typename T>
+__device__ inline errc if_else(T* out, T const* true_value, T const* false_value, bool const* pred)


The nullable if_else overload appears to pass arguments to the non-nullable overload in the wrong order. The callee expects (out, true_value, false_value, pred), but the current call passes (out, pred, true_value, false_value)

fixed in #22514

thirtiseven · 2026-06-11T10:12:04Z

+                               optional<T> const* false_value,
+                               optional<bool> const* pred)
+{
+  if (pred->has_value() && true_value->has_value() && false_value->has_value()) {


The nullable if_else overload writes through out->value() before the optional output is engaged, which looks like an implementation bug. Also, should nullable IF_ELSE follow selected-branch null semantics? Requiring both branch inputs to be valid makes IF(true, valid, null) and IF(false, null, valid) null. If strict null propagation is intended here, can we document that so consumers like Spark can avoid this operator for nullable branches?

fixed in #22514

thirtiseven · 2026-06-11T10:15:54Z

-                  .instantiate(is_null_aware, has_user_data, ins, outs);
+                  .instantiate(error_handling_mode, is_null_aware, has_user_data, ins, outs);

  return jit::get_udf_kernel(


We found a stale-but-loadable Jitify file-cache entry that caused cudaErrorLaunchFailure for a transform JIT kernel. Clearing/disabling the cache or using a fresh LIBCUDF_KERNEL_CACHE_PATH fixed it, and copying the stale file into a clean cache reproduced it. This is not recoverable through the current deserialization retry because get_kernel() succeeds and the failure happens at launch. Could the file cache key or directory include a stronger libcudf/JIT-header build identity, or otherwise avoid reusing stale linked-program files across local rebuilds with the same CUDF_VERSION?

Operational avoidance:

Prefer a build- or deployment-specific LIBCUDF_KERNEL_CACHE_PATH.

Clear the cache when swapping libcudf builds in place.

Use LIBCUDF_KERNEL_CACHE_DISABLED=1 only as a debugging or conservative workaround; it avoids stale files but gives up disk-cache benefits.

The new RTCX library we introduced fixed this headache. It uses the build hash of the headers as a subdirectory in the LIBCUDF_KERNEL_CACHE_PATH directory.

lamarrr · 2026-06-22T16:45:37Z

Closed.
Being upstreamed via #22598

support ANSI SQL operators

8919728

github-actions Bot assigned lamarrr Apr 20, 2026

github-actions Bot added the libcudf Affects libcudf (C++/CUDA) code. label Apr 20, 2026

lamarrr changed the title ~~[FEA] Support ANSI SQL operators~~ [FEA] Support ANSI SQL operators in Transform Apr 20, 2026

lamarrr changed the title ~~[FEA] Support ANSI SQL operators in Transform~~ [WIP][FEA] Support ANSI SQL operators in Transform Apr 20, 2026

lamarrr and others added 3 commits April 21, 2026 23:12

checkpoint

fbf1cec

checkpoint

602390f

Co-authored-by: Copilot <copilot@github.com>

Remove deprecated opcode definitions from opcode.hpp

c52c6b8

lamarrr added feature request New feature or request non-breaking Non-breaking change labels Apr 26, 2026

lamarrr and others added 9 commits April 27, 2026 17:28

update

87b1d7a

Co-authored-by: Copilot <copilot@github.com>

Merge remote-tracking branch 'upstream/main' into ansi-jit

ec0b3d6

add decimal rescale implementation

b2c3ac8

Merge remote-tracking branch 'upstream/main' into ansi-jit

3e19bbf

update

a62cdfd

Merge remote-tracking branch 'upstream/main' into ansi-jit

4adcdaa

Merge remote-tracking branch 'upstream/main' into ansi-jit

e12417f

Merge remote-tracking branch 'upstream/main' into ansi-jit

938205c

Implement JIT expressions and operations for ANSI SQL compatibility

dc47b85

github-actions Bot added the CMake CMake build issue label May 1, 2026

lamarrr added 4 commits May 3, 2026 01:07

add more tests and bug fixes

e90392d

bug fixes and more tests

02878c3

bug fixes and more tests

8292fc4

GregoryKimball added this to libcudf May 4, 2026

GregoryKimball moved this to Burndown in libcudf May 4, 2026

lamarrr added 2 commits May 5, 2026 00:40

bug fixes and more tests

923bd2c

bug fixes and more tests

0423a7d

revans2 reviewed May 5, 2026

View reviewed changes

parameterize tests

59f7f0b

igorpeshansky reviewed May 8, 2026

View reviewed changes

Comment thread cpp/include/cudf/operators/ansi_arithmetic.cuh

Comment thread cpp/include/cudf/operators/ansi_arithmetic.cuh

Comment thread cpp/include/cudf/operators/ansi_arithmetic.cuh

This was referenced May 14, 2026

[FEA] ANSI SQL Operator JIT Support (1) : Refactor ROW IR #22511

Merged

[FEA] ANSI SQL Operator JIT Support (2) : Implement Operator Library #22514

Merged

vuule moved this from Burndown to Slip in libcudf May 14, 2026

lamarrr mentioned this pull request May 20, 2026

[FEA][STORY] Support ANSI Operators using JIT #22598

Closed

28 tasks

thirtiseven reviewed Jun 11, 2026

View reviewed changes

This was referenced Jun 12, 2026

Review ANSI row IR follow-up fixes thirtiseven/cudf#3

Open

Support ANSI JIT AST project expressions NVIDIA/cudf-spark#15069

Closed

[FEA] ANSI SQL Operator JIT Support (4) : Implement IR for ANSI and other JIT Extensions #22602

Merged

lamarrr closed this Jun 22, 2026

GregoryKimball removed this from libcudf Jul 13, 2026


		auto finalized = finalize_outputs(is_null_aware, row_size, std::move(output_columns), stream, mr);

		switch (error_handling_mode) {

Uh oh!

Conversation

lamarrr commented Apr 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

New Operators

Checklist

Uh oh!

copy-pr-bot Bot commented Apr 20, 2026

Uh oh!

revans2 left a comment

Choose a reason for hiding this comment

Uh oh!

lamarrr commented May 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

revans2 commented May 5, 2026

Uh oh!

igorpeshansky left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

thirtiseven left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lamarrr Jun 22, 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!

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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lamarrr commented Jun 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

lamarrr commented Apr 20, 2026 •

edited

Loading

lamarrr commented May 5, 2026 •

edited

Loading

lamarrr Jun 22, 2026 •

edited

Loading