ggml-opencl: port FA barrier fix to f16/f32_f16 siblings + guard upscale zero dims

PR-review follow-ups:

- Port the flash-attn tile-loop barrier/divergence fix (previously only in
  flash_attn_f32.cl) to the f16 and f32_f16 sibling kernels: guard the score
  loop with `if (my_query_row < n_q)` instead of an early `continue`, and add a
  trailing barrier at the end of the K/V tile loop so out-of-range lanes cannot
  race ahead into the next tile's load while active lanes still read l_k/l_v.
  flash_attn_f32_f16 is the kernel the Qwen3-VL vision tower actually dispatches
  (K/V cast to F16), so the fix was missing from the validated workload.
  Score-loop bodies are unchanged.

- Guard zero source dimensions in ggml_cl_upscale: the sf* scale factors divide
  by the source dims, so a zero source dim yields +inf; the existing early-exit
  only covered zero destination dims.

Validated on-device (S25 Ultra / Adreno 830 OpenCL, QVAC-21297): the GPU vision
projector matches CPU exactly (Delta 0.0 pp, task-for-task) and is ~26% faster
than the shipping CPU projector on encode; 0 device-loss.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Author: iancris

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -9573,6 +9573,12 @@ static void ggml_cl_upscale(ggml_backend_t backend, const ggml_tensor * src0, gg
     const int ne2 = dst->ne[2];
     const int ne3 = dst->ne[3];
 
+    // Zero source dims would make the sf* scale factors below divide by zero (+inf);
+    // the later dst-zero early-exit does not cover this.
+    if (ne00 == 0 || ne01 == 0 || ne02 == 0 || ne03 == 0) {
+        return;
+    }
+
     float sf0 = (float)ne0 / ne00;
     float sf1 = (float)ne1 / ne01;
     float sf2 = (float)ne2 / ne02;

ggml/src/ggml-opencl/kernels/flash_attn_f16.cl

@@ -121,10 +121,15 @@ __kernel void flash_attn_f16(
         }
         barrier(CLK_LOCAL_MEM_FENCE);
 
-        if (my_query_row >= n_q) {
-            continue;
-        }
-
+        // NOTE: do NOT `continue` for out-of-range query rows here. Every
+        // work-item must reach the trailing barrier at the end of this loop,
+        // otherwise the extra lanes (my_query_row >= n_q in the last partial
+        // BLOCK_M block) race ahead to the next tile's load and overwrite
+        // l_k/l_v while active lanes are still reading them. That shared-memory
+        // race silently corrupts the K/V tiles for any sequence spanning more
+        // than one BLOCK_N tile (e.g. the bidirectional Qwen3-VL vision tower,
+        // n_kv=247), degrading the encode. Guard the score loop instead.
+        if (my_query_row < n_q) {
         for (int j = 0; j < BLOCK_N; j += 2) {
             const int k_row0 = k_start + j;
             const int k_row1 = k_start + j + 1;
@@ -170,6 +175,11 @@ __kernel void flash_attn_f16(
             l_i = l_i * scale_prev + p0 + p1;
             m_i = m_new;
         }
+        } // end if (my_query_row < n_q)
+
+        // Ensure every work-item has finished reading l_k/l_v before the next
+        // iteration overwrites the shared K/V tiles.
+        barrier(CLK_LOCAL_MEM_FENCE);
     }
 
     if (my_query_row < n_q) {

ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl

@@ -196,10 +196,15 @@ __kernel void flash_attn_f32_f16(
         }
         barrier(CLK_LOCAL_MEM_FENCE);
 
-        if (my_query_row >= n_q) {
-            continue;
-        }
-
+        // NOTE: do NOT `continue` for out-of-range query rows here. Every
+        // work-item must reach the trailing barrier at the end of this loop,
+        // otherwise the extra lanes (my_query_row >= n_q in the last partial
+        // BLOCK_M block) race ahead to the next tile's load and overwrite
+        // l_k/l_v while active lanes are still reading them. That shared-memory
+        // race silently corrupts the K/V tiles for any sequence spanning more
+        // than one BLOCK_N tile (e.g. the bidirectional Qwen3-VL vision tower,
+        // n_kv=247), degrading the encode. Guard the score loop instead.
+        if (my_query_row < n_q) {
         for (int j = 0; j < BLOCK_N; j += 2) {
             const int k_row0 = k_start + j;
             const int k_row1 = k_start + j + 1;
@@ -245,6 +250,11 @@ __kernel void flash_attn_f32_f16(
             l_i = l_i * scale_prev + p0 + p1;
             m_i = m_new;
         }
+        } // end if (my_query_row < n_q)
+
+        // Ensure every work-item has finished reading l_k/l_v before the next
+        // iteration overwrites the shared K/V tiles.
+        barrier(CLK_LOCAL_MEM_FENCE);
     }
 
     if (my_query_row < n_q) {