[pre-commit.ci] pre-commit autoupdate

.pre-commit-config.yaml

@@ -11,7 +11,7 @@ ci:
 
 repos:
   - repo: https://github.com/pre-commit/pre-commit-hooks
-    rev: v5.0.0
+    rev: v6.0.0
     hooks:
       - id: end-of-file-fixer
       - id: mixed-line-ending
@@ -40,7 +40,7 @@ repos:
           )
 
   - repo: https://github.com/pre-commit/mirrors-clang-format
-    rev: v20.1.7
+    rev: v22.1.5
     hooks:
       - id: clang-format
         types_or: [file]
@@ -62,7 +62,7 @@ repos:
   # TODO/REMINDER: add the Ruff vscode extension to the devcontainers
   # Ruff, the Python auto-correcting linter/formatter written in Rust
   - repo: https://github.com/astral-sh/ruff-pre-commit
-    rev: v0.12.2
+    rev: v0.15.20
     hooks:
     - id: ruff  # linter
     - id: ruff-format  # formatter
@@ -74,7 +74,7 @@ repos:
     - id: gersemi
 
   - repo: https://github.com/codespell-project/codespell
-    rev: v2.4.1
+    rev: v2.4.2
     hooks:
       - id: codespell
         additional_dependencies: [tomli]
@@ -86,7 +86,7 @@ repos:
           )
 
   - repo: https://github.com/pre-commit/mirrors-mypy
-    rev: 'v1.16.1'
+    rev: 'v2.1.0'
     hooks:
       - id: mypy
         # numpy is intentionally not installed here: numpy>=2.3 ships PEP 695

benchmarks/scripts/cccl/bench/search.py

@@ -80,12 +80,14 @@ def filter_benchmark_space_for_p0(algname, ct_space, rt_values):
     ]:
         ct_space = list(
             filter(
-                lambda variant: not (
-                    ("OffsetT{ct}=I64" in variant)
-                    or ("KeyT{ct}=I16" in variant)
-                    or ("ValueT{ct}=I16" in variant)
-                    or ("KeyT{ct}=I128" in variant)
-                    or ("ValueT{ct}=I128" in variant)
+                lambda variant: (
+                    not (
+                        ("OffsetT{ct}=I64" in variant)
+                        or ("KeyT{ct}=I16" in variant)
+                        or ("ValueT{ct}=I16" in variant)
+                        or ("KeyT{ct}=I128" in variant)
+                        or ("ValueT{ct}=I128" in variant)
+                    )
                 ),
                 ct_space,
             )

benchmarks/scripts/sol.py

@@ -103,9 +103,7 @@ def alg_bws(dfs, verbose):
     medians["hue"] = medians["hue"] + gpuname + " "
     if medians["variant"].unique().size > 1:
         variant = (
-            medians["variant"]
-            .astype(str)
-            .map(lambda x: (" " + x if x != "base" else ""))
+            medians["variant"].astype(str).map(lambda x: " " + x if x != "base" else "")
         )
         medians["hue"] = medians["hue"] + variant + " "
     if verbose or medians["ctk"].unique().size > 1:

c/parallel/test/test_merge_sort.cpp

@@ -252,7 +252,7 @@ C2H_TEST("DeviceMergeSort:SortPairsCopy works with custom types", "[merge_sort]"
 {
   const size_t num_items      = GENERATE_COPY(take(2, random(1, 100000)), values({5, 10000, 100000}));
   operation_t op              = make_operation("op",
-                                  R"(struct key_pair { short a; size_t b; };
+                                               R"(struct key_pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) {
   key_pair* lhs = static_cast<key_pair*>(lhs_ptr);
   key_pair* rhs = static_cast<key_pair*>(rhs_ptr);
@@ -308,7 +308,7 @@ C2H_TEST("DeviceMergeSort:SortPairsCopy works with custom types with well-known
 {
   const size_t num_items      = GENERATE_COPY(take(2, random(1, 100000)), values({5, 10000, 100000}));
   operation_t op_state        = make_operation("op",
-                                        R"(struct key_pair { short a; size_t b; };
+                                               R"(struct key_pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) {
   key_pair* lhs = static_cast<key_pair*>(lhs_ptr);
   key_pair* rhs = static_cast<key_pair*>(rhs_ptr);

c/parallel/test/test_reduce.cpp

@@ -229,7 +229,7 @@ C2H_TEST("Reduce works with custom types", "[reduce]")
   const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24)));
 
   operation_t op              = make_operation("op",
-                                  R"(struct pair { short a; size_t b; };
+                                               R"(struct pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) {
   pair* lhs = static_cast<pair*>(lhs_ptr);
   pair* rhs = static_cast<pair*>(rhs_ptr);
@@ -266,7 +266,7 @@ C2H_TEST("Reduce works with custom types with well-known operations", "[reduce][
   const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24)));
 
   operation_t op_state        = make_operation("op",
-                                        R"(struct pair { short a; size_t b; };
+                                               R"(struct pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) {
   pair* lhs = static_cast<pair*>(lhs_ptr);
   pair* rhs = static_cast<pair*>(rhs_ptr);

c/parallel/test/test_scan.cpp

@@ -399,7 +399,7 @@ C2H_TEST("Scan works with custom types", "[scan]")
   const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24)));
 
   operation_t op              = make_operation("op",
-                                  R"(struct pair { short a; size_t b; };
+                                               R"(struct pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) {
   pair* lhs = static_cast<pair*>(lhs_ptr);
   pair* rhs = static_cast<pair*>(rhs_ptr);
@@ -439,7 +439,7 @@ C2H_TEST("Scan works with custom types with well-known operations", "[scan][well
   const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24)));
 
   operation_t op_state        = make_operation("op",
-                                        R"(struct pair { short a; size_t b; };
+                                               R"(struct pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, void* out_ptr) {
   pair* lhs = static_cast<pair*>(lhs_ptr);
   pair* rhs = static_cast<pair*>(rhs_ptr);

c/parallel/test/test_transform.cpp

@@ -355,7 +355,7 @@ C2H_TEST("Transform works with output of different type", "[transform]")
   const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24)));
 
   operation_t op               = make_operation("op",
-                                  R"(struct pair { short a; size_t b; };
+                                                R"(struct pair { short a; size_t b; };
 extern "C" __device__ void op(void* x_ptr, void* out_ptr) {
   int* x = static_cast<int*>(x_ptr);
   pair* out = static_cast<pair*>(out_ptr);
@@ -443,7 +443,7 @@ C2H_TEST("Transform works with custom types", "[transform]")
   const std::size_t num_items = GENERATE(0, 42, take(4, random(1 << 12, 1 << 24)));
 
   operation_t op              = make_operation("op",
-                                  R"(struct pair { short a; size_t b; };
+                                               R"(struct pair { short a; size_t b; };
 extern "C" __device__ void op(void* x_ptr, void* out_ptr) {
   pair* x = static_cast<pair*>(x_ptr);
   pair* out = static_cast<pair*>(out_ptr);

c/parallel/test/test_unique_by_key.cpp

@@ -468,7 +468,7 @@ C2H_TEST("DeviceSelect::UniqueByKey works with custom types", "[device][select_u
   const int num_items = GENERATE_COPY(take(2, random(1, 1000000)));
 
   operation_t op              = make_operation("op",
-                                  R"(struct key_pair { short a; size_t b; };
+                                               R"(struct key_pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) {
   key_pair* lhs = static_cast<key_pair*>(lhs_ptr);
   key_pair* rhs = static_cast<key_pair*>(rhs_ptr);
@@ -541,7 +541,7 @@ C2H_TEST("DeviceSelect::UniqueByKey works with custom types with well-known oper
   const int num_items = GENERATE_COPY(take(2, random(1, 1000000)));
 
   operation_t op_state        = make_operation("op",
-                                        R"(struct key_pair { short a; size_t b; };
+                                               R"(struct key_pair { short a; size_t b; };
 extern "C" __device__ void op(void* lhs_ptr, void* rhs_ptr, bool* out_ptr) {
   key_pair* lhs = static_cast<key_pair*>(lhs_ptr);
   key_pair* rhs = static_cast<key_pair*>(rhs_ptr);
@@ -706,7 +706,7 @@ C2H_TEST("DeviceSelect::UniqueByKey fails to build for large types due to no vsm
   const int num_items = 1;
 
   operation_t op           = make_operation("op",
-                                  R"(struct large_key_pair { int a; char c[500]; };
+                                            R"(struct large_key_pair { int a; char c[500]; };
 extern "C" __device__ bool op(large_key_pair lhs, large_key_pair rhs) {
   return lhs.a == rhs.a;
 })");

c/parallel/test/test_util.h

@@ -1099,10 +1099,10 @@ inline std::tuple<std::string, std::string, std::string> make_counting_iterator_
   std::string iterator_state_def_src = std::format("struct {0} {{ {1} value; }};\n", iterator_state_name, value_type);
   std::string advance_fn_def_src     = std::format(
     "extern \"C\" __device__ void {0}(void* state, const void* offset) {{\n"
-        "  auto* typed_state = static_cast<{1}*>(state);\n"
-        "  auto offset_val = *static_cast<const unsigned long long*>(offset);\n"
-        "  typed_state->value += offset_val;\n"
-        "}}",
+    "  auto* typed_state = static_cast<{1}*>(state);\n"
+    "  auto offset_val = *static_cast<const unsigned long long*>(offset);\n"
+    "  typed_state->value += offset_val;\n"
+    "}}",
     advance_fn_name,
     iterator_state_name);
 
@@ -1186,10 +1186,10 @@ inline std::tuple<std::string, std::string, std::string> make_reverse_iterator_s
   std::string iterator_state_src = std::format("struct {0} {{ {1}* data; }};\n", iterator_state_name, value_type);
   std::string advance_fn_src     = std::format(
     "extern \"C\" __device__ void {0}(void* state, const void* offset) {{\n"
-        "  auto* typed_state = static_cast<{1}*>(state);\n"
-        "  auto offset_val = *static_cast<const unsigned long long*>(offset);\n"
-        "  typed_state->data -= offset_val;\n"
-        "}}",
+    "  auto* typed_state = static_cast<{1}*>(state);\n"
+    "  auto offset_val = *static_cast<const unsigned long long*>(offset);\n"
+    "  typed_state->data -= offset_val;\n"
+    "}}",
     advance_fn_name,
     iterator_state_name);
   std::string dereference_fn_src;

cub/cub/agent/agent_batched_topk.cuh

@@ -216,7 +216,7 @@ struct agent_batched_topk_worker_per_segment
     else
     {
       // Process small segment
-      const auto k         = (::cuda::std::min) (params::get_param(k_param, segment_id),
+      const auto k = (::cuda::std::min) (params::get_param(k_param, segment_id),
                                          static_cast<decltype(params::get_param(k_param, segment_id))>(segment_size));
       const auto direction = select_directions.get_param(segment_id);
 

cub/cub/agent/agent_select_if.cuh

@@ -893,7 +893,7 @@ struct AgentSelectIf
       0,
       0,
       num_tile_selections,
-      bool_constant_v < SelectionOpt == SelectImpl::Partition >);
+      bool_constant_v<SelectionOpt == SelectImpl::Partition>);
 
     return num_tile_selections;
   }
@@ -976,7 +976,7 @@ struct AgentSelectIf
       num_selections_prefix,
       num_rejected_prefix,
       num_selections,
-      bool_constant_v < SelectionOpt == SelectImpl::Partition >);
+      bool_constant_v<SelectionOpt == SelectImpl::Partition>);
 
     return num_selections;
   }

cub/cub/agent/agent_three_way_partition.cuh

@@ -335,7 +335,7 @@ struct AgentThreeWayPartition
           // Medium item
           int local_selection_idx = (first_items_selection_indices - num_first_selections_prefix)
                                   + (second_items_selection_indices - num_second_selections_prefix);
-          local_scatter_offset = second_item_end + item_idx - local_selection_idx;
+          local_scatter_offset    = second_item_end + item_idx - local_selection_idx;
         }
 
         temp_storage.raw_exchange.Alias()[local_scatter_offset] = items[ITEM];

cub/cub/block/block_radix_sort.cuh

@@ -741,8 +741,7 @@ public:
   template <class DecomposerT>
   _CCCL_DEVICE _CCCL_FORCEINLINE //
   ::cuda::std::enable_if_t< //
-    !::cuda::std::is_convertible_v<DecomposerT, int>>
-  Sort(KeyT (&keys)[ItemsPerThread], DecomposerT decomposer)
+    !::cuda::std::is_convertible_v<DecomposerT, int>> Sort(KeyT (&keys)[ItemsPerThread], DecomposerT decomposer)
   {
     Sort(keys, decomposer, 0, detail::radix::traits_t<KeyT>::default_end_bit(decomposer));
   }

cub/cub/block/radix_rank_sort_operations.cuh

@@ -544,8 +544,7 @@ private:
 public:
   template <class DecomposerT = detail::identity_decomposer_t>
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE //
-  bit_ordered_type
-  In(bit_ordered_type key, DecomposerT decomposer = {})
+  bit_ordered_type In(bit_ordered_type key, DecomposerT decomposer = {})
   {
     key = bit_ordered_conversion_policy::to_bit_ordered(decomposer, key);
     if constexpr (IS_DESCENDING)
@@ -557,8 +556,7 @@ public:
 
   template <class DecomposerT = detail::identity_decomposer_t>
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE //
-  bit_ordered_type
-  Out(bit_ordered_type key, DecomposerT decomposer = {})
+  bit_ordered_type Out(bit_ordered_type key, DecomposerT decomposer = {})
   {
     if constexpr (IS_DESCENDING)
     {
@@ -570,8 +568,7 @@ public:
 
   template <class DecomposerT = detail::identity_decomposer_t>
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE //
-  bit_ordered_type
-  DefaultKey(DecomposerT decomposer = {})
+  bit_ordered_type DefaultKey(DecomposerT decomposer = {})
   {
     return IS_DESCENDING ? traits::min_raw_binary_key(decomposer) : traits::max_raw_binary_key(decomposer);
   }

cub/cub/detail/fast_modulo_division.cuh

@@ -160,7 +160,7 @@ public:
     _CCCL_ASSERT(static_cast<size_t>(num_bits + BitSize - BitOffset) < sizeof(larger_t) * CHAR_BIT, "overflow error");
     // without explicit power-of-two check, num_bits needs to replace +1 with !::cuda::is_power_of_two(udivisor)
     _multiplier  = static_cast<unsigned_t>(::cuda::ceil_div(larger_t{1} << (num_bits + BitSize - BitOffset), //
-                                                           static_cast<larger_t>(divisor)));
+                                                            static_cast<larger_t>(divisor)));
     _shift_right = num_bits - BitOffset;
     _CCCL_ASSERT(_multiplier != 0, "overflow error");
   }

cub/cub/device/device_scan.cuh

@@ -249,10 +249,10 @@ struct DeviceScan
   {
     using offset_t = detail::choose_offset_t<NumItemsT>;
     using accum_t  = ::cuda::std::__accumulator_t<
-       ScanOpT,
-       cub::detail::it_value_t<ValuesInputIteratorT>,
-       ::cuda::std::
-         _If<::cuda::std::is_same_v<InitValueT, NullType>, cub::detail::it_value_t<ValuesInputIteratorT>, InitValueT>>;
+      ScanOpT,
+      cub::detail::it_value_t<ValuesInputIteratorT>,
+      ::cuda::std::
+        _If<::cuda::std::is_same_v<InitValueT, NullType>, cub::detail::it_value_t<ValuesInputIteratorT>, InitValueT>>;
 
     using default_policy_selector_t =
       detail::scan_by_key::policy_selector_from_types<detail::it_value_t<KeysInputIteratorT>,

cub/cub/device/dispatch/dispatch_adjacent_difference.cuh

@@ -249,14 +249,15 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceAdjacentDifference"
       error                      = CubDebug(
         THRUST_NS_QUALIFIER::cuda_cub::detail::triple_chevron(
           num_tiles, AdjacentDifferencePolicyT::BLOCK_THREADS, 0, stream)
-          .doit(detail::adjacent_difference::DeviceAdjacentDifferenceDifferenceKernel < KernelPolicySelector,
-                InputIteratorT,
-                OutputIteratorT,
-                DifferenceOpT,
-                OffsetT,
-                InputT,
-                AliasOpt == MayAlias::Yes,
-                ReadOpt == ReadOption::Left >,
+          .doit(detail::adjacent_difference::DeviceAdjacentDifferenceDifferenceKernel<
+                  KernelPolicySelector,
+                  InputIteratorT,
+                  OutputIteratorT,
+                  DifferenceOpT,
+                  OffsetT,
+                  InputT,
+                  AliasOpt == MayAlias::Yes,
+                  ReadOpt == ReadOption::Left>,
                 d_input,
                 first_tile_previous,
                 d_output,
@@ -441,14 +442,15 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch(
 
   if (const auto error = CubDebug(
         THRUST_NS_QUALIFIER::cuda_cub::detail::triple_chevron(num_tiles, active_policy.threads_per_block, 0, stream)
-          .doit(DeviceAdjacentDifferenceDifferenceKernel < policy_selector_t,
-                InputIteratorT,
-                OutputIteratorT,
-                DifferenceOpT,
-                offset_t,
-                input_t,
-                AliasOpt == MayAlias::Yes,
-                ReadOpt == ReadOption::Left >,
+          .doit(DeviceAdjacentDifferenceDifferenceKernel<
+                  policy_selector_t,
+                  InputIteratorT,
+                  OutputIteratorT,
+                  DifferenceOpT,
+                  offset_t,
+                  input_t,
+                  AliasOpt == MayAlias::Yes,
+                  ReadOpt == ReadOption::Left>,
                 d_input,
                 first_tile_previous,
                 d_output,

cub/cub/device/dispatch/dispatch_batch_memcpy.cuh

@@ -152,8 +152,8 @@ __launch_bounds__(int(current_policy<PolicySelector>().large_buffer.threads_per_
       {
         if (thread_offset < buffer_sizes[buffer_id])
         {
-          const auto value = read_item < MemcpyOpt == CopyAlg::Memcpy, AliasT,
-                     InputBufferT > (input_buffer_it[buffer_id], thread_offset);
+          const auto value =
+            read_item<MemcpyOpt == CopyAlg::Memcpy, AliasT, InputBufferT>(input_buffer_it[buffer_id], thread_offset);
           write_item<MemcpyOpt == CopyAlg::Memcpy, AliasT, OutputBufferT>(
             output_buffer_it[buffer_id], thread_offset, value);
         }
@@ -346,7 +346,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch(
 
   constexpr BlockOffsetT init_kernel_threads = 128U;
   const auto tile_size                       = static_cast<uint32_t>(active_policy.small_buffer.threads_per_block)
-                       * static_cast<uint32_t>(active_policy.small_buffer.buffers_per_thread);
+                                             * static_cast<uint32_t>(active_policy.small_buffer.buffers_per_thread);
 
   constexpr auto max_num_buffers_per_invocation = ::cuda::std::int64_t{512 * 1024 * 1024};
   static_assert(max_num_buffers_per_invocation <= ::cuda::std::numeric_limits<per_invocation_buffer_offset_t>::max());

cub/cub/device/dispatch/dispatch_histogram.cuh

@@ -216,12 +216,12 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch(
     if constexpr (IsDeviceInit)
     {
       return kernel_source.template HistogramSweepKernelDeviceInit<
-             PolicySelector,
-             PRIVATIZED_SMEM_BINS,
-             FirstLevelArrayT,
-             SecondLevelArrayT,
-             IsEven,
-             IsByteSample>();
+        PolicySelector,
+        PRIVATIZED_SMEM_BINS,
+        FirstLevelArrayT,
+        SecondLevelArrayT,
+        IsEven,
+        IsByteSample>();
     }
     else
     {