[pre-commit.ci] pre-commit autoupdate (#3248)

* [pre-commit.ci] pre-commit autoupdate

updates:
- [github.com/pre-commit/mirrors-clang-format: v18.1.8 → v19.1.6](pre-commit/mirrors-clang-format@v18.1.8...v19.1.6)
- [github.com/astral-sh/ruff-pre-commit: v0.8.3 → v0.8.6](astral-sh/ruff-pre-commit@v0.8.3...v0.8.6)
- [github.com/pre-commit/mirrors-mypy: v1.13.0 → v1.14.1](pre-commit/mirrors-mypy@v1.13.0...v1.14.1)

Co-authored-by: Michael Schellenberger Costa <[email protected]>

.pre-commit-config.yaml

@@ -17,7 +17,7 @@ repos:
       - id: mixed-line-ending
       - id: trailing-whitespace
   - repo: https://github.com/pre-commit/mirrors-clang-format
-    rev: v18.1.8
+    rev: v19.1.6
     hooks:
       - id: clang-format
         types_or: [file]
@@ -39,7 +39,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.8.3
+    rev: v0.8.6
     hooks:
     - id: ruff  # linter
     - id: ruff-format  # formatter
@@ -57,7 +57,7 @@ repos:
 
 
   - repo: https://github.com/pre-commit/mirrors-mypy
-    rev: 'v1.13.0'
+    rev: 'v1.14.1'
     hooks:
       - id: mypy
         additional_dependencies: [types-cachetools, numpy]

cub/cub/agent/agent_histogram.cuh

@@ -629,7 +629,7 @@ struct AgentHistogram
 
     // Set valid flags
     MarkValid<IS_FULL_TILE>(
-      is_valid, valid_samples, Int2Type<AgentHistogramPolicyT::LOAD_ALGORITHM == BLOCK_LOAD_STRIPED>{});
+      is_valid, valid_samples, Int2Type < AgentHistogramPolicyT::LOAD_ALGORITHM == BLOCK_LOAD_STRIPED > {});
 
     // Accumulate samples
     if (prefer_smem)

cub/cub/agent/agent_reduce.cuh

@@ -382,8 +382,8 @@ struct AgentReduce
     even_share.template BlockInit<TILE_ITEMS>(block_offset, block_end);
 
     return (IsAligned(d_in + block_offset, Int2Type<ATTEMPT_VECTORIZATION>()))
-           ? ConsumeRange(even_share, Int2Type < true && ATTEMPT_VECTORIZATION > ())
-           : ConsumeRange(even_share, Int2Type < false && ATTEMPT_VECTORIZATION > ());
+           ? ConsumeRange(even_share, Int2Type<true && ATTEMPT_VECTORIZATION>())
+           : ConsumeRange(even_share, Int2Type<false && ATTEMPT_VECTORIZATION>());
   }
 
   /**
@@ -396,8 +396,8 @@ struct AgentReduce
     even_share.template BlockInit<TILE_ITEMS, GRID_MAPPING_STRIP_MINE>();
 
     return (IsAligned(d_in, Int2Type<ATTEMPT_VECTORIZATION>()))
-           ? ConsumeRange(even_share, Int2Type < true && ATTEMPT_VECTORIZATION > ())
-           : ConsumeRange(even_share, Int2Type < false && ATTEMPT_VECTORIZATION > ());
+           ? ConsumeRange(even_share, Int2Type<true && ATTEMPT_VECTORIZATION>())
+           : ConsumeRange(even_share, Int2Type<false && ATTEMPT_VECTORIZATION>());
   }
 
 private:

cub/cub/block/block_radix_rank.cuh

@@ -606,8 +606,7 @@ private:
     {
       volatile DigitCounterT warp_digit_counters[RADIX_DIGITS][PADDED_WARPS];
       DigitCounterT raking_grid[BLOCK_THREADS][PADDED_RAKING_SEGMENT];
-    }
-    aliasable;
+    } aliasable;
   };
 #endif // !_CCCL_DOXYGEN_INVOKED
 

cub/cub/detail/strong_load.cuh

@@ -59,14 +59,14 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE uint4 load_relaxed(uint4 const* ptr)
   uint4 retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.v4.u32 {%0, %1, %2, %3}, [%4];"
-                  : "=r"(retval.x), "=r"(retval.y), "=r"(retval.z), "=r"(retval.w)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.v4.u32 {%0, %1, %2, %3}, [%4];"
-                  : "=r"(retval.x), "=r"(retval.y), "=r"(retval.z), "=r"(retval.w)
-                  : "l"(ptr)
-                  : "memory");));
+    (asm volatile("ld.relaxed.gpu.v4.u32 {%0, %1, %2, %3}, [%4];" : "=r"(retval.x),
+                  "=r"(retval.y),
+                  "=r"(retval.z),
+                  "=r"(retval.w) : "l"(ptr) : "memory");),
+    (asm volatile("ld.cg.v4.u32 {%0, %1, %2, %3}, [%4];" : "=r"(retval.x),
+                  "=r"(retval.y),
+                  "=r"(retval.z),
+                  "=r"(retval.w) : "l"(ptr) : "memory");));
   return retval;
 }
 
@@ -75,14 +75,8 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE ulonglong2 load_relaxed(ulonglong2 const*
   ulonglong2 retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.v2.u64 {%0, %1}, [%2];"
-                  : "=l"(retval.x), "=l"(retval.y)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];"
-                  : "=l"(retval.x), "=l"(retval.y)
-                  : "l"(ptr)
-                  : "memory");));
+    (asm volatile("ld.relaxed.gpu.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");),
+    (asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");));
   return retval;
 }
 
@@ -91,14 +85,14 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE ushort4 load_relaxed(ushort4 const* ptr)
   ushort4 retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.v4.u16 {%0, %1, %2, %3}, [%4];"
-                  : "=h"(retval.x), "=h"(retval.y), "=h"(retval.z), "=h"(retval.w)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.v4.u16 {%0, %1, %2, %3}, [%4];"
-                  : "=h"(retval.x), "=h"(retval.y), "=h"(retval.z), "=h"(retval.w)
-                  : "l"(ptr)
-                  : "memory");));
+    (asm volatile("ld.relaxed.gpu.v4.u16 {%0, %1, %2, %3}, [%4];" : "=h"(retval.x),
+                  "=h"(retval.y),
+                  "=h"(retval.z),
+                  "=h"(retval.w) : "l"(ptr) : "memory");),
+    (asm volatile("ld.cg.v4.u16 {%0, %1, %2, %3}, [%4];" : "=h"(retval.x),
+                  "=h"(retval.y),
+                  "=h"(retval.z),
+                  "=h"(retval.w) : "l"(ptr) : "memory");));
   return retval;
 }
 
@@ -107,63 +101,36 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE uint2 load_relaxed(uint2 const* ptr)
   uint2 retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.v2.u32 {%0, %1}, [%2];"
-                  : "=r"(retval.x), "=r"(retval.y)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];"
-                  : "=r"(retval.x), "=r"(retval.y)
-                  : "l"(ptr)
-                  : "memory");));
+    (asm volatile("ld.relaxed.gpu.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");),
+    (asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");));
   return retval;
 }
 
 static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned long long load_relaxed(unsigned long long const* ptr)
 {
   unsigned long long retval;
-  NV_IF_TARGET(
-    NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.u64 %0, [%1];"
-                  : "=l"(retval)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.u64 %0, [%1];"
-                  : "=l"(retval)
-                  : "l"(ptr)
-                  : "memory");));
+  NV_IF_TARGET(NV_PROVIDES_SM_70,
+               (asm volatile("ld.relaxed.gpu.u64 %0, [%1];" : "=l"(retval) : "l"(ptr) : "memory");),
+               (asm volatile("ld.cg.u64 %0, [%1];" : "=l"(retval) : "l"(ptr) : "memory");));
   return retval;
 }
 
 static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int load_relaxed(unsigned int const* ptr)
 {
   unsigned int retval;
-  NV_IF_TARGET(
-    NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.u32 %0, [%1];"
-                  : "=r"(retval)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.u32 %0, [%1];"
-                  : "=r"(retval)
-                  : "l"(ptr)
-                  : "memory");));
+  NV_IF_TARGET(NV_PROVIDES_SM_70,
+               (asm volatile("ld.relaxed.gpu.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory");),
+               (asm volatile("ld.cg.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory");));
 
   return retval;
 }
 
 static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned short load_relaxed(unsigned short const* ptr)
 {
   unsigned short retval;
-  NV_IF_TARGET(
-    NV_PROVIDES_SM_70,
-    (asm volatile("ld.relaxed.gpu.u16 %0, [%1];"
-                  : "=h"(retval)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.u16 %0, [%1];"
-                  : "=h"(retval)
-                  : "l"(ptr)
-                  : "memory");));
+  NV_IF_TARGET(NV_PROVIDES_SM_70,
+               (asm volatile("ld.relaxed.gpu.u16 %0, [%1];" : "=h"(retval) : "l"(ptr) : "memory");),
+               (asm volatile("ld.cg.u16 %0, [%1];" : "=h"(retval) : "l"(ptr) : "memory");));
   return retval;
 }
 
@@ -172,24 +139,16 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned char load_relaxed(unsigned char c
   unsigned short retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile(
-       "{"
-       "  .reg .u8 datum;"
-       "  ld.relaxed.gpu.u8 datum, [%1];"
-       "  cvt.u16.u8 %0, datum;"
-       "}"
-       : "=h"(retval)
-       : "l"(ptr)
-       : "memory");),
-    (asm volatile(
-       "{"
-       "  .reg .u8 datum;"
-       "  ld.cg.u8 datum, [%1];"
-       "  cvt.u16.u8 %0, datum;"
-       "}"
-       : "=h"(retval)
-       : "l"(ptr)
-       : "memory");));
+    (asm volatile("{"
+                  "  .reg .u8 datum;"
+                  "  ld.relaxed.gpu.u8 datum, [%1];"
+                  "  cvt.u16.u8 %0, datum;"
+                  "}" : "=h"(retval) : "l"(ptr) : "memory");),
+    (asm volatile("{"
+                  "  .reg .u8 datum;"
+                  "  ld.cg.u8 datum, [%1];"
+                  "  cvt.u16.u8 %0, datum;"
+                  "}" : "=h"(retval) : "l"(ptr) : "memory");));
   return (unsigned char) retval;
 }
 
@@ -198,14 +157,8 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE ulonglong2 load_acquire(ulonglong2 const*
   ulonglong2 retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile("ld.acquire.gpu.v2.u64 {%0, %1}, [%2];"
-                  : "=l"(retval.x), "=l"(retval.y)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];"
-                  : "=l"(retval.x), "=l"(retval.y)
-                  : "l"(ptr)
-                  : "memory");
+    (asm volatile("ld.acquire.gpu.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");),
+    (asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");
      __threadfence();));
   return retval;
 }
@@ -215,32 +168,18 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE uint2 load_acquire(uint2 const* ptr)
   uint2 retval;
   NV_IF_TARGET(
     NV_PROVIDES_SM_70,
-    (asm volatile("ld.acquire.gpu.v2.u32 {%0, %1}, [%2];"
-                  : "=r"(retval.x), "=r"(retval.y)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];"
-                  : "=r"(retval.x), "=r"(retval.y)
-                  : "l"(ptr)
-                  : "memory");
+    (asm volatile("ld.acquire.gpu.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");),
+    (asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");
      __threadfence();));
   return retval;
 }
 
 static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int load_acquire(unsigned int const* ptr)
 {
   unsigned int retval;
-  NV_IF_TARGET(
-    NV_PROVIDES_SM_70,
-    (asm volatile("ld.acquire.gpu.u32 %0, [%1];"
-                  : "=r"(retval)
-                  : "l"(ptr)
-                  : "memory");),
-    (asm volatile("ld.cg.u32 %0, [%1];"
-                  : "=r"(retval)
-                  : "l"(ptr)
-                  : "memory");
-     __threadfence();));
+  NV_IF_TARGET(NV_PROVIDES_SM_70,
+               (asm volatile("ld.acquire.gpu.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory");),
+               (asm volatile("ld.cg.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory"); __threadfence();));
 
   return retval;
 }