[DRAFT] Add AOT JIT+LTO capability - #22390
Closed
divyegala wants to merge 16 commits into
Closed
Conversation
github-actions
Bot
added
libcudf
Contributor
|
Is this overlapping #22209 or meant as an alternative approach? |
Contributor
Author
|
@devavret it is meant as an alternative approach. This is the approach we are using in cuVS in production currently, designed by @robertmaynard, @KyleFromNVIDIA, and myself. |
…/nvjitlink_kernels
divyegala
changed the title
|
ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (3)
✅ Files skipped from review due to trivial changes (1)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughSummary by CodeRabbit
WalkthroughThis PR implements a complete JIT-LTO (Just-In-Time Link-Time Optimization) system for MurmurHash3 x86_32 hashing in CUDF. It adds build infrastructure, CMake code generation tooling, JIT kernel execution abstractions (FragmentEntry, AlgorithmLauncher, AlgorithmPlanner), CUDA device code for type-specialized hashing, and refactors the existing murmurhash3 implementation to use dynamic JIT-linked kernels instead of compile-time template instantiation. ChangesJIT-LTO MurmurHash3 Implementation
🎯 4 (Complex) | ⏱️ ~45 minutes 🚥 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)
Tip 💬 Introducing Slack Agent: The best way for teams to turn conversations into code.Slack Agent is built on CodeRabbit's deep understanding of your code, so your team can collaborate across the entire SDLC without losing context.
Built for teams:
One agent for your entire SDLC. Right inside Slack. 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: 9
🧹 Nitpick comments (7)
cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in (1)
15-25: ⚡ Quick winFail fast if the noop specialization is ever dispatched.
This path should be unreachable for a correctly planned fragment set. Returning
0here turns a fragment-selection bug into silently wrong hashes instead of an immediate failure.Suggested change
+#include <cudf/detail/utilities/assert.cuh> + template <> __device__ hash_value_type murmur_jit_hasher<@storage_cpp@>(column_device_view col, uint32_t seed, bool nullable, size_type row_index) { (void)col; (void)seed; (void)nullable; (void)row_index; - return hash_value_type{0}; + CUDF_UNREACHABLE("Unexpected dispatch to noop murmur_jit_hasher specialization."); }🤖 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/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in` around lines 15 - 25, The noop specialization murmur_jit_hasher<@storage_cpp@> should fail fast instead of returning 0; replace the silent return in murmur_jit_hasher<@storage_cpp@> with a device-side immediate failure (e.g., an assert(false) or device trap/abort) so any dispatch to this unreachable path aborts execution and surfaces the bug; ensure the failure message includes context like "noop specialization dispatched for murmur_jit_hasher" to aid debugging.cpp/src/jit_lto/AlgorithmLauncher.cpp (1)
40-48: 💤 Low valueUse
nullptrinstead ofNULLfor theattrspointer.🔧 Proposed fix
- config.attrs = NULL; + config.attrs = nullptr;🤖 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_lto/AlgorithmLauncher.cpp` around lines 40 - 48, Replace the C-style NULL with C++ nullptr for the cudaLaunchConfig_t attrs member in AlgorithmLauncher.cpp: locate the config variable initialization (config.gridDim, config.blockDim, config.stream, config.dynamicSmemBytes, config.numAttrs, config.attrs) and change config.attrs from NULL to nullptr so it uses modern C++ null pointer semantics before calling cudaLaunchKernelExC.cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp (1)
12-16: 💤 Low value
<string>and<unordered_map>are unused in this header.Neither
std::stringnorstd::unordered_mapappear anywhere inAlgorithmLauncher. Removing them reduces transitive include cost for every TU that includes this header.🔧 Proposed fix
`#include` <cstdint> `#include` <memory> -#include <string> `#include` <unordered_map>Actually,
<unordered_map>is also unused — remove both:`#include` <cstdint> `#include` <memory> -#include <string> -#include <unordered_map>As per coding guidelines, avoid unnecessary includes in headers that add compilation overhead.
🤖 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/detail/jit_lto/AlgorithmLauncher.hpp` around lines 12 - 16, The header AlgorithmLauncher.hpp unnecessarily includes <string> and <unordered_map>; remove those two include lines from the top of the file and rebuild to ensure nothing in AlgorithmLauncher (and its declarations) depends on std::string or std::unordered_map; if any use appears, replace with a forward declaration or move the include into the corresponding .cpp where the types are actually needed.cpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hpp (1)
36-40: ⚡ Quick winPrefer C++20
requiresclause overstd::enable_if_tSFINAE.♻️ Proposed refactor
- template <typename T, typename = std::enable_if_t<std::is_convertible_v<T*, FragmentEntry*>>> - void add_fragment(std::unique_ptr<T> fragment) + template <typename T> + requires std::is_convertible_v<T*, FragmentEntry*> + void add_fragment(std::unique_ptr<T> fragment)As per coding guidelines, use C++20
requiresclauses for type-gating instead ofCUDF_ENABLE_IFor its equivalentstd::enable_if_t.🤖 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/detail/jit_lto/AlgorithmPlanner.hpp` around lines 36 - 40, The template for add_fragment uses SFINAE via std::enable_if_t; change it to a C++20 constrained template using a requires clause to follow the codebase guideline. Update the function template signature for add_fragment to remove the default enable_if parameter and add a requires condition like requires(std::is_convertible_v<T*, FragmentEntry*>) (referencing the add_fragment template and the FragmentEntry type) so the function only accepts types convertible to FragmentEntry*, leaving the implementation that pushes the unique_ptr into fragments unchanged.cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh (1)
130-133: 💤 Low valueUse
size_typeinstead ofintfor the loop variable.🔧 Proposed fix
- for (int i = 0; i < curr_col.size(); ++i) { + for (size_type i = 0; i < curr_col.size(); ++i) {As per coding guidelines,
cudf::size_type(signed 32-bit) should be used for sizes.🤖 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/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh` around lines 130 - 133, The for-loop uses a plain int for indexing curr_col.size(); change the loop variable to cudf::size_type to match cudf sizing conventions and avoid type-mismatch: replace "int i" with "cudf::size_type i" in the loop that updates hash by calling murmur_jit_hash_dispatcher(curr_col, _element_hasher.seed(), _check_nulls, i), ensuring all references (curr_col, murmur_jit_hash_dispatcher, _element_hasher, _check_nulls, hash) remain unchanged.cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp (2)
7-7: ⚡ Quick winDrop the unused project include from this header.
At Line 7, this header is not referenced by any declaration here, so it adds avoidable transitive compile overhead.
Proposed fix
-#include <cudf/detail/jit_lto/nvjitlink_checker.hpp>As per coding guidelines, "Avoid unnecessary includes in headers that add compilation overhead."
🤖 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/detail/jit_lto/FragmentEntry.hpp` at line 7, The header FragmentEntry.hpp unnecessarily includes cudf/detail/jit_lto/nvjitlink_checker.hpp which is not referenced by any declaration in FragmentEntry.hpp; remove the include directive for nvjitlink_checker.hpp from FragmentEntry.hpp to eliminate the transitive compile dependency and rely on source files that need nvjitlink_checker.hpp to include it instead.
24-30: ⚡ Quick winMark pure getter APIs
[[nodiscard]].At Line 24, Line 28, Line 30, Line 37, Line 39, and Line 41, these side-effect-free non-void getters should be
[[nodiscard]]to prevent accidental ignored results.Proposed fix
- virtual const char* get_key() const = 0; + [[nodiscard]] virtual const char* get_key() const = 0; @@ - virtual const uint8_t* get_data() const = 0; + [[nodiscard]] virtual const uint8_t* get_data() const = 0; @@ - virtual size_t get_length() const = 0; + [[nodiscard]] virtual size_t get_length() const = 0; @@ - const uint8_t* get_data() const override { return StaticFatbinFragmentEntry<FragmentTag>::data; } + [[nodiscard]] const uint8_t* get_data() const override + { + return StaticFatbinFragmentEntry<FragmentTag>::data; + } @@ - size_t get_length() const override { return StaticFatbinFragmentEntry<FragmentTag>::length; } + [[nodiscard]] size_t get_length() const override + { + return StaticFatbinFragmentEntry<FragmentTag>::length; + } @@ - const char* get_key() const override + [[nodiscard]] const char* get_key() const override { return typeid(StaticFatbinFragmentEntry<FragmentTag>).name(); }As per coding guidelines, "Add [[nodiscard]] attribute to side-effect-free functions with non-void return types."
Also applies to: 37-43
🤖 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/detail/jit_lto/FragmentEntry.hpp` around lines 24 - 30, Mark the pure getter APIs as [[nodiscard]] so callers can't accidentally ignore their return values: add the [[nodiscard]] attribute to the virtual getters in FragmentEntry and FatbinFragmentEntry (specifically FragmentEntry::get_key, FatbinFragmentEntry::get_data, and FatbinFragmentEntry::get_length) and any other side-effect-free non-void virtual getters in the file (the ones reported around lines 37–43). Update the function declarations to include [[nodiscard]] on the getter signatures (apply the attribute to the function declaration), rebuild and run tests to ensure no call sites break from now-obligatory use of returned values.
🤖 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/cmake/Modules/compute_matrix_product.cmake`:
- Around line 20-33: The current logic in compute_matrix_product.cmake prefers
_JIT_LTO_MATRIX_JSON_FILE when both inputs are present and falls back to stdin
when neither is provided; change this to require exactly one input by adding a
validation block that checks the presence of _JIT_LTO_MATRIX_JSON_FILE and
_JIT_LTO_MATRIX_JSON_STRING (use the same variable names) and calls
message(FATAL_ERROR ...) if both are set or both are unset, otherwise proceed
with the existing execute_process branches (the branches that call
"${Python3_EXECUTABLE}"
"${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" with either the
file path or "-" for stdin).
In `@cpp/cmake/Modules/compute_matrix_product.py`:
- Around line 155-173: The warnings.warn calls inside iterate_next_dimension
(the blocks that emit UsedKeyWarning and UnusedKeyWarning using variables
warn_used/warn_unused) need an explicit stacklevel so the warning points to the
caller; update both warnings.warn invocations to include stacklevel=2 (or a
higher value if you determine deeper call depth from iterate_matrix_product)
while keeping the existing message and category arguments intact.
In `@cpp/cmake/Modules/generate_jit_lto_kernels.cmake`:
- Around line 94-99: The find_program invocation that looks up bin_to_c (NAMES
bin2c, PATHS ${CUDAToolkit_BIN_DIR}) should be made required so configuration
fails fast if the tool is missing; update the find_program call for variable
bin_to_c to include the REQUIRED keyword (in generate_jit_lto_kernels.cmake) so
CMake will error at configure time rather than allowing a later build-time
failure when ${bin_to_c} is used.
In `@cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh`:
- Around line 109-133: The current STRUCT handling uses
structs_column_device_view(curr_col).get_sliced_child(0) which skips children
1..n and yields incorrect hashes; change the code so when curr_col.type().id()
== type_id::STRUCT you iterate over all child columns (for j from 0 to
curr_col.num_child_columns()-1), obtain each child via
structs_column_device_view(curr_col).get_sliced_child(j) and combine each
child’s hash into the running hash (using the same combine logic and
murmur_jit_hash_dispatcher(seed(), _check_nulls, index) or by recursing/pushing
onto the processing stack), rather than unconditionally descending into child 0;
update the loop around curr_col/for i to ensure every struct child contributes
to the final hash.
In `@cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp`:
- Around line 26-32: The anonymous namespace function
murmur_jit_launcher_cache() defines a function-local static LauncherJitCache in
a header which can produce one-cache-per-TU fragmentation; move the cache
accessor and the related helper free functions (collect_*, add_*) out of the
header into a new or existing .cpp implementation file so there is a single
process-wide cache, and in the header expose only a non-anonymous
internal-linkage declaration (or a single inline thin accessor that forwards to
the .cpp implementation) or add a clear comment requiring the header be included
by exactly one TU; update all callers to use the new non-anonymous accessor
(murmur_jit_launcher_cache) or the relocated collect_* / add_* symbols.
- Around line 38-67: The static array murmur_jit_hasher_type_ids currently
declares size 30 but only contains 28 initializers, causing two implicit EMPTY
entries that silently hit the default branch in add_strong_hasher_fragment and
add_noop_hasher_fragment; fix this by making the declaration match the actual
initializers (change std::array<type_id, 30> to std::array<type_id, 28>) or, if
two additional types were intended, add the explicit type_id entries (or
explicit type_id::EMPTY markers) so the array length and contents are correct
and unambiguous.
In `@cpp/src/jit_lto/AlgorithmLauncher.cpp`:
- Around line 25-34: The move-assignment operator
AlgorithmLauncher::operator=(AlgorithmLauncher&& other) noexcept should suppress
the return value from cudaLibraryUnload just like the destructor does; change
the call that currently reads cudaLibraryUnload(library) to explicitly discard
its result (e.g., cast to (void)) before proceeding to assign kernel and library
and null out other.kernel/other.library so the noexcept semantics are
consistent.
In `@cpp/src/jit_lto/AlgorithmPlanner.cpp`:
- Around line 39-45: The current get_fragments_key() concatenates
fragment->get_key() strings directly causing ambiguous keys (e.g., ["ab","c"] vs
["a","bc"]); update AlgorithmPlanner::get_fragments_key to produce an
unambiguous encoding by either adding a delimiter with escaping or, better,
length-prefixing each fragment key (e.g., append the length then a separator
then the key) when iterating over this->fragments so each fragment boundary is
preserved; ensure you reference get_fragments_key and fragment->get_key when
making the change and keep the resulting string deterministic.
- Around line 87-120: Create RAII wrappers so nvJitLinkHandle and cudaLibrary_t
are always cleaned on exceptions: wrap nvJitLinkHandle (created by
nvJitLinkCreate) in a small guard class (e.g., NvJitLinkGuard) that holds the
handle and calls nvJitLinkDestroy(handle) in its destructor and provides
release() to relinquish ownership before returning the AlgorithmLauncher;
similarly, wrap the cudaLibrary_t returned by cudaLibraryLoadData in a
CudaLibraryGuard that calls the appropriate unload function (cudaLibraryUnload
or equivalent) in its destructor and also provides release() to transfer
ownership to AlgorithmLauncher; then replace the raw nvJitLinkHandle and
cudaLibrary_t locals with these guards, call frag->add_to(handle) via
guard.get(), and call guard.release() when constructing and returning
std::make_shared<AlgorithmLauncher>(kernel, library) so the destructor cleanup
is suppressed on success.
---
Nitpick comments:
In `@cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp`:
- Around line 12-16: The header AlgorithmLauncher.hpp unnecessarily includes
<string> and <unordered_map>; remove those two include lines from the top of the
file and rebuild to ensure nothing in AlgorithmLauncher (and its declarations)
depends on std::string or std::unordered_map; if any use appears, replace with a
forward declaration or move the include into the corresponding .cpp where the
types are actually needed.
In `@cpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hpp`:
- Around line 36-40: The template for add_fragment uses SFINAE via
std::enable_if_t; change it to a C++20 constrained template using a requires
clause to follow the codebase guideline. Update the function template signature
for add_fragment to remove the default enable_if parameter and add a requires
condition like requires(std::is_convertible_v<T*, FragmentEntry*>) (referencing
the add_fragment template and the FragmentEntry type) so the function only
accepts types convertible to FragmentEntry*, leaving the implementation that
pushes the unique_ptr into fragments unchanged.
In `@cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp`:
- Line 7: The header FragmentEntry.hpp unnecessarily includes
cudf/detail/jit_lto/nvjitlink_checker.hpp which is not referenced by any
declaration in FragmentEntry.hpp; remove the include directive for
nvjitlink_checker.hpp from FragmentEntry.hpp to eliminate the transitive compile
dependency and rely on source files that need nvjitlink_checker.hpp to include
it instead.
- Around line 24-30: Mark the pure getter APIs as [[nodiscard]] so callers can't
accidentally ignore their return values: add the [[nodiscard]] attribute to the
virtual getters in FragmentEntry and FatbinFragmentEntry (specifically
FragmentEntry::get_key, FatbinFragmentEntry::get_data, and
FatbinFragmentEntry::get_length) and any other side-effect-free non-void virtual
getters in the file (the ones reported around lines 37–43). Update the function
declarations to include [[nodiscard]] on the getter signatures (apply the
attribute to the function declaration), rebuild and run tests to ensure no call
sites break from now-obligatory use of returned values.
In `@cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in`:
- Around line 15-25: The noop specialization murmur_jit_hasher<@storage_cpp@>
should fail fast instead of returning 0; replace the silent return in
murmur_jit_hasher<@storage_cpp@> with a device-side immediate failure (e.g., an
assert(false) or device trap/abort) so any dispatch to this unreachable path
aborts execution and surfaces the bug; ensure the failure message includes
context like "noop specialization dispatched for murmur_jit_hasher" to aid
debugging.
In `@cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh`:
- Around line 130-133: The for-loop uses a plain int for indexing
curr_col.size(); change the loop variable to cudf::size_type to match cudf
sizing conventions and avoid type-mismatch: replace "int i" with
"cudf::size_type i" in the loop that updates hash by calling
murmur_jit_hash_dispatcher(curr_col, _element_hasher.seed(), _check_nulls, i),
ensuring all references (curr_col, murmur_jit_hash_dispatcher, _element_hasher,
_check_nulls, hash) remain unchanged.
In `@cpp/src/jit_lto/AlgorithmLauncher.cpp`:
- Around line 40-48: Replace the C-style NULL with C++ nullptr for the
cudaLaunchConfig_t attrs member in AlgorithmLauncher.cpp: locate the config
variable initialization (config.gridDim, config.blockDim, config.stream,
config.dynamicSmemBytes, config.numAttrs, config.attrs) and change config.attrs
from NULL to nullptr so it uses modern C++ null pointer semantics before calling
cudaLaunchKernelExC.
🪄 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: a6092813-a1ad-410e-bdc1-430c12c93cc0
📒 Files selected for processing (27)
ci/build_wheel_libcudf.shcpp/CMakeLists.txtcpp/cmake/Modules/compute_matrix_product.cmakecpp/cmake/Modules/compute_matrix_product.pycpp/cmake/Modules/generate_jit_lto_kernels.cmakecpp/cmake/Modules/register_fatbin.cpp.incpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hppcpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hppcpp/include/cudf/detail/jit_lto/FragmentEntry.hppcpp/include/cudf/detail/jit_lto/nvjitlink_checker.hppcpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hppcpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuhcpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuhcpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuhcpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.incpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.jsoncpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.incpp/src/hash/jit_lto_kernels/murmurhash_entry_matrix.jsoncpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.incpp/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.jsoncpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.incpp/src/hash/murmurhash3_x86_32.cucpp/src/hash/murmurhash3_x86_32_jit_launch.hppcpp/src/jit_lto/AlgorithmLauncher.cppcpp/src/jit_lto/AlgorithmPlanner.cppcpp/src/jit_lto/FragmentEntry.cppcpp/src/jit_lto/nvjitlink_checker.cpp
Comment on lines
+20
to
+33
| if(_JIT_LTO_MATRIX_JSON_FILE) | ||
| execute_process( | ||
| COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" | ||
| "${_JIT_LTO_MATRIX_JSON_FILE}" # | ||
| OUTPUT_VARIABLE output COMMAND_ERROR_IS_FATAL ANY | ||
| ) | ||
| else() | ||
| execute_process( | ||
| COMMAND ${CMAKE_COMMAND} -E echo "${_JIT_LTO_MATRIX_JSON_STRING}" | ||
| COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" | ||
| - | ||
| OUTPUT_VARIABLE output COMMAND_ERROR_IS_FATAL ANY | ||
| ) | ||
| endif() |
There was a problem hiding this comment.
Require exactly one matrix input source.
compute_matrix_product(...) silently prefers MATRIX_JSON_FILE when both inputs are set, and it falls through to the stdin path when neither is set. Both cases make configuration failures harder to diagnose and can generate the wrong matrix product.
Suggested change
+ if((_JIT_LTO_MATRIX_JSON_FILE AND _JIT_LTO_MATRIX_JSON_STRING) OR
+ (NOT _JIT_LTO_MATRIX_JSON_FILE AND NOT _JIT_LTO_MATRIX_JSON_STRING))
+ message(FATAL_ERROR
+ "compute_matrix_product requires exactly one of MATRIX_JSON_FILE or MATRIX_JSON_STRING")
+ endif()
+
if(_JIT_LTO_MATRIX_JSON_FILE)
execute_process(
COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py"
"${_JIT_LTO_MATRIX_JSON_FILE}" #📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| if(_JIT_LTO_MATRIX_JSON_FILE) | |
| execute_process( | |
| COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" | |
| "${_JIT_LTO_MATRIX_JSON_FILE}" # | |
| OUTPUT_VARIABLE output COMMAND_ERROR_IS_FATAL ANY | |
| ) | |
| else() | |
| execute_process( | |
| COMMAND ${CMAKE_COMMAND} -E echo "${_JIT_LTO_MATRIX_JSON_STRING}" | |
| COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" | |
| - | |
| OUTPUT_VARIABLE output COMMAND_ERROR_IS_FATAL ANY | |
| ) | |
| endif() | |
| if((_JIT_LTO_MATRIX_JSON_FILE AND _JIT_LTO_MATRIX_JSON_STRING) OR | |
| (NOT _JIT_LTO_MATRIX_JSON_FILE AND NOT _JIT_LTO_MATRIX_JSON_STRING)) | |
| message(FATAL_ERROR | |
| "compute_matrix_product requires exactly one of MATRIX_JSON_FILE or MATRIX_JSON_STRING") | |
| endif() | |
| if(_JIT_LTO_MATRIX_JSON_FILE) | |
| execute_process( | |
| COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" | |
| "${_JIT_LTO_MATRIX_JSON_FILE}" # | |
| OUTPUT_VARIABLE output COMMAND_ERROR_IS_FATAL ANY | |
| ) | |
| else() | |
| execute_process( | |
| COMMAND ${CMAKE_COMMAND} -E echo "${_JIT_LTO_MATRIX_JSON_STRING}" | |
| COMMAND "${Python3_EXECUTABLE}" "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" | |
| - | |
| OUTPUT_VARIABLE output COMMAND_ERROR_IS_FATAL ANY | |
| ) | |
| endif() |
🤖 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/cmake/Modules/compute_matrix_product.cmake` around lines 20 - 33, The
current logic in compute_matrix_product.cmake prefers _JIT_LTO_MATRIX_JSON_FILE
when both inputs are present and falls back to stdin when neither is provided;
change this to require exactly one input by adding a validation block that
checks the presence of _JIT_LTO_MATRIX_JSON_FILE and _JIT_LTO_MATRIX_JSON_STRING
(use the same variable names) and calls message(FATAL_ERROR ...) if both are set
or both are unset, otherwise proceed with the existing execute_process branches
(the branches that call "${Python3_EXECUTABLE}"
"${CMAKE_CURRENT_FUNCTION_LIST_DIR}/compute_matrix_product.py" with either the
file path or "-" for stdin).
Comment on lines
+155
to
+173
| f"[{json.dumps(i)}]" for i in path[:-1] | ||
| ) | ||
|
|
||
| if warn_used and used and underscores: | ||
| warnings.warn( | ||
| f"Key {json.dumps(last)} at root{path_repr} " | ||
| f"is used in a matrix product entry even though it " | ||
| f"begins with {json.dumps(underscores)}. Consider " | ||
| f"renaming it to {json.dumps(rest)} to indicate this.", | ||
| category=UsedKeyWarning, | ||
| ) | ||
| elif warn_unused and not used and not underscores: | ||
| warnings.warn( | ||
| f"Key {json.dumps(last)} at root{path_repr} " | ||
| f"is never used in a matrix product entry and is used " | ||
| f"only for grouping. Consider renaming it to " | ||
| f"{json.dumps(f'_{last}')} to indicate this.", | ||
| category=UnusedKeyWarning, | ||
| ) |
There was a problem hiding this comment.
warnings.warn() calls are missing an explicit stacklevel argument (Ruff B028).
Without stacklevel, the warning points into the closure rather than the caller's frame, making the diagnostic difficult to act on. Since iterate_next_dimension is one level below iterate_matrix_product, stacklevel=2 is a reasonable baseline (though a higher value may be needed depending on the actual call depth from user code).
🔧 Proposed fix
warnings.warn(
f"Key {json.dumps(last)} at root{path_repr} "
f"is used in a matrix product entry even though it "
f"begins with {json.dumps(underscores)}. Consider "
f"renaming it to {json.dumps(rest)} to indicate this.",
category=UsedKeyWarning,
+ stacklevel=2,
)
elif warn_unused and not used and not underscores:
warnings.warn(
f"Key {json.dumps(last)} at root{path_repr} "
f"is never used in a matrix product entry and is used "
f"only for grouping. Consider renaming it to "
f"{json.dumps(f'_{last}')} to indicate this.",
category=UnusedKeyWarning,
+ stacklevel=2,
)📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| f"[{json.dumps(i)}]" for i in path[:-1] | |
| ) | |
| if warn_used and used and underscores: | |
| warnings.warn( | |
| f"Key {json.dumps(last)} at root{path_repr} " | |
| f"is used in a matrix product entry even though it " | |
| f"begins with {json.dumps(underscores)}. Consider " | |
| f"renaming it to {json.dumps(rest)} to indicate this.", | |
| category=UsedKeyWarning, | |
| ) | |
| elif warn_unused and not used and not underscores: | |
| warnings.warn( | |
| f"Key {json.dumps(last)} at root{path_repr} " | |
| f"is never used in a matrix product entry and is used " | |
| f"only for grouping. Consider renaming it to " | |
| f"{json.dumps(f'_{last}')} to indicate this.", | |
| category=UnusedKeyWarning, | |
| ) | |
| f"[{json.dumps(i)}]" for i in path[:-1] | |
| ) | |
| if warn_used and used and underscores: | |
| warnings.warn( | |
| f"Key {json.dumps(last)} at root{path_repr} " | |
| f"is used in a matrix product entry even though it " | |
| f"begins with {json.dumps(underscores)}. Consider " | |
| f"renaming it to {json.dumps(rest)} to indicate this.", | |
| category=UsedKeyWarning, | |
| stacklevel=2, | |
| ) | |
| elif warn_unused and not used and not underscores: | |
| warnings.warn( | |
| f"Key {json.dumps(last)} at root{path_repr} " | |
| f"is never used in a matrix product entry and is used " | |
| f"only for grouping. Consider renaming it to " | |
| f"{json.dumps(f'_{last}')} to indicate this.", | |
| category=UnusedKeyWarning, | |
| stacklevel=2, | |
| ) |
🧰 Tools
🪛 Ruff (0.15.12)
[warning] 159-159: No explicit stacklevel keyword argument found
Set stacklevel=2
(B028)
[warning] 167-167: No explicit stacklevel keyword argument found
Set stacklevel=2
(B028)
🤖 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/cmake/Modules/compute_matrix_product.py` around lines 155 - 173, The
warnings.warn calls inside iterate_next_dimension (the blocks that emit
UsedKeyWarning and UnusedKeyWarning using variables warn_used/warn_unused) need
an explicit stacklevel so the warning points to the caller; update both
warnings.warn invocations to include stacklevel=2 (or a higher value if you
determine deeper call depth from iterate_matrix_product) while keeping the
existing message and category arguments intact.
Comment on lines
+94
to
+99
| find_program( | ||
| bin_to_c | ||
| NAMES bin2c | ||
| PATHS ${CUDAToolkit_BIN_DIR} | ||
| ) | ||
|
|
There was a problem hiding this comment.
🧩 Analysis chain
🏁 Script executed:
find /cpp -name "generate_jit_lto_kernels.cmake" -type fRepository: rapidsai/cudf
Length of output: 97
🏁 Script executed:
cat -n cpp/cmake/Modules/generate_jit_lto_kernels.cmake | head -110Repository: rapidsai/cudf
Length of output: 5285
🏁 Script executed:
cat -n cpp/cmake/Modules/generate_jit_lto_kernels.cmake | tail -n +90Repository: rapidsai/cudf
Length of output: 1803
Add REQUIRED to find_program call for bin2c.
The find_program call for bin2c lacks the REQUIRED keyword. If the tool is unavailable, configuration succeeds but the build fails later when the custom command at line 37 attempts to execute ${bin_to_c}. Adding REQUIRED ensures a clear configure-time error instead of a confusing build-time failure.
Proposed fix
find_program(
bin_to_c
NAMES bin2c
PATHS ${CUDAToolkit_BIN_DIR}
+ REQUIRED
)🤖 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/cmake/Modules/generate_jit_lto_kernels.cmake` around lines 94 - 99, The
find_program invocation that looks up bin_to_c (NAMES bin2c, PATHS
${CUDAToolkit_BIN_DIR}) should be made required so configuration fails fast if
the tool is missing; update the find_program call for variable bin_to_c to
include the REQUIRED keyword (in generate_jit_lto_kernels.cmake) so CMake will
error at configure time rather than allowing a later build-time failure when
${bin_to_c} is used.
Comment on lines
+109
to
+133
| while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) { | ||
| if (_check_nulls) { | ||
| auto validity_it = cudf::detail::make_validity_iterator<true>(curr_col); | ||
| hash = cudf::detail::accumulate( | ||
| validity_it, validity_it + curr_col.size(), hash, [](auto h, auto is_valid) { | ||
| return cudf::hashing::detail::hash_combine(h, is_valid ? NON_NULL_HASH : NULL_HASH); | ||
| }); | ||
| } | ||
| if (curr_col.type().id() == type_id::STRUCT) { | ||
| if (curr_col.num_child_columns() == 0) { return hash; } | ||
| curr_col = cudf::detail::structs_column_device_view(curr_col).get_sliced_child(0); | ||
| } else if (curr_col.type().id() == type_id::LIST) { | ||
| auto list_col = cudf::detail::lists_column_device_view(curr_col); | ||
| auto list_sizes = cudf::make_list_size_iterator(list_col); | ||
| hash = cudf::detail::accumulate( | ||
| list_sizes, list_sizes + list_col.size(), hash, [](auto h, auto size) { | ||
| return cudf::hashing::detail::hash_combine(h, MurmurHash3_x86_32<size_type>{}(size)); | ||
| }); | ||
| curr_col = list_col.get_sliced_child(); | ||
| } | ||
| } | ||
| for (int i = 0; i < curr_col.size(); ++i) { | ||
| hash = cudf::hashing::detail::hash_combine( | ||
| hash, murmur_jit_hash_dispatcher(curr_col, _element_hasher.seed(), _check_nulls, i)); | ||
| } |
There was a problem hiding this comment.
STRUCT hashing silently skips all children beyond child 0 — incorrect hash for multi-field structs.
get_sliced_child(0) is unconditionally used on every STRUCT iteration, so children at index 1, 2, … are never reached. For a STRUCT(INT32, STRING), only the INT32 child contributes to the hash. This deviates from the non-JIT murmurhash3_x86_32 implementation, which combines hashes across all children, and would fail any benchmark comparison with the reference path for multi-child structs.
The while-loop approach needs to be changed to iterate over every child of a STRUCT and combine their hashes:
🐛 Proposed structural fix for STRUCT child iteration
if (curr_col.type().id() == type_id::STRUCT) {
if (curr_col.num_child_columns() == 0) { return hash; }
- curr_col = cudf::detail::structs_column_device_view(curr_col).get_sliced_child(0);
+ auto struct_cv = cudf::detail::structs_column_device_view(curr_col);
+ for (size_type ci = 1; ci < curr_col.num_child_columns(); ++ci) {
+ auto sibling = struct_cv.get_sliced_child(ci);
+ for (size_type si = 0; si < sibling.size(); ++si) {
+ hash = cudf::hashing::detail::hash_combine(
+ hash, murmur_jit_hash_dispatcher(sibling, _element_hasher.seed(), _check_nulls, si));
+ }
+ }
+ curr_col = struct_cv.get_sliced_child(0);
}Note: this still only fully recurses into child 0 via the while loop; a fully general fix requires replacing the iterative scheme with recursion or an explicit stack over all children.
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) { | |
| if (_check_nulls) { | |
| auto validity_it = cudf::detail::make_validity_iterator<true>(curr_col); | |
| hash = cudf::detail::accumulate( | |
| validity_it, validity_it + curr_col.size(), hash, [](auto h, auto is_valid) { | |
| return cudf::hashing::detail::hash_combine(h, is_valid ? NON_NULL_HASH : NULL_HASH); | |
| }); | |
| } | |
| if (curr_col.type().id() == type_id::STRUCT) { | |
| if (curr_col.num_child_columns() == 0) { return hash; } | |
| curr_col = cudf::detail::structs_column_device_view(curr_col).get_sliced_child(0); | |
| } else if (curr_col.type().id() == type_id::LIST) { | |
| auto list_col = cudf::detail::lists_column_device_view(curr_col); | |
| auto list_sizes = cudf::make_list_size_iterator(list_col); | |
| hash = cudf::detail::accumulate( | |
| list_sizes, list_sizes + list_col.size(), hash, [](auto h, auto size) { | |
| return cudf::hashing::detail::hash_combine(h, MurmurHash3_x86_32<size_type>{}(size)); | |
| }); | |
| curr_col = list_col.get_sliced_child(); | |
| } | |
| } | |
| for (int i = 0; i < curr_col.size(); ++i) { | |
| hash = cudf::hashing::detail::hash_combine( | |
| hash, murmur_jit_hash_dispatcher(curr_col, _element_hasher.seed(), _check_nulls, i)); | |
| } | |
| while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) { | |
| if (_check_nulls) { | |
| auto validity_it = cudf::detail::make_validity_iterator<true>(curr_col); | |
| hash = cudf::detail::accumulate( | |
| validity_it, validity_it + curr_col.size(), hash, [](auto h, auto is_valid) { | |
| return cudf::hashing::detail::hash_combine(h, is_valid ? NON_NULL_HASH : NULL_HASH); | |
| }); | |
| } | |
| if (curr_col.type().id() == type_id::STRUCT) { | |
| if (curr_col.num_child_columns() == 0) { return hash; } | |
| auto struct_cv = cudf::detail::structs_column_device_view(curr_col); | |
| for (size_type ci = 1; ci < curr_col.num_child_columns(); ++ci) { | |
| auto sibling = struct_cv.get_sliced_child(ci); | |
| for (size_type si = 0; si < sibling.size(); ++si) { | |
| hash = cudf::hashing::detail::hash_combine( | |
| hash, murmur_jit_hash_dispatcher(sibling, _element_hasher.seed(), _check_nulls, si)); | |
| } | |
| } | |
| curr_col = struct_cv.get_sliced_child(0); | |
| } else if (curr_col.type().id() == type_id::LIST) { | |
| auto list_col = cudf::detail::lists_column_device_view(curr_col); | |
| auto list_sizes = cudf::make_list_size_iterator(list_col); | |
| hash = cudf::detail::accumulate( | |
| list_sizes, list_sizes + list_col.size(), hash, [](auto h, auto size) { | |
| return cudf::hashing::detail::hash_combine(h, MurmurHash3_x86_32<size_type>{}(size)); | |
| }); | |
| curr_col = list_col.get_sliced_child(); | |
| } | |
| } | |
| for (int i = 0; i < curr_col.size(); ++i) { | |
| hash = cudf::hashing::detail::hash_combine( | |
| hash, murmur_jit_hash_dispatcher(curr_col, _element_hasher.seed(), _check_nulls, i)); | |
| } |
🤖 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/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh` around lines 109 -
133, The current STRUCT handling uses
structs_column_device_view(curr_col).get_sliced_child(0) which skips children
1..n and yields incorrect hashes; change the code so when curr_col.type().id()
== type_id::STRUCT you iterate over all child columns (for j from 0 to
curr_col.num_child_columns()-1), obtain each child via
structs_column_device_view(curr_col).get_sliced_child(j) and combine each
child’s hash into the running hash (using the same combine logic and
murmur_jit_hash_dispatcher(seed(), _check_nulls, index) or by recursing/pushing
onto the processing stack), rather than unconditionally descending into child 0;
update the loop around curr_col/for i to ensure every struct child contributes
to the final hash.
Comment on lines
+26
to
+32
| namespace { | ||
|
|
||
| inline cudf::detail::jit_lto::LauncherJitCache& murmur_jit_launcher_cache() | ||
| { | ||
| static cudf::detail::jit_lto::LauncherJitCache cache; | ||
| return cache; | ||
| } |
There was a problem hiding this comment.
Anonymous namespace with a function-local static in a .hpp header risks cache fragmentation.
murmur_jit_launcher_cache() contains a static LauncherJitCache cache;. If this header is ever included by more than one translation unit (e.g. a benchmark or test TU), each TU obtains its own copy of the static, breaking the intended per-process cache and causing redundant nvJitLink builds. Even though today it may only be included from murmurhash3_x86_32.cu, the pattern is fragile.
Move the cache accessor and the helper free functions (collect_*, add_*) into a .cpp file, and expose only what is needed via a non-anonymous internal-linkage declaration, or guard with a comment that this header must be included by exactly one TU.
🤖 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/hash/murmurhash3_x86_32_jit_launch.hpp` around lines 26 - 32, The
anonymous namespace function murmur_jit_launcher_cache() defines a
function-local static LauncherJitCache in a header which can produce
one-cache-per-TU fragmentation; move the cache accessor and the related helper
free functions (collect_*, add_*) out of the header into a new or existing .cpp
implementation file so there is a single process-wide cache, and in the header
expose only a non-anonymous internal-linkage declaration (or a single inline
thin accessor that forwards to the .cpp implementation) or add a clear comment
requiring the header be included by exactly one TU; update all callers to use
the new non-anonymous accessor (murmur_jit_launcher_cache) or the relocated
collect_* / add_* symbols.
Comment on lines
+38
to
+67
| static constexpr std::array<type_id, 30> murmur_jit_hasher_type_ids{{ | ||
| type_id::INT8, | ||
| type_id::INT16, | ||
| type_id::INT32, | ||
| type_id::INT64, | ||
| type_id::UINT8, | ||
| type_id::UINT16, | ||
| type_id::UINT32, | ||
| type_id::UINT64, | ||
| type_id::FLOAT32, | ||
| type_id::FLOAT64, | ||
| type_id::BOOL8, | ||
| type_id::TIMESTAMP_DAYS, | ||
| type_id::TIMESTAMP_SECONDS, | ||
| type_id::TIMESTAMP_MILLISECONDS, | ||
| type_id::TIMESTAMP_MICROSECONDS, | ||
| type_id::TIMESTAMP_NANOSECONDS, | ||
| type_id::DURATION_DAYS, | ||
| type_id::DURATION_SECONDS, | ||
| type_id::DURATION_MILLISECONDS, | ||
| type_id::DURATION_MICROSECONDS, | ||
| type_id::DURATION_NANOSECONDS, | ||
| type_id::DICTIONARY32, | ||
| type_id::STRING, | ||
| type_id::LIST, | ||
| type_id::DECIMAL32, | ||
| type_id::DECIMAL64, | ||
| type_id::DECIMAL128, | ||
| type_id::STRUCT, | ||
| }}; |
There was a problem hiding this comment.
std::array<type_id, 30> has only 28 initializers — the declared size is wrong.
Counting the entries in lines 39–66 gives 28 type IDs. The trailing 2 elements default to type_id(0) (EMPTY), which silently hits default: break; in both add_strong_hasher_fragment and add_noop_hasher_fragment. The wrong size could mask a future off-by-one when adding or removing types, and makes it unclear which 2 types are intentionally omitted.
🔧 Proposed fix
-static constexpr std::array<type_id, 30> murmur_jit_hasher_type_ids{{
+static constexpr std::array<type_id, 28> murmur_jit_hasher_type_ids{{📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| static constexpr std::array<type_id, 30> murmur_jit_hasher_type_ids{{ | |
| type_id::INT8, | |
| type_id::INT16, | |
| type_id::INT32, | |
| type_id::INT64, | |
| type_id::UINT8, | |
| type_id::UINT16, | |
| type_id::UINT32, | |
| type_id::UINT64, | |
| type_id::FLOAT32, | |
| type_id::FLOAT64, | |
| type_id::BOOL8, | |
| type_id::TIMESTAMP_DAYS, | |
| type_id::TIMESTAMP_SECONDS, | |
| type_id::TIMESTAMP_MILLISECONDS, | |
| type_id::TIMESTAMP_MICROSECONDS, | |
| type_id::TIMESTAMP_NANOSECONDS, | |
| type_id::DURATION_DAYS, | |
| type_id::DURATION_SECONDS, | |
| type_id::DURATION_MILLISECONDS, | |
| type_id::DURATION_MICROSECONDS, | |
| type_id::DURATION_NANOSECONDS, | |
| type_id::DICTIONARY32, | |
| type_id::STRING, | |
| type_id::LIST, | |
| type_id::DECIMAL32, | |
| type_id::DECIMAL64, | |
| type_id::DECIMAL128, | |
| type_id::STRUCT, | |
| }}; | |
| static constexpr std::array<type_id, 28> murmur_jit_hasher_type_ids{{ | |
| type_id::INT8, | |
| type_id::INT16, | |
| type_id::INT32, | |
| type_id::INT64, | |
| type_id::UINT8, | |
| type_id::UINT16, | |
| type_id::UINT32, | |
| type_id::UINT64, | |
| type_id::FLOAT32, | |
| type_id::FLOAT64, | |
| type_id::BOOL8, | |
| type_id::TIMESTAMP_DAYS, | |
| type_id::TIMESTAMP_SECONDS, | |
| type_id::TIMESTAMP_MILLISECONDS, | |
| type_id::TIMESTAMP_MICROSECONDS, | |
| type_id::TIMESTAMP_NANOSECONDS, | |
| type_id::DURATION_DAYS, | |
| type_id::DURATION_SECONDS, | |
| type_id::DURATION_MILLISECONDS, | |
| type_id::DURATION_MICROSECONDS, | |
| type_id::DURATION_NANOSECONDS, | |
| type_id::DICTIONARY32, | |
| type_id::STRING, | |
| type_id::LIST, | |
| type_id::DECIMAL32, | |
| type_id::DECIMAL64, | |
| type_id::DECIMAL128, | |
| type_id::STRUCT, | |
| }}; |
🤖 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/hash/murmurhash3_x86_32_jit_launch.hpp` around lines 38 - 67, The
static array murmur_jit_hasher_type_ids currently declares size 30 but only
contains 28 initializers, causing two implicit EMPTY entries that silently hit
the default branch in add_strong_hasher_fragment and add_noop_hasher_fragment;
fix this by making the declaration match the actual initializers (change
std::array<type_id, 30> to std::array<type_id, 28>) or, if two additional types
were intended, add the explicit type_id entries (or explicit type_id::EMPTY
markers) so the array length and contents are correct and unambiguous.
Comment on lines
+25
to
+34
| AlgorithmLauncher& AlgorithmLauncher::operator=(AlgorithmLauncher&& other) noexcept | ||
| { | ||
| if (this != &other) { | ||
| if (library != nullptr) { cudaLibraryUnload(library); } | ||
| kernel = other.kernel; | ||
| library = other.library; | ||
| other.kernel = nullptr; | ||
| other.library = nullptr; | ||
| } | ||
| return *this; |
There was a problem hiding this comment.
Move-assignment operator doesn't suppress the cudaLibraryUnload return value — inconsistent with the destructor.
The destructor casts the return to (void), which is the conventional way to intentionally discard the result in a noexcept context. The move-assignment operator should do the same.
🔧 Proposed fix
if (this != &other) {
- if (library != nullptr) { cudaLibraryUnload(library); }
+ if (library != nullptr) { (void)cudaLibraryUnload(library); }🤖 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_lto/AlgorithmLauncher.cpp` around lines 25 - 34, The
move-assignment operator AlgorithmLauncher::operator=(AlgorithmLauncher&& other)
noexcept should suppress the return value from cudaLibraryUnload just like the
destructor does; change the call that currently reads cudaLibraryUnload(library)
to explicitly discard its result (e.g., cast to (void)) before proceeding to
assign kernel and library and null out other.kernel/other.library so the
noexcept semantics are consistent.
Comment on lines
+39
to
+45
| std::string AlgorithmPlanner::get_fragments_key() const | ||
| { | ||
| std::string key = ""; | ||
| for (const auto& fragment : this->fragments) { | ||
| key += fragment->get_key(); | ||
| } | ||
| return key; |
There was a problem hiding this comment.
Encode fragment keys unambiguously before caching launchers.
Concatenating raw fragment keys can alias distinct fragment sets to the same cache entry. For example, ["ab", "c"] and ["a", "bc"] both produce "abc", so a later lookup can reuse the wrong linked kernel.
Suggested change
std::string AlgorithmPlanner::get_fragments_key() const
{
- std::string key = "";
+ std::string key;
for (const auto& fragment : this->fragments) {
- key += fragment->get_key();
+ auto const fragment_key = fragment->get_key();
+ key += std::to_string(fragment_key.size());
+ key += ':';
+ key += fragment_key;
+ key += ';';
}
return key;
}🤖 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_lto/AlgorithmPlanner.cpp` around lines 39 - 45, The current
get_fragments_key() concatenates fragment->get_key() strings directly causing
ambiguous keys (e.g., ["ab","c"] vs ["a","bc"]); update
AlgorithmPlanner::get_fragments_key to produce an unambiguous encoding by either
adding a delimiter with escaping or, better, length-prefixing each fragment key
(e.g., append the length then a separator then the key) when iterating over
this->fragments so each fragment boundary is preserved; ensure you reference
get_fragments_key and fragment->get_key when making the change and keep the
resulting string deterministic.
Comment on lines
+87
to
+120
| nvJitLinkHandle handle; | ||
| const char* lopts[] = {"-lto", archs.c_str()}; | ||
| auto result = nvJitLinkCreate(&handle, 2, lopts); | ||
| check_nvjitlink_result(handle, result); | ||
|
|
||
| for (const auto& frag : this->fragments) { | ||
| frag->add_to(handle); | ||
| } | ||
|
|
||
| result = nvJitLinkComplete(handle); | ||
| check_nvjitlink_result(handle, result); | ||
|
|
||
| size_t cubin_size; | ||
| result = nvJitLinkGetLinkedCubinSize(handle, &cubin_size); | ||
| check_nvjitlink_result(handle, result); | ||
|
|
||
| std::unique_ptr<char[]> cubin{new char[cubin_size]}; | ||
| result = nvJitLinkGetLinkedCubin(handle, cubin.get()); | ||
| check_nvjitlink_result(handle, result); | ||
|
|
||
| result = nvJitLinkDestroy(&handle); | ||
| CUDF_EXPECTS(result == NVJITLINK_SUCCESS, "nvJitLinkDestroy failed"); | ||
|
|
||
| cudaLibrary_t library; | ||
| CUDF_CUDA_TRY( | ||
| cudaLibraryLoadData(&library, cubin.get(), nullptr, nullptr, 0, nullptr, nullptr, 0)); | ||
|
|
||
| cudaKernel_t kernel; | ||
| CUDF_CUDA_TRY(cudaLibraryGetKernel(&kernel, library, this->entrypoint.c_str())); | ||
|
|
||
| double const build_ms = duration_ms(clock::now() - t_build_start).count(); | ||
| emit_jit_lto_build_timing(build_ms); | ||
|
|
||
| return std::make_shared<AlgorithmLauncher>(kernel, library); |
There was a problem hiding this comment.
🧩 Analysis chain
🏁 Script executed:
cat -n cpp/src/jit_lto/AlgorithmPlanner.cpp | head -150 | tail -70Repository: rapidsai/cudf
Length of output: 2791
🏁 Script executed:
# Search for check_nvjitlink_result definition
rg -A 5 "check_nvjitlink_result" cpp/src/jit_lto/ --include="*.cpp" --include="*.h" --include="*.hpp"Repository: rapidsai/cudf
Length of output: 499
🏁 Script executed:
# Look for AlgorithmLauncher class definition and constructor
rg -B 5 -A 15 "class AlgorithmLauncher" cpp/src/jit_lto/ --include="*.cpp" --include="*.h" --include="*.hpp"Repository: rapidsai/cudf
Length of output: 499
🏁 Script executed:
# Check for any existing RAII wrappers or smart pointers in this file
rg "unique_ptr|shared_ptr|reference_wrapper" cpp/src/jit_lto/AlgorithmPlanner.cppRepository: rapidsai/cudf
Length of output: 342
🏁 Script executed:
# Search for nvJitLinkDestroy usage across the codebase to see patterns
rg "nvJitLinkDestroy|nvJitLinkCreate" cpp/ --include="*.cpp" --include="*.h" --include="*.hpp" -B 2 -A 2Repository: rapidsai/cudf
Length of output: 499
🏁 Script executed:
rg "check_nvjitlink_result" cpp/src/jit_lto/Repository: rapidsai/cudf
Length of output: 563
🏁 Script executed:
rg "class AlgorithmLauncher" cpp/src/jit_lto/Repository: rapidsai/cudf
Length of output: 39
🏁 Script executed:
rg -A 10 "CUDF_CUDA_TRY" cpp/src/jit_lto/AlgorithmPlanner.cpp | head -20Repository: rapidsai/cudf
Length of output: 826
🏁 Script executed:
fd "AlgorithmLauncher" cpp/src/jit_lto/Repository: rapidsai/cudf
Length of output: 95
🏁 Script executed:
cat cpp/src/jit_lto/nvjitlink_checker.cppRepository: rapidsai/cudf
Length of output: 996
🏁 Script executed:
cat cpp/src/jit_lto/AlgorithmLauncher.cppRepository: rapidsai/cudf
Length of output: 1486
🏁 Script executed:
# Find CUDF_CUDA_TRY definition
rg -B 2 -A 5 "define CUDF_CUDA_TRY" cpp/Repository: rapidsai/cudf
Length of output: 813
Wrap the nvJitLink and CUDA library handles in RAII wrappers.
The nvJitLinkHandle created at line 87 is not destroyed if any exception occurs during operations at lines 90–105 (including frag->add_to(handle) and calls to check_nvjitlink_result, which throws on failure via CUDF_FAIL). Additionally, if cudaLibraryGetKernel at line 115 throws an exception after cudaLibraryLoadData successfully loads the library at line 112, the loaded cudaLibrary_t is leaked because the AlgorithmLauncher constructor is never reached to register it for cleanup in the destructor.
Per coding guidelines, ensure GPU resource cleanup occurs in all exception paths using RAII patterns.
🤖 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_lto/AlgorithmPlanner.cpp` around lines 87 - 120, Create RAII
wrappers so nvJitLinkHandle and cudaLibrary_t are always cleaned on exceptions:
wrap nvJitLinkHandle (created by nvJitLinkCreate) in a small guard class (e.g.,
NvJitLinkGuard) that holds the handle and calls nvJitLinkDestroy(handle) in its
destructor and provides release() to relinquish ownership before returning the
AlgorithmLauncher; similarly, wrap the cudaLibrary_t returned by
cudaLibraryLoadData in a CudaLibraryGuard that calls the appropriate unload
function (cudaLibraryUnload or equivalent) in its destructor and also provides
release() to transfer ownership to AlgorithmLauncher; then replace the raw
nvJitLinkHandle and cudaLibrary_t locals with these guards, call
frag->add_to(handle) via guard.get(), and call guard.release() when constructing
and returning std::make_shared<AlgorithmLauncher>(kernel, library) so the
destructor cleanup is suppressed on success.
There was a problem hiding this comment.
🧹 Nitpick comments (2)
cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh (1)
14-17: 💤 Low valueMissing declaration or include for
murmur_jit_hasherandhash_value_type.The function uses
murmur_jit_hasher<...>and returnshash_value_type, but there's no include or forward declaration for these. This header relies on being included after the hasher declaration header, which creates an implicit ordering dependency.Consider adding an include for the hasher declaration header (e.g.,
murmurhash3_x86_32_jit_hasher_decl.cuh) or adding a comment documenting the required inclusion order.🤖 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/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh` around lines 14 - 17, The header murmurhash3_x86_32_jit_dispatch_impl.cuh uses symbols murmur_jit_hasher and hash_value_type in murmur_jit_hash_dispatcher_all_types but does not declare or include them; add an explicit include of the hasher declaration header (e.g., murmurhash3_x86_32_jit_hasher_decl.cuh) at the top of this file or add forward declarations for template<class T> struct murmur_jit_hasher and using hash_value_type = ... so the types are resolved without relying on include order; alternatively add a clear comment documenting the required inclusion order and prefer the explicit include to remove the implicit dependency.cpp/CMakeLists.txt (1)
239-239: 💤 Low valueClarify the need for duplicate
find_package(CUDAToolkit).Line 234-238 already calls
rapids_find_package(CUDAToolkit REQUIRED ...). This additionalfind_packageon line 239 specifically requests thenvJitLinkcomponent. If this is intentional (e.g., becauserapids_find_packagedoesn't handleCOMPONENTSthe same way), consider adding a brief comment explaining why both calls are needed.🤖 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/CMakeLists.txt` at line 239, There are two package calls—rapids_find_package(CUDAToolkit REQUIRED ...) and a subsequent find_package(CUDAToolkit REQUIRED COMPONENTS nvJitLink)—which looks duplicated; either consolidate by adding the nvJitLink component to the prior rapids_find_package if it supports COMPONENTS, or keep the separate find_package but add a concise comment above find_package(CUDAToolkit REQUIRED COMPONENTS nvJitLink) explaining why rapids_find_package cannot/does not request nvJitLink (e.g., rapids_find_package omits COMPONENTS handling) so reviewers understand the intentional duplication.
🤖 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/CMakeLists.txt`:
- Line 239: There are two package calls—rapids_find_package(CUDAToolkit REQUIRED
...) and a subsequent find_package(CUDAToolkit REQUIRED COMPONENTS
nvJitLink)—which looks duplicated; either consolidate by adding the nvJitLink
component to the prior rapids_find_package if it supports COMPONENTS, or keep
the separate find_package but add a concise comment above
find_package(CUDAToolkit REQUIRED COMPONENTS nvJitLink) explaining why
rapids_find_package cannot/does not request nvJitLink (e.g., rapids_find_package
omits COMPONENTS handling) so reviewers understand the intentional duplication.
In `@cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh`:
- Around line 14-17: The header murmurhash3_x86_32_jit_dispatch_impl.cuh uses
symbols murmur_jit_hasher and hash_value_type in
murmur_jit_hash_dispatcher_all_types but does not declare or include them; add
an explicit include of the hasher declaration header (e.g.,
murmurhash3_x86_32_jit_hasher_decl.cuh) at the top of this file or add forward
declarations for template<class T> struct murmur_jit_hasher and using
hash_value_type = ... so the types are resolved without relying on include
order; alternatively add a clear comment documenting the required inclusion
order and prefer the explicit include to remove the implicit dependency.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 7ef33748-1642-4006-9fba-e95b97fce80e
📒 Files selected for processing (6)
cpp/CMakeLists.txtcpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hppcpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuhcpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.incpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.jsoncpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
🚧 Files skipped from review as they are similar to previous changes (2)
- cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
- cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
github-actions
Bot
added
Python
This was referenced May 28, 2026
Contributor
|
At this point I feel comfortable closing this PR in favor of #22680. We'll probably pull some parts of this PR into a new one just for adding murmurhash support once librtcx + LTO-IR is fully working in libcudf. |
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
3 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 PR is a POC to add AOT JIT+LTO capability by using the cuVS architecture. The initial POC is on
murmurhash_x86_32functionality, which uses a device-side type dispatcher. The idea behind LTO for this functionality is that we generate a hashing device function fragment for each cudf type, and then only link the fragments for the types present in the input table. For the other types, a no-op fragment is linked instead.Benchmarks:
Checklist