fix(cuda): pin turbo4 cooperative warp-collectives to WARP_SIZE (gfx803 corruption)

The mt_paged_attention turbo4 scatter/decode/tile kernels and the set-rows
turbo4 quant kernel are written for 32-lane warps (WARP_SIZE=32,
DECODE_NUM_WARPS=NUM_THREADS/32, lane=tid%32). They issued bare 3-arg
__shfl_*_sync, which the HIP shim maps to width=warpSize. On gfx803 (GCN)
warpSize=64, so two logical 32-lane warps share one 64-wide hardware wave:

  * scatter/set-rows nibble pack: __shfl_sync(.., my_nibble, lane^1) read the
    partner from the wrong half of the wave for threads 32-63/96-127,
    corrupting ~25% of packed K/V nibbles.
  * decode/tile dequant: __shfl_sync(.., norm_f, 0) broadcast lane-0s norm

Author: kmbandy

ggml/src/ggml-cuda/mt_pagedattn.cu

@@ -421,7 +421,7 @@ __global__ void mt_scatter_kv_turbo4_0_kernel(
     // ---- Step 6: Pack qs (nibble packed, warp-cooperative) ----
     const int      lane            = j % WARP_SIZE;
     const uint8_t  my_nibble       = idx & 0xF;
-    const uint8_t  partner_nibble  = __shfl_sync(0xffffffffu, my_nibble, lane ^ 1);
+    const uint8_t  partner_nibble  = __shfl_sync(0xffffffffu, my_nibble, lane ^ 1, WARP_SIZE);
     if ((j & 1) == 0) {
         blk->qs[j / 2] = my_nibble | (partner_nibble << 4);
     }

ggml/src/ggml-cuda/mt_pagedattn_aiter.cu

@@ -322,7 +322,7 @@ __global__ void mt_scatter_kv_turbo4_aiter_kernel(
     {
         const int      lane            = j % WARP_SIZE;
         const uint8_t  my_nibble       = idx & 0xF;
-        const uint8_t  partner_nibble  = __shfl_sync(0xffffffffu, my_nibble, lane ^ 1);
+        const uint8_t  partner_nibble  = __shfl_sync(0xffffffffu, my_nibble, lane ^ 1, WARP_SIZE);
         if ((j & 1) == 0) {
             blk->qs[j / 2] = my_nibble | (partner_nibble << 4);
         }

ggml/src/ggml-cuda/mt_pagedattn_decode.cu

@@ -238,7 +238,7 @@ static __device__ __forceinline__ void decode_coop_stage_turbo4(
                 }
             }
         }
-        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0);
+        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0, WARP_SIZE);
 
         uint16_t packed = 0;
         if (blk != nullptr) {
@@ -302,7 +302,7 @@ static __device__ __forceinline__ void decode_coop_stage_turbo3(
                 }
             }
         }
-        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0);
+        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0, WARP_SIZE);
 
         uint8_t qs_byte    = 0;
         uint8_t signs_byte = 0;
@@ -386,7 +386,7 @@ template <typename T>
 __device__ __forceinline__ T decode_warp_reduce_sum(T v) {
     #pragma unroll
     for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
-        v += __shfl_xor_sync(0xFFFFFFFF, v, mask);
+        v += __shfl_xor_sync(0xFFFFFFFF, v, mask, WARP_SIZE);
     }
     return v;
 }
@@ -395,7 +395,7 @@ template <typename T>
 __device__ __forceinline__ T decode_warp_reduce_max(T v) {
     #pragma unroll
     for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
-        v = max(v, __shfl_xor_sync(0xFFFFFFFF, v, mask));
+        v = max(v, __shfl_xor_sync(0xFFFFFFFF, v, mask, WARP_SIZE));
     }
     return v;
 }
@@ -938,7 +938,7 @@ __global__ void mt_paged_attention_decode_kernel_wmma(
             float local_max = -INFINITY;
             #pragma unroll
             for (int l = 0; l < scores.ne; ++l) local_max = max(local_max, scores.x[l]);
-            const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16));
+            const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16, WARP_SIZE));
 
             const float new_max = max(running_max, row_max);
             float rescale = 1.0f;
@@ -959,7 +959,7 @@ __global__ void mt_paged_attention_decode_kernel_wmma(
                 scores.x[l] = e;
                 local_sum  += e;
             }
-            const float row_sum = local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16);
+            const float row_sum = local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16, WARP_SIZE);
             running_sum += row_sum;
             running_max  = new_max;
 

ggml/src/ggml-cuda/mt_pagedattn_tile.cu

@@ -230,7 +230,7 @@ static __device__ __forceinline__ void coop_stage_turbo4_tile(
         }
 
         // Broadcast norm from lane 0 to all 32 lanes of this warp.
-        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0);
+        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0, WARP_SIZE);
 
         // Each lane reads 2 bytes (4 nibbles = 4 elements). qs is uint8_t[64],
         // 2-byte aligned at qs[2*lane_id] for any lane.
@@ -308,7 +308,7 @@ static __device__ __forceinline__ void coop_stage_turbo3_tile(
         }
 
         // Broadcast norm from lane 0 to all 32 lanes of this warp.
-        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0);
+        norm_f = __shfl_sync(0xFFFFFFFF, norm_f, 0, WARP_SIZE);
 
         // Each lane reads 1 byte of qs (its 4 elements' low2 bits) and 1
         // byte of signs (covers this lane's 4 elements' high1 bits, plus
@@ -476,7 +476,7 @@ __global__ void mt_paged_attention_tile_kernel(
         for (int l = 0; l < scores.ne; ++l) {
             local_max = max(local_max, scores.x[l]);
         }
-        const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16));
+        const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16, WARP_SIZE));
 
         const float new_max = max(running_max, row_max);
 
@@ -502,7 +502,7 @@ __global__ void mt_paged_attention_tile_kernel(
             scores.x[l]   = e;
             local_sum   += e;
         }
-        const float row_sum = local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16);
+        const float row_sum = local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16, WARP_SIZE);
         running_sum += row_sum;
         running_max  = new_max;
 
@@ -861,7 +861,7 @@ __global__ void mt_paged_attention_tile_mw_kernel(
             for (int l = 0; l < scores.ne; ++l) {
                 local_max = max(local_max, scores.x[l]);
             }
-            const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16));
+            const float row_max = max(local_max, __shfl_xor_sync(0xFFFFFFFF, local_max, 16, WARP_SIZE));
 
             const float new_max = max(running_max, row_max);
 
@@ -887,7 +887,7 @@ __global__ void mt_paged_attention_tile_mw_kernel(
                 scores.x[l]   = e;
                 local_sum   += e;
             }
-            const float row_sum = local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16);
+            const float row_sum = local_sum + __shfl_xor_sync(0xFFFFFFFF, local_sum, 16, WARP_SIZE);
             running_sum += row_sum;
             running_max  = new_max;
 

ggml/src/ggml-cuda/set-rows.cu

@@ -1068,7 +1068,7 @@ static __global__ void k_set_rows_turbo4(
     const uint8_t my_nibble = idx & 0xF;
     uint8_t qs_byte = 0;
     // Gather nibble from partner thread
-    uint8_t partner_nibble = __shfl_sync(0xffffffff, my_nibble, lane ^ 1);
+    uint8_t partner_nibble = __shfl_sync(0xffffffff, my_nibble, lane ^ 1, WARP_SIZE);
     if (j % 2 == 0) {
         qs_byte = my_nibble | (partner_nibble << 4);
         blk->qs[j / 2] = qs_byte;