cuda: disable sparse V skip (warp divergence regression)

Per-lane branching in the VEC FA kernel causes warp divergence that
costs more than the skipped dequants save. Benchmarked at -0.3% to
-2.8% on RTX 3090/4090 across all context lengths.

Metal path unaffected (remains enabled, +4% to +23%).

TODO: revisit with warp-level ballot skip (__ballot_sync + early
exit when entire warp is below threshold).

Data: @sztlink (Qwen3-30B-A3B Q4_K_M, CUDA SM86/SM89)

Author: TheTom

ggml/src/ggml-cuda/fattn-vec.cuh

@@ -333,6 +333,10 @@ static __global__ void flash_attn_ext_vec(
             }
 
             // Sparse V: skip V dequant if all attention weights for this position are negligible
+            // Disabled — per-lane branching causes warp divergence that costs more than the
+            // skipped dequants save (-0.3% to -2.8% on RTX 3090/4090).
+            // TODO: revisit with warp-level ballot skip.
+#if 0
             {
                 bool dominated = true;
 #pragma unroll
@@ -341,6 +345,7 @@ static __global__ void flash_attn_ext_vec(
                 }
                 if (dominated) { continue; }
             }
+#endif
 
 #pragma unroll
             for (int i_VKQ_0 = 0; i_VKQ_0 < D/2; i_VKQ_0 += nthreads_V*V_rows_per_thread/2) {
@@ -373,6 +378,10 @@ static __global__ void flash_attn_ext_vec(
             }
 
             // Sparse V: skip V dequant if all attention weights for this position are negligible
+            // Disabled — per-lane branching causes warp divergence that costs more than the
+            // skipped dequants save (-0.3% to -2.8% on RTX 3090/4090).
+            // TODO: revisit with warp-level ballot skip.
+#if 0
             {
                 bool dominated = true;
 #pragma unroll
@@ -381,6 +390,7 @@ static __global__ void flash_attn_ext_vec(
                 }
                 if (dominated) { continue; }
             }
+#endif
 
 #pragma unroll
             for (int i_VKQ_0 = 0; i_VKQ_0 < D/2; i_VKQ_0 += nthreads_V*V_rows_per_thread/2) {