From 82c4bf151f6328e474814491b063ce94531607d8 Mon Sep 17 00:00:00 2001
From: Vyas Ramasubramani <vyasr@nvidia.com>
Date: Wed, 22 Apr 2026 18:40:26 -0700
Subject: [PATCH 01/11] First pass

---
 cpp/src/hash/murmurhash3_x86_32.cu      | 137 ++++++++++++++++++-
 cpp/src/hash/murmurhash3_x86_32_lto.cuh | 171 ++++++++++++++++++++++++
 2 files changed, 303 insertions(+), 5 deletions(-)
 create mode 100644 cpp/src/hash/murmurhash3_x86_32_lto.cuh

diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu
index bc0f6eb99e09..97c276ad1b2b 100644
--- a/cpp/src/hash/murmurhash3_x86_32.cu
+++ b/cpp/src/hash/murmurhash3_x86_32.cu
@@ -13,10 +13,133 @@
 
 #include <thrust/tabulate.h>
 
+#include "murmurhash3_x86_32_lto.cuh"
+
 namespace cudf {
 namespace hashing {
 namespace detail {
 
+//
+// 1. Build a device-side dispatcher that 
+
+//template <typename T>
+//hash_value_type hasher(cudf::column_device_view col, uint32_t seed);
+
+__device__  __forceinline__ constexpr decltype(auto) hash_dispatcher(cudf::column_device_view col, uint32_t seed, bool const nullable)
+{
+  switch (col.type().id()) {
+    case type_id::INT8:
+      return hasher<id_to_type<type_id::INT8>>(
+        col, seed, nullable);
+    case type_id::INT16:
+      return hasher<id_to_type<type_id::INT16>>(
+        col, seed, nullable);
+    case type_id::INT32:
+      return hasher<id_to_type<type_id::INT32>>(
+        col, seed, nullable);
+    case type_id::INT64:
+      return hasher<id_to_type<type_id::INT64>>(
+        col, seed, nullable);
+    case type_id::UINT8:
+      return hasher<id_to_type<type_id::UINT8>>(
+        col, seed, nullable);
+    case type_id::UINT16:
+      return hasher<id_to_type<type_id::UINT16>>(
+        col, seed, nullable);
+    case type_id::UINT32:
+      return hasher<id_to_type<type_id::UINT32>>(
+        col, seed, nullable);
+    case type_id::UINT64:
+      return hasher<id_to_type<type_id::UINT64>>(
+        col, seed, nullable);
+    case type_id::FLOAT32:
+      return hasher<id_to_type<type_id::FLOAT32>>(
+        col, seed, nullable);
+    case type_id::FLOAT64:
+      return hasher<id_to_type<type_id::FLOAT64>>(
+        col, seed, nullable);
+    case type_id::BOOL8:
+      return hasher<id_to_type<type_id::BOOL8>>(
+        col, seed, nullable);
+    case type_id::TIMESTAMP_DAYS:
+      return hasher<id_to_type<type_id::TIMESTAMP_DAYS>>(
+        col, seed, nullable);
+    case type_id::TIMESTAMP_SECONDS:
+      return hasher<id_to_type<type_id::TIMESTAMP_SECONDS>>(
+        col, seed, nullable);
+    case type_id::TIMESTAMP_MILLISECONDS:
+      return hasher<id_to_type<type_id::TIMESTAMP_MILLISECONDS>>(
+        col, seed, nullable);
+    case type_id::TIMESTAMP_MICROSECONDS:
+      return hasher<id_to_type<type_id::TIMESTAMP_MICROSECONDS>>(
+        col, seed, nullable);
+    case type_id::TIMESTAMP_NANOSECONDS:
+      return hasher<id_to_type<type_id::TIMESTAMP_NANOSECONDS>>(
+        col, seed, nullable);
+    case type_id::DURATION_DAYS:
+      return hasher<id_to_type<type_id::DURATION_DAYS>>(
+        col, seed, nullable);
+    case type_id::DURATION_SECONDS:
+      return hasher<id_to_type<type_id::DURATION_SECONDS>>(
+        col, seed, nullable);
+    case type_id::DURATION_MILLISECONDS:
+      return hasher<id_to_type<type_id::DURATION_MILLISECONDS>>(
+        col, seed, nullable);
+    case type_id::DURATION_MICROSECONDS:
+      return hasher<id_to_type<type_id::DURATION_MICROSECONDS>>(
+        col, seed, nullable);
+    case type_id::DURATION_NANOSECONDS:
+      return hasher<id_to_type<type_id::DURATION_NANOSECONDS>>(
+        col, seed, nullable);
+    case type_id::DICTIONARY32:
+      return hasher<id_to_type<type_id::DICTIONARY32>>(
+        col, seed, nullable);
+    case type_id::STRING:
+      return hasher<id_to_type<type_id::STRING>>(
+        col, seed, nullable);
+    case type_id::LIST:
+      return hasher<id_to_type<type_id::LIST>>(
+        col, seed, nullable);
+    case type_id::DECIMAL32:
+      return hasher<id_to_type<type_id::DECIMAL32>>(
+        col, seed, nullable);
+    case type_id::DECIMAL64:
+      return hasher<id_to_type<type_id::DECIMAL64>>(
+        col, seed, nullable);
+    case type_id::DECIMAL128:
+      return hasher<id_to_type<type_id::DECIMAL128>>(
+        col, seed, nullable);
+    case type_id::STRUCT:
+      return hasher<id_to_type<type_id::STRUCT>>(
+        col, seed, nullable);
+    default: {
+#ifndef __CUDA_ARCH__
+      CUDF_FAIL("Invalid type_id.");
+#else
+      CUDF_UNREACHABLE("Invalid type_id.");
+#endif
+    }
+  }
+}
+
+__global__ void murmurhash3_x86_32_kernel(mutable_column_device_view output,
+                                  uint32_t seed,
+                                  table_device_view const input,
+                                  bool const nullable)
+{
+  cudf::size_type idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < input.num_rows()) {
+      auto const num_cols = input.num_columns();
+      if (num_cols == 0) return;
+      hash_value_type hash_value = hash_dispatcher(input.column(0), seed, nullable);
+      for (int i = 1; i < num_cols; ++i) {
+          hash_value = cudf::hashing::detail::hash_combine(
+            hash_value, hash_dispatcher(input.column(i), seed, nullable));
+      }
+      output.element<hash_value_type>(idx) = hash_value;
+  }
+}
+
 std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
                                            uint32_t seed,
                                            rmm::cuda_stream_view stream,
@@ -36,11 +159,15 @@ std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
   auto output_view      = output->mutable_view();
 
   // Compute the hash value for each row
-  thrust::tabulate(rmm::exec_policy_nosync(stream),
-                   output_view.begin<hash_value_type>(),
-                   output_view.end<hash_value_type>(),
-                   row_hasher.device_hasher<MurmurHash3_x86_32>(nullable, seed));
-
+  //thrust::tabulate(rmm::exec_policy_nosync(stream),
+  //                 output_view.begin<hash_value_type>(),
+  //                 output_view.end<hash_value_type>(),
+  //                 row_hasher.device_hasher<MurmurHash3_x86_32>(nullable, seed));
+  //
+  auto d_output = mutable_column_device_view::create(output_view, stream);
+  auto d_input  = table_device_view::create(input, stream);
+  murmurhash3_x86_32_kernel<<<1, input.num_rows(), 0, stream.value()>>>(
+    *d_output, seed, *d_input, nullable);
   return output;
 }
 
diff --git a/cpp/src/hash/murmurhash3_x86_32_lto.cuh b/cpp/src/hash/murmurhash3_x86_32_lto.cuh
new file mode 100644
index 000000000000..98d6c636eaca
--- /dev/null
+++ b/cpp/src/hash/murmurhash3_x86_32_lto.cuh
@@ -0,0 +1,171 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "cuda/std/__type_traits/is_same.h"
+#include <cudf/detail/row_operator/hashing.cuh>
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/types.hpp>
+#include <cudf/detail/row_operator/hashing.cuh>
+#include <cudf/hashing.hpp>
+#include <cudf/structs/structs_column_device_view.cuh>
+#include <cudf/lists/lists_column_device_view.cuh>
+
+namespace cudf::hashing::detail {
+using result_type = hash_value_type;
+
+template <typename Nullate>
+class element_hasher {
+ public:
+
+  /**
+   * @brief Constructs an element_hasher object.
+   *
+   * @param nulls Indicates whether to check for nulls
+   * @param seed  The seed to use for the hash function
+   * @param null_hash The hash value to use for nulls
+   */
+  __device__ element_hasher(
+    Nullate nulls,
+    result_type seed      = DEFAULT_HASH_SEED,
+    result_type null_hash = cuda::std::numeric_limits<result_type>::max()) noexcept
+    : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
+  {
+  }
+
+  /**
+   * @brief Returns the hash value of the given element.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(column_device_view::has_element_accessor<T>())
+  {
+    if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
+    return MurmurHash3_x86_32<T>{_seed}(col.element<T>(row_index));
+  }
+
+  /**
+   * @brief Returns the hash value of the given element.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(not column_device_view::has_element_accessor<T>())
+  {
+    CUDF_UNREACHABLE("Unsupported type in hash.");
+  }
+
+  Nullate _check_nulls;
+  // Assumes seeds are the same as the result type of the hash function
+  result_type _seed;
+  result_type _null_hash;
+};
+
+template <typename Nullate>
+class element_hasher_adaptor {
+  static constexpr result_type NULL_HASH     = cuda::std::numeric_limits<result_type>::max();
+  static constexpr result_type NON_NULL_HASH = 0;
+
+ public:
+  __device__ element_hasher_adaptor(Nullate _check_nulls, result_type seed) noexcept
+    : _element_hasher(_check_nulls, seed)
+  {
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(not cudf::is_nested<T>() and not cudf::is_dictionary<T>())
+  {
+    return _element_hasher.template operator()<T>(col, row_index);
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(cudf::is_nested<T>() or cudf::is_dictionary<T>())
+  {
+    CUDF_UNREACHABLE("Can't get here yet");
+  }
+
+  //template <typename T>
+  //__device__ result_type operator()(column_device_view const& col,
+  //                                  size_type row_index) const noexcept
+  //  requires(cudf::is_dictionary<T>())
+  //{
+  //  if constexpr (Nullate) {
+  //    if (col.is_null(row_index)) { return NULL_HASH; }
+  //  }
+  //
+  //  auto const keys = col.child(dictionary_column_view::keys_column_index);
+  //  return type_dispatcher<dispatch_storage_type>(
+  //    keys.type(),
+  //    _element_hasher,
+  //    keys,
+  //    static_cast<size_type>(col.element<dictionary32>(row_index)));
+  //}
+  //
+  //template <typename T>
+  //__device__ result_type operator()(column_device_view const& col,
+  //                                  size_type row_index) const noexcept
+  //  requires(cudf::is_nested<T>())
+  //{
+  //  auto hash                   = result_type{0};
+  //  column_device_view curr_col = col.slice(row_index, 1);
+  //  while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) {
+  //    if constexpr (Nullate) {
+  //      auto validity_it = detail::make_validity_iterator<true>(curr_col);
+  //      hash             = detail::accumulate(
+  //        validity_it, validity_it + curr_col.size(), hash, [](auto hash, auto is_valid) {
+  //          return cudf::hashing::detail::hash_combine(hash,
+  //                                                     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 = detail::structs_column_device_view(curr_col).get_sliced_child(0);
+  //    } else if (curr_col.type().id() == type_id::LIST) {
+  //      auto list_col   = detail::lists_column_device_view(curr_col);
+  //      auto list_sizes = make_list_size_iterator(list_col);
+  //      hash            = detail::accumulate(
+  //        list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
+  //          return cudf::hashing::detail::hash_combine(hash, 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,
+  //      type_dispatcher<dispatch_void_if_nested>(curr_col.type(), _element_hasher, curr_col, i));
+  //  }
+  //  return hash;
+  //}
+
+  element_hasher<Nullate> const _element_hasher;
+};
+
+
+template <typename Nullate, typename T>
+__device__ hash_value_type hasher_impl(Nullate check_nulls, cudf::column_device_view col, uint32_t seed) {
+  auto hasher = element_hasher_adaptor<Nullate>{check_nulls, seed};
+  return hasher.template operator()<T>(col, threadIdx.x);
+}
+
+template <typename T>
+__device__ hash_value_type hasher(cudf::column_device_view col, uint32_t seed, bool const nullable) {
+  return hasher_impl<nullate::DYNAMIC, T>(nullate::DYNAMIC{nullable}, col, seed);
+}
+} // namespace cudf::hashing::detail

From 77c483c0624b028cbb1079704d4816f0ec6efd13 Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Wed, 6 May 2026 02:05:26 +0000
Subject: [PATCH 02/11] add jit+lto, tests pass

---
 cpp/CMakeLists.txt                            |  76 ++++-
 .../Modules/compute_matrix_product.cmake      |  53 ++++
 cpp/cmake/Modules/compute_matrix_product.py   | 243 +++++++++++++++
 .../Modules/generate_jit_lto_kernels.cmake    | 135 ++++++++
 cpp/cmake/Modules/register_fatbin.cpp.in      |  22 ++
 .../cudf/detail/jit_lto/AlgorithmLauncher.hpp |  49 +++
 .../cudf/detail/jit_lto/AlgorithmPlanner.hpp  |  57 ++++
 .../cudf/detail/jit_lto/FragmentEntry.hpp     |  67 ++++
 .../cudf/detail/jit_lto/nvjitlink_checker.hpp |  13 +
 .../detail/murmurhash3_x86_32_jit_tags.hpp    |  50 +++
 .../murmurhash3_x86_32_jit_device.cuh         |  82 +++++
 .../murmurhash3_x86_32_jit_hasher_decl.cuh    |  27 ++
 .../murmurhash3_x86_32_lto.cuh                | 162 ++++++++++
 .../murmurhash_entry_kernel.cu.in             |  33 ++
 .../murmurhash_entry_matrix.json              |   7 +
 .../murmurhash_hasher_fragment.cu.in          |  27 ++
 .../murmurhash_hasher_matrix.json             |  32 ++
 .../murmurhash_hasher_noop_fragment.cu.in     |  28 ++
 cpp/src/hash/murmurhash3_x86_32.cu            | 154 +---------
 .../hash/murmurhash3_x86_32_jit_launch.hpp    | 287 ++++++++++++++++++
 cpp/src/hash/murmurhash3_x86_32_lto.cuh       | 171 -----------
 cpp/src/jit_lto/AlgorithmLauncher.cpp         |  51 ++++
 cpp/src/jit_lto/AlgorithmPlanner.cpp          |  98 ++++++
 cpp/src/jit_lto/FragmentEntry.cpp             |  18 ++
 cpp/src/jit_lto/nvjitlink_checker.cpp         |  30 ++
 25 files changed, 1649 insertions(+), 323 deletions(-)
 create mode 100644 cpp/cmake/Modules/compute_matrix_product.cmake
 create mode 100644 cpp/cmake/Modules/compute_matrix_product.py
 create mode 100644 cpp/cmake/Modules/generate_jit_lto_kernels.cmake
 create mode 100644 cpp/cmake/Modules/register_fatbin.cpp.in
 create mode 100644 cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp
 create mode 100644 cpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hpp
 create mode 100644 cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp
 create mode 100644 cpp/include/cudf/detail/jit_lto/nvjitlink_checker.hpp
 create mode 100644 cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in
 create mode 100644 cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
 delete mode 100644 cpp/src/hash/murmurhash3_x86_32_lto.cuh
 create mode 100644 cpp/src/jit_lto/AlgorithmLauncher.cpp
 create mode 100644 cpp/src/jit_lto/AlgorithmPlanner.cpp
 create mode 100644 cpp/src/jit_lto/FragmentEntry.cpp
 create mode 100644 cpp/src/jit_lto/nvjitlink_checker.cpp

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 5a0b2f95e830..7afd47436d41 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -240,6 +240,7 @@ rapids_find_package(
   BUILD_EXPORT_SET cudf-exports
   INSTALL_EXPORT_SET cudf-exports
 )
+find_package(CUDAToolkit REQUIRED COMPONENTS nvJitLink)
 include(cmake/Modules/ConfigureCUDA.cmake) # set other CUDA compilation flags
 
 # ##################################################################################################
@@ -328,6 +329,74 @@ if(NOT BUILD_SHARED_LIBS)
   endif()
 endif()
 
+# ##################################################################################################
+# * JIT+LTO (MurmurHash3 x86 32 device fragments + nvjitlink) --------------------------------------
+# Same CMake pattern as cuVS (`cpp/CMakeLists.txt` + `cmake/modules/generate_jit_lto_kernels.cmake`).
+include(cmake/Modules/generate_jit_lto_kernels.cmake)
+
+add_library(cudf_jit_lto_kernel_usage_requirements INTERFACE)
+target_include_directories(
+  cudf_jit_lto_kernel_usage_requirements
+  INTERFACE "${CMAKE_CURRENT_SOURCE_DIR}/include"
+            "${CMAKE_CURRENT_SOURCE_DIR}/src"
+            "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels"
+)
+target_compile_options(
+  cudf_jit_lto_kernel_usage_requirements
+  INTERFACE "$<$<COMPILE_LANGUAGE:CXX>:${CUDF_CXX_FLAGS}>"
+            "$<$<COMPILE_LANGUAGE:CUDA>:${CUDF_CUDA_FLAGS}>"
+)
+target_compile_features(cudf_jit_lto_kernel_usage_requirements INTERFACE cuda_std_20)
+target_link_libraries(
+  cudf_jit_lto_kernel_usage_requirements INTERFACE CCCL::CCCL rmm::rmm
+                                                    $<BUILD_LOCAL_INTERFACE:cuco::cuco>
+)
+
+# Same arch ladder as cuVS JIT fragments (single real arch for fatbin link).
+set(JIT_LTO_TARGET_ARCHITECTURE "70-real")
+if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
+  set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
+endif()
+
+block(PROPAGATE cudf_jit_lto_generated_sources)
+set(cudf_jit_lto_generated_sources)
+set(CMAKE_CUDA_ARCHITECTURES ${JIT_LTO_TARGET_ARCHITECTURE})
+
+set(murmur_jit_ns "cudf::hashing::detail::jit_lto")
+generate_jit_lto_kernels(
+  cudf_jit_lto_generated_sources
+  NAME_FORMAT "murmurhash_jit_hasher_@abbrev@"
+  MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json"
+  KERNEL_INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in"
+  FRAGMENT_TAG_FORMAT
+    "${murmur_jit_ns}::fragment_tag_murmur_hasher<${murmur_jit_ns}::tag_@abbrev@>"
+  FRAGMENT_TAG_HEADER_FILES "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/hasher"
+  KERNEL_LINK_LIBRARIES cudf_jit_lto_kernel_usage_requirements
+)
+generate_jit_lto_kernels(
+  cudf_jit_lto_generated_sources
+  NAME_FORMAT "murmurhash_jit_hasher_noop_@abbrev@"
+  MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json"
+  KERNEL_INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in"
+  FRAGMENT_TAG_FORMAT
+    "${murmur_jit_ns}::fragment_tag_murmur_hasher_noop<${murmur_jit_ns}::tag_@abbrev@>"
+  FRAGMENT_TAG_HEADER_FILES "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/hasher_noop"
+  KERNEL_LINK_LIBRARIES cudf_jit_lto_kernel_usage_requirements
+)
+generate_jit_lto_kernels(
+  cudf_jit_lto_generated_sources
+  NAME_FORMAT "murmurhash_jit_@k@"
+  MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json"
+  KERNEL_INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in"
+  FRAGMENT_TAG_FORMAT "${murmur_jit_ns}::fragment_tag_murmur_entry"
+  FRAGMENT_TAG_HEADER_FILES "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/entry"
+  KERNEL_LINK_LIBRARIES cudf_jit_lto_kernel_usage_requirements
+)
+endblock()
+
 # ##################################################################################################
 # * library targets -------------------------------------------------------------------------------
 add_library(
@@ -460,6 +529,11 @@ add_library(
   src/groupby/sort/sort_helper.cu
   src/hash/md5_hash.cu
   src/hash/murmurhash3_x86_32.cu
+  src/jit_lto/AlgorithmLauncher.cpp
+  src/jit_lto/AlgorithmPlanner.cpp
+  src/jit_lto/FragmentEntry.cpp
+  src/jit_lto/nvjitlink_checker.cpp
+  ${cudf_jit_lto_generated_sources}
   src/hash/murmurhash3_x64_128.cu
   src/hash/sha1_hash.cu
   src/hash/sha224_hash.cu
@@ -1004,7 +1078,7 @@ target_link_libraries(
   cudf
   PUBLIC CCCL::CCCL rapids_logger::rapids_logger rmm::rmm $<BUILD_LOCAL_INTERFACE:BS::thread_pool>
   PRIVATE $<BUILD_LOCAL_INTERFACE:nvtx3::nvtx3-cpp> $<BUILD_LOCAL_INTERFACE:cuco::cuco> ZLIB::ZLIB
-          nvcomp::nvcomp kvikio::kvikio nanoarrow::nanoarrow zstd
+          nvcomp::nvcomp kvikio::kvikio nanoarrow::nanoarrow zstd CUDA::nvJitLink
 )
 
 # Add Conda library, and include paths if specified
diff --git a/cpp/cmake/Modules/compute_matrix_product.cmake b/cpp/cmake/Modules/compute_matrix_product.cmake
new file mode 100644
index 000000000000..3192bee98989
--- /dev/null
+++ b/cpp/cmake/Modules/compute_matrix_product.cmake
@@ -0,0 +1,53 @@
+# =============================================================================
+# cmake-format: off
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+# SPDX-License-Identifier: Apache-2.0
+# cmake-format: on
+# =============================================================================
+
+include_guard(GLOBAL)
+
+function(compute_matrix_product output_var)
+  set(options)
+  set(one_value MATRIX_JSON_FILE MATRIX_JSON_STRING)
+  set(multi_value)
+
+  cmake_parse_arguments(_JIT_LTO "${options}" "${one_value}" "${multi_value}" ${ARGN})
+
+  find_package(Python3 REQUIRED COMPONENTS Interpreter)
+
+  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()
+
+  set(${output_var}
+      "${output}"
+      PARENT_SCOPE
+  )
+endfunction()
+
+function(populate_matrix_variables matrix_json_entry)
+  string(JSON len LENGTH "${matrix_json_entry}")
+  math(EXPR last "${len} - 1")
+
+  # cmake-lint: disable=C0103,E1120
+  foreach(i RANGE "${last}")
+    string(JSON key MEMBER "${matrix_json_entry}" "${i}")
+    string(JSON value GET "${matrix_json_entry}" "${key}")
+    set(${key}
+        "${value}"
+        PARENT_SCOPE
+    )
+  endforeach()
+endfunction()
diff --git a/cpp/cmake/Modules/compute_matrix_product.py b/cpp/cmake/Modules/compute_matrix_product.py
new file mode 100644
index 000000000000..76262f6597ae
--- /dev/null
+++ b/cpp/cmake/Modules/compute_matrix_product.py
@@ -0,0 +1,243 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+# This algorithm takes a JSON dictionary and computes a matrix product of all
+# of its arrays. We use this to compute all matrix combinations for kernel
+# generation. We *could* write this in CMake with `string(JSON)`, but writing
+# it in Python is much easier. Once we have a version of CMake that has
+# https://gitlab.kitware.com/cmake/cmake/-/merge_requests/11516, we may be able
+# to port the algorithm to CMake script and use it in other RAPIDS projects.
+
+import argparse
+import json
+import re
+import sys
+import warnings
+from itertools import chain
+from typing import TYPE_CHECKING
+
+if TYPE_CHECKING:
+    from collections.abc import Generator, Iterator
+
+    MatrixValue = None | bool | int | float | str
+    Matrix = MatrixValue | list["Matrix"] | dict[str, "Matrix"]
+
+
+class NoKeyError(ValueError):
+    pass
+
+
+class UnusedKeyWarning(UserWarning):
+    pass
+
+
+class UsedKeyWarning(UserWarning):
+    pass
+
+
+IDENTIFIER_RE: re.Pattern = re.compile(r"^(?P<underscores>_*)(?P<rest>.*)$")
+
+
+def iterate_matrix_product(
+    *,
+    matrix: "Matrix",
+    warn_unused=True,
+    warn_used=True,
+) -> "Generator[dict[str, MatrixValue]]":
+    """Computes a matrix product of a JSON document
+
+    This algorithm computes the product of a matrix in a more sophisticated
+    way than can be done with itertools.product(). Multiple related values
+    can be grouped together, and a dimension can have sub-dimensions. Given
+    the following JSON document:
+
+    .. code-block:: json
+
+      {
+        "value": ["one", "two"],
+        "_group": [
+          {
+            "subgroup_value": "three"
+            "subgroup_subdim": ["four", "five"]
+          },
+          {
+            "subgroup_value": "six",
+            "subgroup_subdim": ["seven", "eight"]
+          }
+        ]
+      }
+
+    The following matrix product entries will be produced:
+
+    .. code-block:: json
+
+      {"subgroup_subdim": "four", "subgroup_value": "three", "value": "one"}
+      {"subgroup_subdim": "five", "subgroup_value": "three", "value": "one"}
+      {"subgroup_subdim": "seven", "subgroup_value": "six", "value": "one"}
+      {"subgroup_subdim": "eight", "subgroup_value": "six", "value": "one"}
+      {"subgroup_subdim": "four", "subgroup_value": "three", "value": "two"}
+      {"subgroup_subdim": "five", "subgroup_value": "three", "value": "two"}
+      {"subgroup_subdim": "seven", "subgroup_value": "six", "value": "two"}
+      {"subgroup_subdim": "eight", "subgroup_value": "six", "value": "two"}
+
+    Notice that the name ``_group`` does not appear in any of the matrix
+    product entries. This is because all leaf nodes beneath it are under
+    a different key that's closer to them in the hierarchy, which is used in
+    the final product. If a key is used only for grouping and does not appear
+    in the final product, it should be prefixed with an underscore (``_``) to
+    indicate that they are hidden. Likewise, keys that appear in the final
+    product should not be prefixed with an underscore. Failure to follow this
+    convention will not affect the proper functioning of the algorithm, but a
+    warning will be emitted unless the respective ``warn_unused`` and/or
+    ``warn_used`` parameters are set to ``False``.
+
+    Every leaf node in the input document must have at least one dictionary key
+    in its path, or else a ``NoKeyError`` will be thrown. For example, a
+    document consisting only of an array of strings is invalid.
+
+    Parameters
+    ----------
+    matrix : Matrix
+        JSON document on which to compute the product.
+    warn_unused : bool
+        Whether or not to warn if unused keys are not prefixed with underscores
+        (true by default).
+    warn_used : bool
+        Whether or not to warn if used keys are prefixed with underscores (true
+        by default).
+
+    Returns
+    -------
+    Generator[dict[str, MatrixValue]]
+        Iterator of matrix product entries.
+
+    Raises
+    ------
+    NoKeyError
+        If a leaf node in the document does not have any key leading to it.
+
+    Warns
+    -----
+    UnusedKeyWarning
+        If an unused key is not prefixed with an underscore.
+    UsedKeyWarning
+        If a used key is prefixed with an underscore.
+    """
+
+    def iterate_next_dimension(
+        queue: "Iterator[tuple[tuple[str | int, ...], str | None, Matrix]]",
+        entry: "dict[str, MatrixValue]",
+    ) -> "Generator[tuple[dict[str, MatrixValue], bool]]":
+        try:
+            path, key, matrix = next(queue)
+        except StopIteration:
+            yield (entry, False)
+        else:
+            used = False
+            for e, u in iterate_impl(path, key, matrix, queue, entry):
+                if u:
+                    used = True
+                yield e, u
+
+            try:
+                last = path[-1]
+            except IndexError:
+                pass
+            else:
+                if isinstance(last, str):
+                    match = IDENTIFIER_RE.search(last)
+                    assert match
+                    underscores = match.group("underscores")
+                    rest = match.group("rest")
+                    path_repr = "".join(
+                        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,
+                        )
+
+    def iterate_impl(
+        path: tuple[str | int, ...],
+        key: str | None,
+        matrix: "Matrix",
+        queue: "Iterator[tuple[tuple[str | int, ...], str | None, Matrix]]",
+        entry: "dict[str, MatrixValue]",
+    ) -> "Generator[tuple[dict[str, MatrixValue], bool]]":
+        if isinstance(matrix, dict):
+            yield from (
+                (e, False)
+                for e, _ in iterate_next_dimension(
+                    chain(
+                        (
+                            ((*path, k), k, v)
+                            for (k, v) in sorted(matrix.items())
+                        ),
+                        queue,
+                    ),
+                    entry,
+                )
+            )
+        elif isinstance(matrix, list):
+            queue_list = list(queue)
+            for i, v in enumerate(matrix):
+                yield from iterate_next_dimension(
+                    chain([((*path, i), key, v)], queue_list), entry
+                )
+        else:
+            if key is None:
+                raise NoKeyError
+            yield from (
+                (e, True)
+                for e, _ in iterate_next_dimension(
+                    queue, {**entry, key: matrix}
+                )
+            )
+
+    yield from (
+        entry for entry, _ in iterate_impl((), None, matrix, chain(), {})
+    )
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser()
+    parser.add_argument(
+        "--warn-unused", action=argparse.BooleanOptionalAction, default=True
+    )
+    parser.add_argument(
+        "--warn-used", action=argparse.BooleanOptionalAction, default=True
+    )
+    parser.add_argument("filename", nargs="?", default="-")
+    namespace = parser.parse_args()
+    with (
+        sys.stdin
+        if namespace.filename == "-"
+        else open(namespace.filename) as f
+    ):
+        matrix: "Matrix" = json.load(f)
+
+    json.dump(
+        list(
+            iterate_matrix_product(
+                matrix=matrix,
+                warn_unused=namespace.warn_unused,
+                warn_used=namespace.warn_used,
+            )
+        ),
+        sys.stdout,
+        indent=2,
+        sort_keys=True,
+    )
diff --git a/cpp/cmake/Modules/generate_jit_lto_kernels.cmake b/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
new file mode 100644
index 000000000000..0f45ab884a76
--- /dev/null
+++ b/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
@@ -0,0 +1,135 @@
+# =============================================================================
+# cmake-format: off
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+# SPDX-License-Identifier: Apache-2.0
+# cmake-format: on
+# =============================================================================
+# cuVS uses DEPENDS $<TARGET_OBJECTS:${kernel_target}>. With CUDA_FATBIN_COMPILATION that genex
+# expands to .fatbin paths; add_custom_command then requires those files as separate Makefile
+# prerequisites, but no rule exists (gmake: "No rule to make target ... .fatbin"). Depending on
+# the OBJECT target keeps ordering correct; COMMAND still passes $<TARGET_OBJECTS:...> to bin2c.
+
+include_guard(GLOBAL)
+
+include(${CMAKE_CURRENT_LIST_DIR}/compute_matrix_product.cmake)
+
+function(add_jit_lto_kernel kernel_target)
+  set(options)
+  set(one_value KERNEL_FILE FATBIN_HEADER_FILE)
+  set(multi_value LINK_LIBRARIES)
+
+  cmake_parse_arguments(_JIT_LTO "${options}" "${one_value}" "${multi_value}" ${ARGN})
+
+  add_library(${kernel_target} OBJECT EXCLUDE_FROM_ALL "${_JIT_LTO_KERNEL_FILE}")
+  # Do not modify these properties, options, and libraries. Usage requirements (including CUDA
+  # version, etc.) should be propagated to the kernel targets via INTERFACE libraries passed in
+  # through the LINK_LIBRARIES argument.
+  target_link_libraries(${kernel_target} PRIVATE ${_JIT_LTO_LINK_LIBRARIES})
+  target_compile_options(${kernel_target} PRIVATE -Xfatbin=--compress-all --compress-mode=size)
+  set_target_properties(
+    ${kernel_target}
+    PROPERTIES CUDA_SEPARABLE_COMPILATION ON
+               CUDA_FATBIN_COMPILATION ON
+               POSITION_INDEPENDENT_CODE ON
+               INTERPROCEDURAL_OPTIMIZATION ON
+  )
+
+  add_custom_command(
+    OUTPUT "${_JIT_LTO_FATBIN_HEADER_FILE}"
+    COMMAND "${bin_to_c}" --const --name embedded_fatbin --static $<TARGET_OBJECTS:${kernel_target}>
+            > "${_JIT_LTO_FATBIN_HEADER_FILE}"
+    DEPENDS ${kernel_target}
+  )
+endfunction()
+
+function(process_jit_lto_matrix_entry source_list_var)
+  set(options)
+  set(one_value NAME_FORMAT KERNEL_INPUT_FILE OUTPUT_DIRECTORY FRAGMENT_TAG_FORMAT
+                MATRIX_JSON_ENTRY
+  )
+  set(multi_value KERNEL_LINK_LIBRARIES FRAGMENT_TAG_HEADER_FILES)
+
+  cmake_parse_arguments(_JIT_LTO "${options}" "${one_value}" "${multi_value}" ${ARGN})
+
+  populate_matrix_variables("${_JIT_LTO_MATRIX_JSON_ENTRY}")
+  string(CONFIGURE "${_JIT_LTO_NAME_FORMAT}" kernel_name @ONLY)
+  string(CONFIGURE "${_JIT_LTO_FRAGMENT_TAG_FORMAT}" fragment_tag @ONLY)
+
+  set(fragment_tag_header_files "")
+  foreach(header_file IN LISTS _JIT_LTO_FRAGMENT_TAG_HEADER_FILES)
+    if(NOT header_file MATCHES "^(\".*\"|<.*>)$")
+      set(header_file "\"${header_file}\"")
+    endif()
+    string(APPEND fragment_tag_header_files "#include ${header_file}\n")
+  endforeach()
+
+  set(kernel_file "${_JIT_LTO_OUTPUT_DIRECTORY}/${kernel_name}_kernel.cu")
+  set(kernel_target "${kernel_name}_kernel")
+  set(fatbin_header_file "${_JIT_LTO_OUTPUT_DIRECTORY}/${kernel_name}_fatbin.h")
+  set(fatbin_file "${_JIT_LTO_OUTPUT_DIRECTORY}/${kernel_name}_fatbin.cpp")
+  configure_file("${_JIT_LTO_KERNEL_INPUT_FILE}" "${kernel_file}" @ONLY)
+  configure_file("${CMAKE_CURRENT_FUNCTION_LIST_DIR}/register_fatbin.cpp.in" "${fatbin_file}" @ONLY)
+
+  add_jit_lto_kernel(
+    ${kernel_target}
+    KERNEL_FILE "${kernel_file}"
+    FATBIN_HEADER_FILE "${fatbin_header_file}"
+    LINK_LIBRARIES ${_JIT_LTO_KERNEL_LINK_LIBRARIES}
+  )
+  list(APPEND ${source_list_var} "${fatbin_header_file}" "${fatbin_file}")
+  set(${source_list_var}
+      "${${source_list_var}}"
+      PARENT_SCOPE
+  )
+endfunction()
+
+function(generate_jit_lto_kernels source_list_var)
+  set(options)
+  set(one_value NAME_FORMAT MATRIX_JSON_FILE MATRIX_JSON_STRING KERNEL_INPUT_FILE
+                FRAGMENT_TAG_FORMAT OUTPUT_DIRECTORY
+  )
+  set(multi_value KERNEL_LINK_LIBRARIES FRAGMENT_TAG_HEADER_FILES)
+
+  cmake_parse_arguments(_JIT_LTO "${options}" "${one_value}" "${multi_value}" ${ARGN})
+
+  find_package(CUDAToolkit REQUIRED)
+  find_program(
+    bin_to_c
+    NAMES bin2c
+    PATHS ${CUDAToolkit_BIN_DIR}
+  )
+
+  if(_JIT_LTO_MATRIX_JSON_FILE)
+    set_property(
+      DIRECTORY
+      PROPERTY CMAKE_CONFIGURE_DEPENDS "${_JIT_LTO_MATRIX_JSON_FILE}"
+      APPEND
+    )
+    compute_matrix_product(matrix_product MATRIX_JSON_FILE "${_JIT_LTO_MATRIX_JSON_FILE}")
+  else()
+    compute_matrix_product(matrix_product MATRIX_JSON_STRING "${_JIT_LTO_MATRIX_JSON_STRING}")
+  endif()
+
+  string(JSON len LENGTH "${matrix_product}")
+  math(EXPR last "${len} - 1")
+
+  # cmake-lint: disable=C0103,E1120
+  foreach(i RANGE "${last}")
+    string(JSON matrix_json_entry GET "${matrix_product}" "${i}")
+    process_jit_lto_matrix_entry(
+      "${source_list_var}"
+      NAME_FORMAT "${_JIT_LTO_NAME_FORMAT}"
+      KERNEL_INPUT_FILE "${_JIT_LTO_KERNEL_INPUT_FILE}"
+      FRAGMENT_TAG_FORMAT "${_JIT_LTO_FRAGMENT_TAG_FORMAT}"
+      FRAGMENT_TAG_HEADER_FILES ${_JIT_LTO_FRAGMENT_TAG_HEADER_FILES}
+      OUTPUT_DIRECTORY "${_JIT_LTO_OUTPUT_DIRECTORY}"
+      MATRIX_JSON_ENTRY "${matrix_json_entry}"
+      KERNEL_LINK_LIBRARIES ${_JIT_LTO_KERNEL_LINK_LIBRARIES}
+    )
+  endforeach()
+
+  set(${source_list_var}
+      "${${source_list_var}}"
+      PARENT_SCOPE
+  )
+endfunction()
diff --git a/cpp/cmake/Modules/register_fatbin.cpp.in b/cpp/cmake/Modules/register_fatbin.cpp.in
new file mode 100644
index 000000000000..6f11ce4bad77
--- /dev/null
+++ b/cpp/cmake/Modules/register_fatbin.cpp.in
@@ -0,0 +1,22 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include "@fatbin_header_file@"
+#include <cudf/detail/jit_lto/FragmentEntry.hpp>
+
+@fragment_tag_header_files@
+
+namespace {
+
+using fragment_tag   = @fragment_tag@;
+using fragment_entry = cudf::detail::jit_lto::StaticFatbinFragmentEntry<fragment_tag>;
+
+}  // namespace
+
+template <>
+const uint8_t* const fragment_entry::data = embedded_fatbin;
+
+template <>
+const size_t fragment_entry::length = sizeof(embedded_fatbin);
diff --git a/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp b/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp
new file mode 100644
index 000000000000..c68cd70e185e
--- /dev/null
+++ b/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp
@@ -0,0 +1,49 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cstdint>
+#include <driver_types.h>
+#include <string>
+#include <unordered_map>
+#include <vector_types.h>
+
+#include <cuda_runtime.h>
+#include <memory>
+
+namespace cudf::detail::jit_lto {
+
+struct AlgorithmLauncher {
+  AlgorithmLauncher() : kernel{nullptr}, library{nullptr} {}
+
+  AlgorithmLauncher(cudaKernel_t k, cudaLibrary_t lib);
+
+  ~AlgorithmLauncher();
+
+  AlgorithmLauncher(const AlgorithmLauncher&)            = delete;
+  AlgorithmLauncher& operator=(const AlgorithmLauncher&) = delete;
+
+  AlgorithmLauncher(AlgorithmLauncher&& other) noexcept;
+  AlgorithmLauncher& operator=(AlgorithmLauncher&& other) noexcept;
+
+  template <typename FuncT, typename... Args>
+  void dispatch(cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, Args&&... args)
+  {
+    static_assert(std::is_same_v<FuncT, void(std::remove_reference_t<Args>...)>,
+                  "dispatch() argument types do not match the kernel function signature FuncT");
+
+    void* kernel_args[] = {const_cast<void*>(static_cast<void const*>(&args))...};
+    this->call(stream, grid, block, shared_mem, kernel_args);
+  }
+
+  cudaKernel_t get_kernel() { return this->kernel; }
+
+ private:
+  void call(cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, void** args);
+  cudaKernel_t kernel;
+  cudaLibrary_t library;
+};
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hpp b/cpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hpp
new file mode 100644
index 000000000000..e0e71421254c
--- /dev/null
+++ b/cpp/include/cudf/detail/jit_lto/AlgorithmPlanner.hpp
@@ -0,0 +1,57 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/detail/jit_lto/AlgorithmLauncher.hpp>
+#include <cudf/detail/jit_lto/FragmentEntry.hpp>
+
+#include <memory>
+#include <shared_mutex>
+#include <string>
+#include <type_traits>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+namespace cudf::detail::jit_lto {
+
+struct LauncherJitCache {
+  std::shared_mutex mutex;
+  std::unordered_map<std::string, std::shared_ptr<AlgorithmLauncher>> launchers;
+};
+
+struct AlgorithmPlanner {
+  AlgorithmPlanner(std::string entrypoint, LauncherJitCache& jit_cache)
+    : entrypoint(std::move(entrypoint)), jit_cache_(jit_cache)
+  {
+  }
+
+  std::shared_ptr<AlgorithmLauncher> get_launcher();
+
+  std::string entrypoint;
+  std::vector<std::unique_ptr<FragmentEntry>> fragments;
+
+  template <typename T, typename = std::enable_if_t<std::is_convertible_v<T*, FragmentEntry*>>>
+  void add_fragment(std::unique_ptr<T> fragment)
+  {
+    fragments.push_back(std::unique_ptr<FragmentEntry>(std::move(fragment)));
+  }
+
+  template <typename FragmentTag>
+  void add_static_fragment()
+  {
+    add_fragment(std::make_unique<StaticFatbinFragmentEntry<FragmentTag>>());
+  }
+
+ private:
+  std::string get_fragments_key() const;
+  std::shared_ptr<AlgorithmLauncher> build();
+
+  std::shared_ptr<AlgorithmLauncher> read_cache(std::string const& launch_key) const;
+
+  LauncherJitCache& jit_cache_;
+};
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp b/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp
new file mode 100644
index 000000000000..d4269bd2e7f2
--- /dev/null
+++ b/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp
@@ -0,0 +1,67 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+#include <string>
+#include <typeinfo>
+#include <vector>
+
+#include <nvJitLink.h>
+
+#include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
+
+namespace cudf::detail::jit_lto {
+
+struct FragmentEntry {
+  virtual ~FragmentEntry() = default;
+
+  virtual bool add_to(nvJitLinkHandle& handle) const = 0;
+
+  virtual const char* get_key() const = 0;
+};
+
+struct FatbinFragmentEntry : FragmentEntry {
+  virtual const uint8_t* get_data() const = 0;
+
+  virtual size_t get_length() const = 0;
+
+  bool add_to(nvJitLinkHandle& handle) const override final;
+};
+
+template <typename FragmentTag>
+struct StaticFatbinFragmentEntry final : FatbinFragmentEntry {
+  const uint8_t* get_data() const override { return StaticFatbinFragmentEntry<FragmentTag>::data; }
+
+  size_t get_length() const override { return StaticFatbinFragmentEntry<FragmentTag>::length; }
+
+  const char* get_key() const override
+  {
+    return typeid(StaticFatbinFragmentEntry<FragmentTag>).name();
+  }
+
+  static const uint8_t* const data;
+  static const size_t length;
+};
+
+struct UDFFatbinFragment final : FatbinFragmentEntry {
+  UDFFatbinFragment(std::string key, std::vector<uint8_t> bytes)
+    : key_(std::move(key)), bytes_(std::move(bytes))
+  {
+  }
+
+  const uint8_t* get_data() const override { return bytes_.data(); }
+
+  size_t get_length() const override { return bytes_.size(); }
+
+  const char* get_key() const override { return key_.c_str(); }
+
+ private:
+  std::string key_;
+  std::vector<uint8_t> bytes_;
+};
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/include/cudf/detail/jit_lto/nvjitlink_checker.hpp b/cpp/include/cudf/detail/jit_lto/nvjitlink_checker.hpp
new file mode 100644
index 000000000000..70755a9adf96
--- /dev/null
+++ b/cpp/include/cudf/detail/jit_lto/nvjitlink_checker.hpp
@@ -0,0 +1,13 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <nvJitLink.h>
+
+namespace cudf::detail::jit_lto {
+
+void check_nvjitlink_result(nvJitLinkHandle handle, nvJitLinkResult result);
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
new file mode 100644
index 000000000000..ecacd6de96bb
--- /dev/null
+++ b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
@@ -0,0 +1,50 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+namespace cudf::hashing::detail::jit_lto {
+
+struct tag_murmur_entry {};
+
+struct tag_i8 {};
+struct tag_i16 {};
+struct tag_i32 {};
+struct tag_i64 {};
+struct tag_u8 {};
+struct tag_u16 {};
+struct tag_u32 {};
+struct tag_u64 {};
+struct tag_f32 {};
+struct tag_f64 {};
+struct tag_b8 {};
+struct tag_ts_day {};
+struct tag_ts_s {};
+struct tag_ts_ms {};
+struct tag_ts_us {};
+struct tag_ts_ns {};
+struct tag_du_day {};
+struct tag_du_s {};
+struct tag_du_ms {};
+struct tag_du_us {};
+struct tag_du_ns {};
+struct tag_dict {};
+struct tag_str {};
+struct tag_list {};
+struct tag_dec32 {};
+struct tag_dec64 {};
+struct tag_dec128 {};
+struct tag_struct {};
+
+template <typename StorageTag>
+struct fragment_tag_murmur_hasher {};
+
+/// Strong `murmur_jit_hasher<T>` with a no-op body; used for storage types not present in the
+/// table so nvJitLink still sees exactly one definition per `T` (weak overrides are not supported).
+template <typename StorageTag>
+struct fragment_tag_murmur_hasher_noop {};
+
+struct fragment_tag_murmur_entry {};
+
+}  // namespace cudf::hashing::detail::jit_lto
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
new file mode 100644
index 000000000000..7edcdd95fcb8
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
@@ -0,0 +1,82 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+
+#include <cudf/detail/utilities/grid_1d.cuh>
+#include <cudf/detail/utilities/assert.cuh>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/utilities/type_dispatcher.hpp>
+
+namespace cudf::hashing::detail {
+
+__device__ inline hash_value_type murmur_jit_hash_dispatcher(column_device_view col,
+                                                             uint32_t seed,
+                                                             bool const nullable,
+                                                             size_type row_index)
+{
+  switch (col.type().id()) {
+    case type_id::INT8:
+      return murmur_jit_hasher<id_to_type<type_id::INT8>>(col, seed, nullable, row_index);
+    case type_id::INT16:
+      return murmur_jit_hasher<id_to_type<type_id::INT16>>(col, seed, nullable, row_index);
+    case type_id::INT32:
+      return murmur_jit_hasher<id_to_type<type_id::INT32>>(col, seed, nullable, row_index);
+    case type_id::INT64:
+      return murmur_jit_hasher<id_to_type<type_id::INT64>>(col, seed, nullable, row_index);
+    case type_id::UINT8:
+      return murmur_jit_hasher<id_to_type<type_id::UINT8>>(col, seed, nullable, row_index);
+    case type_id::UINT16:
+      return murmur_jit_hasher<id_to_type<type_id::UINT16>>(col, seed, nullable, row_index);
+    case type_id::UINT32:
+      return murmur_jit_hasher<id_to_type<type_id::UINT32>>(col, seed, nullable, row_index);
+    case type_id::UINT64:
+      return murmur_jit_hasher<id_to_type<type_id::UINT64>>(col, seed, nullable, row_index);
+    case type_id::FLOAT32:
+      return murmur_jit_hasher<id_to_type<type_id::FLOAT32>>(col, seed, nullable, row_index);
+    case type_id::FLOAT64:
+      return murmur_jit_hasher<id_to_type<type_id::FLOAT64>>(col, seed, nullable, row_index);
+    case type_id::BOOL8:
+      return murmur_jit_hasher<id_to_type<type_id::BOOL8>>(col, seed, nullable, row_index);
+    case type_id::TIMESTAMP_DAYS:
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_DAYS>>(col, seed, nullable, row_index);
+    case type_id::TIMESTAMP_SECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_SECONDS>>(col, seed, nullable, row_index);
+    case type_id::TIMESTAMP_MILLISECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_MILLISECONDS>>(col, seed, nullable, row_index);
+    case type_id::TIMESTAMP_MICROSECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_MICROSECONDS>>(col, seed, nullable, row_index);
+    case type_id::TIMESTAMP_NANOSECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_NANOSECONDS>>(col, seed, nullable, row_index);
+    case type_id::DURATION_DAYS:
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_DAYS>>(col, seed, nullable, row_index);
+    case type_id::DURATION_SECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_SECONDS>>(col, seed, nullable, row_index);
+    case type_id::DURATION_MILLISECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_MILLISECONDS>>(col, seed, nullable, row_index);
+    case type_id::DURATION_MICROSECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_MICROSECONDS>>(col, seed, nullable, row_index);
+    case type_id::DURATION_NANOSECONDS:
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_NANOSECONDS>>(col, seed, nullable, row_index);
+    case type_id::DICTIONARY32:
+      return murmur_jit_hasher<id_to_type<type_id::DICTIONARY32>>(col, seed, nullable, row_index);
+    case type_id::STRING:
+      return murmur_jit_hasher<id_to_type<type_id::STRING>>(col, seed, nullable, row_index);
+    case type_id::LIST:
+      return murmur_jit_hasher<id_to_type<type_id::LIST>>(col, seed, nullable, row_index);
+    case type_id::DECIMAL32:
+      return murmur_jit_hasher<id_to_type<type_id::DECIMAL32>>(col, seed, nullable, row_index);
+    case type_id::DECIMAL64:
+      return murmur_jit_hasher<id_to_type<type_id::DECIMAL64>>(col, seed, nullable, row_index);
+    case type_id::DECIMAL128:
+      return murmur_jit_hasher<id_to_type<type_id::DECIMAL128>>(col, seed, nullable, row_index);
+    case type_id::STRUCT:
+      return murmur_jit_hasher<id_to_type<type_id::STRUCT>>(col, seed, nullable, row_index);
+    default: CUDF_UNREACHABLE("Invalid type_id.");
+  }
+}
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
new file mode 100644
index 000000000000..44994a48bcf0
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
@@ -0,0 +1,27 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/hashing.hpp>
+#include <cudf/types.hpp>
+
+namespace cudf::hashing::detail {
+
+/**
+ * @brief Forward declaration only (no `murmur_jit_hash_dispatcher` here).
+ *
+ * Per-type explicit specializations must appear in the TU before any use of that specialization.
+ * Hasher / noop fragment TUs include this header, then define `template <> ... murmur_jit_hasher<T>`.
+ * The entry kernel includes `murmurhash3_x86_32_jit_device.cuh`, which pulls this in and adds the
+ * dispatcher that calls `murmur_jit_hasher` for every storage type.
+ */
+template <typename T>
+extern __device__ hash_value_type murmur_jit_hasher(column_device_view col,
+                                                    uint32_t seed,
+                                                    bool nullable,
+                                                    size_type row_index);
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
new file mode 100644
index 000000000000..c5fdfae95349
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
@@ -0,0 +1,162 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/detail/iterator.cuh>
+#include <cudf/detail/row_operator/common_utils.cuh>
+#include <cudf/detail/utilities/accumulate.cuh>
+#include <cudf/detail/utilities/assert.cuh>
+#include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/hashing/detail/hashing.hpp>
+#include <cudf/hashing/detail/murmurhash3_x86_32.cuh>
+#include <cudf/hashing.hpp>
+#include <cudf/lists/list_device_view.cuh>
+#include <cudf/lists/lists_column_device_view.cuh>
+#include <cudf/structs/structs_column_device_view.cuh>
+#include <cudf/types.hpp>
+#include <cudf/utilities/traits.hpp>
+#include <cudf/utilities/type_dispatcher.hpp>
+
+#include <cuda/std/limits>
+
+namespace cudf::hashing::detail {
+
+using result_type = hash_value_type;
+
+class element_hasher {
+ public:
+  /**
+   * @brief Constructs an element_hasher object.
+   *
+   * @param nulls Indicates whether to check for nulls
+   * @param seed  The seed to use for the hash function
+   * @param null_hash The hash value to use for nulls
+   */
+  __device__ element_hasher(
+    bool nulls,
+    result_type seed      = DEFAULT_HASH_SEED,
+    result_type null_hash = cuda::std::numeric_limits<result_type>::max()) noexcept
+    : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
+  {
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(column_device_view::has_element_accessor<T>())
+  {
+    if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
+    return MurmurHash3_x86_32<T>{_seed}(col.element<T>(row_index));
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(not column_device_view::has_element_accessor<T>())
+  {
+    CUDF_UNREACHABLE("Unsupported type in hash.");
+  }
+
+  bool _check_nulls;
+  result_type _seed;
+  result_type _null_hash;
+};
+
+class element_hasher_adaptor {
+  static constexpr result_type NULL_HASH     = cuda::std::numeric_limits<result_type>::max();
+  static constexpr result_type NON_NULL_HASH = 0;
+
+ public:
+  __device__ element_hasher_adaptor(bool check_nulls, result_type seed) noexcept
+    : _element_hasher(check_nulls, seed), _check_nulls(check_nulls)
+  {
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(not cudf::is_nested<T>() and not cudf::is_dictionary<T>())
+  {
+    return _element_hasher.template operator()<T>(col, row_index);
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(cudf::is_dictionary<T>())
+  {
+    if (_check_nulls && col.is_null(row_index)) { return NULL_HASH; }
+
+    auto const keys = col.child(dictionary_column_view::keys_column_index);
+    return type_dispatcher<dispatch_storage_type>(
+      keys.type(),
+      _element_hasher,
+      keys,
+      static_cast<size_type>(col.element<dictionary32>(row_index)));
+  }
+
+  template <typename T>
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
+    requires(cudf::is_nested<T>())
+  {
+    auto hash                   = result_type{0};
+    column_device_view curr_col = col.slice(row_index, 1);
+    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,
+        type_dispatcher<cudf::detail::dispatch_void_if_nested>(
+          curr_col.type(), _element_hasher, curr_col, i));
+    }
+    return hash;
+  }
+
+  element_hasher const _element_hasher;
+  bool const _check_nulls;
+};
+
+template <typename T>
+__device__ hash_value_type hasher_impl(bool check_nulls,
+                                       cudf::column_device_view col,
+                                       uint32_t seed,
+                                       size_type row_index)
+{
+  auto const hasher = element_hasher_adaptor{check_nulls, seed};
+  return hasher.template operator()<T>(col, row_index);
+}
+
+template <typename T>
+__device__ hash_value_type hasher(cudf::column_device_view col,
+                                  uint32_t seed,
+                                  bool const nullable,
+                                  size_type row_index)
+{
+  return hasher_impl<T>(nullable, col, seed, row_index);
+}
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
new file mode 100644
index 000000000000..c8f406a4f728
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
@@ -0,0 +1,33 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/hashing/detail/hashing.hpp>
+#include "murmurhash3_x86_32_jit_device.cuh"
+#include <cudf/table/table_device_view.cuh>
+
+extern "C" __global__ void cudf_murmurhash3_x86_32_jit_link_kernel(cudf::mutable_column_device_view output,
+                                                                   uint32_t seed,
+                                                                   cudf::table_device_view input,
+                                                                   bool nullable)
+{
+  auto const idx = cudf::detail::grid_1d::global_thread_id();
+  if (idx >= input.num_rows()) { return; }
+
+  auto const num_cols = input.num_columns();
+  if (num_cols == 0) {
+    output.element<cudf::hash_value_type>(idx) = seed;
+    return;
+  }
+
+  cudf::hash_value_type hash_value =
+    cudf::hashing::detail::murmur_jit_hash_dispatcher(input.column(0), seed, nullable, idx);
+  for (cudf::size_type i = 1; i < num_cols; ++i) {
+    hash_value = cudf::hashing::detail::hash_combine(
+      hash_value,
+      cudf::hashing::detail::murmur_jit_hash_dispatcher(input.column(i), seed, nullable, idx));
+  }
+  output.element<cudf::hash_value_type>(idx) = hash_value;
+}
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json
new file mode 100644
index 000000000000..c108798fd88a
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json
@@ -0,0 +1,7 @@
+{
+  "_singleton": [
+    {
+      "k": "entry"
+    }
+  ]
+}
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
new file mode 100644
index 000000000000..3b52c19ba8b1
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
@@ -0,0 +1,27 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+// Strong explicit specialization of `murmur_jit_hasher<T>` (real hasher); linked when logical
+// type T appears in the input table. Other storage types use `murmurhash_jit_hasher_noop_*`
+// (strong no-op) so nvJitLink sees exactly one definition per T.
+
+#include <cudf/column/column_device_view.cuh>
+
+// Decl only (not `jit_device.cuh`): dispatcher would use every `murmur_jit_hasher<T>` before the
+// explicit specialization below, which nvcc rejects ("must precede its first use").
+#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+#include "murmurhash3_x86_32_lto.cuh"
+
+namespace cudf::hashing::detail {
+
+template <>
+__device__ hash_value_type murmur_jit_hasher<@storage_cpp@>(column_device_view col,
+                                                            uint32_t seed,
+                                                            bool nullable,
+                                                            size_type row_index)
+{
+  return hasher<@storage_cpp@>(col, seed, nullable, row_index);
+}
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json
new file mode 100644
index 000000000000..5216fcc72c94
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json
@@ -0,0 +1,32 @@
+{
+  "_row": [
+    {"storage_cpp": "int8_t", "abbrev": "i8"},
+    {"storage_cpp": "int16_t", "abbrev": "i16"},
+    {"storage_cpp": "int32_t", "abbrev": "i32"},
+    {"storage_cpp": "int64_t", "abbrev": "i64"},
+    {"storage_cpp": "uint8_t", "abbrev": "u8"},
+    {"storage_cpp": "uint16_t", "abbrev": "u16"},
+    {"storage_cpp": "uint32_t", "abbrev": "u32"},
+    {"storage_cpp": "uint64_t", "abbrev": "u64"},
+    {"storage_cpp": "float", "abbrev": "f32"},
+    {"storage_cpp": "double", "abbrev": "f64"},
+    {"storage_cpp": "bool", "abbrev": "b8"},
+    {"storage_cpp": "cudf::timestamp_D", "abbrev": "ts_day"},
+    {"storage_cpp": "cudf::timestamp_s", "abbrev": "ts_s"},
+    {"storage_cpp": "cudf::timestamp_ms", "abbrev": "ts_ms"},
+    {"storage_cpp": "cudf::timestamp_us", "abbrev": "ts_us"},
+    {"storage_cpp": "cudf::timestamp_ns", "abbrev": "ts_ns"},
+    {"storage_cpp": "cudf::duration_D", "abbrev": "du_day"},
+    {"storage_cpp": "cudf::duration_s", "abbrev": "du_s"},
+    {"storage_cpp": "cudf::duration_ms", "abbrev": "du_ms"},
+    {"storage_cpp": "cudf::duration_us", "abbrev": "du_us"},
+    {"storage_cpp": "cudf::duration_ns", "abbrev": "du_ns"},
+    {"storage_cpp": "cudf::dictionary32", "abbrev": "dict"},
+    {"storage_cpp": "cudf::string_view", "abbrev": "str"},
+    {"storage_cpp": "cudf::list_view", "abbrev": "list"},
+    {"storage_cpp": "numeric::decimal32", "abbrev": "dec32"},
+    {"storage_cpp": "numeric::decimal64", "abbrev": "dec64"},
+    {"storage_cpp": "numeric::decimal128", "abbrev": "dec128"},
+    {"storage_cpp": "cudf::struct_view", "abbrev": "struct"}
+  ]
+}
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in
new file mode 100644
index 000000000000..65dd6cb03681
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in
@@ -0,0 +1,28 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+// Strong explicit specialization with a no-op body for `T` not used by any column in the table.
+// nvJitLink does not coalesce weak + strong device symbols, so unused types use this fatbin
+// instead of `fragment_tag_murmur_hasher` (real hasher).
+
+#include <cudf/column/column_device_view.cuh>
+
+#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+
+namespace cudf::hashing::detail {
+
+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};
+}
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu
index 97c276ad1b2b..8a3e10b26f63 100644
--- a/cpp/src/hash/murmurhash3_x86_32.cu
+++ b/cpp/src/hash/murmurhash3_x86_32.cu
@@ -2,173 +2,25 @@
  * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
-#include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
-#include <cudf/detail/row_operator/hashing.cuh>
 #include <cudf/hashing/detail/hashing.hpp>
-#include <cudf/hashing/detail/murmurhash3_x86_32.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/exec_policy.hpp>
 
-#include <thrust/tabulate.h>
+#include <memory>
 
-#include "murmurhash3_x86_32_lto.cuh"
+#include "murmurhash3_x86_32_jit_launch.hpp"
 
 namespace cudf {
 namespace hashing {
 namespace detail {
 
-//
-// 1. Build a device-side dispatcher that 
-
-//template <typename T>
-//hash_value_type hasher(cudf::column_device_view col, uint32_t seed);
-
-__device__  __forceinline__ constexpr decltype(auto) hash_dispatcher(cudf::column_device_view col, uint32_t seed, bool const nullable)
-{
-  switch (col.type().id()) {
-    case type_id::INT8:
-      return hasher<id_to_type<type_id::INT8>>(
-        col, seed, nullable);
-    case type_id::INT16:
-      return hasher<id_to_type<type_id::INT16>>(
-        col, seed, nullable);
-    case type_id::INT32:
-      return hasher<id_to_type<type_id::INT32>>(
-        col, seed, nullable);
-    case type_id::INT64:
-      return hasher<id_to_type<type_id::INT64>>(
-        col, seed, nullable);
-    case type_id::UINT8:
-      return hasher<id_to_type<type_id::UINT8>>(
-        col, seed, nullable);
-    case type_id::UINT16:
-      return hasher<id_to_type<type_id::UINT16>>(
-        col, seed, nullable);
-    case type_id::UINT32:
-      return hasher<id_to_type<type_id::UINT32>>(
-        col, seed, nullable);
-    case type_id::UINT64:
-      return hasher<id_to_type<type_id::UINT64>>(
-        col, seed, nullable);
-    case type_id::FLOAT32:
-      return hasher<id_to_type<type_id::FLOAT32>>(
-        col, seed, nullable);
-    case type_id::FLOAT64:
-      return hasher<id_to_type<type_id::FLOAT64>>(
-        col, seed, nullable);
-    case type_id::BOOL8:
-      return hasher<id_to_type<type_id::BOOL8>>(
-        col, seed, nullable);
-    case type_id::TIMESTAMP_DAYS:
-      return hasher<id_to_type<type_id::TIMESTAMP_DAYS>>(
-        col, seed, nullable);
-    case type_id::TIMESTAMP_SECONDS:
-      return hasher<id_to_type<type_id::TIMESTAMP_SECONDS>>(
-        col, seed, nullable);
-    case type_id::TIMESTAMP_MILLISECONDS:
-      return hasher<id_to_type<type_id::TIMESTAMP_MILLISECONDS>>(
-        col, seed, nullable);
-    case type_id::TIMESTAMP_MICROSECONDS:
-      return hasher<id_to_type<type_id::TIMESTAMP_MICROSECONDS>>(
-        col, seed, nullable);
-    case type_id::TIMESTAMP_NANOSECONDS:
-      return hasher<id_to_type<type_id::TIMESTAMP_NANOSECONDS>>(
-        col, seed, nullable);
-    case type_id::DURATION_DAYS:
-      return hasher<id_to_type<type_id::DURATION_DAYS>>(
-        col, seed, nullable);
-    case type_id::DURATION_SECONDS:
-      return hasher<id_to_type<type_id::DURATION_SECONDS>>(
-        col, seed, nullable);
-    case type_id::DURATION_MILLISECONDS:
-      return hasher<id_to_type<type_id::DURATION_MILLISECONDS>>(
-        col, seed, nullable);
-    case type_id::DURATION_MICROSECONDS:
-      return hasher<id_to_type<type_id::DURATION_MICROSECONDS>>(
-        col, seed, nullable);
-    case type_id::DURATION_NANOSECONDS:
-      return hasher<id_to_type<type_id::DURATION_NANOSECONDS>>(
-        col, seed, nullable);
-    case type_id::DICTIONARY32:
-      return hasher<id_to_type<type_id::DICTIONARY32>>(
-        col, seed, nullable);
-    case type_id::STRING:
-      return hasher<id_to_type<type_id::STRING>>(
-        col, seed, nullable);
-    case type_id::LIST:
-      return hasher<id_to_type<type_id::LIST>>(
-        col, seed, nullable);
-    case type_id::DECIMAL32:
-      return hasher<id_to_type<type_id::DECIMAL32>>(
-        col, seed, nullable);
-    case type_id::DECIMAL64:
-      return hasher<id_to_type<type_id::DECIMAL64>>(
-        col, seed, nullable);
-    case type_id::DECIMAL128:
-      return hasher<id_to_type<type_id::DECIMAL128>>(
-        col, seed, nullable);
-    case type_id::STRUCT:
-      return hasher<id_to_type<type_id::STRUCT>>(
-        col, seed, nullable);
-    default: {
-#ifndef __CUDA_ARCH__
-      CUDF_FAIL("Invalid type_id.");
-#else
-      CUDF_UNREACHABLE("Invalid type_id.");
-#endif
-    }
-  }
-}
-
-__global__ void murmurhash3_x86_32_kernel(mutable_column_device_view output,
-                                  uint32_t seed,
-                                  table_device_view const input,
-                                  bool const nullable)
-{
-  cudf::size_type idx = blockIdx.x * blockDim.x + threadIdx.x;
-  if (idx < input.num_rows()) {
-      auto const num_cols = input.num_columns();
-      if (num_cols == 0) return;
-      hash_value_type hash_value = hash_dispatcher(input.column(0), seed, nullable);
-      for (int i = 1; i < num_cols; ++i) {
-          hash_value = cudf::hashing::detail::hash_combine(
-            hash_value, hash_dispatcher(input.column(i), seed, nullable));
-      }
-      output.element<hash_value_type>(idx) = hash_value;
-  }
-}
-
 std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
                                            uint32_t seed,
                                            rmm::cuda_stream_view stream,
                                            rmm::device_async_resource_ref mr)
 {
-  auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),
-                                    input.num_rows(),
-                                    mask_state::UNALLOCATED,
-                                    stream,
-                                    mr);
-
-  // Return early if there's nothing to hash
-  if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
-
-  bool const nullable   = has_nulls(input);
-  auto const row_hasher = cudf::detail::row::hash::row_hasher(input, stream);
-  auto output_view      = output->mutable_view();
-
-  // Compute the hash value for each row
-  //thrust::tabulate(rmm::exec_policy_nosync(stream),
-  //                 output_view.begin<hash_value_type>(),
-  //                 output_view.end<hash_value_type>(),
-  //                 row_hasher.device_hasher<MurmurHash3_x86_32>(nullable, seed));
-  //
-  auto d_output = mutable_column_device_view::create(output_view, stream);
-  auto d_input  = table_device_view::create(input, stream);
-  murmurhash3_x86_32_kernel<<<1, input.num_rows(), 0, stream.value()>>>(
-    *d_output, seed, *d_input, nullable);
-  return output;
+  return murmurhash3_x86_32_jit(input, seed, stream, mr);
 }
 
 }  // namespace detail
diff --git a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
new file mode 100644
index 000000000000..44e2a810d9ea
--- /dev/null
+++ b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
@@ -0,0 +1,287 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/column/column_factories.hpp>
+#include <cudf/detail/jit_lto/AlgorithmPlanner.hpp>
+#include <cudf/detail/row_operator/preprocessed_table.cuh>
+#include <cudf/detail/utilities/grid_1d.cuh>
+#include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/hashing/detail/hashing.hpp>
+#include <cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>
+#include <cudf/table/table_device_view.cuh>
+#include <cudf/table/table_view.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <array>
+#include <memory>
+#include <unordered_set>
+
+namespace cudf::hashing::detail {
+
+namespace {
+
+inline cudf::detail::jit_lto::LauncherJitCache& murmur_jit_launcher_cache()
+{
+  static cudf::detail::jit_lto::LauncherJitCache cache;
+  return cache;
+}
+
+// Every `type_id` handled by `murmur_jit_hash_dispatcher` must have exactly one device
+// definition of `murmur_jit_hasher<T>` in the nvJitLink input. Types present in the table use the
+// real hasher fatbin; all others use a separate strong no-op fatbin (nvJitLink does not merge
+// weak + strong overrides for the same symbol).
+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,
+}};
+
+inline void insert_logical_type(std::unordered_set<type_id>& ids, column_view const& col)
+{
+  ids.insert(col.type().id());
+  if (col.type().id() == type_id::DICTIONARY32) {
+    dictionary_column_view const dcv(col);
+    insert_logical_type(ids, dcv.keys());
+  }
+}
+
+inline void collect_table_logical_types(std::unordered_set<type_id>& ids, table_view const& table)
+{
+  for (size_type c = 0; c < table.num_columns(); ++c) {
+    insert_logical_type(ids, table.column(c));
+  }
+}
+
+inline void add_strong_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& planner, type_id id)
+{
+  using namespace cudf::hashing::detail::jit_lto;
+  switch (id) {
+    case type_id::INT8: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_i8>>(); break;
+    case type_id::INT16: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_i16>>(); break;
+    case type_id::INT32: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_i32>>(); break;
+    case type_id::INT64: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_i64>>(); break;
+    case type_id::UINT8: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u8>>(); break;
+    case type_id::UINT16: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u16>>(); break;
+    case type_id::UINT32: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u32>>(); break;
+    case type_id::UINT64: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u64>>(); break;
+    case type_id::FLOAT32: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_f32>>(); break;
+    case type_id::FLOAT64: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_f64>>(); break;
+    case type_id::BOOL8: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_b8>>(); break;
+    case type_id::TIMESTAMP_DAYS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_ts_day>>();
+      break;
+    case type_id::TIMESTAMP_SECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_ts_s>>();
+      break;
+    case type_id::TIMESTAMP_MILLISECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_ts_ms>>();
+      break;
+    case type_id::TIMESTAMP_MICROSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_ts_us>>();
+      break;
+    case type_id::TIMESTAMP_NANOSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_ts_ns>>();
+      break;
+    case type_id::DURATION_DAYS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_du_day>>();
+      break;
+    case type_id::DURATION_SECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_du_s>>();
+      break;
+    case type_id::DURATION_MILLISECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_du_ms>>();
+      break;
+    case type_id::DURATION_MICROSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_du_us>>();
+      break;
+    case type_id::DURATION_NANOSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_du_ns>>();
+      break;
+    case type_id::DICTIONARY32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_dict>>();
+      break;
+    case type_id::STRING: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_str>>(); break;
+    case type_id::LIST: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_list>>(); break;
+    case type_id::DECIMAL32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_dec32>>();
+      break;
+    case type_id::DECIMAL64:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_dec64>>();
+      break;
+    case type_id::DECIMAL128:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_dec128>>();
+      break;
+    case type_id::STRUCT:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_struct>>();
+      break;
+    default: break;
+  }
+}
+
+inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& planner, type_id id)
+{
+  using namespace cudf::hashing::detail::jit_lto;
+  switch (id) {
+    case type_id::INT8: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i8>>(); break;
+    case type_id::INT16:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i16>>();
+      break;
+    case type_id::INT32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i32>>();
+      break;
+    case type_id::INT64:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i64>>();
+      break;
+    case type_id::UINT8: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u8>>(); break;
+    case type_id::UINT16:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u16>>();
+      break;
+    case type_id::UINT32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u32>>();
+      break;
+    case type_id::UINT64:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u64>>();
+      break;
+    case type_id::FLOAT32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_f32>>();
+      break;
+    case type_id::FLOAT64:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_f64>>();
+      break;
+    case type_id::BOOL8: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_b8>>(); break;
+    case type_id::TIMESTAMP_DAYS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_ts_day>>();
+      break;
+    case type_id::TIMESTAMP_SECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_ts_s>>();
+      break;
+    case type_id::TIMESTAMP_MILLISECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_ts_ms>>();
+      break;
+    case type_id::TIMESTAMP_MICROSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_ts_us>>();
+      break;
+    case type_id::TIMESTAMP_NANOSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_ts_ns>>();
+      break;
+    case type_id::DURATION_DAYS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_du_day>>();
+      break;
+    case type_id::DURATION_SECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_du_s>>();
+      break;
+    case type_id::DURATION_MILLISECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_du_ms>>();
+      break;
+    case type_id::DURATION_MICROSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_du_us>>();
+      break;
+    case type_id::DURATION_NANOSECONDS:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_du_ns>>();
+      break;
+    case type_id::DICTIONARY32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_dict>>();
+      break;
+    case type_id::STRING: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_str>>(); break;
+    case type_id::LIST: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_list>>(); break;
+    case type_id::DECIMAL32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_dec32>>();
+      break;
+    case type_id::DECIMAL64:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_dec64>>();
+      break;
+    case type_id::DECIMAL128:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_dec128>>();
+      break;
+    case type_id::STRUCT:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_struct>>();
+      break;
+    default: break;
+  }
+}
+
+}  // namespace
+
+inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
+                                                     uint32_t seed,
+                                                     rmm::cuda_stream_view stream,
+                                                     rmm::device_async_resource_ref mr)
+{
+  auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),
+                                    input.num_rows(),
+                                    mask_state::UNALLOCATED,
+                                    stream,
+                                    mr);
+
+  if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
+
+  bool const nullable = has_nulls(input);
+  auto const preprocessed =
+    cudf::detail::row::hash::preprocessed_table::create(input, stream);
+  table_device_view const input_dv{*preprocessed};
+
+  auto output_view = output->mutable_view();
+  auto d_output    = mutable_column_device_view::create(output_view, stream);
+
+  cudf::detail::jit_lto::AlgorithmPlanner planner{"cudf_murmurhash3_x86_32_jit_link_kernel",
+                                                  murmur_jit_launcher_cache()};
+
+  using namespace cudf::hashing::detail::jit_lto;
+  planner.add_static_fragment<fragment_tag_murmur_entry>();
+
+  std::unordered_set<type_id> logical_types;
+  collect_table_logical_types(logical_types, input);
+  for (type_id const id : murmur_jit_hasher_type_ids) {
+    if (logical_types.count(id) != 0u) {
+      add_strong_hasher_fragment(planner, id);
+    } else {
+      add_noop_hasher_fragment(planner, id);
+    }
+  }
+
+  auto launcher = planner.get_launcher();
+
+  cudf::detail::grid_1d const grid{input.num_rows(), 256};
+  launcher->dispatch<void(cudf::mutable_column_device_view, uint32_t, cudf::table_device_view, bool)>(
+    stream.value(),
+    dim3(grid.num_blocks),
+    dim3(grid.num_threads_per_block),
+    0,
+    *d_output,
+    seed,
+    input_dv,
+    nullable);
+
+  return output;
+}
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/murmurhash3_x86_32_lto.cuh b/cpp/src/hash/murmurhash3_x86_32_lto.cuh
deleted file mode 100644
index 98d6c636eaca..000000000000
--- a/cpp/src/hash/murmurhash3_x86_32_lto.cuh
+++ /dev/null
@@ -1,171 +0,0 @@
-/*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
- * SPDX-License-Identifier: Apache-2.0
- */
-
-#include "cuda/std/__type_traits/is_same.h"
-#include <cudf/detail/row_operator/hashing.cuh>
-#include <cudf/column/column_device_view.cuh>
-#include <cudf/types.hpp>
-#include <cudf/detail/row_operator/hashing.cuh>
-#include <cudf/hashing.hpp>
-#include <cudf/structs/structs_column_device_view.cuh>
-#include <cudf/lists/lists_column_device_view.cuh>
-
-namespace cudf::hashing::detail {
-using result_type = hash_value_type;
-
-template <typename Nullate>
-class element_hasher {
- public:
-
-  /**
-   * @brief Constructs an element_hasher object.
-   *
-   * @param nulls Indicates whether to check for nulls
-   * @param seed  The seed to use for the hash function
-   * @param null_hash The hash value to use for nulls
-   */
-  __device__ element_hasher(
-    Nullate nulls,
-    result_type seed      = DEFAULT_HASH_SEED,
-    result_type null_hash = cuda::std::numeric_limits<result_type>::max()) noexcept
-    : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
-  {
-  }
-
-  /**
-   * @brief Returns the hash value of the given element.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T>
-  __device__ result_type operator()(column_device_view const& col,
-                                    size_type row_index) const noexcept
-    requires(column_device_view::has_element_accessor<T>())
-  {
-    if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
-    return MurmurHash3_x86_32<T>{_seed}(col.element<T>(row_index));
-  }
-
-  /**
-   * @brief Returns the hash value of the given element.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T>
-  __device__ result_type operator()(column_device_view const& col,
-                                    size_type row_index) const noexcept
-    requires(not column_device_view::has_element_accessor<T>())
-  {
-    CUDF_UNREACHABLE("Unsupported type in hash.");
-  }
-
-  Nullate _check_nulls;
-  // Assumes seeds are the same as the result type of the hash function
-  result_type _seed;
-  result_type _null_hash;
-};
-
-template <typename Nullate>
-class element_hasher_adaptor {
-  static constexpr result_type NULL_HASH     = cuda::std::numeric_limits<result_type>::max();
-  static constexpr result_type NON_NULL_HASH = 0;
-
- public:
-  __device__ element_hasher_adaptor(Nullate _check_nulls, result_type seed) noexcept
-    : _element_hasher(_check_nulls, seed)
-  {
-  }
-
-  template <typename T>
-  __device__ result_type operator()(column_device_view const& col,
-                                    size_type row_index) const noexcept
-    requires(not cudf::is_nested<T>() and not cudf::is_dictionary<T>())
-  {
-    return _element_hasher.template operator()<T>(col, row_index);
-  }
-
-  template <typename T>
-  __device__ result_type operator()(column_device_view const& col,
-                                    size_type row_index) const noexcept
-    requires(cudf::is_nested<T>() or cudf::is_dictionary<T>())
-  {
-    CUDF_UNREACHABLE("Can't get here yet");
-  }
-
-  //template <typename T>
-  //__device__ result_type operator()(column_device_view const& col,
-  //                                  size_type row_index) const noexcept
-  //  requires(cudf::is_dictionary<T>())
-  //{
-  //  if constexpr (Nullate) {
-  //    if (col.is_null(row_index)) { return NULL_HASH; }
-  //  }
-  //
-  //  auto const keys = col.child(dictionary_column_view::keys_column_index);
-  //  return type_dispatcher<dispatch_storage_type>(
-  //    keys.type(),
-  //    _element_hasher,
-  //    keys,
-  //    static_cast<size_type>(col.element<dictionary32>(row_index)));
-  //}
-  //
-  //template <typename T>
-  //__device__ result_type operator()(column_device_view const& col,
-  //                                  size_type row_index) const noexcept
-  //  requires(cudf::is_nested<T>())
-  //{
-  //  auto hash                   = result_type{0};
-  //  column_device_view curr_col = col.slice(row_index, 1);
-  //  while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) {
-  //    if constexpr (Nullate) {
-  //      auto validity_it = detail::make_validity_iterator<true>(curr_col);
-  //      hash             = detail::accumulate(
-  //        validity_it, validity_it + curr_col.size(), hash, [](auto hash, auto is_valid) {
-  //          return cudf::hashing::detail::hash_combine(hash,
-  //                                                     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 = detail::structs_column_device_view(curr_col).get_sliced_child(0);
-  //    } else if (curr_col.type().id() == type_id::LIST) {
-  //      auto list_col   = detail::lists_column_device_view(curr_col);
-  //      auto list_sizes = make_list_size_iterator(list_col);
-  //      hash            = detail::accumulate(
-  //        list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
-  //          return cudf::hashing::detail::hash_combine(hash, 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,
-  //      type_dispatcher<dispatch_void_if_nested>(curr_col.type(), _element_hasher, curr_col, i));
-  //  }
-  //  return hash;
-  //}
-
-  element_hasher<Nullate> const _element_hasher;
-};
-
-
-template <typename Nullate, typename T>
-__device__ hash_value_type hasher_impl(Nullate check_nulls, cudf::column_device_view col, uint32_t seed) {
-  auto hasher = element_hasher_adaptor<Nullate>{check_nulls, seed};
-  return hasher.template operator()<T>(col, threadIdx.x);
-}
-
-template <typename T>
-__device__ hash_value_type hasher(cudf::column_device_view col, uint32_t seed, bool const nullable) {
-  return hasher_impl<nullate::DYNAMIC, T>(nullate::DYNAMIC{nullable}, col, seed);
-}
-} // namespace cudf::hashing::detail
diff --git a/cpp/src/jit_lto/AlgorithmLauncher.cpp b/cpp/src/jit_lto/AlgorithmLauncher.cpp
new file mode 100644
index 000000000000..def84fb3626d
--- /dev/null
+++ b/cpp/src/jit_lto/AlgorithmLauncher.cpp
@@ -0,0 +1,51 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/detail/jit_lto/AlgorithmLauncher.hpp>
+#include <cudf/utilities/error.hpp>
+
+namespace cudf::detail::jit_lto {
+
+AlgorithmLauncher::AlgorithmLauncher(cudaKernel_t k, cudaLibrary_t lib) : kernel{k}, library{lib} {}
+
+AlgorithmLauncher::~AlgorithmLauncher()
+{
+  if (library != nullptr) { (void)cudaLibraryUnload(library); }
+}
+
+AlgorithmLauncher::AlgorithmLauncher(AlgorithmLauncher&& other) noexcept
+  : kernel{other.kernel}, library{other.library}
+{
+  other.kernel  = nullptr;
+  other.library = nullptr;
+}
+
+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;
+}
+
+void AlgorithmLauncher::call(
+  cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, void** kernel_args)
+{
+  cudaLaunchConfig_t config{};
+  config.gridDim          = grid;
+  config.blockDim         = block;
+  config.stream           = stream;
+  config.dynamicSmemBytes = shared_mem;
+  config.numAttrs         = 0;
+  config.attrs            = NULL;
+
+  CUDF_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
+}
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/src/jit_lto/AlgorithmPlanner.cpp b/cpp/src/jit_lto/AlgorithmPlanner.cpp
new file mode 100644
index 000000000000..79a5caad2ce7
--- /dev/null
+++ b/cpp/src/jit_lto/AlgorithmPlanner.cpp
@@ -0,0 +1,98 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/detail/jit_lto/AlgorithmPlanner.hpp>
+#include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <cuda_runtime.h>
+#include <nvJitLink.h>
+
+#include <memory>
+#include <shared_mutex>
+#include <string>
+#include <vector>
+
+#include <mutex>
+
+namespace cudf::detail::jit_lto {
+
+std::string AlgorithmPlanner::get_fragments_key() const
+{
+  std::string key = "";
+  for (const auto& fragment : this->fragments) {
+    key += fragment->get_key();
+  }
+  return key;
+}
+
+std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::read_cache(std::string const& launch_key) const
+{
+  auto& launchers = jit_cache_.launchers;
+  std::shared_lock<std::shared_mutex> read_lock(jit_cache_.mutex);
+  if (auto it = launchers.find(launch_key); it != launchers.end()) { return it->second; }
+  return nullptr;
+}
+
+std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::get_launcher()
+{
+  auto& launchers = jit_cache_.launchers;
+  auto launch_key = this->get_fragments_key();
+
+  if (auto hit = read_cache(launch_key)) { return hit; }
+
+  std::unique_lock<std::shared_mutex> write_lock(jit_cache_.mutex);
+  if (auto it = launchers.find(launch_key); it != launchers.end()) { return it->second; }
+
+  auto launcher         = this->build();
+  launchers[launch_key] = launcher;
+  return launcher;
+}
+
+std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::build()
+{
+  int device = 0;
+  int major  = 0;
+  int minor  = 0;
+  CUDF_CUDA_TRY(cudaGetDevice(&device));
+  CUDF_CUDA_TRY(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device));
+  CUDF_CUDA_TRY(cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device));
+
+  std::string archs = "-arch=sm_" + std::to_string((major * 10 + minor));
+
+  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()));
+
+  return std::make_shared<AlgorithmLauncher>(kernel, library);
+}
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/src/jit_lto/FragmentEntry.cpp b/cpp/src/jit_lto/FragmentEntry.cpp
new file mode 100644
index 000000000000..7c66c33dfd80
--- /dev/null
+++ b/cpp/src/jit_lto/FragmentEntry.cpp
@@ -0,0 +1,18 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/detail/jit_lto/FragmentEntry.hpp>
+
+namespace cudf::detail::jit_lto {
+
+bool FatbinFragmentEntry::add_to(nvJitLinkHandle& handle) const
+{
+  auto result = nvJitLinkAddData(handle, NVJITLINK_INPUT_ANY, get_data(), get_length(), get_key());
+
+  check_nvjitlink_result(handle, result);
+  return true;
+}
+
+}  // namespace cudf::detail::jit_lto
diff --git a/cpp/src/jit_lto/nvjitlink_checker.cpp b/cpp/src/jit_lto/nvjitlink_checker.cpp
new file mode 100644
index 000000000000..63d458ddf102
--- /dev/null
+++ b/cpp/src/jit_lto/nvjitlink_checker.cpp
@@ -0,0 +1,30 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <memory>
+#include <nvJitLink.h>
+#include <string>
+
+namespace cudf::detail::jit_lto {
+
+void check_nvjitlink_result(nvJitLinkHandle handle, nvJitLinkResult result)
+{
+  if (result != NVJITLINK_SUCCESS) {
+    std::string error_msg = "nvJITLink failed with error " + std::to_string(result);
+    size_t log_size       = 0;
+    result                = nvJitLinkGetErrorLogSize(handle, &log_size);
+    if (result == NVJITLINK_SUCCESS && log_size > 0) {
+      std::unique_ptr<char[]> log{new char[log_size]};
+      result = nvJitLinkGetErrorLog(handle, log.get());
+      if (result == NVJITLINK_SUCCESS) { error_msg += "\n" + std::string(log.get()); }
+    }
+    CUDF_FAIL(error_msg);
+  }
+}
+
+}  // namespace cudf::detail::jit_lto

From 1881bfa2223482c4dce71ee44b33004a7c3794b9 Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Wed, 6 May 2026 23:20:54 +0000
Subject: [PATCH 03/11] remove udf reference

---
 .../cudf/detail/jit_lto/FragmentEntry.hpp     | 25 +++----------------
 1 file changed, 4 insertions(+), 21 deletions(-)

diff --git a/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp b/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp
index d4269bd2e7f2..aed48c242393 100644
--- a/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp
+++ b/cpp/include/cudf/detail/jit_lto/FragmentEntry.hpp
@@ -4,16 +4,16 @@
  */
 #pragma once
 
+#include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
+
+#include <nvJitLink.h>
+
 #include <cstddef>
 #include <cstdint>
 #include <string>
 #include <typeinfo>
 #include <vector>
 
-#include <nvJitLink.h>
-
-#include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
-
 namespace cudf::detail::jit_lto {
 
 struct FragmentEntry {
@@ -47,21 +47,4 @@ struct StaticFatbinFragmentEntry final : FatbinFragmentEntry {
   static const size_t length;
 };
 
-struct UDFFatbinFragment final : FatbinFragmentEntry {
-  UDFFatbinFragment(std::string key, std::vector<uint8_t> bytes)
-    : key_(std::move(key)), bytes_(std::move(bytes))
-  {
-  }
-
-  const uint8_t* get_data() const override { return bytes_.data(); }
-
-  size_t get_length() const override { return bytes_.size(); }
-
-  const char* get_key() const override { return key_.c_str(); }
-
- private:
-  std::string key_;
-  std::vector<uint8_t> bytes_;
-};
-
 }  // namespace cudf::detail::jit_lto

From 5bfce796d4cde8cdb0bd81ba4550011b6b56d994 Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Wed, 6 May 2026 23:26:13 +0000
Subject: [PATCH 04/11] run pre-commit

---
 cpp/CMakeLists.txt                            | 73 +++++++++++--------
 .../Modules/generate_jit_lto_kernels.cmake    | 33 +++++----
 .../cudf/detail/jit_lto/AlgorithmLauncher.hpp |  9 ++-
 .../murmurhash3_x86_32_jit_device.cuh         | 28 ++++---
 .../murmurhash3_x86_32_jit_hasher_decl.cuh    | 14 ++--
 .../murmurhash3_x86_32_lto.cuh                | 29 +++-----
 .../murmurhash_entry_kernel.cu.in             |  2 +-
 .../murmurhash_hasher_fragment.cu.in          |  2 +-
 .../murmurhash_hasher_noop_fragment.cu.in     |  2 +-
 cpp/src/hash/murmurhash3_x86_32.cu            |  4 +-
 .../hash/murmurhash3_x86_32_jit_launch.hpp    | 60 +++++++++------
 cpp/src/jit_lto/AlgorithmPlanner.cpp          |  4 +-
 cpp/src/jit_lto/nvjitlink_checker.cpp         |  3 +-
 13 files changed, 150 insertions(+), 113 deletions(-)

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 7afd47436d41..30de4057c88b 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -331,28 +331,24 @@ endif()
 
 # ##################################################################################################
 # * JIT+LTO (MurmurHash3 x86 32 device fragments + nvjitlink) --------------------------------------
-# Same CMake pattern as cuVS (`cpp/CMakeLists.txt` + `cmake/modules/generate_jit_lto_kernels.cmake`).
 include(cmake/Modules/generate_jit_lto_kernels.cmake)
 
 add_library(cudf_jit_lto_kernel_usage_requirements INTERFACE)
 target_include_directories(
   cudf_jit_lto_kernel_usage_requirements
-  INTERFACE "${CMAKE_CURRENT_SOURCE_DIR}/include"
-            "${CMAKE_CURRENT_SOURCE_DIR}/src"
+  INTERFACE "${CMAKE_CURRENT_SOURCE_DIR}/include" "${CMAKE_CURRENT_SOURCE_DIR}/src"
             "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels"
 )
 target_compile_options(
-  cudf_jit_lto_kernel_usage_requirements
-  INTERFACE "$<$<COMPILE_LANGUAGE:CXX>:${CUDF_CXX_FLAGS}>"
-            "$<$<COMPILE_LANGUAGE:CUDA>:${CUDF_CUDA_FLAGS}>"
+  cudf_jit_lto_kernel_usage_requirements INTERFACE "$<$<COMPILE_LANGUAGE:CXX>:${CUDF_CXX_FLAGS}>"
+                                                   "$<$<COMPILE_LANGUAGE:CUDA>:${CUDF_CUDA_FLAGS}>"
 )
 target_compile_features(cudf_jit_lto_kernel_usage_requirements INTERFACE cuda_std_20)
 target_link_libraries(
   cudf_jit_lto_kernel_usage_requirements INTERFACE CCCL::CCCL rmm::rmm
-                                                    $<BUILD_LOCAL_INTERFACE:cuco::cuco>
+                                                   $<BUILD_LOCAL_INTERFACE:cuco::cuco>
 )
 
-# Same arch ladder as cuVS JIT fragments (single real arch for fatbin link).
 set(JIT_LTO_TARGET_ARCHITECTURE "70-real")
 if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
   set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
@@ -365,35 +361,54 @@ set(CMAKE_CUDA_ARCHITECTURES ${JIT_LTO_TARGET_ARCHITECTURE})
 set(murmur_jit_ns "cudf::hashing::detail::jit_lto")
 generate_jit_lto_kernels(
   cudf_jit_lto_generated_sources
-  NAME_FORMAT "murmurhash_jit_hasher_@abbrev@"
-  MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json"
-  KERNEL_INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in"
+  NAME_FORMAT
+  "murmurhash_jit_hasher_@abbrev@"
+  MATRIX_JSON_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json"
+  KERNEL_INPUT_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in"
   FRAGMENT_TAG_FORMAT
-    "${murmur_jit_ns}::fragment_tag_murmur_hasher<${murmur_jit_ns}::tag_@abbrev@>"
-  FRAGMENT_TAG_HEADER_FILES "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
-  OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/hasher"
-  KERNEL_LINK_LIBRARIES cudf_jit_lto_kernel_usage_requirements
+  "${murmur_jit_ns}::fragment_tag_murmur_hasher<${murmur_jit_ns}::tag_@abbrev@>"
+  FRAGMENT_TAG_HEADER_FILES
+  "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY
+  "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/hasher"
+  KERNEL_LINK_LIBRARIES
+  cudf_jit_lto_kernel_usage_requirements
 )
 generate_jit_lto_kernels(
   cudf_jit_lto_generated_sources
-  NAME_FORMAT "murmurhash_jit_hasher_noop_@abbrev@"
-  MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json"
-  KERNEL_INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in"
+  NAME_FORMAT
+  "murmurhash_jit_hasher_noop_@abbrev@"
+  MATRIX_JSON_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_matrix.json"
+  KERNEL_INPUT_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in"
   FRAGMENT_TAG_FORMAT
-    "${murmur_jit_ns}::fragment_tag_murmur_hasher_noop<${murmur_jit_ns}::tag_@abbrev@>"
-  FRAGMENT_TAG_HEADER_FILES "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
-  OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/hasher_noop"
-  KERNEL_LINK_LIBRARIES cudf_jit_lto_kernel_usage_requirements
+  "${murmur_jit_ns}::fragment_tag_murmur_hasher_noop<${murmur_jit_ns}::tag_@abbrev@>"
+  FRAGMENT_TAG_HEADER_FILES
+  "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY
+  "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/hasher_noop"
+  KERNEL_LINK_LIBRARIES
+  cudf_jit_lto_kernel_usage_requirements
 )
 generate_jit_lto_kernels(
   cudf_jit_lto_generated_sources
-  NAME_FORMAT "murmurhash_jit_@k@"
-  MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json"
-  KERNEL_INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in"
-  FRAGMENT_TAG_FORMAT "${murmur_jit_ns}::fragment_tag_murmur_entry"
-  FRAGMENT_TAG_HEADER_FILES "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
-  OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/entry"
-  KERNEL_LINK_LIBRARIES cudf_jit_lto_kernel_usage_requirements
+  NAME_FORMAT
+  "murmurhash_jit_@k@"
+  MATRIX_JSON_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_entry_matrix.json"
+  KERNEL_INPUT_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in"
+  FRAGMENT_TAG_FORMAT
+  "${murmur_jit_ns}::fragment_tag_murmur_entry"
+  FRAGMENT_TAG_HEADER_FILES
+  "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY
+  "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/entry"
+  KERNEL_LINK_LIBRARIES
+  cudf_jit_lto_kernel_usage_requirements
 )
 endblock()
 
diff --git a/cpp/cmake/Modules/generate_jit_lto_kernels.cmake b/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
index 0f45ab884a76..b074aaa8f27a 100644
--- a/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
+++ b/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
@@ -4,15 +4,12 @@
 # SPDX-License-Identifier: Apache-2.0
 # cmake-format: on
 # =============================================================================
-# cuVS uses DEPENDS $<TARGET_OBJECTS:${kernel_target}>. With CUDA_FATBIN_COMPILATION that genex
-# expands to .fatbin paths; add_custom_command then requires those files as separate Makefile
-# prerequisites, but no rule exists (gmake: "No rule to make target ... .fatbin"). Depending on
-# the OBJECT target keeps ordering correct; COMMAND still passes $<TARGET_OBJECTS:...> to bin2c.
 
 include_guard(GLOBAL)
 
 include(${CMAKE_CURRENT_LIST_DIR}/compute_matrix_product.cmake)
 
+# This function compiles a JIT LTO kernel object target and writes a C header embedding its fatbin.
 function(add_jit_lto_kernel kernel_target)
   set(options)
   set(one_value KERNEL_FILE FATBIN_HEADER_FILE)
@@ -36,12 +33,14 @@ function(add_jit_lto_kernel kernel_target)
 
   add_custom_command(
     OUTPUT "${_JIT_LTO_FATBIN_HEADER_FILE}"
+    COMMENT "Generate embedded fatbin header for ${kernel_target}"
     COMMAND "${bin_to_c}" --const --name embedded_fatbin --static $<TARGET_OBJECTS:${kernel_target}>
             > "${_JIT_LTO_FATBIN_HEADER_FILE}"
     DEPENDS ${kernel_target}
   )
 endfunction()
 
+# This function materializes one matrix entry into configured sources and registers its fatbin.
 function(process_jit_lto_matrix_entry source_list_var)
   set(options)
   set(one_value NAME_FORMAT KERNEL_INPUT_FILE OUTPUT_DIRECTORY FRAGMENT_TAG_FORMAT
@@ -71,9 +70,7 @@ function(process_jit_lto_matrix_entry source_list_var)
   configure_file("${CMAKE_CURRENT_FUNCTION_LIST_DIR}/register_fatbin.cpp.in" "${fatbin_file}" @ONLY)
 
   add_jit_lto_kernel(
-    ${kernel_target}
-    KERNEL_FILE "${kernel_file}"
-    FATBIN_HEADER_FILE "${fatbin_header_file}"
+    ${kernel_target} KERNEL_FILE "${kernel_file}" FATBIN_HEADER_FILE "${fatbin_header_file}"
     LINK_LIBRARIES ${_JIT_LTO_KERNEL_LINK_LIBRARIES}
   )
   list(APPEND ${source_list_var} "${fatbin_header_file}" "${fatbin_file}")
@@ -83,6 +80,7 @@ function(process_jit_lto_matrix_entry source_list_var)
   )
 endfunction()
 
+# This function expands a JIT LTO matrix JSON description into kernel and fatbin source files.
 function(generate_jit_lto_kernels source_list_var)
   set(options)
   set(one_value NAME_FORMAT MATRIX_JSON_FILE MATRIX_JSON_STRING KERNEL_INPUT_FILE
@@ -118,13 +116,20 @@ function(generate_jit_lto_kernels source_list_var)
     string(JSON matrix_json_entry GET "${matrix_product}" "${i}")
     process_jit_lto_matrix_entry(
       "${source_list_var}"
-      NAME_FORMAT "${_JIT_LTO_NAME_FORMAT}"
-      KERNEL_INPUT_FILE "${_JIT_LTO_KERNEL_INPUT_FILE}"
-      FRAGMENT_TAG_FORMAT "${_JIT_LTO_FRAGMENT_TAG_FORMAT}"
-      FRAGMENT_TAG_HEADER_FILES ${_JIT_LTO_FRAGMENT_TAG_HEADER_FILES}
-      OUTPUT_DIRECTORY "${_JIT_LTO_OUTPUT_DIRECTORY}"
-      MATRIX_JSON_ENTRY "${matrix_json_entry}"
-      KERNEL_LINK_LIBRARIES ${_JIT_LTO_KERNEL_LINK_LIBRARIES}
+      NAME_FORMAT
+      "${_JIT_LTO_NAME_FORMAT}"
+      KERNEL_INPUT_FILE
+      "${_JIT_LTO_KERNEL_INPUT_FILE}"
+      FRAGMENT_TAG_FORMAT
+      "${_JIT_LTO_FRAGMENT_TAG_FORMAT}"
+      FRAGMENT_TAG_HEADER_FILES
+      ${_JIT_LTO_FRAGMENT_TAG_HEADER_FILES}
+      OUTPUT_DIRECTORY
+      "${_JIT_LTO_OUTPUT_DIRECTORY}"
+      MATRIX_JSON_ENTRY
+      "${matrix_json_entry}"
+      KERNEL_LINK_LIBRARIES
+      ${_JIT_LTO_KERNEL_LINK_LIBRARIES}
     )
   endforeach()
 
diff --git a/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp b/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp
index c68cd70e185e..c09592608c56 100644
--- a/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp
+++ b/cpp/include/cudf/detail/jit_lto/AlgorithmLauncher.hpp
@@ -4,14 +4,15 @@
  */
 #pragma once
 
-#include <cstdint>
+#include <cuda_runtime.h>
+
 #include <driver_types.h>
-#include <string>
-#include <unordered_map>
 #include <vector_types.h>
 
-#include <cuda_runtime.h>
+#include <cstdint>
 #include <memory>
+#include <string>
+#include <unordered_map>
 
 namespace cudf::detail::jit_lto {
 
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
index 7edcdd95fcb8..8b9d061b9993 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
@@ -1,13 +1,13 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
 
 #include "murmurhash3_x86_32_jit_hasher_decl.cuh"
 
-#include <cudf/detail/utilities/grid_1d.cuh>
 #include <cudf/detail/utilities/assert.cuh>
+#include <cudf/detail/utilities/grid_1d.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/utilities/type_dispatcher.hpp>
 
@@ -44,23 +44,31 @@ __device__ inline hash_value_type murmur_jit_hash_dispatcher(column_device_view
     case type_id::TIMESTAMP_DAYS:
       return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_DAYS>>(col, seed, nullable, row_index);
     case type_id::TIMESTAMP_SECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_SECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_SECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::TIMESTAMP_MILLISECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_MILLISECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_MILLISECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::TIMESTAMP_MICROSECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_MICROSECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_MICROSECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::TIMESTAMP_NANOSECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_NANOSECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::TIMESTAMP_NANOSECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::DURATION_DAYS:
       return murmur_jit_hasher<id_to_type<type_id::DURATION_DAYS>>(col, seed, nullable, row_index);
     case type_id::DURATION_SECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::DURATION_SECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_SECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::DURATION_MILLISECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::DURATION_MILLISECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_MILLISECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::DURATION_MICROSECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::DURATION_MICROSECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_MICROSECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::DURATION_NANOSECONDS:
-      return murmur_jit_hasher<id_to_type<type_id::DURATION_NANOSECONDS>>(col, seed, nullable, row_index);
+      return murmur_jit_hasher<id_to_type<type_id::DURATION_NANOSECONDS>>(
+        col, seed, nullable, row_index);
     case type_id::DICTIONARY32:
       return murmur_jit_hasher<id_to_type<type_id::DICTIONARY32>>(col, seed, nullable, row_index);
     case type_id::STRING:
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
index 44994a48bcf0..2be689ef3c56 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -14,14 +14,12 @@ namespace cudf::hashing::detail {
  * @brief Forward declaration only (no `murmur_jit_hash_dispatcher` here).
  *
  * Per-type explicit specializations must appear in the TU before any use of that specialization.
- * Hasher / noop fragment TUs include this header, then define `template <> ... murmur_jit_hasher<T>`.
- * The entry kernel includes `murmurhash3_x86_32_jit_device.cuh`, which pulls this in and adds the
- * dispatcher that calls `murmur_jit_hasher` for every storage type.
+ * Hasher / noop fragment TUs include this header, then define `template <> ...
+ * murmur_jit_hasher<T>`. The entry kernel includes `murmurhash3_x86_32_jit_device.cuh`, which pulls
+ * this in and adds the dispatcher that calls `murmur_jit_hasher` for every storage type.
  */
 template <typename T>
-extern __device__ hash_value_type murmur_jit_hasher(column_device_view col,
-                                                    uint32_t seed,
-                                                    bool nullable,
-                                                    size_type row_index);
+extern __device__ hash_value_type
+murmur_jit_hasher(column_device_view col, uint32_t seed, bool nullable, size_type row_index);
 
 }  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
index c5fdfae95349..27c40f4cd245 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -10,9 +10,9 @@
 #include <cudf/detail/utilities/accumulate.cuh>
 #include <cudf/detail/utilities/assert.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/hashing.hpp>
 #include <cudf/hashing/detail/hashing.hpp>
 #include <cudf/hashing/detail/murmurhash3_x86_32.cuh>
-#include <cudf/hashing.hpp>
 #include <cudf/lists/list_device_view.cuh>
 #include <cudf/lists/lists_column_device_view.cuh>
 #include <cudf/structs/structs_column_device_view.cuh>
@@ -108,10 +108,9 @@ class element_hasher_adaptor {
     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(
+        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);
+            return cudf::hashing::detail::hash_combine(h, is_valid ? NON_NULL_HASH : NULL_HASH);
           });
       }
       if (curr_col.type().id() == type_id::STRUCT) {
@@ -128,10 +127,10 @@ class element_hasher_adaptor {
       }
     }
     for (int i = 0; i < curr_col.size(); ++i) {
-      hash = cudf::hashing::detail::hash_combine(
-        hash,
-        type_dispatcher<cudf::detail::dispatch_void_if_nested>(
-          curr_col.type(), _element_hasher, curr_col, i));
+      hash =
+        cudf::hashing::detail::hash_combine(hash,
+                                            type_dispatcher<cudf::detail::dispatch_void_if_nested>(
+                                              curr_col.type(), _element_hasher, curr_col, i));
     }
     return hash;
   }
@@ -141,20 +140,16 @@ class element_hasher_adaptor {
 };
 
 template <typename T>
-__device__ hash_value_type hasher_impl(bool check_nulls,
-                                       cudf::column_device_view col,
-                                       uint32_t seed,
-                                       size_type row_index)
+__device__ hash_value_type
+hasher_impl(bool check_nulls, cudf::column_device_view col, uint32_t seed, size_type row_index)
 {
   auto const hasher = element_hasher_adaptor{check_nulls, seed};
   return hasher.template operator()<T>(col, row_index);
 }
 
 template <typename T>
-__device__ hash_value_type hasher(cudf::column_device_view col,
-                                  uint32_t seed,
-                                  bool const nullable,
-                                  size_type row_index)
+__device__ hash_value_type
+hasher(cudf::column_device_view col, uint32_t seed, bool const nullable, size_type row_index)
 {
   return hasher_impl<T>(nullable, col, seed, row_index);
 }
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
index c8f406a4f728..1ff7a1c343d1 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
index 3b52c19ba8b1..af8e6c239380 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 // Strong explicit specialization of `murmur_jit_hasher<T>` (real hasher); linked when logical
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in
index 65dd6cb03681..f6284bd55aa8 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_noop_fragment.cu.in
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 // Strong explicit specialization with a no-op body for `T` not used by any column in the table.
diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu
index 8a3e10b26f63..b493204a1cdd 100644
--- a/cpp/src/hash/murmurhash3_x86_32.cu
+++ b/cpp/src/hash/murmurhash3_x86_32.cu
@@ -2,6 +2,8 @@
  * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
+#include "murmurhash3_x86_32_jit_launch.hpp"
+
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/hashing/detail/hashing.hpp>
 
@@ -9,8 +11,6 @@
 
 #include <memory>
 
-#include "murmurhash3_x86_32_jit_launch.hpp"
-
 namespace cudf {
 namespace hashing {
 namespace detail {
diff --git a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
index 44e2a810d9ea..ad6d490cbe68 100644
--- a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
+++ b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -93,8 +93,12 @@ inline void add_strong_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner&
     case type_id::UINT16: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u16>>(); break;
     case type_id::UINT32: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u32>>(); break;
     case type_id::UINT64: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_u64>>(); break;
-    case type_id::FLOAT32: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_f32>>(); break;
-    case type_id::FLOAT64: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_f64>>(); break;
+    case type_id::FLOAT32:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_f32>>();
+      break;
+    case type_id::FLOAT64:
+      planner.add_static_fragment<fragment_tag_murmur_hasher<tag_f64>>();
+      break;
     case type_id::BOOL8: planner.add_static_fragment<fragment_tag_murmur_hasher<tag_b8>>(); break;
     case type_id::TIMESTAMP_DAYS:
       planner.add_static_fragment<fragment_tag_murmur_hasher<tag_ts_day>>();
@@ -151,7 +155,9 @@ inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& pl
 {
   using namespace cudf::hashing::detail::jit_lto;
   switch (id) {
-    case type_id::INT8: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i8>>(); break;
+    case type_id::INT8:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i8>>();
+      break;
     case type_id::INT16:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i16>>();
       break;
@@ -161,7 +167,9 @@ inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& pl
     case type_id::INT64:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_i64>>();
       break;
-    case type_id::UINT8: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u8>>(); break;
+    case type_id::UINT8:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u8>>();
+      break;
     case type_id::UINT16:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_u16>>();
       break;
@@ -177,7 +185,9 @@ inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& pl
     case type_id::FLOAT64:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_f64>>();
       break;
-    case type_id::BOOL8: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_b8>>(); break;
+    case type_id::BOOL8:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_b8>>();
+      break;
     case type_id::TIMESTAMP_DAYS:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_ts_day>>();
       break;
@@ -211,8 +221,12 @@ inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& pl
     case type_id::DICTIONARY32:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_dict>>();
       break;
-    case type_id::STRING: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_str>>(); break;
-    case type_id::LIST: planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_list>>(); break;
+    case type_id::STRING:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_str>>();
+      break;
+    case type_id::LIST:
+      planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_list>>();
+      break;
     case type_id::DECIMAL32:
       planner.add_static_fragment<fragment_tag_murmur_hasher_noop<tag_dec32>>();
       break;
@@ -232,9 +246,9 @@ inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& pl
 }  // namespace
 
 inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
-                                                     uint32_t seed,
-                                                     rmm::cuda_stream_view stream,
-                                                     rmm::device_async_resource_ref mr)
+                                                      uint32_t seed,
+                                                      rmm::cuda_stream_view stream,
+                                                      rmm::device_async_resource_ref mr)
 {
   auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),
                                     input.num_rows(),
@@ -244,9 +258,8 @@ inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
 
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
-  bool const nullable = has_nulls(input);
-  auto const preprocessed =
-    cudf::detail::row::hash::preprocessed_table::create(input, stream);
+  bool const nullable     = has_nulls(input);
+  auto const preprocessed = cudf::detail::row::hash::preprocessed_table::create(input, stream);
   table_device_view const input_dv{*preprocessed};
 
   auto output_view = output->mutable_view();
@@ -271,15 +284,16 @@ inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
   auto launcher = planner.get_launcher();
 
   cudf::detail::grid_1d const grid{input.num_rows(), 256};
-  launcher->dispatch<void(cudf::mutable_column_device_view, uint32_t, cudf::table_device_view, bool)>(
-    stream.value(),
-    dim3(grid.num_blocks),
-    dim3(grid.num_threads_per_block),
-    0,
-    *d_output,
-    seed,
-    input_dv,
-    nullable);
+  launcher
+    ->dispatch<void(cudf::mutable_column_device_view, uint32_t, cudf::table_device_view, bool)>(
+      stream.value(),
+      dim3(grid.num_blocks),
+      dim3(grid.num_threads_per_block),
+      0,
+      *d_output,
+      seed,
+      input_dv,
+      nullable);
 
   return output;
 }
diff --git a/cpp/src/jit_lto/AlgorithmPlanner.cpp b/cpp/src/jit_lto/AlgorithmPlanner.cpp
index 79a5caad2ce7..8be6e638d79f 100644
--- a/cpp/src/jit_lto/AlgorithmPlanner.cpp
+++ b/cpp/src/jit_lto/AlgorithmPlanner.cpp
@@ -8,15 +8,15 @@
 #include <cudf/utilities/error.hpp>
 
 #include <cuda_runtime.h>
+
 #include <nvJitLink.h>
 
 #include <memory>
+#include <mutex>
 #include <shared_mutex>
 #include <string>
 #include <vector>
 
-#include <mutex>
-
 namespace cudf::detail::jit_lto {
 
 std::string AlgorithmPlanner::get_fragments_key() const
diff --git a/cpp/src/jit_lto/nvjitlink_checker.cpp b/cpp/src/jit_lto/nvjitlink_checker.cpp
index 63d458ddf102..25bdc27cd5b9 100644
--- a/cpp/src/jit_lto/nvjitlink_checker.cpp
+++ b/cpp/src/jit_lto/nvjitlink_checker.cpp
@@ -6,8 +6,9 @@
 #include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
 #include <cudf/utilities/error.hpp>
 
-#include <memory>
 #include <nvJitLink.h>
+
+#include <memory>
 #include <string>
 
 namespace cudf::detail::jit_lto {

From 7a0c3f1d9f2ed9781044b0fcbf53f87fbbc6550f Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Thu, 7 May 2026 00:15:40 +0000
Subject: [PATCH 05/11] fix style check

---
 cpp/cmake/Modules/compute_matrix_product.cmake | 2 ++
 cpp/cmake/Modules/compute_matrix_product.py    | 2 ++
 2 files changed, 4 insertions(+)

diff --git a/cpp/cmake/Modules/compute_matrix_product.cmake b/cpp/cmake/Modules/compute_matrix_product.cmake
index 3192bee98989..a00070a6018d 100644
--- a/cpp/cmake/Modules/compute_matrix_product.cmake
+++ b/cpp/cmake/Modules/compute_matrix_product.cmake
@@ -7,6 +7,7 @@
 
 include_guard(GLOBAL)
 
+# This function runs compute_matrix_product.py and stores the JSON matrix product in output_var.
 function(compute_matrix_product output_var)
   set(options)
   set(one_value MATRIX_JSON_FILE MATRIX_JSON_STRING)
@@ -37,6 +38,7 @@ function(compute_matrix_product output_var)
   )
 endfunction()
 
+# This function unpacks a JSON object into CMake variables in the caller scope.
 function(populate_matrix_variables matrix_json_entry)
   string(JSON len LENGTH "${matrix_json_entry}")
   math(EXPR last "${len} - 1")
diff --git a/cpp/cmake/Modules/compute_matrix_product.py b/cpp/cmake/Modules/compute_matrix_product.py
index 76262f6597ae..1d88942b2216 100644
--- a/cpp/cmake/Modules/compute_matrix_product.py
+++ b/cpp/cmake/Modules/compute_matrix_product.py
@@ -8,6 +8,8 @@
 # https://gitlab.kitware.com/cmake/cmake/-/merge_requests/11516, we may be able
 # to port the algorithm to CMake script and use it in other RAPIDS projects.
 
+from __future__ import annotations
+
 import argparse
 import json
 import re

From 22f1761e0b74b90175347aae207745f206dfbe6e Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Thu, 7 May 2026 00:49:30 +0000
Subject: [PATCH 06/11] exclude nvjitlink from wheel install path

---
 ci/build_wheel_libcudf.sh | 1 +
 1 file changed, 1 insertion(+)

diff --git a/ci/build_wheel_libcudf.sh b/ci/build_wheel_libcudf.sh
index 3680f1ebe10a..b468400aa28b 100755
--- a/ci/build_wheel_libcudf.sh
+++ b/ci/build_wheel_libcudf.sh
@@ -35,6 +35,7 @@ export SKBUILD_CMAKE_ARGS="-DUSE_NVCOMP_RUNTIME_WHEEL=ON"
 python -m auditwheel repair \
     --exclude libkvikio.so \
     --exclude libnvcomp.so.5 \
+    --exclude libnvJitLink.so.* \
     --exclude librapids_logger.so \
     --exclude librmm.so \
     -w "${RAPIDS_WHEEL_BLD_OUTPUT_DIR}" \

From 551693f0bdd85c8a260c4de6efe477b7329d6a6b Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Thu, 7 May 2026 01:38:08 +0000
Subject: [PATCH 07/11] remove double compilation

---
 cpp/cmake/Modules/generate_jit_lto_kernels.cmake | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/cmake/Modules/generate_jit_lto_kernels.cmake b/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
index b074aaa8f27a..95e8c314e3c6 100644
--- a/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
+++ b/cpp/cmake/Modules/generate_jit_lto_kernels.cmake
@@ -36,7 +36,7 @@ function(add_jit_lto_kernel kernel_target)
     COMMENT "Generate embedded fatbin header for ${kernel_target}"
     COMMAND "${bin_to_c}" --const --name embedded_fatbin --static $<TARGET_OBJECTS:${kernel_target}>
             > "${_JIT_LTO_FATBIN_HEADER_FILE}"
-    DEPENDS ${kernel_target}
+    DEPENDS $<TARGET_OBJECTS:${kernel_target}>
   )
 endfunction()
 

From bab9125dcace7bad42dc85519a1e39057872ed50 Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Thu, 7 May 2026 21:01:57 +0000
Subject: [PATCH 08/11] make dispatcher its own fragment

---
 cpp/CMakeLists.txt                            | 17 ++++++++++
 .../detail/murmurhash3_x86_32_jit_tags.hpp    |  3 ++
 ... murmurhash3_x86_32_jit_dispatch_impl.cuh} | 15 ++++----
 .../murmurhash3_x86_32_jit_hasher_decl.cuh    | 15 ++++++--
 .../murmurhash3_x86_32_lto.cuh                | 21 ++++++------
 .../murmurhash_dispatch_fragment.cu.in        | 11 ++++++
 .../murmurhash_dispatch_matrix.json           |  7 ++++
 .../murmurhash_entry_kernel.cu.in             |  4 ++-
 .../murmurhash_hasher_fragment.cu.in          |  5 ++-
 .../hash/murmurhash3_x86_32_jit_launch.hpp    | 34 ++++++++++++++++---
 cpp/src/jit_lto/AlgorithmPlanner.cpp          | 25 ++++++++++++++
 11 files changed, 126 insertions(+), 31 deletions(-)
 rename cpp/src/hash/jit_lto_kernels/{murmurhash3_x86_32_jit_device.cuh => murmurhash3_x86_32_jit_dispatch_impl.cuh} (88%)
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in
 create mode 100644 cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 6178c8be370a..16aa72d87779 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -406,6 +406,23 @@ generate_jit_lto_kernels(
   KERNEL_LINK_LIBRARIES
   cudf_jit_lto_kernel_usage_requirements
 )
+generate_jit_lto_kernels(
+  cudf_jit_lto_generated_sources
+  NAME_FORMAT
+  "murmurhash_jit_@k@"
+  MATRIX_JSON_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json"
+  KERNEL_INPUT_FILE
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in"
+  FRAGMENT_TAG_FORMAT
+  "${murmur_jit_ns}::fragment_tag_murmur_dispatch"
+  FRAGMENT_TAG_HEADER_FILES
+  "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
+  OUTPUT_DIRECTORY
+  "${CMAKE_CURRENT_BINARY_DIR}/generated_jit_lto/murmurhash/dispatch"
+  KERNEL_LINK_LIBRARIES
+  cudf_jit_lto_kernel_usage_requirements
+)
 endblock()
 
 # ##################################################################################################
diff --git a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
index ecacd6de96bb..e0f2f90fc40b 100644
--- a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
+++ b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
@@ -47,4 +47,7 @@ struct fragment_tag_murmur_hasher_noop {};
 
 struct fragment_tag_murmur_entry {};
 
+/// Owns ``murmur_jit_hash_dispatcher`` (shared by entry + hasher fragments).
+struct fragment_tag_murmur_dispatch {};
+
 }  // namespace cudf::hashing::detail::jit_lto
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh
similarity index 88%
rename from cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
rename to cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh
index 8b9d061b9993..27cf0de3546e 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_device.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh
@@ -4,19 +4,18 @@
  */
 #pragma once
 
-#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
-
+#include <cudf/column/column_device_view.cuh>
 #include <cudf/detail/utilities/assert.cuh>
-#include <cudf/detail/utilities/grid_1d.cuh>
-#include <cudf/table/table_device_view.cuh>
 #include <cudf/utilities/type_dispatcher.hpp>
 
 namespace cudf::hashing::detail {
 
-__device__ inline hash_value_type murmur_jit_hash_dispatcher(column_device_view col,
-                                                             uint32_t seed,
-                                                             bool const nullable,
-                                                             size_type row_index)
+// Strong definition lives in the ``murmurhash_jit_dispatch`` fatbin only; declared in
+// ``murmurhash3_x86_32_jit_hasher_decl.cuh`` for entry / hasher / LTO device code.
+__device__ hash_value_type murmur_jit_hash_dispatcher(column_device_view col,
+                                                      uint32_t seed,
+                                                      bool const nullable,
+                                                      size_type row_index)
 {
   switch (col.type().id()) {
     case type_id::INT8:
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
index 2be689ef3c56..f05b427b089c 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_hasher_decl.cuh
@@ -11,13 +11,22 @@
 namespace cudf::hashing::detail {
 
 /**
- * @brief Forward declaration only (no `murmur_jit_hash_dispatcher` here).
+ * @brief Forward declarations for the Murmur JIT link graph.
  *
  * Per-type explicit specializations must appear in the TU before any use of that specialization.
  * Hasher / noop fragment TUs include this header, then define `template <> ...
- * murmur_jit_hasher<T>`. The entry kernel includes `murmurhash3_x86_32_jit_device.cuh`, which pulls
- * this in and adds the dispatcher that calls `murmur_jit_hasher` for every storage type.
+ * murmur_jit_hasher<T>`. The ``murmurhash_jit_dispatch`` fatbin defines
+ * ``murmur_jit_hash_dispatcher``; the entry kernel and hasher fragments only declare it here.
+ *
+ * `murmur_jit_hash_dispatcher` is declared here so `murmurhash3_x86_32_lto.cuh` can call it from
+ * nested/dictionary paths without including the dispatcher body before `murmur_jit_hasher`
+ * specializations in hasher fragment TUs.
  */
+extern __device__ hash_value_type murmur_jit_hash_dispatcher(column_device_view col,
+                                                             uint32_t seed,
+                                                             bool nullable,
+                                                             size_type row_index);
+
 template <typename T>
 extern __device__ hash_value_type
 murmur_jit_hasher(column_device_view col, uint32_t seed, bool nullable, size_type row_index);
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
index 27c40f4cd245..cc3184f62a58 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_lto.cuh
@@ -4,9 +4,10 @@
  */
 #pragma once
 
+#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/detail/iterator.cuh>
-#include <cudf/detail/row_operator/common_utils.cuh>
 #include <cudf/detail/utilities/accumulate.cuh>
 #include <cudf/detail/utilities/assert.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
@@ -18,7 +19,6 @@
 #include <cudf/structs/structs_column_device_view.cuh>
 #include <cudf/types.hpp>
 #include <cudf/utilities/traits.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
 
 #include <cuda/std/limits>
 
@@ -43,6 +43,8 @@ class element_hasher {
   {
   }
 
+  __device__ result_type seed() const noexcept { return _seed; }
+
   template <typename T>
   __device__ result_type operator()(column_device_view const& col,
                                     size_type row_index) const noexcept
@@ -91,11 +93,10 @@ class element_hasher_adaptor {
     if (_check_nulls && col.is_null(row_index)) { return NULL_HASH; }
 
     auto const keys = col.child(dictionary_column_view::keys_column_index);
-    return type_dispatcher<dispatch_storage_type>(
-      keys.type(),
-      _element_hasher,
-      keys,
-      static_cast<size_type>(col.element<dictionary32>(row_index)));
+    return murmur_jit_hash_dispatcher(keys,
+                                      _element_hasher.seed(),
+                                      _check_nulls,
+                                      static_cast<size_type>(col.element<dictionary32>(row_index)));
   }
 
   template <typename T>
@@ -127,10 +128,8 @@ class element_hasher_adaptor {
       }
     }
     for (int i = 0; i < curr_col.size(); ++i) {
-      hash =
-        cudf::hashing::detail::hash_combine(hash,
-                                            type_dispatcher<cudf::detail::dispatch_void_if_nested>(
-                                              curr_col.type(), _element_hasher, curr_col, i));
+      hash = cudf::hashing::detail::hash_combine(
+        hash, murmur_jit_hash_dispatcher(curr_col, _element_hasher.seed(), _check_nulls, i));
     }
     return hash;
   }
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in
new file mode 100644
index 000000000000..bdd3f4226ba9
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in
@@ -0,0 +1,11 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+// Owns the single device definition of ``murmur_jit_hash_dispatcher`` so the entry kernel and
+// hasher fragments can all reference it and resolve at nvJitLink time.
+
+#include <cudf/column/column_device_view.cuh>
+
+#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+#include "murmurhash3_x86_32_jit_dispatch_impl.cuh"
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json
new file mode 100644
index 000000000000..d7566537e8bb
--- /dev/null
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json
@@ -0,0 +1,7 @@
+{
+  "_singleton": [
+    {
+      "k": "dispatch"
+    }
+  ]
+}
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
index 1ff7a1c343d1..e66c042cacc7 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_entry_kernel.cu.in
@@ -4,10 +4,12 @@
  */
 
 #include <cudf/column/column_device_view.cuh>
+#include <cudf/detail/utilities/grid_1d.cuh>
 #include <cudf/hashing/detail/hashing.hpp>
-#include "murmurhash3_x86_32_jit_device.cuh"
 #include <cudf/table/table_device_view.cuh>
 
+#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+
 extern "C" __global__ void cudf_murmurhash3_x86_32_jit_link_kernel(cudf::mutable_column_device_view output,
                                                                    uint32_t seed,
                                                                    cudf::table_device_view input,
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
index af8e6c239380..5956f8010e08 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_hasher_fragment.cu.in
@@ -8,9 +8,8 @@
 
 #include <cudf/column/column_device_view.cuh>
 
-// Decl only (not `jit_device.cuh`): dispatcher would use every `murmur_jit_hasher<T>` before the
-// explicit specialization below, which nvcc rejects ("must precede its first use").
-#include "murmurhash3_x86_32_jit_hasher_decl.cuh"
+// `murmurhash3_x86_32_lto.cuh` pulls in `murmurhash3_x86_32_jit_hasher_decl.cuh` first (dispatcher
+// decl + `murmur_jit_hasher` template decl). Dispatcher body lives only in the entry fatbin.
 #include "murmurhash3_x86_32_lto.cuh"
 
 namespace cudf::hashing::detail {
diff --git a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
index ad6d490cbe68..45386ee53851 100644
--- a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
+++ b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
@@ -11,6 +11,7 @@
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/hashing/detail/hashing.hpp>
 #include <cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>
+#include <cudf/lists/lists_column_view.hpp>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/table/table_view.hpp>
 
@@ -65,19 +66,41 @@ static constexpr std::array<type_id, 30> murmur_jit_hasher_type_ids{{
   type_id::STRUCT,
 }};
 
-inline void insert_logical_type(std::unordered_set<type_id>& ids, column_view const& col)
+/**
+ * @brief Collect logical `type_id`s for fragment selection, including descendants of nested
+ * columns.
+ *
+ * Nested hashing dispatches primitives via `murmur_jit_hash_dispatcher`; strong `murmur_jit_hasher`
+ * fragments must be linked for every storage type that can appear under struct/list/dictionary
+ * columns, not only top-level columns.
+ */
+inline void collect_nested_logical_types(std::unordered_set<type_id>& ids, column_view const& col)
 {
   ids.insert(col.type().id());
-  if (col.type().id() == type_id::DICTIONARY32) {
-    dictionary_column_view const dcv(col);
-    insert_logical_type(ids, dcv.keys());
+  switch (col.type().id()) {
+    case type_id::STRUCT:
+      for (size_type i = 0; i < col.num_children(); ++i) {
+        collect_nested_logical_types(ids, col.child(i));
+      }
+      break;
+    case type_id::LIST: {
+      lists_column_view const lcv(col);
+      collect_nested_logical_types(ids, lcv.child());
+      break;
+    }
+    case type_id::DICTIONARY32: {
+      dictionary_column_view const dcv(col);
+      collect_nested_logical_types(ids, dcv.keys());
+      break;
+    }
+    default: break;
   }
 }
 
 inline void collect_table_logical_types(std::unordered_set<type_id>& ids, table_view const& table)
 {
   for (size_type c = 0; c < table.num_columns(); ++c) {
-    insert_logical_type(ids, table.column(c));
+    collect_nested_logical_types(ids, table.column(c));
   }
 }
 
@@ -270,6 +293,7 @@ inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
 
   using namespace cudf::hashing::detail::jit_lto;
   planner.add_static_fragment<fragment_tag_murmur_entry>();
+  planner.add_static_fragment<fragment_tag_murmur_dispatch>();
 
   std::unordered_set<type_id> logical_types;
   collect_table_logical_types(logical_types, input);
diff --git a/cpp/src/jit_lto/AlgorithmPlanner.cpp b/cpp/src/jit_lto/AlgorithmPlanner.cpp
index 8be6e638d79f..8d1b204485c5 100644
--- a/cpp/src/jit_lto/AlgorithmPlanner.cpp
+++ b/cpp/src/jit_lto/AlgorithmPlanner.cpp
@@ -5,12 +5,16 @@
 
 #include <cudf/detail/jit_lto/AlgorithmPlanner.hpp>
 #include <cudf/detail/jit_lto/nvjitlink_checker.hpp>
+#include <cudf/logger.hpp>
 #include <cudf/utilities/error.hpp>
 
 #include <cuda_runtime.h>
 
 #include <nvJitLink.h>
 
+#include <chrono>
+#include <cstdio>
+#include <cstdlib>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
@@ -19,6 +23,19 @@
 
 namespace cudf::detail::jit_lto {
 
+namespace {
+
+void emit_jit_lto_build_timing(double const build_ms)
+{
+  CUDF_LOG_INFO("jit_lto: AlgorithmPlanner::build %.6f ms", build_ms);
+  if (std::getenv("CUDF_JIT_LTO_LINK_TIMING") != nullptr) {
+    std::fprintf(stderr, "CUDF_JIT_LTO_LINK_TIMING build_ms=%.6f\n", build_ms);
+    std::fflush(stderr);
+  }
+}
+
+}  // namespace
+
 std::string AlgorithmPlanner::get_fragments_key() const
 {
   std::string key = "";
@@ -53,6 +70,11 @@ std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::get_launcher()
 
 std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::build()
 {
+  using clock       = std::chrono::steady_clock;
+  using duration_ms = std::chrono::duration<double, std::milli>;
+
+  auto const t_build_start = clock::now();
+
   int device = 0;
   int major  = 0;
   int minor  = 0;
@@ -92,6 +114,9 @@ std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::build()
   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);
 }
 

From 0caa15d0c72cf588e7a358d23fbb7d7707b8193a Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Fri, 8 May 2026 21:21:10 +0000
Subject: [PATCH 09/11] specialized dispatch and lto measurement

---
 cpp/CMakeLists.txt                            |  4 +--
 .../detail/murmurhash3_x86_32_jit_tags.hpp    |  7 ++--
 .../murmurhash3_x86_32_jit_dispatch_impl.cuh  | 11 +++----
 .../murmurhash_dispatch_fragment.cu.in        | 26 +++++++++++++--
 .../murmurhash_dispatch_matrix.json           |  9 +++--
 .../hash/murmurhash3_x86_32_jit_launch.hpp    | 33 ++++++++++++++-----
 6 files changed, 67 insertions(+), 23 deletions(-)

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 16aa72d87779..0e26f419f7c3 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -409,13 +409,13 @@ generate_jit_lto_kernels(
 generate_jit_lto_kernels(
   cudf_jit_lto_generated_sources
   NAME_FORMAT
-  "murmurhash_jit_@k@"
+  "murmurhash_jit_dispatch_@suffix@"
   MATRIX_JSON_FILE
   "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json"
   KERNEL_INPUT_FILE
   "${CMAKE_CURRENT_SOURCE_DIR}/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in"
   FRAGMENT_TAG_FORMAT
-  "${murmur_jit_ns}::fragment_tag_murmur_dispatch"
+  "${murmur_jit_ns}::fragment_tag_murmur_dispatch_@suffix@"
   FRAGMENT_TAG_HEADER_FILES
   "<cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp>"
   OUTPUT_DIRECTORY
diff --git a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
index e0f2f90fc40b..7def96d96bde 100644
--- a/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
+++ b/cpp/include/cudf/hashing/detail/murmurhash3_x86_32_jit_tags.hpp
@@ -47,7 +47,10 @@ struct fragment_tag_murmur_hasher_noop {};
 
 struct fragment_tag_murmur_entry {};
 
-/// Owns ``murmur_jit_hash_dispatcher`` (shared by entry + hasher fragments).
-struct fragment_tag_murmur_dispatch {};
+/// Owns ``murmur_jit_hash_dispatcher`` with the full type-id switch (matrix ``suffix`` = ``all``).
+struct fragment_tag_murmur_dispatch_all {};
+
+/// Owns ``murmur_jit_hash_dispatcher`` with an INT32-only branch (matrix ``suffix`` = ``int32``).
+struct fragment_tag_murmur_dispatch_int32 {};
 
 }  // namespace cudf::hashing::detail::jit_lto
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh
index 27cf0de3546e..385dcdf541cf 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash3_x86_32_jit_dispatch_impl.cuh
@@ -10,12 +10,11 @@
 
 namespace cudf::hashing::detail {
 
-// Strong definition lives in the ``murmurhash_jit_dispatch`` fatbin only; declared in
-// ``murmurhash3_x86_32_jit_hasher_decl.cuh`` for entry / hasher / LTO device code.
-__device__ hash_value_type murmur_jit_hash_dispatcher(column_device_view col,
-                                                      uint32_t seed,
-                                                      bool const nullable,
-                                                      size_type row_index)
+/// Full type-id switch used when the dispatch fatbin is built with ``dispatch_int32_only == 0``.
+__device__ inline hash_value_type murmur_jit_hash_dispatcher_all_types(column_device_view col,
+                                                                       uint32_t seed,
+                                                                       bool const nullable,
+                                                                       size_type row_index)
 {
   switch (col.type().id()) {
     case type_id::INT8:
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in
index bdd3f4226ba9..f640a807ccd1 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_fragment.cu.in
@@ -2,10 +2,32 @@
  * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
-// Owns the single device definition of ``murmur_jit_hash_dispatcher`` so the entry kernel and
-// hasher fragments can all reference it and resolve at nvJitLink time.
+// Owns the device definition of ``murmur_jit_hash_dispatcher``. Matrix row injects
+// ``dispatch_int32_only`` (0 = full switch, 1 = INT32-only) for ``if constexpr`` so the unused
+// branch is not codegen'd in this fatbin.
 
 #include <cudf/column/column_device_view.cuh>
 
 #include "murmurhash3_x86_32_jit_hasher_decl.cuh"
 #include "murmurhash3_x86_32_jit_dispatch_impl.cuh"
+
+namespace cudf::hashing::detail {
+
+static constexpr bool dispatch_int32_only_v = static_cast<bool>(@dispatch_int32_only@);
+
+__device__ hash_value_type murmur_jit_hash_dispatcher(column_device_view col,
+                                                      uint32_t seed,
+                                                      bool const nullable,
+                                                      size_type row_index)
+{
+  if constexpr (dispatch_int32_only_v) {
+    if (col.type().id() == type_id::INT32) {
+      return murmur_jit_hasher<id_to_type<type_id::INT32>>(col, seed, nullable, row_index);
+    }
+    CUDF_UNREACHABLE("Invalid type_id.");
+  } else {
+    return murmur_jit_hash_dispatcher_all_types(col, seed, nullable, row_index);
+  }
+}
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json
index d7566537e8bb..756f858393c1 100644
--- a/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json
+++ b/cpp/src/hash/jit_lto_kernels/murmurhash_dispatch_matrix.json
@@ -1,7 +1,12 @@
 {
-  "_singleton": [
+  "_row": [
     {
-      "k": "dispatch"
+      "suffix": "all",
+      "dispatch_int32_only": 0
+    },
+    {
+      "suffix": "int32",
+      "dispatch_int32_only": 1
     }
   ]
 }
diff --git a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
index 45386ee53851..cc7170a87a0b 100644
--- a/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
+++ b/cpp/src/hash/murmurhash3_x86_32_jit_launch.hpp
@@ -31,10 +31,11 @@ inline cudf::detail::jit_lto::LauncherJitCache& murmur_jit_launcher_cache()
   return cache;
 }
 
-// Every `type_id` handled by `murmur_jit_hash_dispatcher` must have exactly one device
-// definition of `murmur_jit_hasher<T>` in the nvJitLink input. Types present in the table use the
-// real hasher fatbin; all others use a separate strong no-op fatbin (nvJitLink does not merge
-// weak + strong overrides for the same symbol).
+// Every `type_id` handled by `murmur_jit_hash_dispatcher` must have exactly one device definition
+// of `murmur_jit_hasher<T>` in the nvJitLink input for that link. Types present in the table use
+// the real hasher fatbin; unused types use a strong no-op fatbin when linking the full dispatcher.
+// Flat INT32-only tables use the INT32-only dispatcher and link only the INT32 strong hasher (no
+// no-op stubs).
 static constexpr std::array<type_id, 30> murmur_jit_hasher_type_ids{{
   type_id::INT8,
   type_id::INT16,
@@ -268,6 +269,15 @@ inline void add_noop_hasher_fragment(cudf::detail::jit_lto::AlgorithmPlanner& pl
 
 }  // namespace
 
+/// True for a flat ``INT32``-only table: ``collect_table_logical_types`` yields exactly
+/// ``{INT32}``. Then we link ``fragment_tag_murmur_dispatch_int32`` and **omit** noop hasher
+/// fatbins. Nested or mixed types keep the full dispatcher and noop stubs for unreferenced
+/// ``murmur_jit_hasher`` symbols.
+inline bool use_murmur_jit_int32_only_dispatch(std::unordered_set<type_id> const& logical_types)
+{
+  return logical_types.size() == 1u && logical_types.count(type_id::INT32) == 1u;
+}
+
 inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
                                                       uint32_t seed,
                                                       rmm::cuda_stream_view stream,
@@ -291,16 +301,21 @@ inline std::unique_ptr<column> murmurhash3_x86_32_jit(table_view const& input,
   cudf::detail::jit_lto::AlgorithmPlanner planner{"cudf_murmurhash3_x86_32_jit_link_kernel",
                                                   murmur_jit_launcher_cache()};
 
-  using namespace cudf::hashing::detail::jit_lto;
-  planner.add_static_fragment<fragment_tag_murmur_entry>();
-  planner.add_static_fragment<fragment_tag_murmur_dispatch>();
-
   std::unordered_set<type_id> logical_types;
   collect_table_logical_types(logical_types, input);
+
+  using namespace cudf::hashing::detail::jit_lto;
+  planner.add_static_fragment<fragment_tag_murmur_entry>();
+  bool const int32_only_jit_link = use_murmur_jit_int32_only_dispatch(logical_types);
+  if (int32_only_jit_link) {
+    planner.add_static_fragment<fragment_tag_murmur_dispatch_int32>();
+  } else {
+    planner.add_static_fragment<fragment_tag_murmur_dispatch_all>();
+  }
   for (type_id const id : murmur_jit_hasher_type_ids) {
     if (logical_types.count(id) != 0u) {
       add_strong_hasher_fragment(planner, id);
-    } else {
+    } else if (not int32_only_jit_link) {
       add_noop_hasher_fragment(planner, id);
     }
   }

From c9c5de0a7f886d8090b3e260bffa3d6803c695ad Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Mon, 11 May 2026 22:35:30 +0000
Subject: [PATCH 10/11] bench options

---
 cpp/src/jit_lto/AlgorithmPlanner.cpp | 54 ++++++++++++++++++++++++++--
 1 file changed, 51 insertions(+), 3 deletions(-)

diff --git a/cpp/src/jit_lto/AlgorithmPlanner.cpp b/cpp/src/jit_lto/AlgorithmPlanner.cpp
index 8d1b204485c5..357bdab6cb19 100644
--- a/cpp/src/jit_lto/AlgorithmPlanner.cpp
+++ b/cpp/src/jit_lto/AlgorithmPlanner.cpp
@@ -13,12 +13,15 @@
 #include <nvJitLink.h>
 
 #include <chrono>
+#include <cstdint>
 #include <cstdio>
 #include <cstdlib>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
+#include <sstream>
 #include <string>
+#include <utility>
 #include <vector>
 
 namespace cudf::detail::jit_lto {
@@ -29,11 +32,50 @@ void emit_jit_lto_build_timing(double const build_ms)
 {
   CUDF_LOG_INFO("jit_lto: AlgorithmPlanner::build %.6f ms", build_ms);
   if (std::getenv("CUDF_JIT_LTO_LINK_TIMING") != nullptr) {
-    std::fprintf(stderr, "CUDF_JIT_LTO_LINK_TIMING build_ms=%.6f\n", build_ms);
+    if (auto const* meta = std::getenv("CUDF_JIT_LTO_LINK_TIMING_META");
+        meta != nullptr && meta[0] != '\0') {
+      std::fprintf(stderr, "CUDF_JIT_LTO_LINK_TIMING build_ms=%.6f %s\n", build_ms, meta);
+    } else {
+      std::fprintf(stderr, "CUDF_JIT_LTO_LINK_TIMING build_ms=%.6f\n", build_ms);
+    }
     std::fflush(stderr);
   }
 }
 
+std::vector<std::string> make_nvjitlink_option_strings(std::string arch_sm)
+{
+  std::vector<std::string> opts;
+  opts.reserve(8);
+  opts.emplace_back("-lto");
+  opts.push_back(std::move(arch_sm));
+
+  // Whitespace-separated extra nvJitLink flags (e.g. "-O3", "-split-compile=0"). When non-empty,
+  // this replaces compile-time CUDF_JIT_LTO_NVJITLINK_OPTSET extras so one build can sweep flags.
+  // Note: "-O0" with "-lto" has triggered cudaErrorLaunchFailure on some toolchains; axis
+  // benchmarks skip that pair (see benchmark_murmur_jit_lto_axes.sh). JIT fragments are LTO-IR;
+  // -lto is required.
+  if (auto const* env = std::getenv("CUDF_JIT_LTO_NVJITLINK_OPTIONS");
+      env != nullptr && env[0] != '\0') {
+    std::istringstream iss(env);
+    std::string tok;
+    while (iss >> tok) {
+      opts.push_back(std::move(tok));
+    }
+    return opts;
+  }
+
+#if defined(CUDF_JIT_LTO_NVJITLINK_PROFILE_O0)
+  opts.emplace_back("-O0");
+#elif defined(CUDF_JIT_LTO_NVJITLINK_PROFILE_O3)
+  opts.emplace_back("-O3");
+#elif defined(CUDF_JIT_LTO_NVJITLINK_PROFILE_SPLIT_COMPILE)
+  opts.emplace_back("-split-compile=0");
+#elif defined(CUDF_JIT_LTO_NVJITLINK_PROFILE_SPLIT_COMPILE_EXTENDED)
+  opts.emplace_back("-split-compile-extended=0");
+#endif
+  return opts;
+}
+
 }  // namespace
 
 std::string AlgorithmPlanner::get_fragments_key() const
@@ -84,9 +126,15 @@ std::shared_ptr<AlgorithmLauncher> AlgorithmPlanner::build()
 
   std::string archs = "-arch=sm_" + std::to_string((major * 10 + minor));
 
+  auto const opt_strings = make_nvjitlink_option_strings(std::move(archs));
+  std::vector<char const*> opt_ptrs;
+  opt_ptrs.reserve(opt_strings.size());
+  for (auto const& s : opt_strings) {
+    opt_ptrs.push_back(s.c_str());
+  }
+
   nvJitLinkHandle handle;
-  const char* lopts[] = {"-lto", archs.c_str()};
-  auto result         = nvJitLinkCreate(&handle, 2, lopts);
+  auto result = nvJitLinkCreate(&handle, static_cast<uint32_t>(opt_ptrs.size()), opt_ptrs.data());
   check_nvjitlink_result(handle, result);
 
   for (const auto& frag : this->fragments) {

From fa9181426c3695b2ddf54a6a299239b414f91385 Mon Sep 17 00:00:00 2001
From: Divye Gala <divyegala@gmail.com>
Date: Mon, 11 May 2026 22:35:48 +0000
Subject: [PATCH 11/11] bench scripts

---
 .../benchmark_murmur_jit_lto_axes.sh          | 123 ++++
 .../benchmark_murmur_jit_lto_link.py          | 538 ++++++++++++++++++
 2 files changed, 661 insertions(+)
 create mode 100755 python/pylibcudf/benchmark_murmur_jit_lto_axes.sh
 create mode 100644 python/pylibcudf/benchmark_murmur_jit_lto_link.py

diff --git a/python/pylibcudf/benchmark_murmur_jit_lto_axes.sh b/python/pylibcudf/benchmark_murmur_jit_lto_axes.sh
new file mode 100755
index 000000000000..03d4e9289e60
--- /dev/null
+++ b/python/pylibcudf/benchmark_murmur_jit_lto_axes.sh
@@ -0,0 +1,123 @@
+#!/usr/bin/env bash
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+# SPDX-License-Identifier: Apache-2.0
+#
+# Murmur JIT+LTO link benchmark: sweep nvJitLink *extra* flags via CUDF_JIT_LTO_NVJITLINK_OPTIONS.
+# libcudf always passes -lto; embedded fragments are LTO-IR, so omitting -lto is not supported.
+#
+# Skipped preset:
+#   - O0  (nvJitLink ``-lto`` + ``-O0`` → ``cudaErrorLaunchFailure`` on observed stacks)
+#
+# Each profile clears ~/.nv/ComputeCache once, then runs two Python processes (cold/warm disk);
+# each run does one in-process table sweep. All timings append to one CSV.
+#
+# Usage:
+#   conda activate cudf_2606   # or set CONDA_ENV
+#   ./python/pylibcudf/benchmark_murmur_jit_lto_axes.sh [PROFILE ...]
+#
+# No arguments: lto, split-compile, split-compile-extended. (-O3 is nvJitLink default with -lto;
+# pass ``O3`` explicitly if needed.) ``O0`` is not run (see above).
+#
+# Environment:
+#   CONDA_ENV   conda env name (default: cudf_2606)
+#   REPO_ROOT   cuDF git root (default: inferred from this script)
+#   OUTDIR      output directory (default: $REPO_ROOT/jit_lto_bench_out)
+#   CSV_OUT      raw per-link CSV (default: $OUTDIR/murmur_jit_lto_bench.csv); removed at start
+#   SUMMARY_CSV  pivot summary for analysis (default: $OUTDIR/murmur_jit_lto_bench_summary.csv)
+
+set -euo pipefail
+
+REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)"
+CONDA_ENV="${CONDA_ENV:-cudf_2606}"
+OUTDIR="${OUTDIR:-${REPO_ROOT}/jit_lto_bench_out}"
+CSV_OUT="${CSV_OUT:-${OUTDIR}/murmur_jit_lto_bench.csv}"
+SUMMARY_CSV="${SUMMARY_CSV:-${OUTDIR}/murmur_jit_lto_bench_summary.csv}"
+BENCH_PY="${REPO_ROOT}/python/pylibcudf/benchmark_murmur_jit_lto_link.py"
+
+mkdir -p "${OUTDIR}"
+rm -f "${CSV_OUT}"
+
+declare -A NVJITLINK_PROFILE_OPTS=(
+  [lto]=""
+  [O0]="-O0"
+  [O3]="-O3"
+  [split-compile]="-split-compile=0"
+  [split-compile-extended]="-split-compile-extended=0"
+)
+
+DEFAULT_PROFILES=(lto split-compile split-compile-extended)
+
+filter_profiles() {
+  local -a in=( "$@" )
+  RUN_PROFILES=()
+  local p
+  for p in "${in[@]}"; do
+    if [[ "${p}" == O0 ]]; then
+      echo "Skipping profile O0: nvJitLink -lto -O0 fails at kernel launch on observed stacks." >&2
+      continue
+    fi
+    RUN_PROFILES+=( "${p}" )
+  done
+}
+
+if [[ ${#} -gt 0 ]]; then
+  filter_profiles "$@"
+else
+  RUN_PROFILES=( "${DEFAULT_PROFILES[@]}" )
+fi
+
+if [[ ${#RUN_PROFILES[@]} -eq 0 ]]; then
+  echo "No profiles left to run after filtering." >&2
+  exit 1
+fi
+
+for p in "${RUN_PROFILES[@]}"; do
+  if [[ ! -v NVJITLINK_PROFILE_OPTS[$p] ]]; then
+    echo "Unknown profile '${p}'. Valid: ${!NVJITLINK_PROFILE_OPTS[*]}" >&2
+    exit 1
+  fi
+done
+
+unset CUDF_JIT_LTO_DISABLE_LTO
+
+# Conda's cuda-nvcc activate hook uses NVCC_PREPEND_FLAGS without a default; ``set -u`` makes that
+# an error. Allow unset variables only while initializing conda.
+set +u
+# shellcheck source=/dev/null
+source "$(conda info --base)/etc/profile.d/conda.sh"
+conda activate "${CONDA_ENV}"
+set -u
+
+echo "REPO_ROOT=${REPO_ROOT}"
+echo "CONDA_DEFAULT_ENV=${CONDA_DEFAULT_ENV}"
+echo "OUTDIR=${OUTDIR}"
+echo "CSV_OUT=${CSV_OUT}"
+echo "SUMMARY_CSV=${SUMMARY_CSV}"
+echo "PROFILES=${RUN_PROFILES[*]}"
+echo ""
+
+for tag in "${RUN_PROFILES[@]}"; do
+  extra="${NVJITLINK_PROFILE_OPTS[$tag]}"
+  if [[ -n "${extra}" ]]; then
+    export CUDF_JIT_LTO_NVJITLINK_OPTIONS="${extra}"
+  else
+    unset CUDF_JIT_LTO_NVJITLINK_OPTIONS
+  fi
+
+  echo "======== Profile ${tag} (CUDF_JIT_LTO_NVJITLINK_OPTIONS=${extra:-<unset>}, always -lto) ========"
+  echo "======== Clear CUDA compute cache ========"
+  rm -rf "${HOME}/.nv/ComputeCache"
+
+  for inv in 1 2; do
+    echo "-------- Python: profile=${tag} script_invocation=${inv} (append ${CSV_OUT}) --------"
+    python "${BENCH_PY}" \
+      --nvjitlink-optset "${tag}" \
+      --script-invocation "${inv}" \
+      --output-csv "${CSV_OUT}"
+  done
+  echo ""
+done
+
+python "${BENCH_PY}" --summarize-bench-csv "${CSV_OUT}" --summary-output "${SUMMARY_CSV}"
+
+echo "Done. Raw: ${CSV_OUT}  Summary: ${SUMMARY_CSV}"
diff --git a/python/pylibcudf/benchmark_murmur_jit_lto_link.py b/python/pylibcudf/benchmark_murmur_jit_lto_link.py
new file mode 100644
index 000000000000..5528d7e22386
--- /dev/null
+++ b/python/pylibcudf/benchmark_murmur_jit_lto_link.py
@@ -0,0 +1,538 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+# SPDX-License-Identifier: Apache-2.0
+"""
+Benchmark ``AlgorithmPlanner::build()`` (MurmurHash3 x86_32 JIT+LTO / nvJitLink).
+
+**Axes (how this script fits in)**
+
+- **CUDA / nvJitLink disk cache** (~``~/.nv/ComputeCache``): clear before the *first* process
+  invocation; the *second* process invocation (separate ``python`` run) sees a warm disk cache.
+- **In-process sweep**: one full A→E table pass per Python run (``--process-passes`` defaults to
+  **1**; increase only if you explicitly want repeated sweeps in the same process).
+- **nvJitLink link flags**: **``CUDF_JIT_LTO_NVJITLINK_OPTIONS``** — whitespace-separated extras after
+  ``-lto`` and ``-arch=sm_XX``. Murmur JIT fragments are LTO-IR; **``-lto`` is always required** (the
+  axis driver does not attempt no-LTO links). The axis script does **not** sweep **``-O0``** with
+  ``-lto`` (observed ``cudaErrorLaunchFailure``). No redundant default **``-O3``** row (O3 is the
+  default optimization level with LTO).
+
+Set ``CUDF_JIT_LTO_LINK_TIMING=1`` (set automatically here) so ``build()`` emits one stderr line per
+link. Optional ``CUDF_JIT_LTO_LINK_TIMING_META`` is set per pass for parsing (``script_inv``,
+``process_pass``).
+
+**Suggested sweep** (one libcudf build, then)::
+
+    ./python/pylibcudf/benchmark_murmur_jit_lto_axes.sh
+
+Use ``--output-csv PATH`` to append one row per ``build_ms`` event (header written if the file is
+new or empty). The axis shell driver writes a single CSV for the full sweep, then a **summary CSV**
+(``--summarize-bench-csv``) with cold/warm columns and percent deltas vs ``lto`` baseline per
+scenario and totals.
+
+Table order: **A** ``float32`` (full dispatcher first), **B** ``int32``, then struct, list, int64.
+"""
+
+from __future__ import annotations
+
+import argparse
+import csv
+import io
+import json
+import os
+import re
+import sys
+import threading
+from pathlib import Path
+from typing import Any
+
+# (letter, description) aligned with table order in ``_benchmark_tables``.
+_BENCHMARK_TYPE_LABELS: list[tuple[str, str]] = [
+    ("A", "float32 column (first build: full dispatcher)"),
+    ("B", "int32 column"),
+    ("C", "struct<int32, int32> column (nested)"),
+    ("D", "list<int32> column (nested)"),
+    ("E", "int64 column"),
+]
+
+
+def _drop_script_dir_from_sys_path() -> None:
+    """If this file lives under ``python/pylibcudf/``, running it prepends that directory to
+    ``sys.path`` so ``import pylibcudf`` loads the **source** tree (no built Cython extensions)
+    instead of the conda/site-packages install. Drop that entry so the installed package wins.
+    """
+    script_dir = os.path.dirname(os.path.abspath(__file__))
+    if sys.path and os.path.realpath(sys.path[0]) == os.path.realpath(
+        script_dir
+    ):
+        del sys.path[0]
+
+
+_drop_script_dir_from_sys_path()
+
+# Optional trailing meta: space-separated key=value tokens (from CUDF_JIT_LTO_LINK_TIMING_META).
+_TIMING_LINE = re.compile(
+    r"^CUDF_JIT_LTO_LINK_TIMING build_ms=(?P<build>[\d.]+)(?:\s+(?P<meta>.+))?\s*$"
+)
+
+
+def _parse_timing_meta(meta: str | None) -> dict[str, str]:
+    if not meta:
+        return {}
+    out: dict[str, str] = {}
+    for part in meta.split():
+        if "=" in part:
+            k, _, v = part.partition("=")
+            out[k] = v
+    return out
+
+
+def _parse_build_timing_lines(captured_stderr: str) -> list[dict[str, Any]]:
+    events: list[dict[str, Any]] = []
+    for line in captured_stderr.splitlines():
+        m = _TIMING_LINE.match(line.strip())
+        if m:
+            ev: dict[str, Any] = {"build_ms": float(m.group("build"))}
+            meta = _parse_timing_meta(m.group("meta"))
+            if meta:
+                ev["meta"] = meta
+            events.append(ev)
+    return events
+
+
+def _attach_scenario_labels(
+    events: list[dict[str, Any]],
+) -> list[dict[str, Any]]:
+    """First ``build`` line in a pass is scenario A, second is B, ..."""
+    out = []
+    for i, ev in enumerate(events):
+        letter, desc = _BENCHMARK_TYPE_LABELS[i] if i < len(_BENCHMARK_TYPE_LABELS) else (
+            "?",
+            "unknown",
+        )
+        row = {**ev, "type_letter": letter, "type_desc": desc}
+        out.append(row)
+    return out
+
+
+def _capture_cpp_stderr(work):
+    """Redirect FD 2 so libcudf ``fprintf(stderr, ...)`` is captured."""
+    read_fd, write_fd = os.pipe()
+    saved = os.dup(2)
+    os.dup2(write_fd, 2)
+    os.close(write_fd)
+    buf = io.StringIO()
+
+    def reader():
+        with os.fdopen(read_fd, "r") as pipe:
+            buf.write(pipe.read())
+
+    th = threading.Thread(target=reader)
+    th.start()
+    try:
+        work()
+    finally:
+        os.dup2(saved, 2)
+        os.close(saved)
+    th.join()
+    return buf.getvalue()
+
+
+def _benchmark_tables():
+    """Return ``pylibcudf.Table`` for types A–E (float32 before int32 for full-dispatch first link)."""
+    import pyarrow as pa
+
+    import pylibcudf as plc
+
+    pa_float_first = pa.table(
+        {"x": pa.array([1.0, 2.0, 3.0, 4.0, 5.0], type=pa.float32())}
+    )
+    pa_int32 = pa.table({"x": pa.array([1, 2, 3, 4, 5], type=pa.int32())})
+    pa_c = pa.table(
+        {
+            "s": pa.array(
+                [{"a": 1, "b": 2}, {"a": 3, "b": 4}, {"a": 5, "b": 6}],
+                type=pa.struct([("a", pa.int32()), ("b", pa.int32())]),
+            )
+        }
+    )
+    pa_d = pa.table(
+        {
+            "lst": pa.array(
+                [[1, 2], [3, 4, 5], []],
+                type=pa.list_(pa.int32()),
+            )
+        }
+    )
+    pa_e = pa.table({"x": pa.array([1, 2, 3, 4, 5], type=pa.int64())})
+
+    return [
+        plc.Table.from_arrow(pa_float_first),
+        plc.Table.from_arrow(pa_int32),
+        plc.Table.from_arrow(pa_c),
+        plc.Table.from_arrow(pa_d),
+        plc.Table.from_arrow(pa_e),
+    ]
+
+
+def _cache_axis_interpretation(
+    *, script_invocation: int, process_pass: int
+) -> dict[str, str]:
+    """Human-readable labels for the benchmark grid (not exhaustive of every subsystem)."""
+    disk = "warm" if script_invocation >= 2 else "cold"
+    mem = "warm" if process_pass >= 2 else "cold"
+    return {
+        "cuda_compute_cache_disk": disk,
+        "libcudf_jit_lto_in_process": mem,
+    }
+
+
+def run_benchmark(
+    *,
+    script_invocation: int = 1,
+    process_passes: int = 1,
+    nvjitlink_optset: str = "lto",
+) -> dict[str, Any]:
+    """Run ``process_passes`` full table sweeps; capture timing lines per pass."""
+    os.environ["CUDF_JIT_LTO_LINK_TIMING"] = "1"
+
+    import pylibcudf as plc
+
+    seed = plc.hashing.LIBCUDF_DEFAULT_HASH_SEED
+    tables = _benchmark_tables()
+
+    pass_results: list[dict[str, Any]] = []
+
+    for process_pass in range(1, process_passes + 1):
+        os.environ["CUDF_JIT_LTO_LINK_TIMING_META"] = (
+            f"script_inv={script_invocation} process_pass={process_pass} "
+            f"nvjitlink_optset={nvjitlink_optset}"
+        )
+
+        def work():
+            for tbl in tables:
+                plc.hashing.murmurhash3_x86_32(tbl, seed)
+
+        captured = _capture_cpp_stderr(work)
+        raw_events = _parse_build_timing_lines(captured)
+        labeled = _attach_scenario_labels(raw_events)
+        pass_results.append(
+            {
+                "process_pass": process_pass,
+                "axes_interpretation": _cache_axis_interpretation(
+                    script_invocation=script_invocation,
+                    process_pass=process_pass,
+                ),
+                "link_events": labeled,
+                "expected_scenarios": len(_BENCHMARK_TYPE_LABELS),
+            }
+        )
+
+    return {
+        "nvjitlink_optset": nvjitlink_optset,
+        "script_invocation": script_invocation,
+        "process_passes": process_passes,
+        "passes": pass_results,
+    }
+
+
+_CSV_COLUMNS = (
+    "nvjitlink_optset",
+    "script_invocation",
+    "process_pass",
+    "cuda_compute_cache_disk",
+    "libcudf_jit_lto_in_process",
+    "scenario",
+    "type_desc",
+    "build_ms",
+)
+
+
+_SUMMARY_CSV_COLUMNS = (
+    "nvjitlink_optset",
+    "scenario",
+    "type_desc",
+    "cold_ms",
+    "warm_ms",
+    "warm_pct_of_cold",
+    "cold_pct_vs_lto_baseline",
+    "warm_pct_vs_lto_baseline",
+)
+
+
+def summarize_bench_csv(raw_path: Path, summary_path: Path) -> None:
+    """Read raw benchmark CSV (per-link rows); write a pivot-style summary for analysis.
+
+    Expects ``script_invocation`` 1 = cold disk, 2 = warm disk (as produced by the axis shell).
+    Percent columns: ``warm_pct_of_cold`` = 100×warm/cold; ``*_pct_vs_lto_baseline`` = 100×(x/lto−1)
+    for the same scenario and invocation class. Baseline rows (``lto``) use 0 for those deltas.
+    Appends a **TOTAL** row per ``nvjitlink_optset`` (sum of A–E present in the raw file).
+    """
+    with raw_path.open(newline="", encoding="utf-8") as f:
+        rows = list(csv.DictReader(f))
+    if not rows:
+        raise SystemExit(f"empty or missing benchmark CSV: {raw_path}")
+
+    # ms[optset][scenario][inv] = float
+    ms: dict[str, dict[str, dict[int, float]]] = {}
+    type_desc: dict[tuple[str, str], str] = {}
+    opt_order: list[str] = []
+
+    for r in rows:
+        opt = r["nvjitlink_optset"]
+        scen = r["scenario"]
+        inv = int(r["script_invocation"])
+        val = float(r["build_ms"])
+        if opt not in ms:
+            ms[opt] = {}
+            opt_order.append(opt)
+        ms[opt].setdefault(scen, {})[inv] = val
+        type_desc[(opt, scen)] = r.get("type_desc", "")
+
+    known_scen = ["A", "B", "C", "D", "E"]
+    all_scen: set[str] = set()
+    for d in ms.values():
+        all_scen |= set(d.keys())
+    scen_order = [s for s in known_scen if s in all_scen] + sorted(
+        all_scen - set(known_scen)
+    )
+
+    def lto_ms(scen: str, inv: int) -> float | None:
+        if "lto" not in ms or scen not in ms["lto"]:
+            return None
+        return ms["lto"][scen].get(inv)
+
+    def pct_vs_baseline(
+        opt: str, scen: str, inv: int, value: float | None
+    ) -> str:
+        if value is None:
+            return ""
+        if opt == "lto":
+            return "0"
+        base = lto_ms(scen, inv)
+        if base is None or base <= 0:
+            return ""
+        return f"{100.0 * (value / base - 1.0):.2f}"
+
+    out_rows: list[dict[str, str]] = []
+
+    for opt in opt_order:
+        cold_sum = 0.0
+        warm_sum = 0.0
+        for scen in scen_order:
+            if scen not in ms[opt]:
+                continue
+            d = ms[opt][scen]
+            cold = d.get(1)
+            warm = d.get(2)
+            if cold is not None:
+                cold_sum += cold
+            if warm is not None:
+                warm_sum += warm
+
+            warm_pct_of_cold = ""
+            if cold is not None and warm is not None and cold > 0:
+                warm_pct_of_cold = f"{100.0 * warm / cold:.2f}"
+
+            out_rows.append(
+                {
+                    "nvjitlink_optset": opt,
+                    "scenario": scen,
+                    "type_desc": type_desc.get((opt, scen), ""),
+                    "cold_ms": f"{cold:.6f}" if cold is not None else "",
+                    "warm_ms": f"{warm:.6f}" if warm is not None else "",
+                    "warm_pct_of_cold": warm_pct_of_cold,
+                    "cold_pct_vs_lto_baseline": pct_vs_baseline(opt, scen, 1, cold),
+                    "warm_pct_vs_lto_baseline": pct_vs_baseline(opt, scen, 2, warm),
+                }
+            )
+
+        warm_pct_tot = ""
+        if cold_sum > 0 and warm_sum > 0:
+            warm_pct_tot = f"{100.0 * warm_sum / cold_sum:.2f}"
+        lto_cold_tot = sum(
+            ms["lto"][s][1]
+            for s in scen_order
+            if "lto" in ms and s in ms["lto"] and 1 in ms["lto"][s]
+        )
+        lto_warm_tot = sum(
+            ms["lto"][s][2]
+            for s in scen_order
+            if "lto" in ms and s in ms["lto"] and 2 in ms["lto"][s]
+        )
+        cold_pct_tot = ""
+        warm_pct_lto_tot = ""
+        if opt != "lto" and lto_cold_tot > 0 and cold_sum > 0:
+            cold_pct_tot = f"{100.0 * (cold_sum / lto_cold_tot - 1.0):.2f}"
+        if opt == "lto":
+            cold_pct_tot = "0"
+            warm_pct_lto_tot = "0"
+        elif lto_warm_tot > 0 and warm_sum > 0:
+            warm_pct_lto_tot = f"{100.0 * (warm_sum / lto_warm_tot - 1.0):.2f}"
+
+        out_rows.append(
+            {
+                "nvjitlink_optset": opt,
+                "scenario": "TOTAL",
+                "type_desc": "",
+                "cold_ms": f"{cold_sum:.6f}" if cold_sum > 0 else "",
+                "warm_ms": f"{warm_sum:.6f}" if warm_sum > 0 else "",
+                "warm_pct_of_cold": warm_pct_tot,
+                "cold_pct_vs_lto_baseline": cold_pct_tot,
+                "warm_pct_vs_lto_baseline": warm_pct_lto_tot,
+            }
+        )
+
+    summary_path.parent.mkdir(parents=True, exist_ok=True)
+    with summary_path.open("w", newline="", encoding="utf-8") as f:
+        w = csv.DictWriter(f, fieldnames=_SUMMARY_CSV_COLUMNS)
+        w.writeheader()
+        w.writerows(out_rows)
+
+
+def append_report_to_csv(path: Path, report: dict[str, Any]) -> None:
+    """Append one CSV row per link timing event; write header if *path* is missing or empty."""
+    path.parent.mkdir(parents=True, exist_ok=True)
+    new_file = (not path.exists()) or path.stat().st_size == 0
+    with path.open("a", newline="", encoding="utf-8") as f:
+        w = csv.DictWriter(f, fieldnames=_CSV_COLUMNS)
+        if new_file:
+            w.writeheader()
+        for p in report["passes"]:
+            ax = p["axes_interpretation"]
+            for ev in p["link_events"]:
+                w.writerow(
+                    {
+                        "nvjitlink_optset": report["nvjitlink_optset"],
+                        "script_invocation": report["script_invocation"],
+                        "process_pass": p["process_pass"],
+                        "cuda_compute_cache_disk": ax["cuda_compute_cache_disk"],
+                        "libcudf_jit_lto_in_process": ax["libcudf_jit_lto_in_process"],
+                        "scenario": ev["type_letter"],
+                        "type_desc": ev["type_desc"],
+                        "build_ms": ev["build_ms"],
+                    }
+                )
+
+
+def _print_human_report(report: dict[str, Any]) -> None:
+    opt = report["nvjitlink_optset"]
+    inv = report["script_invocation"]
+    print(
+        "AlgorithmPlanner::build() (Murmur JIT+LTO), CUDF_JIT_LTO_LINK_TIMING=1\n"
+        f"nvjitlink_optset={opt} script_invocation={inv}\n"
+    )
+    for p in report["passes"]:
+        pn = p["process_pass"]
+        ax = p["axes_interpretation"]
+        evs = p["link_events"]
+        nexp = p["expected_scenarios"]
+        print(
+            f"--- process_pass={pn} "
+            f"(cuda_compute_cache_disk~{ax['cuda_compute_cache_disk']}, "
+            f"libcudf_jit_lto_in_process~{ax['libcudf_jit_lto_in_process']}) ---"
+        )
+        if len(evs) != nexp:
+            print(
+                f"  Expected {nexp} timing lines, got {len(evs)} "
+                "(fragment collapse, wrong build, or cache hit)."
+            )
+        if not evs:
+            print("  (no nvJitLink build lines — expected for pass 2 if in-memory cache hit)\n")
+            continue
+        for row in evs:
+            letter = row["type_letter"]
+            desc = row["type_desc"]
+            ms = row["build_ms"]
+            print(f"  Type {letter} ({desc}): {ms:.6f} ms")
+        print("")
+
+
+def main(argv: list[str] | None = None) -> int:
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument(
+        "--summarize-bench-csv",
+        type=Path,
+        metavar="RAW.csv",
+        help="Read raw per-link benchmark CSV and exit after writing summary (see --summary-output)",
+    )
+    p.add_argument(
+        "--summary-output",
+        type=Path,
+        metavar="PATH",
+        help="Summary CSV path (default: RAW stem + _summary.csv next to RAW)",
+    )
+    p.add_argument(
+        "--script-invocation",
+        type=int,
+        default=1,
+        metavar="N",
+        help="1 = first python process after clearing ~/.nv/ComputeCache; 2 = second process (warm disk cache). Default: 1",
+    )
+    p.add_argument(
+        "--process-passes",
+        type=int,
+        default=1,
+        metavar="N",
+        help="In-process repetitions of the full table sweep (default: 1).",
+    )
+    p.add_argument(
+        "--nvjitlink-optset",
+        default="lto",
+        metavar="NAME",
+        help="Label for this run (e.g. matches CUDF_JIT_LTO_NVJITLINK_OPTIONS preset in the shell driver). Default: lto",
+    )
+    p.add_argument(
+        "--output-json",
+        type=Path,
+        metavar="PATH",
+        help="Write full report as JSON to PATH",
+    )
+    p.add_argument(
+        "--output-csv",
+        type=Path,
+        metavar="PATH",
+        help="Append link timing rows to PATH (creates file and header if missing or empty)",
+    )
+    p.add_argument(
+        "-q",
+        "--quiet",
+        action="store_true",
+        help="Suppress human-readable report (use with --output-csv or --output-json)",
+    )
+    args = p.parse_args(argv)
+
+    if args.summarize_bench_csv is not None:
+        raw = args.summarize_bench_csv
+        if not raw.is_file():
+            p.error(f"not a file: {raw}")
+        out = args.summary_output
+        if out is None:
+            out = raw.with_name(f"{raw.stem}_summary.csv")
+        summarize_bench_csv(raw, out)
+        print(f"Wrote summary CSV: {out}")
+        return 0
+
+    if args.process_passes < 1:
+        p.error("--process-passes must be >= 1")
+
+    report = run_benchmark(
+        script_invocation=args.script_invocation,
+        process_passes=args.process_passes,
+        nvjitlink_optset=args.nvjitlink_optset,
+    )
+
+    if args.output_json:
+        args.output_json.write_text(
+            json.dumps(report, indent=2) + "\n", encoding="utf-8"
+        )
+
+    if args.output_csv:
+        append_report_to_csv(args.output_csv, report)
+
+    if not args.quiet:
+        _print_human_report(report)
+
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())