diff --git a/cccl_upstream/cub/benchmarks/bench/bitonic_sort/bitonic_common.cuh b/cccl_upstream/cub/benchmarks/bench/bitonic_sort/bitonic_common.cuh
new file mode 100644
index 00000000..364d4d67
--- /dev/null
+++ b/cccl_upstream/cub/benchmarks/bench/bitonic_sort/bitonic_common.cuh
@@ -0,0 +1,90 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#pragma once
+
+#include <cub/util_arch.cuh>
+
+#include <cuda/std/limits>
+
+#include <string>
+
+#include <device_side_benchmark.cuh>
+#include <nvbench_helper.cuh>
+
+inline constexpr int warp_threads                       = cub::detail::warp_threads;
+inline constexpr int num_iterations_for_latency_mode    = 100;
+inline constexpr int num_iterations_for_throughput_mode = 10;
+inline constexpr int block_dim_for_throughput_mode      = 128;
+inline constexpr int grid_threads_for_throughput_mode   = 1 << 27;
+
+enum class Mode
+{
+  // launch a single warp
+  Latency,
+  // launch grid_threads_for_throughput_mode threads. Measure Elem/s.
+  Throughput
+};
+
+NVBENCH_DECLARE_ENUM_TYPE_STRINGS(
+  Mode,
+  // Callable to generate input strings:
+  [](Mode value) {
+    switch (value)
+    {
+      case Mode::Latency:
+        return "latency";
+      case Mode::Throughput:
+        return "throughput";
+      default:
+        return "Unknown";
+    }
+  },
+  // Callable to generate descriptions:
+  [](auto) {
+    return std::string{};
+  })
+
+template <typename ActionT, Mode mode, typename KeyT, typename ValueT, int Len>
+void run_bench(nvbench::state& state)
+{
+  constexpr int items_per_thread = Len / warp_threads;
+  const auto kernel              = benchmark_kernel<items_per_thread, KeyT, ValueT, ActionT, int>;
+
+  int block_dim;
+  int grid_dim;
+  int num_iterations;
+  if (mode == Mode::Latency)
+  {
+    block_dim      = warp_threads;
+    grid_dim       = 1;
+    num_iterations = num_iterations_for_latency_mode;
+  }
+  else
+  {
+    block_dim      = block_dim_for_throughput_mode;
+    grid_dim       = grid_threads_for_throughput_mode / block_dim;
+    num_iterations = num_iterations_for_throughput_mode;
+  }
+
+  state.add_element_count(static_cast<size_t>(grid_dim) * (block_dim / warp_threads) * Len * num_iterations);
+
+  state.exec([grid_dim, block_dim, kernel, num_iterations](nvbench::launch& launch) {
+    kernel<<<grid_dim, block_dim, 0, launch.get_stream()>>>(num_iterations, ActionT{}, Len);
+  });
+}
+
+struct CustomLess
+{
+  template <typename T>
+  __device__ bool operator()(const T& lhs, const T& rhs) const
+  {
+    return lhs < rhs;
+  }
+
+  template <typename T>
+  static constexpr T oob_default =
+    cuda::std::numeric_limits<T>::has_infinity
+      ? cuda::std::numeric_limits<T>::infinity()
+      : cuda::std::numeric_limits<T>::max();
+};
diff --git a/cccl_upstream/cub/benchmarks/bench/bitonic_sort/warp_keys.cu b/cccl_upstream/cub/benchmarks/bench/bitonic_sort/warp_keys.cu
new file mode 100644
index 00000000..8f8445ff
--- /dev/null
+++ b/cccl_upstream/cub/benchmarks/bench/bitonic_sort/warp_keys.cu
@@ -0,0 +1,80 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <cub/warp/warp_bitonic_sort.cuh>
+
+#include <nvbench_helper.cuh>
+
+#include "bitonic_common.cuh"
+
+using modes      = nvbench::enum_type_list<Mode::Latency, Mode::Throughput>;
+using key_types  = nvbench::type_list<int16_t, float>;
+using len_values = nvbench::enum_type_list<32, 64, 96, 128, 160, 192, 224, 256>;
+
+template <int ItemsPerThread>
+struct full_op_t
+{
+  template <typename KeyT>
+  _CCCL_DEVICE _CCCL_FORCEINLINE void operator()(KeyT (&keys)[ItemsPerThread], int) const
+  {
+    using WarpBitonicSort = cub::detail::WarpBitonicSort<KeyT, ItemsPerThread>;
+    using TempStorage     = typename WarpBitonicSort::TempStorage;
+    __shared__ TempStorage temp_storage[32];
+    const auto warp_id = threadIdx.x / 32;
+    WarpBitonicSort{temp_storage[warp_id]}.Sort(keys, CustomLess{});
+  }
+};
+
+template <Mode mode, typename KeyT, int Len>
+void full(nvbench::state& state, nvbench::type_list<nvbench::enum_type<mode>, KeyT, nvbench::enum_type<Len>>)
+{
+  run_bench<full_op_t<Len / warp_threads>, mode, KeyT, void, Len>(state);
+}
+
+NVBENCH_BENCH_TYPES(full, NVBENCH_TYPE_AXES(modes, key_types, len_values)).set_type_axes_names({"mode", "KeyT", "len"});
+
+template <int ItemsPerThread>
+struct partial_oob_op_t
+{
+  template <typename KeyT>
+  _CCCL_DEVICE _CCCL_FORCEINLINE void operator()(KeyT (&keys)[ItemsPerThread], int len) const
+  {
+    using WarpBitonicSort = cub::detail::WarpBitonicSort<KeyT, ItemsPerThread>;
+    using TempStorage     = typename WarpBitonicSort::TempStorage;
+    __shared__ TempStorage temp_storage[32];
+    const auto warp_id = threadIdx.x / 32;
+    WarpBitonicSort{temp_storage[warp_id]}.Sort(keys, CustomLess{}, len, CustomLess::oob_default<KeyT>);
+  }
+};
+
+template <Mode mode, typename KeyT, int Len>
+void partial_oob(nvbench::state& state, nvbench::type_list<nvbench::enum_type<mode>, KeyT, nvbench::enum_type<Len>>)
+{
+  run_bench<partial_oob_op_t<Len / warp_threads>, mode, KeyT, void, Len>(state);
+}
+
+NVBENCH_BENCH_TYPES(partial_oob, NVBENCH_TYPE_AXES(modes, key_types, len_values))
+  .set_type_axes_names({"mode", "KeyT", "len"});
+
+template <int ItemsPerThread>
+struct partial_op_t
+{
+  template <typename KeyT>
+  _CCCL_DEVICE _CCCL_FORCEINLINE void operator()(KeyT (&keys)[ItemsPerThread], int len) const
+  {
+    using WarpBitonicSort = cub::detail::WarpBitonicSort<KeyT, ItemsPerThread>;
+    using TempStorage     = typename WarpBitonicSort::TempStorage;
+    __shared__ TempStorage temp_storage[32];
+    const auto warp_id = threadIdx.x / 32;
+    WarpBitonicSort{temp_storage[warp_id]}.Sort(keys, CustomLess{}, len);
+  }
+};
+
+template <Mode mode, typename KeyT, int Len>
+void partial(nvbench::state& state, nvbench::type_list<nvbench::enum_type<mode>, KeyT, nvbench::enum_type<Len>>)
+{
+  run_bench<partial_op_t<Len / warp_threads>, mode, KeyT, void, Len>(state);
+}
+
+NVBENCH_BENCH_TYPES(partial, NVBENCH_TYPE_AXES(modes, key_types, len_values))
+  .set_type_axes_names({"mode", "KeyT", "len"});
diff --git a/cccl_upstream/cub/benchmarks/bench/bitonic_sort/warp_pairs.cu b/cccl_upstream/cub/benchmarks/bench/bitonic_sort/warp_pairs.cu
new file mode 100644
index 00000000..5db4ae19
--- /dev/null
+++ b/cccl_upstream/cub/benchmarks/bench/bitonic_sort/warp_pairs.cu
@@ -0,0 +1,86 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <cub/warp/warp_bitonic_sort.cuh>
+
+#include <nvbench_helper.cuh>
+
+#include "bitonic_common.cuh"
+
+using modes       = nvbench::enum_type_list<Mode::Latency, Mode::Throughput>;
+using key_types   = nvbench::type_list<int16_t, float>;
+using value_types = offset_types;
+using len_values  = nvbench::enum_type_list<32, 64, 96, 128, 160, 192, 224, 256>;
+
+template <int ItemsPerThread>
+struct full_op_t
+{
+  template <typename KeyT, typename ValueT>
+  _CCCL_DEVICE _CCCL_FORCEINLINE void
+  operator()(KeyT (&keys)[ItemsPerThread], ValueT (&values)[ItemsPerThread], int) const
+  {
+    using WarpBitonicSort = cub::detail::WarpBitonicSort<KeyT, ItemsPerThread, warp_threads, ValueT>;
+    using TempStorage     = typename WarpBitonicSort::TempStorage;
+    __shared__ TempStorage temp_storage[32];
+    const auto warp_id = threadIdx.x / 32;
+    WarpBitonicSort{temp_storage[warp_id]}.Sort(keys, values, CustomLess{});
+  }
+};
+
+template <Mode mode, typename KeyT, typename ValueT, int Len>
+void full(nvbench::state& state, nvbench::type_list<nvbench::enum_type<mode>, KeyT, ValueT, nvbench::enum_type<Len>>)
+{
+  run_bench<full_op_t<Len / warp_threads>, mode, KeyT, ValueT, Len>(state);
+}
+
+NVBENCH_BENCH_TYPES(full, NVBENCH_TYPE_AXES(modes, key_types, value_types, len_values))
+  .set_type_axes_names({"mode", "KeyT", "ValueT", "len"});
+
+template <int ItemsPerThread>
+struct partial_oob_op_t
+{
+  template <typename KeyT, typename ValueT>
+  _CCCL_DEVICE _CCCL_FORCEINLINE void
+  operator()(KeyT (&keys)[ItemsPerThread], ValueT (&values)[ItemsPerThread], int len) const
+  {
+    using WarpBitonicSort = cub::detail::WarpBitonicSort<KeyT, ItemsPerThread, warp_threads, ValueT>;
+    using TempStorage     = typename WarpBitonicSort::TempStorage;
+    __shared__ TempStorage temp_storage[32];
+    const auto warp_id = threadIdx.x / 32;
+    WarpBitonicSort{temp_storage[warp_id]}.Sort(keys, values, CustomLess{}, len, CustomLess::oob_default<KeyT>);
+  }
+};
+
+template <Mode mode, typename KeyT, typename ValueT, int Len>
+void partial_oob(nvbench::state& state,
+                 nvbench::type_list<nvbench::enum_type<mode>, KeyT, ValueT, nvbench::enum_type<Len>>)
+{
+  run_bench<partial_oob_op_t<Len / warp_threads>, mode, KeyT, ValueT, Len>(state);
+}
+
+NVBENCH_BENCH_TYPES(partial_oob, NVBENCH_TYPE_AXES(modes, key_types, value_types, len_values))
+  .set_type_axes_names({"mode", "KeyT", "ValueT", "len"});
+
+template <int ItemsPerThread>
+struct partial_op_t
+{
+  template <typename KeyT, typename ValueT>
+  _CCCL_DEVICE _CCCL_FORCEINLINE void
+  operator()(KeyT (&keys)[ItemsPerThread], ValueT (&values)[ItemsPerThread], int len) const
+  {
+    using WarpBitonicSort = cub::detail::WarpBitonicSort<KeyT, ItemsPerThread, warp_threads, ValueT>;
+    using TempStorage     = typename WarpBitonicSort::TempStorage;
+    __shared__ TempStorage temp_storage[32];
+    const auto warp_id = threadIdx.x / 32;
+    WarpBitonicSort{temp_storage[warp_id]}.Sort(keys, values, CustomLess{}, len);
+  }
+};
+
+template <Mode mode, typename KeyT, typename ValueT, int Len>
+void partial(nvbench::state& state, nvbench::type_list<nvbench::enum_type<mode>, KeyT, ValueT, nvbench::enum_type<Len>>)
+{
+  run_bench<partial_op_t<Len / warp_threads>, mode, KeyT, ValueT, Len>(state);
+}
+
+NVBENCH_BENCH_TYPES(partial, NVBENCH_TYPE_AXES(modes, key_types, value_types, len_values))
+  .set_type_axes_names({"mode", "KeyT", "ValueT", "len"});
diff --git a/cccl_upstream/cub/benchmarks/bench/for_each/extents.cu b/cccl_upstream/cub/benchmarks/bench/for_each/extents.cu
index 53476386..1003e2c5 100644
--- a/cccl_upstream/cub/benchmarks/bench/for_each/extents.cu
+++ b/cccl_upstream/cub/benchmarks/bench/for_each/extents.cu
@@ -56,7 +56,10 @@ void for_each_in_extents(nvbench::state& state, nvbench::type_list<T, OffsetT>)
   });
 }
 
-NVBENCH_BENCH_TYPES(for_each_in_extents, NVBENCH_TYPE_AXES(fundamental_types, offset_types))
+// in tile there are not 128 bit types, so we cannot use 64 bit types because they get promoted internally
+using offsets = cuda::std::conditional_t<_CCCL_TILE_COMPILATION(), nvbench::type_list<int32_t>, offset_types>;
+
+NVBENCH_BENCH_TYPES(for_each_in_extents, NVBENCH_TYPE_AXES(fundamental_types, offsets))
   .set_name("base")
   .set_type_axes_names({"T{ct}", "OffsetT{ct}"})
   .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4));
diff --git a/cccl_upstream/cub/benchmarks/bench/histogram/histogram_common.cuh b/cccl_upstream/cub/benchmarks/bench/histogram/histogram_common.cuh
index 348c03c3..c8756526 100644
--- a/cccl_upstream/cub/benchmarks/bench/histogram/histogram_common.cuh
+++ b/cccl_upstream/cub/benchmarks/bench/histogram/histogram_common.cuh
@@ -38,7 +38,7 @@ constexpr cub::BlockHistogramMemoryPreference MEM_PREFERENCE = cub::BLEND;
 template <typename SampleT, int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS>
 struct bench_policy_selector
 {
-  _CCCL_API constexpr auto operator()(::cuda::compute_capability) const -> cub::HistogramPolicy
+  _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const -> cub::HistogramPolicy
   {
     constexpr cub::BlockLoadAlgorithm load_algorithm =
       (TUNE_LOAD_ALGORITHM == cub::BLOCK_LOAD_STRIPED)
diff --git a/cccl_upstream/cub/benchmarks/bench/select/flagged.cu b/cccl_upstream/cub/benchmarks/bench/select/flagged.cu
index ac9a6921..7508c49f 100644
--- a/cccl_upstream/cub/benchmarks/bench/select/flagged.cu
+++ b/cccl_upstream/cub/benchmarks/bench/select/flagged.cu
@@ -17,12 +17,13 @@
 // %RANGE% TUNE_MAGIC_NS ns 0:2048:4
 // %RANGE% TUNE_DELAY_CONSTRUCTOR_ID dcid 0:7:1
 // %RANGE% TUNE_L2_WRITE_LATENCY_NS l2w 0:1200:5
+// %RANGE% TUNE_PREFETCH pf 0:3:1
 
 #if !TUNE_BASE
 template <typename InputT>
 struct bench_policy_selector
 {
-  [[nodiscard]] _CCCL_API constexpr auto operator()(cuda::compute_capability) const -> cub::SelectPolicy
+  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::SelectPolicy
   {
     return {cub::SelectAlgorithm::lookback,
             {TUNE_THREADS_PER_BLOCK,
@@ -30,7 +31,8 @@ struct bench_policy_selector
              (TUNE_TRANSPOSE == 0 ? cub::BLOCK_LOAD_DIRECT : cub::BLOCK_LOAD_WARP_TRANSPOSE),
              (TUNE_LOAD == 0 ? cub::LOAD_DEFAULT : cub::LOAD_CA),
              cub::BLOCK_SCAN_WARP_SCANS,
-             lookback_delay_policy}};
+             lookback_delay_policy,
+             static_cast<cub::detail::LoadPrefetch>(TUNE_PREFETCH)}};
   }
 };
 #endif // !TUNE_BASE
diff --git a/cccl_upstream/cub/benchmarks/bench/select/if.cu b/cccl_upstream/cub/benchmarks/bench/select/if.cu
index 06f46318..fd4bfb7f 100644
--- a/cccl_upstream/cub/benchmarks/bench/select/if.cu
+++ b/cccl_upstream/cub/benchmarks/bench/select/if.cu
@@ -24,7 +24,7 @@
 template <typename InputT>
 struct bench_policy_selector
 {
-  [[nodiscard]] _CCCL_API constexpr auto operator()(cuda::compute_capability) const -> cub::SelectPolicy
+  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::SelectPolicy
   {
     return {cub::SelectAlgorithm::lookback,
             {TUNE_THREADS_PER_BLOCK,
diff --git a/cccl_upstream/cub/benchmarks/bench/select/unique.cu b/cccl_upstream/cub/benchmarks/bench/select/unique.cu
index ed54d58c..8e046121 100644
--- a/cccl_upstream/cub/benchmarks/bench/select/unique.cu
+++ b/cccl_upstream/cub/benchmarks/bench/select/unique.cu
@@ -22,7 +22,7 @@
 template <typename InputT>
 struct bench_policy_selector
 {
-  [[nodiscard]] _CCCL_API constexpr auto operator()(cuda::compute_capability) const -> cub::SelectPolicy
+  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::SelectPolicy
   {
     return {cub::SelectAlgorithm::lookback,
             {TUNE_THREADS_PER_BLOCK,