diff --git a/csrc/gemm/gemv.cu b/csrc/gemm/gemv.cu
index 4afb05f..4554f94 100644
--- a/csrc/gemm/gemv.cu
+++ b/csrc/gemm/gemv.cu
@@ -28,18 +28,25 @@ __global__ void gemv_bf16_fused_kernel(
     const int t = threadIdx.x;
     const int col = block_n * GEMV_TILE_N + t;
 
-    if (col >= N) return;
-
     const int k_start = block_k * GEMV_TILE_K;
     const int k_end = min(k_start + GEMV_TILE_K, K);
     const int k_len = k_end - k_start;
 
+    // Cooperative load of x into shared memory uses ALL threads in the block
+    // (indexed by t, independent of col). Threads whose column is out of range
+    // must still help load and reach the barrier — returning early here would
+    // leave part of x_shared uninitialized AND make __syncthreads divergent
+    // (UB). So the col>=N check happens only AFTER the load + barrier. This bug
+    // produced intermittent huge/garbage outputs whenever N % GEMV_TILE_N != 0
+    // (e.g. gpt-oss decode o_proj with N=2880), collapsing the forward pass.
     __shared__ float x_shared[GEMV_TILE_K];
     for (int i = t; i < k_len; i += GEMV_BLOCK) {
         x_shared[i] = __bfloat162float(x[k_start + i]);
     }
     __syncthreads();
 
+    if (col >= N) return;
+
     float sum = 0.0f;
     for (int ki = 0; ki < k_len; ki++) {
         sum += x_shared[ki] * __bfloat162float(W[(long long)(k_start + ki) * N + col]);