[FEA] ANSI SQL Operator JIT Support (4) : Implement IR for ANSI and other JIT Extensions#22602
Merged
Merged
Conversation
github-actions
Bot
added
libcudf
28 tasks
|
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. |
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughAdds public error types and an evaluation_error exception, introduces JIT expression builders and an internal operation node, implements opcode-level fallibility and type validation, threads error sinks through row-IR CUDA emission and kernel launches, and propagates an error_policy through transform/compute APIs and tests. ChangesTransform and Filter Error Handling Integration
Estimated code review effort🎯 4 (Complex) | ⏱️ ~60 minutes Possibly related issues
Possibly related PRs
Suggested labels
Suggested reviewers
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 15
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
cpp/include/cudf/transform.hpp (1)
172-183:⚠️ Potential issue | 🟠 Major | ⚡ Quick winThis signature change is source-breaking for positional callers.
Adding
error_modeat Line 181 beforestream/mrchanges argument positions for existingmulti_transform(..., row_size, stream, mr)call sites. Please keep backward compatibility via an overload (old signature) that forwards withops::error_mode::IGNORE.As per coding guidelines, "C++ API changes in public headers (cpp/include/cudf/, cpp/include/nvtext/) require proper deprecation: [[deprecated]] attribute,
@deprecateddoxygen tag, and PR labels (deprecation/breaking)."🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/include/cudf/transform.hpp` around lines 172 - 183, The new multi_transform signature is source-breaking because inserting the error_mode parameter shifts positional arguments (breaking calls like multi_transform(..., row_size, stream, mr)); restore backward compatibility by adding an overloaded multi_transform with the old parameter order (no error_mode) that is marked [[deprecated]] and has a `@deprecated` doxygen tag, and have that overload simply forward to the new multi_transform passing ops::error_mode::IGNORE for error_mode; ensure the forwarding overload accepts the same parameters (std::optional<size_type> row_size, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) and documents deprecation so callers can migrate.cpp/src/transform/transform.cu (1)
162-197:⚠️ Potential issue | 🔴 Critical | ⚡ Quick winFix function signatures for
instantiateandlaunchto match their callers.The
runfunction callsinstantiate(error_handling_mode, is_null_aware, ...)andlaunch(..., d_error_sink, stream), but the function definitions at lines 162 and 186 have not been updated to accept these parameters. Additionally, the template instantiation at line 182 needs to passerror_handling_modeas the first argument. This is a compile-time blocker.Required changes
-jitify2::Kernel instantiate(null_aware is_null_aware, +jitify2::Kernel instantiate(ops::error_mode error_handling_mode, + null_aware is_null_aware, bool has_user_data, std::string const& ins, std::string const& outs, std::vector<std::string> const& ptx_input_types, std::vector<std::string> const& ptx_output_types, std::string const& udf, udf_source_type source_type) { @@ auto kernel = jitify2::reflection::Template("cudf::jit::transform_kernel") - .instantiate(is_null_aware, has_user_data, ins, outs); + .instantiate(error_handling_mode, is_null_aware, has_user_data, ins, outs); @@ void launch(jitify2::Kernel const& kernel, size_type row_size, bitmask_type const* stencil, void* user_data, column_device_view_core const* input_cols, mutable_column_device_view_core const* output_cols, + jit::error_sink* d_error_sink, rmm::cuda_stream_view stream) { @@ - void* args[] = {&row_size, &stencil, &user_data, &input_cols, &output_cols}; + void* args[] = {&row_size, &stencil, &user_data, &input_cols, &output_cols, &d_error_sink}; kernel->configure_1d_max_occupancy(0, 0, nullptr, stream.value())->launch_raw(args); }🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/src/transform/transform.cu` around lines 162 - 197, The instantiate and launch function signatures do not match their callers: update instantiate(jitify2::Kernel instantiate(...)) to accept the additional first parameter error_handling_mode (same enum/type used by callers) and pass that error_handling_mode as the first template argument when creating jitify2::reflection::Template("cudf::jit::transform_kernel").instantiate(...); also update launch(void launch(...)) to accept the extra trailing parameter d_error_sink (the device error sink pointer passed from run) and ensure it is included in the kernel argument list (add &d_error_sink or equivalent to the void* args array). Adjust parameter order and names to match callers (run) so compilation succeeds, using the symbols instantiate, launch, transform_kernel, error_handling_mode, and d_error_sink to locate changes.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/include/cudf/operators/ansi_arithmetic.cuh`:
- Around line 397-402: In ansi_div(decimal<R>* out, decimal<R> const* a,
decimal<R> const* b) handle divide-by-zero separately: first check if b->value()
== 0 and return errc::DIVISION_BY_ZERO, otherwise check
numeric::division_overflow<R>(a->value(), b->value()) and return errc::OVERFLOW;
update the conditional logic in ansi_div to preserve the ANSI error class for
decimal division instead of lumping zero-divide into OVERFLOW.
- Around line 654-660: The ansi_neg template currently only constrains signed
types but should follow the file's pattern: change the existing ansi_neg<T>
constraint to requires(cuda::std::is_integral_v<T> && cuda::std::is_signed_v<T>)
so it applies only to signed integral types (and still checks overflow for
numeric_limits<T>::min()), and add a second overload ansi_neg<T> constrained
with requires(cuda::std::is_floating_point_v<T>) that performs negation without
the integer overflow check and returns errc::OK; update both
declarations/definitions to match the file's style and ensure unique symbol
names remain ansi_neg for overload resolution.
In `@cpp/include/cudf/operators/arithmetic.cuh`:
- Around line 21-27: In the signed-integer paths of __device__ functions abs(T*
out, T const* a) and neg(...) detect the negation overflow when *a equals the
type's minimum representable value (e.g., std::numeric_limits<T>::min()) and
handle it by returning errc::OVERFLOW (or alternatively saturating to
std::numeric_limits<T>::max() if saturation semantics are desired); update abs
and neg to perform this check only for integral signed types (preserve current
behavior for floating types), and ensure you use the same errc::OVERFLOW enum to
signal the condition rather than producing UB by negating the min value.
- Around line 129-133: The integer arithmetic functions div, floor_div, mod, and
pymod must check the divisor before performing / or % and return
errc::DIVISION_BY_ZERO when the divisor is zero; update each function (div,
floor_div, mod, pymod) to first detect integer types (e.g., via
std::is_integral<T>::value) and if the divisor (*b) == 0 return
errc::DIVISION_BY_ZERO, otherwise proceed with the existing operation and return
errc::OK so floating-point paths remain unchanged.
In `@cpp/include/cudf/operators/bitwise.cuh`:
- Around line 171-175: bit_shift_left and bit_shift_right perform shifts without
validating the shift count; update both functions (bit_shift_left,
bit_shift_right) to guard/clamp the shift amount by first checking the count is
non-negative and less than the type width (e.g., width = sizeof(T)*8), and if
invalid return an appropriate errc (e.g., errc::INVALID_ARGUMENT or
errc::OUT_OF_RANGE), otherwise perform the shift using the validated/clamped
count and return errc::OK; ensure you handle signed/unsigned count types
consistently and use the validated count variable in the shift expression.
In `@cpp/include/cudf/operators/casts.cuh`:
- Around line 461-466: decimal_cast currently static_casts a->value() into To
which can silently wrap/truncate; change decimal_cast(decimal<To>* out,
decimal<From> const* a) to first read the source representation (auto src =
a->value()) and check that src is within the representable range of To (compare
against std::numeric_limits<To>::min()/max() or equivalent numeric::bounds
helpers) before performing static_cast; if the value is out of range return a
non-OK errc (e.g., errc::ARITHMETIC or the appropriate overflow error enum) and
only on success perform the static_cast and assign to *out preserving the scale.
In `@cpp/include/cudf/operators/detail/promote.cuh`:
- Around line 19-20: Add a guarded primary template for promoted_t that triggers
a dependent static_assert with a clear diagnostic when users try to instantiate
promoted_t<T> for unsupported types: modify the primary template struct
promoted_t to contain a static_assert(sizeof(T) == 0, "promoted_t<T>
instantiated for an unsupported type; specialize promoted_t for supported types
like ..."), so instantiations of promoted_t<T> without a matching specialization
produce an actionable error message rather than an opaque incomplete-type
failure. Ensure the static_assert is dependent on T (e.g., using sizeof(T)==0)
so specializations still work.
In `@cpp/include/cudf/operators/error.hpp`:
- Line 13: Change the unscoped enum declaration to a scoped enum by replacing
"enum errc : int { OK = 0, OVERFLOW = 1, DIVISION_BY_ZERO = 2 };" with "enum
class errc : int { OK = 0, OVERFLOW = 1, DIVISION_BY_ZERO = 2 };" so that call
sites using ops::errc::OK, ops::errc::OVERFLOW, and ops::errc::DIVISION_BY_ZERO
compile; after this change, update any code that relied on implicit conversion
to int (e.g., comparisons or assignments) to use explicit casts like
static_cast<int>(ops::errc::OK) where necessary.
In `@cpp/include/cudf/operators/logic.cuh`:
- Around line 249-255: The current code incorrectly requires both true_value and
false_value to be present before evaluating optional branch semantics; instead,
check pred->has_value() first, then branch on pred->value(): if true, require
true_value->has_value() and use true_value (call if_else<T> or directly assign)
to set *out, if false require false_value->has_value() and use false_value to
set *out; if the selected branch is missing or pred is missing set *out =
nullopt. Update the logic around pred, true_value, false_value, if_else, and out
to implement this per-row branch-dependent null handling.
In `@cpp/include/cudf/operators/op_traits.hpp`:
- Around line 9-10: The header op_traits.hpp uses std::min (referenced around
line 310) but doesn't include <algorithm>, which breaks header self-containment;
fix by adding the <algorithm> include to op_traits.hpp (alongside the existing
includes like <span> and <vector>) so std::min is declared when this header is
included.
In `@cpp/include/cudf/operators/types.cuh`:
- Around line 42-43: Add a compile-time guard to the template function ipow10 to
restrict instantiation to integral types: inside the template (or as a template
constraint) use a static_assert checking std::is_integral_v<T> (include
<type_traits> if necessary) with a clear message like "ipow10 requires an
integral type" so non-integral instantiations fail to compile; apply this to the
template declaration/definition for ipow10 to prevent misuse.
In `@cpp/src/jit/row_ir.cpp`:
- Around line 443-447: The emitted code uses `return e;` for fallible ops but
`generate_code()` currently declares the generated device function as `void`,
producing invalid CUDA; fix by changing the generated device function ABI in
`generate_code()` to return `cudf::ops::errc` (update the function signature and
all call sites accordingly) so the `sink.emit` template that emits `return e;`
becomes valid—ensure `sink.emit` places `cudf::ops::errc` as the function return
type and propagate/handle the returned `errc` in callers.
- Around line 67-80: The constructor currently only checks arity for non-RESCALE
ops but allows a stray target_scale_; add a
CUDF_EXPECTS(!target_scale_.has_value(), std::format("Target scale must be
nullopt for operator `{}`. Got a target scale.", get_op_name(op_))) in the
non-RESCALE branch (alongside the existing arity check) so target_scale_ is
enforced exclusive to opcode::RESCALE; keep the existing
target_scale_.has_value() check in the RESCALE branch and update error text
there if needed to reference get_op_name(op_) for consistency.
- Around line 371-374: When handling opcode::RESCALE in row_ir.cpp, avoid
constructing cudf::numeric_scalar<int32_t> with implicit defaults; instead pass
the instance_context's stream_ and mr_ into the scalar so stream ordering and
memory resource are preserved. Update the branch that sets scale_reference_ (the
input_reference created via ctx.add_input) to construct
cudf::numeric_scalar<int32_t> with ctx.stream_ and ctx.mr_ (accessing the
private stream_/mr_ available to node/instance_context) and then pass that
scalar to ctx.add_input so the input uses the correct stream and memory
resource.
In `@cpp/src/transform/transform.cu`:
- Around line 829-835: Replace usage of rmm::device_scalar with
cudf::detail::device_scalar for d_error_sink: change the optional type
declaration std::optional<rmm::device_scalar<jit::error_sink>> d_error_sink and
the construction in the ops::error_mode::ANY_ROW case to use
cudf::detail::device_scalar<jit::error_sink> instead; also add the missing
include <cudf/detail/device_scalar.hpp> to the file so the type is available.
This keeps the variable d_error_sink and the switch on
error_handling_mode/ops::error_mode::ANY_ROW the same while swapping the
concrete device_scalar implementation.
---
Outside diff comments:
In `@cpp/include/cudf/transform.hpp`:
- Around line 172-183: The new multi_transform signature is source-breaking
because inserting the error_mode parameter shifts positional arguments (breaking
calls like multi_transform(..., row_size, stream, mr)); restore backward
compatibility by adding an overloaded multi_transform with the old parameter
order (no error_mode) that is marked [[deprecated]] and has a `@deprecated`
doxygen tag, and have that overload simply forward to the new multi_transform
passing ops::error_mode::IGNORE for error_mode; ensure the forwarding overload
accepts the same parameters (std::optional<size_type> row_size,
rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) and documents
deprecation so callers can migrate.
In `@cpp/src/transform/transform.cu`:
- Around line 162-197: The instantiate and launch function signatures do not
match their callers: update instantiate(jitify2::Kernel instantiate(...)) to
accept the additional first parameter error_handling_mode (same enum/type used
by callers) and pass that error_handling_mode as the first template argument
when creating
jitify2::reflection::Template("cudf::jit::transform_kernel").instantiate(...);
also update launch(void launch(...)) to accept the extra trailing parameter
d_error_sink (the device error sink pointer passed from run) and ensure it is
included in the kernel argument list (add &d_error_sink or equivalent to the
void* args array). Adjust parameter order and names to match callers (run) so
compilation succeeds, using the symbols instantiate, launch, transform_kernel,
error_handling_mode, and d_error_sink to locate changes.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: d4b3a073-e61d-4e84-8032-949584894420
📒 Files selected for processing (27)
cpp/CMakeLists.txtcpp/include/cudf/ast/jit_expressions.hppcpp/include/cudf/operators/ansi_arithmetic.cuhcpp/include/cudf/operators/arithmetic.cuhcpp/include/cudf/operators/bitwise.cuhcpp/include/cudf/operators/casts.cuhcpp/include/cudf/operators/comparison.cuhcpp/include/cudf/operators/detail/promote.cuhcpp/include/cudf/operators/error.hppcpp/include/cudf/operators/logic.cuhcpp/include/cudf/operators/math.cuhcpp/include/cudf/operators/null_handling.cuhcpp/include/cudf/operators/op_traits.hppcpp/include/cudf/operators/opcodes.hppcpp/include/cudf/operators/trigonometric.cuhcpp/include/cudf/operators/types.cuhcpp/include/cudf/transform.hppcpp/src/ast/jit_expressions.cppcpp/src/jit/error_sink.cuhcpp/src/jit/row_ir.cppcpp/src/jit/row_ir.hppcpp/src/transform/jit/kernel.cucpp/src/transform/transform.cucpp/tests/CMakeLists.txtcpp/tests/ast/jit_ast_tests.cppcpp/tests/jit/row_ir.cppcpp/tests/transform/integration/unary_transform_test.cpp
3 tasks
lamarrr
force-pushed
the
ansi-jit-3--ansi-row-ir
branch
from
6c3ad8c to
b05bdc0
Compare
…ons and new ansi operators
github-actions
Bot
added
Python
lamarrr
added
feature request
…ng in compute_column functions
- Introduced a new enum class `error_policy` to specify how functions handle errors (PROPAGATE or NULLIFY). - Updated the `evaluation_error` class to reflect changes in error handling. - Modified JIT operation functions to accept `error_policy` instead of a boolean for nullification on error. - Adjusted the internal representation of operations to store the error policy. - Updated tests to use the new error policy handling, ensuring correct behavior for operations that may fail.
8 tasks
bdice
approved these changes
Jun 30, 2026
bdice
left a comment
bdice
left a comment
Contributor
There was a problem hiding this comment.
Approving, conditional on a satisfactory outcome for the Spark performance tests.
wence-
approved these changes
Jul 1, 2026
wence-
left a comment
wence-
left a comment
Contributor
There was a problem hiding this comment.
Thanks for addressing my comments.
Contributor
Author
|
/ok to test b687474 |
…on older CUDA drivers
…cudf into ansi-jit-3--ansi-row-ir
Contributor
Author
|
/ok to test 309b2c8 |
Contributor
Author
|
/ok to test 0d885f3 |
Contributor
Author
|
/merge |
thirtiseven
reviewed
Jul 2, 2026
thirtiseven
left a comment
thirtiseven
left a comment
Contributor
There was a problem hiding this comment.
Sorry I have a few more comments from spark integration side.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
This pull request is the last of the set in #22598.
It:
transformExample usage of the newly added ANSI-mode operators can be found in
cpp/tests/ast/jit_expressions_tests.cpp.Checklist