[FEA] RTCX: Replace Software SHA256 with SIMD-accelerated xxHash128 for JIT cache - #22920
Closed
lamarrr wants to merge 8 commits into
Closed
[FEA] RTCX: Replace Software SHA256 with SIMD-accelerated xxHash128 for JIT cache#22920lamarrr wants to merge 8 commits into
lamarrr wants to merge 8 commits into
Conversation
- Introduced a new hash128 structure to replace the previous sha256 implementation. - Added hash128_hex_string for hexadecimal representation of hash128. - Updated cache management functions to utilize hash128 instead of sha256. - Removed sha256.hpp and its associated logic, consolidating hashing functionality. - Adjusted file handling and caching mechanisms to accommodate the new hash type. - Ensured compatibility with existing code by updating relevant function signatures and implementations.
|
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:
📝 WalkthroughWalkthroughReplaces the custom in-header SHA-256 implementation ( ChangesSHA-256 → XXH3 128-bit hash migration
Estimated code review effort🎯 4 (Complex) | ⏱️ ~45 minutes 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: 2
🧹 Nitpick comments (1)
cpp/librtcx/rtcx.hpp (1)
478-518: 💤 Low valueRename parameter
shatohashfor consistency.The parameter at line 511 is still named
shaeven though the type has been changed tohash128. This naming inconsistency could cause confusion during future maintenance.- void insert(hash128 const& sha, T&& value, std::uint64_t tick) + void insert(hash128 const& hash, T&& value, std::uint64_t tick) { if (limit_ == 0) { return; } if ((entries_.size() + 1) > limit_) { purge(); } - entries_.emplace(sha, entry{tick, std::move(value)}); + entries_.emplace(hash, entry{tick, std::move(value)}); }🤖 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/librtcx/rtcx.hpp` around lines 478 - 518, The `insert` method in the `lru_memory_cache` class has a parameter named `sha` with type `hash128`, which is inconsistent. Rename the parameter from `sha` to `hash` in the method signature, and update all references to this parameter within the method body, including the call to `entries_.emplace()` where it is used.
🤖 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/librtcx/embed.hpp`:
- Around line 110-117: The XXH3_createState() function call can return NULL on
memory allocation failure, but the code does not check for this condition before
using the state pointer in subsequent XXH3_128bits_update and
XXH3_128bits_digest calls. Add a NULL check immediately after the
XXH3_createState() assignment using the RTCX_EMBED_EXPECTS macro to validate
that the state pointer is not NULL before proceeding with the hash operations.
This will prevent potential null pointer dereferences if memory allocation
fails.
In `@cpp/librtcx/hash.hpp`:
- Around line 70-73: The operator[] method in the hash class has a critical
out-of-bounds memory access bug. When index=0, the calculation 16 - index
results in accessing byte[16], which is invalid for a 16-byte buffer (valid
indices are 0-15). Change the indexing calculation from 16 - index to 15 - index
to properly access the valid byte range while maintaining big-endian byte order
where index=0 returns the most significant byte (at position 15) and index=15
returns the least significant byte (at position 0).
---
Nitpick comments:
In `@cpp/librtcx/rtcx.hpp`:
- Around line 478-518: The `insert` method in the `lru_memory_cache` class has a
parameter named `sha` with type `hash128`, which is inconsistent. Rename the
parameter from `sha` to `hash` in the method signature, and update all
references to this parameter within the method body, including the call to
`entries_.emplace()` where it is used.
🪄 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: 24df94fa-a489-4583-9332-4d68f5750363
📒 Files selected for processing (7)
cpp/librtcx/README.mdcpp/librtcx/embed.hppcpp/librtcx/hash.hppcpp/librtcx/rtcx.cppcpp/librtcx/rtcx.hppcpp/librtcx/sha256.hppcpp/src/jit/cache.cpp
💤 Files with no reviewable changes (1)
- cpp/librtcx/sha256.hpp
lamarrr
added
feature request
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cpp/src/jit/cache.cpp (1)
32-36:⚠️ Potential issue | 🔴 Critical | ⚡ Quick winCRITICAL: Encode header boundaries in the cache key.
Issue:
hash(state, header_include_names)andhash(state, headers)concatenate C strings without counts or per-entry lengths, so inputs like{"a","bc"}and{"ab","c"}hash identically.
Why: Thatcache_keydrives memory/diskcuLibraryreuse, so a collision here can return a library compiled with different headers.Suggested fix
void hash(XXH3_state_t* ctx, std::span<char const> input) { XXH3_128bits_update(ctx, input.data(), input.size()); } +void hash(XXH3_state_t* ctx, std::uint64_t input) +{ + hash(ctx, std::span<char const>{reinterpret_cast<char const*>(&input), sizeof(input)}); +} + void hash(XXH3_state_t* ctx, std::span<char const* const> inputs) { + hash(ctx, static_cast<std::uint64_t>(inputs.size())); for (auto const* input : inputs) { - XXH3_128bits_update(ctx, input, std::strlen(input)); + auto const size = std::strlen(input); + hash(ctx, static_cast<std::uint64_t>(size)); + hash(ctx, std::span<char const>{input, size}); } }Also applies to: 321-323
🤖 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/jit/cache.cpp` around lines 32 - 36, The hash function `hash` is concatenating C string contents without encoding boundary information, causing different input combinations like `{"a","bc"}` and `{"ab","c"}` to produce identical hashes. This creates cache collisions that could lead to incorrect library reuse. Fix this by encoding the length of each string before hashing its content in the loop that iterates through inputs, ensuring that each string's length is hashed as a distinct value using XXH3_128bits_update before hashing the string content itself. Apply this same fix to the other hash calls at lines 321-323 that also pass header data without boundary encoding.
🤖 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.
Outside diff comments:
In `@cpp/src/jit/cache.cpp`:
- Around line 32-36: The hash function `hash` is concatenating C string contents
without encoding boundary information, causing different input combinations like
`{"a","bc"}` and `{"ab","c"}` to produce identical hashes. This creates cache
collisions that could lead to incorrect library reuse. Fix this by encoding the
length of each string before hashing its content in the loop that iterates
through inputs, ensuring that each string's length is hashed as a distinct value
using XXH3_128bits_update before hashing the string content itself. Apply this
same fix to the other hash calls at lines 321-323 that also pass header data
without boundary encoding.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: aa6d7399-5c21-4102-bb50-c39c2110b21d
📒 Files selected for processing (2)
cpp/librtcx/embed.hppcpp/src/jit/cache.cpp
🚧 Files skipped from review as they are similar to previous changes (1)
- cpp/librtcx/embed.hpp
There was a problem hiding this comment.
🧹 Nitpick comments (1)
cpp/src/jit/cache.cpp (1)
300-303: ⚡ Quick winConsider stack-allocating
XXH3_state_tto avoid heap allocation in hot path.
XXH3_state_tcan be declared directly on the stack, eliminating themalloc/freeoverhead fromXXH3_createState()/XXH3_freeState().♻️ Proposed change
- XXH3_state_t* state = XXH3_createState(); - CUDF_EXPECTS(state != nullptr, "Failed to create XXH3 state", std::runtime_error); - XXH3_128bits_reset(state); - RTCX_DEFER([state] { XXH3_freeState(state); }); + XXH3_state_t state; + XXH3_128bits_reset(&state);Then update usages to pass
&stateinstead ofstate.🤖 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/jit/cache.cpp` around lines 300 - 303, The XXH3_state_t object is being heap-allocated using XXH3_createState() in this hot path, which incurs malloc/free overhead. Instead, declare XXH3_state_t directly as a stack-allocated variable, remove the XXH3_createState() call and the null check, update the XXH3_128bits_reset() call to pass the address of the stack-allocated state using the address-of operator, and remove the RTCX_DEFER cleanup block that calls XXH3_freeState() since the stack allocation will be automatically cleaned up when it goes out of scope.
🤖 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.
Nitpick comments:
In `@cpp/src/jit/cache.cpp`:
- Around line 300-303: The XXH3_state_t object is being heap-allocated using
XXH3_createState() in this hot path, which incurs malloc/free overhead. Instead,
declare XXH3_state_t directly as a stack-allocated variable, remove the
XXH3_createState() call and the null check, update the XXH3_128bits_reset() call
to pass the address of the stack-allocated state using the address-of operator,
and remove the RTCX_DEFER cleanup block that calls XXH3_freeState() since the
stack allocation will be automatically cleaned up when it goes out of scope.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: d2b51fab-fcbd-4d24-bde9-1687b915be0b
📒 Files selected for processing (1)
cpp/src/jit/cache.cpp
3 tasks
Contributor
|
How worse it is if we use |
Contributor
Author
|
closed and merged into #22680 |
Contributor
Author
|
|
rapids-bot Bot
pushed a commit
that referenced
this pull request
Jul 3, 2026
This PR adds LTO-based transforms to libcudf by introducing a new `cudf::transform_lto` API that accepts LTO-IR or FATBIN UDF binaries and executes them through the existing transform pipeline. The API supports typed transform inputs/outputs, null-awareness, and optional user data. The change enables libcudf transforms to execute precompiled device UDF fragments instead of relying only on source/PTX-style runtime compilation. This creates a path for lower-overhead, link-time-optimized transform kernels while preserving the existing transform abstraction around input columns, scalar inputs, output specifications, and null policy. <img width="1693" height="929" alt="image" src="https://github.com/user-attachments/assets/36e4850b-3176-4297-942f-e532e93cb378" /> The AOT-compiled transform UDF has a similar ABI signature as NUMBA-CUDA UDFs (https://nvidia.github.io/numba-cuda/user/cuda_ffi.html): ```cpp extern "C" __device__ int transform(Output * ... outputs, Inputs... inputs); ``` The integer return is used for signaling errors (non-zero values) and may be discarded or propagated by the implementation. This pull request also replaces the software SHA256 implementation with a SIMD-accelerated xxHash implementation (XXH3-128, ported from #22920). SHA256 was initially chosen as the obvious choice for a collision-free hash, and it was implemented to keep dependencies minimal. For a hot JIT cache, the previous cryptographic SHA256 hash took ~88.1% (~4727ns) of the `get_kernel` function execution time. With the new non-cryptographic SIMD-accelerated hash, the hash time is reduced by 16.33x to 288ns, ~31.2% of the `get_kernel` function execution time. This pull request also: - Hoists the `cudaGetDeviceProperties` values into the cudf context object; each call takes ~1.8ms per-call (see: #23074) - Refactors librtcx's CMake functions to comply with RAPIDS' naming standards - Makes the `nvrtc`-related flags be dispatched by the `nvrtcVersion` and not `cudaGetRuntimeVersion`. They can be different - Removes `discard_errors` specialization of the transform kernel; this helps reduce the number of kernel instantiations that need to be pre-compiled, and also reduces kernel variance Closes #19578 & #23074 #### Benchmarks Throughput benchmarks are provided in #22680 (comment), and compilation-time benchmarks are provided in #22680 (comment) Authors: - Basit Ayantunde (https://github.com/lamarrr) - Bradley Dice (https://github.com/bdice) - Kyle Edwards (https://github.com/KyleFromNVIDIA) Approvers: - Bradley Dice (https://github.com/bdice) - Nghia Truong (https://github.com/ttnghia) URL: #22680
3 tasks
3 tasks
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
2 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 replaces the software SHA256 implementation with a SIMD-accelerated xxHash implementation (XXH3-128). SHA256 was initially chosen as the obvious choice for a collision-free hash and it was implemented to keep dependencies minimal.
For a hot JIT cache the previous cryptographic sha256 hash took ~88.1% (~4727ns) of the
get_kernelfunction execution time.With the new non-cryptographic SIMD-accelerated hash the hash time is reduced by 16.33x to 288ns ~31.2% of the
get_kernelfunction execution time.Changeset Summary
Checklist