CUTLASS 2.3.0 final.

include/cutlass/gemm/kernel/gemm.h

@@ -224,7 +224,7 @@ struct Gemm {
 
     // Broadcast the warp_id computed by lane 0 to ensure dependent code
     // is compiled as warp-uniform.
-    int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+    int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
     int lane_idx = threadIdx.x % 32;
 
     //

include/cutlass/gemm/kernel/gemm_array.h

@@ -184,7 +184,7 @@ struct GemmArray {
       
       // Broadcast the warp_id computed by lane 0 to ensure dependent code
       // is compiled as warp-uniform.
-      int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+      int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
 
       int lane_idx = threadIdx.x % 32;
       

include/cutlass/gemm/kernel/gemm_batched.h

@@ -196,7 +196,7 @@ struct GemmBatched {
 
       // Broadcast the warp_id computed by lane 0 to ensure dependent code
       // is compiled as warp-uniform.
-      int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+      int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
 
       int lane_idx = threadIdx.x % 32;
       

include/cutlass/gemm/kernel/gemm_planar_complex.h

@@ -512,7 +512,7 @@ public:
 
     // Broadcast the warp_id computed by lane 0 to ensure dependent code
     // is compiled as warp-uniform.
-    int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+    int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
 
     int lane_idx = threadIdx.x % 32;
 

include/cutlass/gemm/kernel/gemm_planar_complex_array.h

@@ -441,7 +441,7 @@ public:
 
         // Broadcast the warp_id computed by lane 0 to ensure dependent code
         // is compiled as warp-uniform.
-        int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+        int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
         int lane_idx = threadIdx.x % 32;
     
         //

include/cutlass/gemm/kernel/gemm_universal.h

@@ -402,7 +402,7 @@ public:
 
     // Broadcast the warp_id computed by lane 0 to ensure dependent code
     // is compiled as warp-uniform.
-    int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+    int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
 
     int lane_idx = threadIdx.x % 32;
 

include/cutlass/gemm/kernel/sparse_gemm.h

@@ -269,7 +269,7 @@ struct SparseGemm {
 
     // Broadcast the warp_id computed by lane 0 to ensure dependent code
     // is compiled as warp-uniform.
-    int warp_idx = __shfl_sync(0x1f, threadIdx.x / 32, 0);
+    int warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0);
     int lane_idx = threadIdx.x % 32;
 
     //