diff --git a/onnxruntime/contrib_ops/webgpu/bert/flash_attention.cc b/onnxruntime/contrib_ops/webgpu/bert/flash_attention.cc
index cfd084d58c2d9..a03f50616cf3a 100644
--- a/onnxruntime/contrib_ops/webgpu/bert/flash_attention.cc
+++ b/onnxruntime/contrib_ops/webgpu/bert/flash_attention.cc
@@ -174,7 +174,8 @@ FlashAttentionProgram::FlashAttentionProgram(const std::string& kernel_name,
                                              bool has_local_window,
                                              uint32_t kv_cache_quantization_bits,
                                              int compressed_head_size_u32,
-                                             bool use_seqlens_q)
+                                             bool use_seqlens_q,
+                                             bool has_softcap)
     : Program{kernel_name},
       has_attention_bias_(has_attention_bias),
       is_qualcomm_(is_qualcomm),
@@ -191,7 +192,8 @@ FlashAttentionProgram::FlashAttentionProgram(const std::string& kernel_name,
       kv_cache_quantization_(kv_cache_quantization_bits != 0),
       kv_cache_quantization_bits_(kv_cache_quantization_bits),
       compressed_head_size_u32_(compressed_head_size_u32),
-      use_seqlens_q_(use_seqlens_q) {
+      use_seqlens_q_(use_seqlens_q),
+      has_softcap_(has_softcap) {
 }
 
 Status SplitPackedQKVWithRotaryEmbeddingAndCopyKVProgram::GenerateShaderCode(ShaderHelper& sh) const {
@@ -404,22 +406,26 @@ Status FlashAttentionProgram::GenerateShaderCode(ShaderHelper& shader) const {
   //  Expectation is that present_key, and present_value contain past key and values since
   //  we are out of storage buffers a shader can have and both past/present cant be passed.
   // The hidden size of each q head should be a multiple of 4 because shader uses vectorized loads.
-  shader.AddInput("q", ShaderUsage::UseUniform | ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias);
-  shader.AddInput("present_key", ShaderUsage::UseUniform);
-  shader.AddInput("present_value", ShaderUsage::UseUniform);
+  const auto& q = shader.AddInput("q", ShaderUsage::UseUniform | ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias);
+  const auto& present_key = shader.AddInput("present_key", ShaderUsage::UseUniform);
+  const auto& present_value = shader.AddInput("present_value", ShaderUsage::UseUniform);
+  const ShaderVariableHelper* attention_bias = nullptr;
   if (has_attention_bias_) {
-    shader.AddInput("attention_bias", ShaderUsage::UseUniform);
+    attention_bias = &shader.AddInput("attention_bias", ShaderUsage::UseUniform);
   }
+  const ShaderVariableHelper* seqlens_k = nullptr;
   if (use_seqlen_k_) {
-    shader.AddInput("seqlens_k", ShaderUsage::None);
+    seqlens_k = &shader.AddInput("seqlens_k", ShaderUsage::None);
   }
+  const ShaderVariableHelper* seqlens_q = nullptr;
   if (use_seqlens_q_) {
-    shader.AddInput("seqlens_q", ShaderUsage::None);
+    seqlens_q = &shader.AddInput("seqlens_q", ShaderUsage::None);
   }
+  const ShaderVariableHelper* head_sink = nullptr;
   if (has_head_sink_) {
-    shader.AddInput("head_sink", ShaderUsage::UseUniform);
+    head_sink = &shader.AddInput("head_sink", ShaderUsage::UseUniform);
   }
-  shader.AddOutput("output", ShaderUsage::UseUniform);
+  const auto& output = shader.AddOutput("output", ShaderUsage::UseUniform);
 
   // Constant private-array indices improve NVIDIA's shared-memory path, but
   // expanding larger heads can regress performance. The selection inputs
@@ -448,18 +454,28 @@ Status FlashAttentionProgram::GenerateShaderCode(ShaderHelper& shader) const {
                              WGSL_TEMPLATE_PARAMETER(has_attention_bias, has_attention_bias_),
                              WGSL_TEMPLATE_PARAMETER(has_head_sink, has_head_sink_),
                              WGSL_TEMPLATE_PARAMETER(has_local_window, has_local_window_),
+                             WGSL_TEMPLATE_PARAMETER(has_softcap, has_softcap_),
                              WGSL_TEMPLATE_PARAMETER(is_qualcomm, is_qualcomm_),
                              WGSL_TEMPLATE_PARAMETER(is_unidirectional, is_unidirectional_),
                              WGSL_TEMPLATE_PARAMETER(kv_cache_quantization, kv_cache_quantization_),
                              WGSL_TEMPLATE_PARAMETER(max_k_step_param, max_k_step_),
                              WGSL_TEMPLATE_PARAMETER(prefer_subgroupshuffle, !is_nvidia_),
                              WGSL_TEMPLATE_PARAMETER(q_BNSH, q_BNSH_),
+                             WGSL_TEMPLATE_PARAMETER(qc_cooperative, is_qualcomm_ && SubgroupSize() == 128),
                              WGSL_TEMPLATE_PARAMETER(qkv_head_size, qkv_head_size_),
                              WGSL_TEMPLATE_PARAMETER(qkv_num_heads, qkv_num_heads_),
                              WGSL_TEMPLATE_PARAMETER(unroll_shm_inner, unroll_shm_inner),
                              WGSL_TEMPLATE_PARAMETER(use_seqlen_k, use_seqlen_k_),
                              WGSL_TEMPLATE_PARAMETER(use_seqlens_q, use_seqlens_q_),
-                             WGSL_TEMPLATE_PARAMETER(use_shm_path, use_shm_path_));
+                             WGSL_TEMPLATE_PARAMETER(use_shm_path, use_shm_path_),
+                             WGSL_TEMPLATE_OPTIONAL_VARIABLE(attention_bias, attention_bias),
+                             WGSL_TEMPLATE_OPTIONAL_VARIABLE(head_sink, head_sink),
+                             WGSL_TEMPLATE_VARIABLE(output, output),
+                             WGSL_TEMPLATE_VARIABLE(present_key, present_key),
+                             WGSL_TEMPLATE_VARIABLE(present_value, present_value),
+                             WGSL_TEMPLATE_VARIABLE(q, q),
+                             WGSL_TEMPLATE_OPTIONAL_VARIABLE(seqlens_k, seqlens_k),
+                             WGSL_TEMPLATE_OPTIONAL_VARIABLE(seqlens_q, seqlens_q));
 }
 
 Status FlashAttentionPagedPrefillProgram::GenerateShaderCode(ShaderHelper& shader) const {
@@ -1228,7 +1244,8 @@ Status ApplyDensePrefill(const Tensor* Q, Tensor* present_key, Tensor* present_v
                                 has_local_window,
                                 kv_cache_quantization_bits,
                                 compressed_head_size_u32,
-                                use_seqlens_q};
+                                use_seqlens_q,
+                                parameters.softcap_ > 0.0f};
   program.AddInputs({{Q, ProgramTensorMetadataDependency::TypeAndRank, 4},
                      {present_key, ProgramTensorMetadataDependency::TypeAndRank,
                       kv_cache_quantization_enabled ? 1 : 4},
@@ -1248,17 +1265,27 @@ Status ApplyDensePrefill(const Tensor* Q, Tensor* present_key, Tensor* present_v
   }
   program.AddOutputs({{attn_output, ProgramTensorMetadataDependency::TypeAndRank, 4}});
 
-  // On Apple GPUs, use a larger workgroup size to reduce barrier overhead.
-  const uint32_t prefill_tile_size = is_apple ? 128 : tile_size;
+  // Share each query across four lanes to reduce register pressure on Qualcomm.
+  const bool qc_cooperative = is_qualcomm && is_fp16 && has_subgroups &&
+                              parameters.head_size_ == 128 && parameters.sequence_length_ >= 128 &&
+                              !kv_cache_quantization_enabled &&
+                              context.HasWGSLLanguageFeature(wgpu::WGSLLanguageFeatureName::SubgroupId) &&
+                              context.HasFeature(wgpu::FeatureName::SubgroupSizeControl) &&
+                              context.AdapterInfo().subgroupMinSize <= 128 && context.AdapterInfo().subgroupMaxSize >= 128;
+  const uint32_t prefill_tile_size = qc_cooperative ? 64 : (is_apple ? 128 : tile_size);
+  if (qc_cooperative) {
+    program.SetSubgroupSize(128);
+    program.SetUseSubgroupId(true);
+  }
   const uint32_t num_seq_tile = (parameters.sequence_length_ + prefill_tile_size - 1) / prefill_tile_size;
 
   program.SetDispatchGroupSize(parameters.batch_size_ * parameters.num_heads_ * num_seq_tile)
-      .SetWorkgroupSize(prefill_tile_size)
+      .SetWorkgroupSize(qc_cooperative ? 256 : prefill_tile_size)
       .CacheHint(has_attention_bias, parameters.head_size_, parameters.num_heads_,
                  parameters.is_unidirectional_, is_qualcomm, is_nvidia, is_apple,
                  has_subgroups, q_BNSH, use_seqlen_k, has_head_sink, has_local_window,
                  kv_cache_quantization_bits,
-                 compressed_head_size_u32, program.max_k_step(), use_seqlens_q)
+                 compressed_head_size_u32, program.max_k_step(), use_seqlens_q, parameters.softcap_ > 0.0f)
       .AddUniformVariables({{static_cast<uint32_t>(parameters.sequence_length_)},
                             {static_cast<uint32_t>(parameters.total_sequence_length_)},
                             {static_cast<uint32_t>(present_sequence_length)},
@@ -1269,7 +1296,8 @@ Status ApplyDensePrefill(const Tensor* Q, Tensor* present_key, Tensor* present_v
                             {attn_bias_dim0},
                             {attn_bias_dim1},
                             {attn_bias_dim3},
-                            {static_cast<uint32_t>(has_local_window ? local_window_size : 0)}});
+                            {static_cast<uint32_t>(has_local_window ? local_window_size : 0)},
+                            {parameters.softcap_}});
 
   return context.RunProgram(program);
 }
diff --git a/onnxruntime/contrib_ops/webgpu/bert/flash_attention.h b/onnxruntime/contrib_ops/webgpu/bert/flash_attention.h
index 1f30813ce63f8..4c1a54e40cd73 100644
--- a/onnxruntime/contrib_ops/webgpu/bert/flash_attention.h
+++ b/onnxruntime/contrib_ops/webgpu/bert/flash_attention.h
@@ -105,7 +105,8 @@ class FlashAttentionProgram final : public Program<FlashAttentionProgram> {
                         bool has_local_window = false,
                         uint32_t kv_cache_quantization_bits = 0,
                         int compressed_head_size_u32 = 0,
-                        bool use_seqlens_q = false);
+                        bool use_seqlens_q = false,
+                        bool has_softcap = false);
 
   Status GenerateShaderCode(ShaderHelper& sh) const override;
 
@@ -121,7 +122,8 @@ class FlashAttentionProgram final : public Program<FlashAttentionProgram> {
                                           {"attn_bias_dim0", ProgramUniformVariableDataType::Uint32},
                                           {"attn_bias_dim1", ProgramUniformVariableDataType::Uint32},
                                           {"attn_bias_dim3", ProgramUniformVariableDataType::Uint32},
-                                          {"local_window_size", ProgramUniformVariableDataType::Uint32});
+                                          {"local_window_size", ProgramUniformVariableDataType::Uint32},
+                                          {"softcap", ProgramUniformVariableDataType::Float32});
 
  private:
   bool has_attention_bias_;
@@ -143,6 +145,7 @@ class FlashAttentionProgram final : public Program<FlashAttentionProgram> {
   // seqlens_q[b] and computes past_sequence_length_b = total_kv_b - q_len_b.
   // When unset (default), callers keep the uniform-q_len clamp path unchanged.
   bool use_seqlens_q_;
+  bool has_softcap_;
 };
 
 // Dedicated single-kernel prefill program for PagedAttention. It mirrors the
diff --git a/onnxruntime/contrib_ops/webgpu/bert/flash_attention.wgsl.template b/onnxruntime/contrib_ops/webgpu/bert/flash_attention.wgsl.template
index 446f5a75dfa5a..eaec9ed125aaa 100644
--- a/onnxruntime/contrib_ops/webgpu/bert/flash_attention.wgsl.template
+++ b/onnxruntime/contrib_ops/webgpu/bert/flash_attention.wgsl.template
@@ -1,7 +1,9 @@
+#param qc_cooperative
 #param bit_width
 #param has_attention_bias
 #param has_head_sink
 #param has_local_window
+#param has_softcap
 #param is_qualcomm
 #param is_unidirectional
 #param prefer_subgroupshuffle
@@ -16,6 +18,8 @@
 #param unroll_shm_inner
 #param max_k_step_param
 
+#use .getByOffset .setByOffset
+
 const head_size : u32 = qkv_head_size;
 const num_heads : u32 = qkv_num_heads;
 
@@ -28,9 +32,7 @@ const COMPRESSED_HEAD_U32 : u32 = compressed_head_size_u32;
 
 // Dequantize 4 consecutive elements from a TQ-packed KV cache into a vec4.
 fn kv_cache_dequant_vec4(
-    kv_cache: ptr<storage, array<u32>, read>, base: u32, elem_base: u32, scale: f32) -> q_value_t {
-  let word_idx = elem_base / KV_CACHE_QUANT_ELEMENTS_PER_WORD;
-  let packed = (*kv_cache)[base + 1u + word_idx];
+    packed: u32, elem_base: u32, scale: f32) -> q_value_t {
   let shift = (elem_base % KV_CACHE_QUANT_ELEMENTS_PER_WORD) * KV_CACHE_QUANT_BITS;
   return kv_cache_quant_dequant_vec4(packed >> shift, scale);
 }
@@ -39,7 +41,7 @@ fn kv_cache_dequant_vec4(
 #if use_seqlen_k
 // When seqlens_k is provided, total_sequence_length is read per batch from the GPU buffer.
 fn get_total_sequence_length(batch_idx: u32) -> u32 {
-  return min(u32(max(seqlens_k[batch_idx], 0)) + 1u, uniforms.present_sequence_length);
+  return min(u32(max(seqlens_k.getByOffset(batch_idx), 0)) + 1u, uniforms.present_sequence_length);
 }
 #else
 // Without seqlens_k, total_sequence_length comes from uniforms (max across batches).
@@ -82,11 +84,11 @@ fn loadq(batch_idx : u32, q_idx_global : u32, head_idx : u32, alpha : q_element_
                head_idx * head_size_vec;
 #endif
   for (var idx : u32 = 0; idx < head_size_vec; idx++) {
-    q_tile[idx] = q[idx + offset] * alpha;
+    q_tile[idx] = q.getByOffset(idx + offset) * alpha;
   }
 }
 
-#if use_shm_path
+#if use_shm_path || qc_cooperative
 
 var<private> qk_scores : array<qk_precision, max_k_step>;
 
@@ -99,9 +101,10 @@ fn loadk(k_start : u32, batch_head_idx : u32, local_idx : u32, total_seq : u32)
     let seq_idx = k_start + slot;
     if (seq_idx < total_seq) {
       let base = kv_head_idx * uniforms.present_sequence_length * COMPRESSED_HEAD_U32 + seq_idx * COMPRESSED_HEAD_U32;
-      let scale = bitcast<f32>(present_key[base]);
+      let scale = bitcast<f32>(present_key.getByOffset(base));
       for (var v : u32 = 0u; v < head_size_vec; v++) {
-        k_tile[slot][v] = kv_cache_dequant_vec4(&present_key, base, v * 4u, scale);
+        let packed = present_key.getByOffset(base + 1u + v * 4u / KV_CACHE_QUANT_ELEMENTS_PER_WORD);
+        k_tile[slot][v] = kv_cache_dequant_vec4(packed, v * 4u, scale);
       }
     } else {
       for (var v : u32 = 0u; v < head_size_vec; v++) {
@@ -114,7 +117,7 @@ fn loadk(k_start : u32, batch_head_idx : u32, local_idx : u32, total_seq : u32)
                k_start * head_size_vec;
   for (var idx : u32 = local_idx; idx < head_size_vec * max_k_step; idx += workgroup_size_x) {
     let slot = u32(idx / head_size_vec);
-    k_tile[slot][idx % head_size_vec] = select(q_value_t(0), present_key[offset + idx], k_start + slot < total_seq);
+    k_tile[slot][idx % head_size_vec] = select(q_value_t(0), present_key.getByOffset(offset + idx), k_start + slot < total_seq);
   }
 #endif
 }
@@ -127,9 +130,10 @@ fn loadv(v_start : u32, batch_head_idx : u32, local_idx : u32, total_seq : u32)
     let seq_idx = v_start + slot;
     if (seq_idx < total_seq) {
       let base = kv_head_idx * uniforms.present_sequence_length * COMPRESSED_HEAD_U32 + seq_idx * COMPRESSED_HEAD_U32;
-      let scale = bitcast<f32>(present_value[base]);
+      let scale = bitcast<f32>(present_value.getByOffset(base));
       for (var v : u32 = 0u; v < head_size_vec; v++) {
-        v_tile[slot][v] = kv_cache_dequant_vec4(&present_value, base, v * 4u, scale);
+        let packed = present_value.getByOffset(base + 1u + v * 4u / KV_CACHE_QUANT_ELEMENTS_PER_WORD);
+        v_tile[slot][v] = kv_cache_dequant_vec4(packed, v * 4u, scale);
       }
     } else {
       for (var v : u32 = 0u; v < head_size_vec; v++) {
@@ -142,7 +146,7 @@ fn loadv(v_start : u32, batch_head_idx : u32, local_idx : u32, total_seq : u32)
                v_start * head_size_vec;
   for (var idx : u32 = local_idx; idx < head_size_vec * max_k_step; idx += workgroup_size_x) {
     let slot = u32(idx / head_size_vec);
-    v_tile[slot][idx % head_size_vec] = select(q_value_t(0), present_value[offset + idx], v_start + slot < total_seq);
+    v_tile[slot][idx % head_size_vec] = select(q_value_t(0), present_value.getByOffset(offset + idx), v_start + slot < total_seq);
   }
 #endif
 }
@@ -152,7 +156,7 @@ fn writeo(batch_idx : u32, o_idx_global : u32, head_idx : u32) {
   let offset = batch_idx * uniforms.new_sequence_length * num_heads * head_size_vec +
                o_idx_global * num_heads * head_size_vec + head_idx * head_size_vec;
   for (var idx : u32 = 0; idx < head_size_vec; idx++) {
-    output[offset + idx] = o_tile[idx];
+    output.setByOffset(offset + idx, o_tile[idx]);
   }
 }
 
@@ -169,7 +173,7 @@ fn loadAttentionBias(batch_idx : u32, q_idx_global : u32, k_idx_global : u32, he
   let stride_total_seq = uniforms.attn_bias_dim3;
   let offset_base = bias_batch_idx * uniforms.attn_bias_dim1 * uniforms.new_sequence_length * stride_total_seq +
                     bias_head_idx * uniforms.new_sequence_length * stride_total_seq + q_idx_global * stride_total_seq;
-  return qk_precision(attention_bias[min(offset_base + k_idx_global, offset_base + stride_total_seq - 1u)]);
+  return qk_precision(attention_bias.getByOffset(min(offset_base + k_idx_global, offset_base + stride_total_seq - 1u)));
 }
  #endif
 
@@ -185,9 +189,10 @@ fn loadk(k_start : u32, batch_head_idx : u32, local_idx : u32, k_step : u32, tot
     let seq_idx = k_start + slot;
     if (seq_idx < total_seq) {
       let base = kv_head_idx * uniforms.present_sequence_length * COMPRESSED_HEAD_U32 + seq_idx * COMPRESSED_HEAD_U32;
-      let scale = bitcast<f32>(present_key[base]);
+      let scale = bitcast<f32>(present_key.getByOffset(base));
       for (var v : u32 = 0u; v < head_size_vec; v++) {
-        k_tile[slot][v] = kv_cache_dequant_vec4(&present_key, base, v * 4u, scale);
+        let packed = present_key.getByOffset(base + 1u + v * 4u / KV_CACHE_QUANT_ELEMENTS_PER_WORD);
+        k_tile[slot][v] = kv_cache_dequant_vec4(packed, v * 4u, scale);
       }
     } else {
       for (var v : u32 = 0u; v < head_size_vec; v++) {
@@ -201,7 +206,7 @@ fn loadk(k_start : u32, batch_head_idx : u32, local_idx : u32, k_step : u32, tot
                k_start * head_size_vec;
   for (var idx : u32 = local_idx; idx < head_size_vec * k_step; idx += workgroup_size_x) {
     let slot = u32(idx / head_size_vec);
-    let val = select(q_value_t(0), present_key[offset + idx], k_start + slot < total_seq);
+    let val = select(q_value_t(0), present_key.getByOffset(offset + idx), k_start + slot < total_seq);
     k_tile[slot][idx % head_size_vec] = val;
   }
 #endif
@@ -215,9 +220,10 @@ fn loadv(v_start : u32, batch_head_idx : u32, local_idx : u32, v_step : u32, tot
     let seq_idx = v_start + slot;
     if (seq_idx < total_seq) {
       let base = kv_head_idx * uniforms.present_sequence_length * COMPRESSED_HEAD_U32 + seq_idx * COMPRESSED_HEAD_U32;
-      let scale = bitcast<f32>(present_value[base]);
+      let scale = bitcast<f32>(present_value.getByOffset(base));
       for (var v : u32 = 0u; v < head_size_vec; v++) {
-        v_tile[slot][v] = kv_cache_dequant_vec4(&present_value, base, v * 4u, scale);
+        let packed = present_value.getByOffset(base + 1u + v * 4u / KV_CACHE_QUANT_ELEMENTS_PER_WORD);
+        v_tile[slot][v] = kv_cache_dequant_vec4(packed, v * 4u, scale);
       }
     } else {
       for (var v : u32 = 0u; v < head_size_vec; v++) {
@@ -231,7 +237,7 @@ fn loadv(v_start : u32, batch_head_idx : u32, local_idx : u32, v_step : u32, tot
                v_start * head_size_vec;
   for (var idx : u32 = local_idx; idx < head_size_vec * v_step; idx += workgroup_size_x) {
     let slot = u32(idx / head_size_vec);
-    let val = select(q_value_t(0), present_value[offset + idx], v_start + slot < total_seq);
+    let val = select(q_value_t(0), present_value.getByOffset(offset + idx), v_start + slot < total_seq);
     v_tile[slot][idx % head_size_vec] = val;
   }
 #endif
@@ -248,8 +254,8 @@ fn writeo(batch_idx : u32, o_idx_global : u32, head_idx : u32, local_idx : u32)
   let offset = batch_idx * uniforms.new_sequence_length * num_heads * head_size_vec +
                o_idx_global * num_heads * head_size_vec + head_idx * head_size_vec;
   for (var idx : u32 = 0; idx < half_head_size_vec; idx++) {
-    output[offset + idx] = o_tile[idx];
-    output[offset + idx + half_head_size_vec] = o_tile_r[local_idx][idx];
+    output.setByOffset(offset + idx, o_tile[idx]);
+    output.setByOffset(offset + idx + half_head_size_vec, o_tile_r[local_idx][idx]);
   }
 }
  #else
@@ -258,7 +264,7 @@ fn writeo(batch_idx : u32, o_idx_global : u32, head_idx : u32) {
   let offset = batch_idx * uniforms.new_sequence_length * num_heads * head_size_vec +
                o_idx_global * num_heads * head_size_vec + head_idx * head_size_vec;
   for (var idx : u32 = 0; idx < head_size_vec; idx++) {
-    output[offset + idx] = o_tile[idx];
+    output.setByOffset(offset + idx, o_tile[idx]);
   }
 }
  #endif
@@ -280,10 +286,10 @@ fn loadAttentionBias(batch_idx : u32, q_idx_global : u32, k_idx_global : u32, he
                     bias_head_idx * uniforms.new_sequence_length * stride_total_seq + q_idx_global * stride_total_seq;
   let offset = offset_base + k_idx_global;
   let offset_max = offset_base + stride_total_seq - 1u;
-  let c1 = qk_precision(attention_bias[min(offset, offset_max)]);
-  let c2 = qk_precision(attention_bias[min(offset + 1, offset_max)]);
-  let c3 = qk_precision(attention_bias[min(offset + 2, offset_max)]);
-  let c4 = qk_precision(attention_bias[min(offset + 3, offset_max)]);
+  let c1 = qk_precision(attention_bias.getByOffset(min(offset, offset_max)));
+  let c2 = qk_precision(attention_bias.getByOffset(min(offset + 1, offset_max)));
+  let c3 = qk_precision(attention_bias.getByOffset(min(offset + 2, offset_max)));
+  let c4 = qk_precision(attention_bias.getByOffset(min(offset + 3, offset_max)));
   return vec4<qk_precision>(c1, c2, c3, c4);
 }
  #else
@@ -311,6 +317,121 @@ fn fetchVTile(k_idx: u32, vec_idx: u32, v_val: q_value_t) -> q_value_t {
 #endif
 
 $MAIN {
+#if qc_cooperative
+  // Four adjacent lanes own disjoint dimensions of one query and output row.
+  let bh = workgroup_idx / uniforms.num_seq_tile;
+  let batch = bh / num_heads;
+  if (batch >= uniforms.batch_size) {
+    return;
+  }
+  let head = bh % num_heads;
+  let rows_per_group = workgroup_size_x / 4u;
+  let row_start = (workgroup_idx % uniforms.num_seq_tile) * rows_per_group;
+  let row = row_start + subgroup_index * (sg_size / 4u) + sg_id / 4u;
+  let d_lane = sg_id % 4u;
+  let valid = row < uniforms.new_sequence_length;
+  let total = get_total_sequence_length(batch);
+#if use_seqlens_q
+  let past = u32(max(i32(total) - seqlens_q.getByOffset(batch), 0));
+#else
+  let past = u32(max(i32(total) - i32(uniforms.new_sequence_length), 0));
+#endif
+#if is_unidirectional
+  let end = min(total, past + row_start + rows_per_group);
+  let causal_end = min(total, past + row + 1u);
+#else
+  let end = total;
+  let causal_end = total;
+#endif
+#if has_local_window
+  let first_query_end = past + row_start + 1u;
+  let query_end = past + row + 1u;
+  let first_window_start = first_query_end - min(first_query_end, uniforms.local_window_size);
+  let window_start = query_end - min(query_end, uniforms.local_window_size);
+#else
+  let first_window_start = 0u;
+  let window_start = 0u;
+#endif
+  var qv : array<q_value_t, 8>;
+  var ov : array<vec4<f32>, 8>;
+#if q_BNSH
+  let q_base = (bh * uniforms.new_sequence_length + row) * 32u;
+#else
+  let q_base = ((batch * uniforms.new_sequence_length + row) * num_heads + head) * 32u;
+#endif
+  if (valid) {
+    for (var d = 0u; d < 8u; d++) {
+      qv[d] = q.getByOffset(q_base + d * 4u + d_lane);
+    }
+  }
+  // Accumulate the unnormalized output in f32; normalize once at the store.
+#if has_head_sink
+  var previous_max = f32(head_sink.getByOffset(head));
+  var denominator = 1.0;
+#else
+  var previous_max = qk_min_value;
+  var denominator = 0.0;
+#endif
+  for (var start = first_window_start / 16u * 16u; start < end; start += 16u) {
+    for (var idx = local_idx; idx < 512u; idx += workgroup_size_x) {
+      let key_row = start + idx / 32u;
+      let offset = (bh / uniforms.n_reps * uniforms.present_sequence_length + key_row) * 32u + idx % 32u;
+      var kval = q_value_t(0);
+      var vval = q_value_t(0);
+      if (key_row < total) {
+        kval = present_key.getByOffset(offset);
+        vval = present_value.getByOffset(offset);
+      }
+      k_tile[idx / 32u][idx % 32u] = kval;
+      v_tile[idx / 32u][idx % 32u] = vval;
+    }
+    workgroupBarrier();
+    var scores : array<f32, 16>;
+    var block_max = previous_max;
+    for (var key = 0u; key < 16u; key++) {
+      var score = 0.0;
+      for (var d = 0u; d < 8u; d++) {
+        score += dot(vec4<f32>(qv[d]), vec4<f32>(k_tile[key][d * 4u + d_lane]));
+      }
+      score += subgroupShuffleXor(score, 1u);
+      score += subgroupShuffleXor(score, 2u);
+      score *= uniforms.alpha;
+#if has_softcap
+      score = tanh(score / uniforms.softcap) * uniforms.softcap;
+#endif
+#if has_attention_bias
+      if (valid) {
+        score += loadAttentionBias(batch, row, start + key, head, total);
+      }
+#endif
+      score = select(qk_min_value, score, is_key_visible(start + key, window_start, causal_end));
+      scores[key] = score;
+      block_max = max(block_max, score);
+    }
+    let correction = exp(previous_max - block_max);
+    denominator *= correction;
+    for (var key = 0u; key < 16u; key++) {
+      scores[key] = select(0.0, exp(scores[key] - block_max), is_key_visible(start + key, window_start, causal_end));
+      denominator += scores[key];
+    }
+    for (var d = 0u; d < 8u; d++) {
+      var acc = ov[d] * correction;
+      for (var key = 0u; key < 16u; key++) {
+        acc += scores[key] * vec4<f32>(v_tile[key][d * 4u + d_lane]);
+      }
+      ov[d] = acc;
+    }
+    previous_max = block_max;
+    workgroupBarrier();
+  }
+  if (valid) {
+    let base = ((batch * uniforms.new_sequence_length + row) * num_heads + head) * 32u;
+    for (var d = 0u; d < 8u; d++) {
+      output.setByOffset(base + d * 4u + d_lane, q_value_t(ov[d] / max(denominator, 1e-20)));
+    }
+  }
+
+#else
   let batch_head_idx = u32(workgroup_idx / uniforms.num_seq_tile);
   let head_idx = batch_head_idx % num_heads;
   let batch_idx = batch_head_idx / num_heads;
@@ -335,7 +456,7 @@ $MAIN {
   }
 
 #if has_head_sink
-  let sink_value = q_element_t(head_sink[head_idx]);
+  let sink_value = q_element_t(head_sink.getByOffset(head_idx));
   var previous_max : qk_precision = qk_precision(sink_value);
   var previous_denom : qk_precision = 1;
 #else
@@ -346,7 +467,7 @@ $MAIN {
 
 #if use_seqlens_q
   // LEFT-aligned per-batch Q length: past_len_b = total_kv_b - q_len_b (>= 0).
-  let q_len_b = u32(seqlens_q[batch_idx]);
+  let q_len_b = u32(seqlens_q.getByOffset(batch_idx));
   let past_sequence_length = total_sequence_length - q_len_b;
 #else
   // Uniform q_len path. total_kv < new_sequence_length can only happen for
@@ -715,4 +836,5 @@ $MAIN {
   }
  #endif
 #endif
+#endif
 }  // MAIN
diff --git a/onnxruntime/test/contrib_ops/group_query_attention_op_test.cc b/onnxruntime/test/contrib_ops/group_query_attention_op_test.cc
index 0cdb4a1f82e10..606b96422cd26 100644
--- a/onnxruntime/test/contrib_ops/group_query_attention_op_test.cc
+++ b/onnxruntime/test/contrib_ops/group_query_attention_op_test.cc
@@ -43,6 +43,7 @@
 #include "core/framework/session_options.h"
 #include "core/providers/webgpu/webgpu_provider_options.h"
 #include "core/session/onnxruntime_session_options_config_keys.h"
+#include "test/contrib_ops/group_query_attention_op_test_helper.h"
 #endif
 
 namespace onnxruntime {
@@ -6300,24 +6301,6 @@ static std::unique_ptr<IExecutionProvider> WebGpuEPWithKVCacheQuantization(
                                .kv_cache_quant_bits = bit_width});
 }
 
-static std::vector<float> RunGQAReference(
-    int batch_size,
-    int sequence_length,
-    int num_heads,
-    int kv_num_heads,
-    int head_size,
-    const std::vector<float>& query_data,
-    const std::vector<float>& key_data,
-    const std::vector<float>& value_data,
-    bool do_rotary,
-    bool use_fp16 = false,
-    bool rotary_interleaved = false);
-
-static void ExpectBlockQuantInt8Close(const std::vector<float>& reference,
-                                      const std::vector<float>& actual,
-                                      float max_relative_rmse,
-                                      float max_absolute_error);
-
 // Graph capture requires the indirect-dispatch dimensions to be prepared on the GPU.
 // Verify that static-cache preprocessing uses the batch-wide total_sequence_length input
 // instead of deriving the dispatch width from batch 0's (possibly shorter) seqlens_k value. The
@@ -7662,7 +7645,7 @@ TEST(GroupQueryAttentionTest, WebGPU_TurboQuant_Prefill_MultiBatch_RightPadding_
 // ---------------------------------------------------------------------------
 
 // Helper: runs GQA without TurboQuant (standard uncompressed KV cache) and returns the output.
-static std::vector<float> RunGQAReference(
+std::vector<float> RunGQAReference(
     int batch_size,
     int sequence_length,
     int num_heads,
@@ -7763,9 +7746,9 @@ static std::vector<float> RunGQAReference(
   tester.SetOutputTolerance(1e6f);
   tester.SetCustomOutputVerifier([](const std::vector<OrtValue>&, const std::string&) {});
 
-  std::vector<std::unique_ptr<IExecutionProvider>> execution_providers;
-  execution_providers.push_back(DefaultWebGpuExecutionProvider());
-  tester.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers);
+  SessionOptions options;
+  ORT_THROW_IF_ERROR(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1"));
+  tester.Config(options).ConfigEp(DefaultWebGpuExecutionProvider()).RunWithConfig();
 
   auto fetches = tester.GetFetches();
   if (use_fp16) {
@@ -8098,10 +8081,10 @@ TEST(GroupQueryAttentionTest, WebGPU_TurboQuant_CrossValidate_Prefill8_MultiKVHe
   EXPECT_LT(max_abs_err, 0.3f) << "TurboQuant 8-token multi-head max absolute error too large: " << max_abs_err;
 }
 
-static void ExpectBlockQuantInt8Close(const std::vector<float>& reference,
-                                      const std::vector<float>& actual,
-                                      float max_relative_rmse,
-                                      float max_absolute_error) {
+void ExpectBlockQuantInt8Close(const std::vector<float>& reference,
+                               const std::vector<float>& actual,
+                               float max_relative_rmse,
+                               float max_absolute_error) {
   ASSERT_EQ(reference.size(), actual.size());
   float max_abs_err = 0.0f;
   float sum_sq_err = 0.0f;
@@ -8180,7 +8163,7 @@ TEST(GroupQueryAttentionTest, WebGPU_BlockQuantInt8_CrossValidate_FlashPrefill)
                             /*max_absolute_error=*/0.03f);
 }
 
-static void RunFp16HighMagnitudeAttention(int sequence_length, uint32_t bit_width, int head_size = 128) {
+void RunFp16HighMagnitudeAttention(int sequence_length, uint32_t bit_width, int head_size) {
   auto ep = bit_width == 0 ? DefaultWebGpuExecutionProvider() : WebGpuEPWithKVCacheQuantization(bit_width);
   if (!ep) {
     GTEST_SKIP() << "WebGPU EP not available";
diff --git a/onnxruntime/test/contrib_ops/group_query_attention_op_test_helper.h b/onnxruntime/test/contrib_ops/group_query_attention_op_test_helper.h
new file mode 100644
index 0000000000000..d0205ddd80eb6
--- /dev/null
+++ b/onnxruntime/test/contrib_ops/group_query_attention_op_test_helper.h
@@ -0,0 +1,31 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include <cstdint>
+#include <vector>
+
+namespace onnxruntime::test {
+
+std::vector<float> RunGQAReference(
+    int batch_size,
+    int sequence_length,
+    int num_heads,
+    int kv_num_heads,
+    int head_size,
+    const std::vector<float>& query_data,
+    const std::vector<float>& key_data,
+    const std::vector<float>& value_data,
+    bool do_rotary,
+    bool use_fp16 = false,
+    bool rotary_interleaved = false);
+
+void ExpectBlockQuantInt8Close(const std::vector<float>& reference,
+                               const std::vector<float>& actual,
+                               float max_relative_rmse,
+                               float max_absolute_error);
+
+void RunFp16HighMagnitudeAttention(int sequence_length, uint32_t bit_width, int head_size = 128);
+
+}  // namespace onnxruntime::test
diff --git a/onnxruntime/test/contrib_ops/paged_attention_op_test.cc b/onnxruntime/test/contrib_ops/paged_attention_op_test.cc
index c5fba3d7b88e0..e15eeea38b083 100644
--- a/onnxruntime/test/contrib_ops/paged_attention_op_test.cc
+++ b/onnxruntime/test/contrib_ops/paged_attention_op_test.cc
@@ -29,6 +29,7 @@
 #include "default_providers.h"
 #include "test/common/cuda_op_test_utils.h"
 #include "test/common/tensor_op_test_utils.h"
+#include "test/contrib_ops/paged_attention_op_test_helper.h"
 #include "test/providers/provider_test_utils.h"
 #include "test/unittest_util/framework_test_utils.h"
 #include "test/util/include/scoped_env_vars.h"
@@ -49,22 +50,6 @@ int CacheIndex(int block_id, int slot_in_block, int kv_head, int dim,
 
 namespace {
 
-struct EndToEndCase {
-  int batch_size = 1;
-  int token_count = 1;
-  int num_heads = 1;
-  int kv_num_heads = 1;
-  int head_size = 8;
-  int block_size = 256;
-  int num_blocks = 2;
-  int max_num_blocks_per_seq = 1;
-  float scale = 0.0f;  // 0 means kernel default = 1/sqrt(head_size)
-  bool is_causal = true;
-  std::vector<int32_t> cumulative_seqlens_q;
-  std::vector<int32_t> past_seqlens;
-  std::vector<int32_t> block_table;
-};
-
 struct IoBindingCase {
   int batch_size = 1;
   int token_count = 0;  // 0 means one token per batch entry
@@ -106,15 +91,15 @@ struct IoBindingCase {
 
 // Masked positions get zero probability. Uses fp32 throughout to establish a
 // reference the fp16 kernel is compared against with a loose tolerance.
-void MaskedSoftmax(std::vector<float>& scores, int valid_key_end) {
+void MaskedSoftmax(std::vector<float>& scores, int valid_key_end, int valid_key_start) {
   const int len = static_cast<int>(scores.size());
   float max_val = -std::numeric_limits<float>::infinity();
-  for (int i = 0; i < valid_key_end && i < len; ++i) {
+  for (int i = valid_key_start; i < valid_key_end && i < len; ++i) {
     max_val = std::max(max_val, scores[i]);
   }
   float sum = 0.0f;
   for (int i = 0; i < len; ++i) {
-    if (i >= valid_key_end) {
+    if (i < valid_key_start || i >= valid_key_end) {
       scores[i] = 0.0f;
     } else {
       scores[i] = std::exp(scores[i] - max_val);
@@ -127,7 +112,10 @@ void MaskedSoftmax(std::vector<float>& scores, int valid_key_end) {
   }
 }
 
-void RunEndToEndCase(const EndToEndCase& c, std::unique_ptr<IExecutionProvider> execution_provider) {
+}  // namespace
+
+void RunPagedAttentionTest(const PagedAttentionTestCase& c, std::unique_ptr<IExecutionProvider> execution_provider,
+                           float output_abs_error) {
   ASSERT_EQ(c.cumulative_seqlens_q.size(), static_cast<size_t>(c.batch_size + 1));
   ASSERT_EQ(c.past_seqlens.size(), static_cast<size_t>(c.batch_size));
   ASSERT_EQ(c.block_table.size(), static_cast<size_t>(c.batch_size * c.max_num_blocks_per_seq));
@@ -223,6 +211,7 @@ void RunEndToEndCase(const EndToEndCase& c, std::unique_ptr<IExecutionProvider>
       const int t = cum_lo + local_tok;
       const int q_pos = past + local_tok;  // Absolute position in the KV window.
       const int valid_key_end = c.is_causal ? q_pos + 1 : total_kv_len;
+      const int valid_key_start = c.local_window_size > 0 ? std::max(0, q_pos + 1 - c.local_window_size) : 0;
       for (int n_q = 0; n_q < c.num_heads; ++n_q) {
         const int h_kv = n_q / gqa_factor;
 
@@ -237,7 +226,7 @@ void RunEndToEndCase(const EndToEndCase& c, std::unique_ptr<IExecutionProvider>
           }
           scores[s] = dot * scale;
         }
-        MaskedSoftmax(scores, valid_key_end);
+        MaskedSoftmax(scores, valid_key_end, valid_key_start);
 
         // out[t, n_q * head_size + d] = sum_s scores[s] * v[s, d]
         for (int d = 0; d < c.head_size; ++d) {
@@ -258,6 +247,9 @@ void RunEndToEndCase(const EndToEndCase& c, std::unique_ptr<IExecutionProvider>
   test.AddAttribute<float>("scale", c.scale);
   test.AddAttribute<int64_t>("is_causal", c.is_causal ? 1 : 0);
   test.AddAttribute<int64_t>("do_rotary", 0);
+  if (c.local_window_size > 0) {
+    test.AddAttribute<int64_t>("local_window_size", c.local_window_size);
+  }
 
   test.AddInput<MLFloat16>("query", {c.token_count, hidden_size}, FloatsToMLFloat16s(query_f));
   test.AddInput<MLFloat16>("key", {c.token_count, kv_hidden_size}, FloatsToMLFloat16s(key_f));
@@ -285,13 +277,15 @@ void RunEndToEndCase(const EndToEndCase& c, std::unique_ptr<IExecutionProvider>
   // fp16 softmax + dot-product accumulation over a short KV window.
   // Reference is fp32; kernel is fp16. 2e-2 is a comfortable envelope for
   // these value ranges and short sequences.
-  test.SetOutputAbsErr("output", 2e-2f);
+  test.SetOutputAbsErr("output", output_abs_error);
   test.SetOutputAbsErr("key_cache_out", 1e-3f);
   test.SetOutputAbsErr("value_cache_out", 1e-3f);
 
   test.ConfigEp(std::move(execution_provider)).RunWithConfig();
 }
 
+namespace {
+
 void RunIoBindingCase(std::unique_ptr<IExecutionProvider> execution_provider,
                       const char* provider_type,
                       bool alias_cache_outputs,
@@ -1093,7 +1087,7 @@ void RunIoBindingCase(std::unique_ptr<IExecutionProvider> execution_provider,
   }
 }
 
-void RunEndToEndCaseOnAvailableProviders(const EndToEndCase& c) {
+void RunEndToEndCaseOnAvailableProviders(const PagedAttentionTestCase& c) {
   bool ran = false;
 
   // CUDA PagedAttention requires key_cache/value_cache outputs to alias the
@@ -1103,7 +1097,7 @@ void RunEndToEndCaseOnAvailableProviders(const EndToEndCase& c) {
   // output buffers are explicitly supported by the kernel fallback path.
 
   if (auto webgpu_ep = DefaultWebGpuExecutionProvider()) {
-    RunEndToEndCase(c, std::move(webgpu_ep));
+    RunPagedAttentionTest(c, std::move(webgpu_ep));
     ran = true;
   }
 
@@ -3186,7 +3180,7 @@ TEST(PagedAttention, WebGpu_RejectsShortRotaryCache) {
 // Decode tier: single batch, single new Q token, non-zero past. The FA
 // tier selector uses `sequence_length_ < 32` → split-reduce path.
 TEST(PagedAttention, EndToEnd_Decode_SingleBatch_WithPast) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 1;
   c.cumulative_seqlens_q = {0, 1};
@@ -3198,7 +3192,7 @@ TEST(PagedAttention, EndToEnd_Decode_SingleBatch_WithPast) {
 // Prefill tier: single batch, multiple new Q tokens, zero past. Exercises
 // causal masking across new tokens.
 TEST(PagedAttention, EndToEnd_Prefill_SingleBatch_NoPast) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 4;
   c.cumulative_seqlens_q = {0, 4};
@@ -3208,7 +3202,7 @@ TEST(PagedAttention, EndToEnd_Prefill_SingleBatch_NoPast) {
 }
 
 TEST(PagedAttention, EndToEnd_Prefill_PagedFlashAttention) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 32;
   c.cumulative_seqlens_q = {0, 32};
@@ -3220,7 +3214,7 @@ TEST(PagedAttention, EndToEnd_Prefill_PagedFlashAttention) {
 // Multi-batch decode with differing past lengths — exercises variable-length
 // packing across batches.
 TEST(PagedAttention, EndToEnd_Decode_MultiBatch_VariablePast) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 2;
   c.token_count = 2;  // one new Q per batch
   c.num_blocks = 3;
@@ -3233,7 +3227,7 @@ TEST(PagedAttention, EndToEnd_Decode_MultiBatch_VariablePast) {
 // GQA (num_heads > kv_num_heads): broadcasts each KV head across gqa_factor
 // query heads. Verifies the head-index mapping matches FA's convention.
 TEST(PagedAttention, EndToEnd_Decode_GQA) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 1;
   c.num_heads = 2;
@@ -3247,7 +3241,7 @@ TEST(PagedAttention, EndToEnd_Decode_GQA) {
 // Empty-token fast path: the output is empty and cache outputs must still
 // preserve the input cache contents in OpTester's non-aliased allocation path.
 TEST(PagedAttention, EndToEnd_EmptyTokens_CopyCaches) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 0;
   c.cumulative_seqlens_q = {0, 0};
@@ -3259,7 +3253,7 @@ TEST(PagedAttention, EndToEnd_EmptyTokens_CopyCaches) {
 // Mixed variable-length prefill (seq 0: 3 tokens, seq 1: 2 tokens) with
 // non-zero past on both — the most realistic paged-attention scenario.
 TEST(PagedAttention, EndToEnd_MixedPrefillDecode_MultiBatch_VariablePast) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 2;
   c.token_count = 5;
   c.num_heads = 2;
@@ -3276,7 +3270,7 @@ TEST(PagedAttention, EndToEnd_MixedPrefillDecode_MultiBatch_VariablePast) {
 // q_varlen=1 code path (raw-packed Q read via cumulative_seqlens_q). Uses
 // head_size = 128 to match the shared-memory tile the shader is sized for.
 TEST(PagedAttention, EndToEnd_Prefill_MultiBatch_Varlen_Fused) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 2;
   c.token_count = 48;  // seq 0: 32 tokens, seq 1: 16 tokens
   c.num_heads = 2;
@@ -3297,7 +3291,7 @@ TEST(PagedAttention, EndToEnd_Prefill_MultiBatch_Varlen_Fused) {
 // adjacency of consecutive logical pages will fail here (see the alignment
 // invariant note at the top of flash_attention_paged_prefill.wgsl.template).
 TEST(PagedAttention, EndToEnd_Prefill_FusedPrefill_BlockBoundaryCrossing) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 64;  // >= 32 for fused prefill; spans 2 blocks of 32
   c.num_heads = 2;
@@ -3318,7 +3312,7 @@ TEST(PagedAttention, EndToEnd_Prefill_FusedPrefill_BlockBoundaryCrossing) {
 // any indexing bug in the per-slot block_table lookup inside
 // flash_attention_paged_decode_qkv.wgsl.template.
 TEST(PagedAttention, EndToEnd_Decode_BlockBoundaryCrossing) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 1;  // decode: one new Q token
   c.num_heads = 1;
@@ -3340,7 +3334,7 @@ TEST(PagedAttention, EndToEnd_Decode_BlockBoundaryCrossing) {
 // the fallback path, locking down the reject-boundary so a future refactor
 // of CanApplyFlashAttentionPagedPrefill can't silently change it.
 TEST(PagedAttention, EndToEnd_Prefill_ForcedFallback_BlockSizeBelowMaxKStep) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 64;  // prefill (>= 32)
   c.num_heads = 2;
@@ -3363,7 +3357,7 @@ TEST(PagedAttention, EndToEnd_Prefill_ForcedFallback_BlockSizeBelowMaxKStep) {
 // head_idx / uniforms.n_reps mapping and the causal-mask past-offset
 // derivation on the same shape.
 TEST(PagedAttention, EndToEnd_Prefill_FusedPrefill_GQA_WithPast) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 1;
   c.token_count = 32;  // prefill: max_seqlen_q >= 32 fires fused shader
   c.num_heads = 4;
@@ -3379,7 +3373,7 @@ TEST(PagedAttention, EndToEnd_Prefill_FusedPrefill_GQA_WithPast) {
 }
 
 TEST(PagedAttention, WebGpu_NonCausal_DirectPagedDecode) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.token_count = 4;
   c.is_causal = false;
   c.cumulative_seqlens_q = {0, 4};
@@ -3389,7 +3383,7 @@ TEST(PagedAttention, WebGpu_NonCausal_DirectPagedDecode) {
 }
 
 TEST(PagedAttention, WebGpu_NonCausal_FusedPagedPrefill) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.token_count = 32;
   c.num_heads = 2;
   c.kv_num_heads = 2;
@@ -3402,7 +3396,7 @@ TEST(PagedAttention, WebGpu_NonCausal_FusedPagedPrefill) {
 }
 
 TEST(PagedAttention, WebGpu_NonCausal_ForcedFallback) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.token_count = 64;
   c.num_heads = 2;
   c.kv_num_heads = 2;
@@ -3418,7 +3412,7 @@ TEST(PagedAttention, WebGpu_NonCausal_ForcedFallback) {
 }
 
 TEST(PagedAttention, WebGpu_NonCausal_RaggedSequences) {
-  EndToEndCase c{};
+  PagedAttentionTestCase c{};
   c.batch_size = 2;
   c.token_count = 4;
   c.num_heads = 2;
diff --git a/onnxruntime/test/contrib_ops/paged_attention_op_test_helper.h b/onnxruntime/test/contrib_ops/paged_attention_op_test_helper.h
new file mode 100644
index 0000000000000..ea0793cfc006d
--- /dev/null
+++ b/onnxruntime/test/contrib_ops/paged_attention_op_test_helper.h
@@ -0,0 +1,37 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include <cstdint>
+#include <memory>
+#include <vector>
+
+namespace onnxruntime {
+class IExecutionProvider;
+
+namespace test {
+
+struct PagedAttentionTestCase {
+  int batch_size = 1;
+  int token_count = 1;
+  int num_heads = 1;
+  int kv_num_heads = 1;
+  int head_size = 8;
+  int block_size = 256;
+  int num_blocks = 2;
+  int max_num_blocks_per_seq = 1;
+  float scale = 0.0f;  // 0 means kernel default = 1/sqrt(head_size)
+  bool is_causal = true;
+  int local_window_size = -1;
+  std::vector<int32_t> cumulative_seqlens_q;
+  std::vector<int32_t> past_seqlens;
+  std::vector<int32_t> block_table;
+};
+
+void RunPagedAttentionTest(const PagedAttentionTestCase& c,
+                           std::unique_ptr<IExecutionProvider> execution_provider,
+                           float output_abs_error = 2e-2f);
+
+}  // namespace test
+}  // namespace onnxruntime
diff --git a/onnxruntime/test/providers/webgpu/attention_test_utils.h b/onnxruntime/test/providers/webgpu/attention_test_utils.h
new file mode 100644
index 0000000000000..6cc7e5b16c609
--- /dev/null
+++ b/onnxruntime/test/providers/webgpu/attention_test_utils.h
@@ -0,0 +1,101 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include <chrono>
+#include <string>
+#include <string_view>
+
+#include "gtest/gtest.h"
+
+#include "test/util/include/default_providers.h"
+
+#if defined(USE_WEBGPU) && !defined(ORT_USE_EP_API_ADAPTERS)
+#include "core/providers/webgpu/webgpu_context.h"
+#endif
+
+namespace onnxruntime::test {
+
+inline bool WebGpuSupportsFp16() {
+#if defined(USE_WEBGPU) && !defined(ORT_USE_EP_API_ADAPTERS)
+  auto ep = DefaultWebGpuExecutionProvider();
+  return ep && webgpu::WebGpuContextFactory::GetContext(ep->GetDeviceId())
+                   .DeviceHasFeature(wgpu::FeatureName::ShaderF16);
+#else
+  // The plugin EP interface cannot report Dawn device features.
+  return false;
+#endif
+}
+
+inline bool SupportsQualcommCooperativePrefill() {
+#if defined(USE_WEBGPU) && !defined(ORT_USE_EP_API_ADAPTERS)
+  auto ep = DefaultWebGpuExecutionProvider();
+  if (!ep) {
+    return false;
+  }
+  const auto& context = webgpu::WebGpuContextFactory::GetContext(ep->GetDeviceId());
+  return context.AdapterInfo().vendor == std::string_view{"qualcomm"} &&
+         context.DeviceHasFeature(wgpu::FeatureName::ShaderF16) &&
+         context.DeviceHasFeature(wgpu::FeatureName::Subgroups) &&
+         context.DeviceHasFeature(wgpu::FeatureName::SubgroupSizeControl) &&
+         context.Instance().HasWGSLLanguageFeature(wgpu::WGSLLanguageFeatureName::SubgroupId) &&
+         context.AdapterInfo().subgroupMinSize <= 128 && context.AdapterInfo().subgroupMaxSize >= 128;
+#else
+  return false;
+#endif
+}
+
+inline bool SupportsCooperativeDispatchObservation() {
+#if defined(USE_WEBGPU) && !defined(ORT_USE_EP_API_ADAPTERS)
+  if (SupportsQualcommCooperativePrefill()) {
+    auto ep = DefaultWebGpuExecutionProvider();
+    return webgpu::WebGpuContextFactory::GetContext(ep->GetDeviceId())
+        .DeviceHasFeature(wgpu::FeatureName::TimestampQuery);
+  }
+#endif
+  return false;
+}
+
+#if defined(USE_WEBGPU) && !defined(ORT_USE_EP_API_ADAPTERS)
+inline void ExpectCooperativeDispatch(const profiling::Events& events) {
+  int attention_dispatches = 0;
+  for (const auto& event : events) {
+    if (!std::string_view{event.name}.ends_with("&FlashAttention")) {
+      continue;
+    }
+    ++attention_dispatches;
+    const auto key = event.args.find("cache_key");
+    ASSERT_NE(key, event.args.end());
+    EXPECT_EQ(key->second.find("FlashAttention[SubgroupId]"), 0u) << key->second;
+    // Cache keys include labels in Debug builds.
+    EXPECT_TRUE(key->second.find(":256,1,1:128:") != std::string::npos ||
+                key->second.find(":WorkgroupSize=256,1,1:SubgroupSize=128:") != std::string::npos)
+        << key->second;
+  }
+  EXPECT_GT(attention_dispatches, 0) << "No cooperative attention dispatch was observed";
+}
+#endif
+
+template <typename Run>
+void RunWithCooperativeDispatchCheck(bool verify_dispatch, Run&& run) {
+  if (!verify_dispatch) {
+    run();
+    return;
+  }
+  ASSERT_TRUE(SupportsCooperativeDispatchObservation());
+#if defined(USE_WEBGPU) && !defined(ORT_USE_EP_API_ADAPTERS)
+  auto ep = DefaultWebGpuExecutionProvider();
+  auto& context = webgpu::WebGpuContextFactory::GetContext(ep->GetDeviceId());
+  profiling::Events events;
+  context.SetProfilingStartTime(std::chrono::high_resolution_clock::now());
+  context.StartProfiling();
+  {
+    auto stop_profiling = gsl::finally([&] { context.CollectProfilingData(events); });
+    run();
+  }
+  ExpectCooperativeDispatch(events);
+#endif
+}
+
+}  // namespace onnxruntime::test
diff --git a/onnxruntime/test/providers/webgpu/contrib_ops/group_query_attention_test.cc b/onnxruntime/test/providers/webgpu/contrib_ops/group_query_attention_test.cc
new file mode 100644
index 0000000000000..8f0883fb560b7
--- /dev/null
+++ b/onnxruntime/test/providers/webgpu/contrib_ops/group_query_attention_test.cc
@@ -0,0 +1,579 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#if !defined(ORT_MINIMAL_BUILD) && !defined(DISABLE_CONTRIB_OPS)
+
+#include <algorithm>
+#include <array>
+#include <chrono>
+#include <cmath>
+#include <limits>
+#include <string>
+#include <string_view>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+#include "gtest/gtest.h"
+
+#include "core/common/inlined_containers.h"
+#include "core/framework/allocator.h"
+#include "core/framework/tensor.h"
+#include "core/providers/webgpu/webgpu_execution_provider.h"
+#include "core/providers/webgpu/webgpu_provider_options.h"
+#include "core/session/IOBinding.h"
+#include "core/session/inference_session.h"
+#include "core/session/onnxruntime_session_options_config_keys.h"
+#include "test/common/tensor_op_test_utils.h"
+#include "test/contrib_ops/group_query_attention_op_test_helper.h"
+#include "test/providers/provider_test_utils.h"
+#include "test/providers/webgpu/attention_test_utils.h"
+#include "test/test_environment.h"
+#include "test/util/include/asserts.h"
+
+namespace onnxruntime::test {
+namespace {
+
+constexpr int kBatchSize = 1;
+constexpr int kSequenceLength = 129;
+constexpr int kNumHeads = 2;
+constexpr int kHeadSize = 128;
+constexpr int kCapacity = 208;
+
+InlinedVector<MLFloat16> MakePrefillData(size_t size, int seed) {
+  InlinedVector<MLFloat16> data(size);
+  for (size_t i = 0; i < size; ++i) {
+    const int value = static_cast<int>((i * seed + i / kHeadSize) % 29) - 14;
+    data[i] = MLFloat16(static_cast<float>(value) * 0.0625f);
+  }
+  return data;
+}
+
+InlinedVector<MLFloat16> PrefillReference(gsl::span<const MLFloat16> query,
+                                          gsl::span<const MLFloat16> key_cache,
+                                          gsl::span<const MLFloat16> value_cache,
+                                          const std::array<int, kBatchSize>& past_lengths,
+                                          int window, int capacity) {
+  InlinedVector<MLFloat16> expected(query.size());
+  for (int batch = 0; batch < kBatchSize; ++batch) {
+    for (int row = 0; row < kSequenceLength; ++row) {
+      const int end = past_lengths[batch] + row + 1;
+      const int start = window > 0 ? std::max(0, end - window) : 0;
+      for (int head = 0; head < kNumHeads; ++head) {
+        const int q_offset = ((batch * kSequenceLength + row) * kNumHeads + head) * kHeadSize;
+        InlinedVector<double> scores(end);
+        double maximum = -std::numeric_limits<double>::infinity();
+        for (int key = start; key < end; ++key) {
+          double score = 0.0;
+          for (int dim = 0; dim < kHeadSize; ++dim) {
+            score += static_cast<double>(query[q_offset + dim].ToFloat()) *
+                     key_cache[(batch * capacity + key) * kHeadSize + dim].ToFloat();
+          }
+          scores[key] = score / std::sqrt(static_cast<double>(kHeadSize));
+          maximum = std::max(maximum, scores[key]);
+        }
+        double denominator = 0.0;
+        for (int key = start; key < end; ++key) {
+          scores[key] = std::exp(scores[key] - maximum);
+          denominator += scores[key];
+        }
+        for (int dim = 0; dim < kHeadSize; ++dim) {
+          double sum = 0.0;
+          for (int key = start; key < end; ++key) {
+            sum += scores[key] * value_cache[(batch * capacity + key) * kHeadSize + dim].ToFloat();
+          }
+          expected[q_offset + dim] = MLFloat16(static_cast<float>(sum / denominator));
+        }
+      }
+    }
+  }
+  return expected;
+}
+
+class CooperativePrefillTester final : public OpTester {
+ public:
+  CooperativePrefillTester() : OpTester("GroupQueryAttention", 1, kMSDomain) {}
+  ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CooperativePrefillTester);
+
+  void RunPrefill(bool capture, int window, uint64_t binding_limit = 0) {
+    const int capacity = binding_limit > 0 ? 512 : kCapacity;
+    const int first_past = binding_limit > 0 ? 193 : 17;
+    SetTestFunctionCalled();
+    const auto query = MakePrefillData(kBatchSize * kSequenceLength * kNumHeads * kHeadSize, 7);
+    const auto key = MakePrefillData(kBatchSize * kSequenceLength * kHeadSize, 11);
+    const auto value = MakePrefillData(key.size(), 13);
+    const auto initial_key = MakePrefillData(kBatchSize * capacity * kHeadSize, 17);
+    const auto initial_value = MakePrefillData(initial_key.size(), 19);
+    const TensorShape output_shape{kBatchSize, kSequenceLength, kNumHeads * kHeadSize};
+    const TensorShape cache_shape{kBatchSize, 1, capacity, kHeadSize};
+    AddAttribute<int64_t>("num_heads", kNumHeads);
+    AddAttribute<int64_t>("kv_num_heads", 1);
+    AddAttribute<int64_t>("local_window_size", window);
+    AddInput<MLFloat16>("query", {kBatchSize, kSequenceLength, kNumHeads * kHeadSize}, query.data(), query.size());
+    AddInput<MLFloat16>("key", {kBatchSize, kSequenceLength, kHeadSize}, key.data(), key.size());
+    AddInput<MLFloat16>("value", {kBatchSize, kSequenceLength, kHeadSize}, value.data(), value.size());
+    AddInput<MLFloat16>("past_key", {kBatchSize, 1, capacity, kHeadSize}, initial_key.data(), initial_key.size());
+    AddInput<MLFloat16>("past_value", {kBatchSize, 1, capacity, kHeadSize}, initial_value.data(), initial_value.size());
+    AddInput<int32_t>("seqlens_k", {kBatchSize}, {first_past + kSequenceLength - 1});
+    AddInput<int32_t>("total_sequence_length", {1}, {first_past + kSequenceLength});
+    AddOutput<MLFloat16>("output", {kBatchSize, kSequenceLength, kNumHeads * kHeadSize}, query.data(), query.size());
+    AddOutput<MLFloat16>("present_key", {kBatchSize, 1, capacity, kHeadSize}, initial_key.data(), initial_key.size());
+    AddOutput<MLFloat16>("present_value", {kBatchSize, 1, capacity, kHeadSize}, initial_value.data(), initial_value.size());
+    auto& model = BuildModel();
+    ASSERT_STATUS_OK(model.MainGraph().Resolve(GetRunContext().resolve_options));
+    const auto serialized = model.ToProto().SerializeAsString();
+
+    ConfigOptions provider_options;
+    if (capture) {
+      ASSERT_STATUS_OK(provider_options.AddConfigEntry(webgpu::options::kEnableGraphCapture,
+                                                       webgpu::options::kEnableGraphCapture_ON));
+    }
+    WebGpuExecutionProviderTestOptions test_options{};
+    test_options.max_storage_buffer_binding_size = binding_limit;
+    auto provider = binding_limit > 0 ? WebGpuExecutionProviderWithTestOptions(provider_options, test_options)
+                                      : WebGpuExecutionProviderWithOptions(provider_options);
+    ASSERT_NE(provider, nullptr);
+    auto* ep = provider.get();
+    auto& context = webgpu::WebGpuContextFactory::GetContext(ep->GetDeviceId());
+    if (binding_limit > 0) {
+      ASSERT_EQ(context.DeviceLimits().maxStorageBufferBindingSize, binding_limit);
+    }
+    SessionOptions options;
+    ASSERT_STATUS_OK(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1"));
+    InferenceSession session(options, GetEnvironment());
+    ASSERT_STATUS_OK(session.RegisterExecutionProvider(std::move(provider)));
+    ASSERT_STATUS_OK(session.Load(serialized.data(), static_cast<int>(serialized.size())));
+    ASSERT_STATUS_OK(session.Initialize());
+    const auto allocators = ep->CreatePreferredAllocators();
+    ASSERT_FALSE(allocators.empty());
+    const auto allocator = allocators.back();
+    ASSERT_NE(allocator->Info().device.Type(), OrtDevice::CPU);
+    std::unordered_map<std::string, OrtValue> feeds;
+    std::vector<std::string> output_names;
+    FillFeedsAndOutputNames(feeds, output_names);
+    std::unique_ptr<IOBinding> binding;
+    ASSERT_STATUS_OK(session.NewIOBinding(&binding));
+    for (auto& [name, input] : feeds) {
+      if (capture || name != "total_sequence_length") {
+        OrtValue gpu_input;
+        const auto& source = input.Get<Tensor>();
+        Tensor::InitOrtValue(source.DataType(), source.Shape(), allocator, gpu_input);
+        ASSERT_STATUS_OK(session.GetDataTransferManager().CopyTensor(source, *gpu_input.GetMutable<Tensor>()));
+        input = std::move(gpu_input);
+      }
+      ASSERT_STATUS_OK(binding->BindInput(name, input));
+    }
+    OrtValue output;
+    Tensor::InitOrtValue(DataTypeImpl::GetType<MLFloat16>(), output_shape, allocator, output);
+    ASSERT_STATUS_OK(binding->BindOutput("output", output));
+    ASSERT_STATUS_OK(binding->BindOutput("present_key", feeds.at("past_key")));
+    ASSERT_STATUS_OK(binding->BindOutput("present_value", feeds.at("past_value")));
+    ASSERT_STATUS_OK(binding->SynchronizeInputs());
+
+    auto upload = [&](const char* name, const void* data) {
+      auto& target = *feeds.at(name).GetMutable<Tensor>();
+      const Tensor source(target.DataType(), target.Shape(), const_cast<void*>(data),
+                          CPUAllocator::DefaultInstance()->Info());
+      ORT_THROW_IF_ERROR(session.GetDataTransferManager().CopyTensor(source, target));
+    };
+    const std::array<std::array<int, kBatchSize>, 3> past_lengths =
+        binding_limit > 0 ? std::array<std::array<int, kBatchSize>, 3>{{{193}, {257}, {129}}}
+                          : std::array<std::array<int, kBatchSize>, 3>{{{17}, {41}, {3}}};
+    for (size_t run = 0; run < (capture ? past_lengths.size() : 1u); ++run) {
+      SCOPED_TRACE(testing::Message() << "capture=" << capture << " window=" << window << " run=" << run);
+      const auto& past = past_lengths[run];
+      std::array<int32_t, kBatchSize> lengths{};
+      for (int batch = 0; batch < kBatchSize; ++batch) {
+        lengths[batch] = past[batch] + kSequenceLength - 1;
+      }
+      const int32_t total = *std::max_element(lengths.begin(), lengths.end()) + 1;
+      upload("seqlens_k", lengths.data());
+      if (capture) {
+        // A device-resident total leaves the host total-length parameter at zero.
+        ASSERT_NE(feeds.at("total_sequence_length").Get<Tensor>().Location().device.Type(), OrtDevice::CPU);
+        upload("total_sequence_length", &total);
+      }
+      // Start each replay with known past caches; update the same bound buffers.
+      upload("past_key", initial_key.data());
+      upload("past_value", initial_value.data());
+      Tensor cpu_output(DataTypeImpl::GetType<MLFloat16>(), output_shape, CPUAllocator::DefaultInstance());
+      std::fill_n(cpu_output.MutableData<MLFloat16>(), query.size(), MLFloat16(-99.0f));
+      ASSERT_STATUS_OK(session.GetDataTransferManager().CopyTensor(cpu_output, *output.GetMutable<Tensor>()));
+      profiling::Events events;
+      context.SetProfilingStartTime(std::chrono::high_resolution_clock::now());
+      context.StartProfiling();
+      {
+        auto stop_profiling = gsl::finally([&] { context.CollectProfilingData(events); });
+        if (run == 0) {
+          ASSERT_STATUS_OK(session.Run(RunOptions{}, *binding));
+        } else {
+          ASSERT_TRUE(ep->IsGraphCaptured(0));
+          ASSERT_STATUS_OK(ep->ReplayGraph(0));
+        }
+        ASSERT_STATUS_OK(binding->SynchronizeOutputs());
+      }
+      ExpectCooperativeDispatch(events);
+      ASSERT_FALSE(testing::Test::HasFailure());
+      if (capture) {
+        ASSERT_TRUE(ep->IsGraphCaptured(0));
+      }
+      auto expected_key = initial_key;
+      auto expected_value = initial_value;
+      for (int batch = 0; batch < kBatchSize; ++batch) {
+        const int source = batch * kSequenceLength * kHeadSize;
+        const int target = (batch * capacity + past[batch]) * kHeadSize;
+        std::copy_n(key.begin() + source, kSequenceLength * kHeadSize, expected_key.begin() + target);
+        std::copy_n(value.begin() + source, kSequenceLength * kHeadSize, expected_value.begin() + target);
+      }
+      const auto expected = PrefillReference(query, expected_key, expected_value, past, window, capacity);
+      ASSERT_STATUS_OK(session.GetDataTransferManager().CopyTensor(output.Get<Tensor>(), cpu_output));
+      for (size_t i = 0; i < expected.size(); ++i) {
+        ASSERT_NEAR(cpu_output.Data<MLFloat16>()[i].ToFloat(), expected[i].ToFloat(), 0.001f) << "index=" << i;
+      }
+      for (const auto& [name, expected_cache] : {std::pair{"past_key", &expected_key},
+                                                 std::pair{"past_value", &expected_value}}) {
+        Tensor cache(DataTypeImpl::GetType<MLFloat16>(), cache_shape, CPUAllocator::DefaultInstance());
+        ASSERT_STATUS_OK(session.GetDataTransferManager().CopyTensor(feeds.at(name).Get<Tensor>(), cache));
+        ASSERT_TRUE(std::equal(expected_cache->begin(), expected_cache->end(), cache.Data<MLFloat16>())) << name;
+      }
+    }
+  }
+};
+
+void RunCooperativePrefillTest(bool capture) {
+  if (!SupportsQualcommCooperativePrefill()) {
+    GTEST_SKIP() << "Requires Qualcomm FP16 and 128-lane subgroup control";
+  }
+  auto provider = DefaultWebGpuExecutionProvider();
+  ASSERT_NE(provider, nullptr);
+  const auto& context = webgpu::WebGpuContextFactory::GetContext(provider->GetDeviceId());
+  if (!context.DeviceHasFeature(wgpu::FeatureName::TimestampQuery)) {
+    GTEST_SKIP() << "Dispatch observation requires WebGPU timestamp queries";
+  }
+  for (int window : {-1, 17}) {
+    CooperativePrefillTester test;
+    test.RunPrefill(capture, window);
+  }
+}
+
+}  // namespace
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill) {
+  if (!SupportsQualcommCooperativePrefill()) {
+    GTEST_SKIP() << "Requires Qualcomm FP16 and 128-lane subgroup control";
+  }
+  for (int sequence_length : {128, 129, 257}) {
+    for (int batch_size : {1, 2}) {
+      SCOPED_TRACE(testing::Message() << "sequence_length=" << sequence_length << ", batch=" << batch_size);
+      constexpr int num_heads = 6;
+      constexpr int kv_num_heads = 2;
+      constexpr int head_size = 128;
+      RandomValueGenerator random{1234};
+      const auto query = random.Gaussian<float>(std::array<int64_t, 3>{batch_size, sequence_length, num_heads * head_size}, 0.0f, 1.0f);
+      const auto key = random.Gaussian<float>(std::array<int64_t, 3>{batch_size, sequence_length, kv_num_heads * head_size}, 0.0f, 1.0f);
+      const auto value = random.Gaussian<float>(std::array<int64_t, 3>{batch_size, sequence_length, kv_num_heads * head_size}, 0.0f, 1.0f);
+      const bool rotary = batch_size == 2;
+      const auto reference = RunGQAReference(batch_size, sequence_length, num_heads, kv_num_heads, head_size,
+                                             query, key, value, rotary);
+      const auto actual = RunGQAReference(batch_size, sequence_length, num_heads, kv_num_heads, head_size,
+                                          query, key, value, rotary, /*use_fp16=*/true);
+      ExpectBlockQuantInt8Close(reference, actual, /*max_relative_rmse=*/0.003f,
+                                /*max_absolute_error=*/0.005f);
+    }
+  }
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_HighMagnitude) {
+  if (!SupportsQualcommCooperativePrefill()) {
+    GTEST_SKIP() << "Requires Qualcomm FP16 and 128-lane subgroup control";
+  }
+  RunFp16HighMagnitudeAttention(/*sequence_length=*/129, /*bit_width=*/0);
+}
+
+static void RunWebGpuPrefillFeatures(bool verify_dispatch) {
+  if (!SupportsQualcommCooperativePrefill()) {
+    GTEST_SKIP() << "Requires Qualcomm FP16 and 128-lane subgroup control";
+  }
+  constexpr int batch_size = 2;
+  constexpr int sequence_length = 129;
+  constexpr int num_heads = 6;
+  constexpr int kv_num_heads = 2;
+  constexpr int head_size = 128;
+  const float scale = 1.0f / std::sqrt(static_cast<float>(head_size));
+  RandomValueGenerator random{4321};
+  const auto query = ToFloat16(random.Gaussian<float>(
+      std::array<int64_t, 3>{batch_size, sequence_length, num_heads * head_size}, 0.0f, 1.0f));
+  const auto key = ToFloat16(random.Gaussian<float>(
+      std::array<int64_t, 3>{batch_size, sequence_length, kv_num_heads * head_size}, 0.0f, 1.0f));
+  const auto value = ToFloat16(random.Gaussian<float>(
+      std::array<int64_t, 3>{batch_size, sequence_length, kv_num_heads * head_size}, 0.0f, 1.0f));
+  const auto sinks = ToFloat16({-2.0f, 0.0f, 2.0f, 4.0f, 6.0f, 8.0f});
+  struct FeatureCase {
+    float softcap;
+    int bias_batches;
+    int bias_heads;
+    bool sink;
+    int window;
+  };
+  for (const auto& c : {FeatureCase{0.5f, 0, 0, false, -1},
+                        FeatureCase{0.0f, 1, 1, false, -1},
+                        FeatureCase{0.0f, 1, num_heads, false, -1},
+                        FeatureCase{0.0f, batch_size, 1, false, -1},
+                        FeatureCase{0.0f, batch_size, num_heads, false, -1},
+                        FeatureCase{0.0f, 0, 0, true, -1},
+                        FeatureCase{0.0f, 0, 0, false, 1},
+                        FeatureCase{0.0f, 0, 0, false, 17},
+                        FeatureCase{0.0f, 0, 0, false, 2048},
+                        FeatureCase{0.5f, batch_size, num_heads, true, 17}}) {
+    SCOPED_TRACE(testing::Message() << "softcap=" << c.softcap << " bias=" << c.bias_batches << ","
+                                    << c.bias_heads << " sink=" << c.sink << " window=" << c.window);
+    const auto bias = ToFloat16(random.Gaussian<float>(
+        std::array<int64_t, 4>{c.bias_batches, c.bias_heads, sequence_length, sequence_length}, 0.0f, 1.0f));
+    InlinedVector<MLFloat16> expected(query.size());
+    for (int batch = 0; batch < batch_size; ++batch) {
+      for (int row = 0; row < sequence_length; ++row) {
+        const int first = c.window > 0 ? std::max(0, row + 1 - c.window) : 0;
+        for (int head = 0; head < num_heads; ++head) {
+          const int kv_head = head / (num_heads / kv_num_heads);
+          const int q_offset = ((batch * sequence_length + row) * num_heads + head) * head_size;
+          InlinedVector<double> scores(row + 1);
+          double maximum = c.sink ? sinks[head].ToFloat() : -std::numeric_limits<double>::infinity();
+          for (int k = first; k <= row; ++k) {
+            const int k_offset = ((batch * sequence_length + k) * kv_num_heads + kv_head) * head_size;
+            double score = 0.0;
+            for (int d = 0; d < head_size; ++d) {
+              score += static_cast<double>(query[q_offset + d].ToFloat()) * key[k_offset + d].ToFloat();
+            }
+            score *= scale;
+            if (c.softcap > 0.0f) {
+              score = std::tanh(score / c.softcap) * c.softcap;
+            }
+            if (!bias.empty()) {
+              const int bias_batch = c.bias_batches == 1 ? 0 : batch;
+              const int bias_head = c.bias_heads == 1 ? 0 : head;
+              const int offset =
+                  ((bias_batch * c.bias_heads + bias_head) * sequence_length + row) * sequence_length + k;
+              score += bias[offset].ToFloat();
+            }
+            scores[k] = score;
+            maximum = std::max(maximum, score);
+          }
+          double denominator = c.sink ? std::exp(sinks[head].ToFloat() - maximum) : 0.0;
+          for (int k = first; k <= row; ++k) {
+            scores[k] = std::exp(scores[k] - maximum);
+            denominator += scores[k];
+          }
+          for (int d = 0; d < head_size; ++d) {
+            double sum = 0.0;
+            for (int k = first; k <= row; ++k) {
+              const int offset = ((batch * sequence_length + k) * kv_num_heads + kv_head) * head_size + d;
+              sum += scores[k] * value[offset].ToFloat();
+            }
+            expected[q_offset + d] = MLFloat16(static_cast<float>(sum / denominator));
+          }
+        }
+      }
+    }
+    OpTester test("GroupQueryAttention", 1, kMSDomain);
+    test.AddAttribute<int64_t>("num_heads", num_heads);
+    test.AddAttribute<int64_t>("kv_num_heads", kv_num_heads);
+    test.AddAttribute<float>("softcap", c.softcap);
+    test.AddAttribute<int64_t>("local_window_size", c.window);
+    test.AddInput<MLFloat16>("query", {batch_size, sequence_length, num_heads * head_size}, query);
+    test.AddInput<MLFloat16>("key", {batch_size, sequence_length, kv_num_heads * head_size}, key);
+    test.AddInput<MLFloat16>("value", {batch_size, sequence_length, kv_num_heads * head_size}, value);
+    test.AddOptionalInputEdge<MLFloat16>();  // past_key
+    test.AddOptionalInputEdge<MLFloat16>();  // past_value
+    test.AddInput<int32_t>("seqlens_k", {batch_size}, {sequence_length - 1, sequence_length - 1});
+    test.AddInput<int32_t>("total_sequence_length", {1}, {sequence_length}, true);
+    test.AddOptionalInputEdge<MLFloat16>();  // cos_cache
+    test.AddOptionalInputEdge<MLFloat16>();  // sin_cache
+    test.AddOptionalInputEdge<int64_t>();    // position_ids
+    if (!bias.empty()) {
+      test.AddInput<MLFloat16>("attention_bias", {c.bias_batches, c.bias_heads, sequence_length, sequence_length},
+                               bias);
+    } else {
+      test.AddOptionalInputEdge<MLFloat16>();
+    }
+    if (c.sink) {
+      test.AddInput<MLFloat16>("head_sink", {num_heads}, sinks);
+    }
+    test.AddOutput<MLFloat16>("output", {batch_size, sequence_length, num_heads * head_size},
+                              expected.data(), expected.size());
+    test.SetOutputAbsErr("output", 0.003f);
+    SessionOptions options;
+    ASSERT_STATUS_OK(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1"));
+    RunWithCooperativeDispatchCheck(verify_dispatch, [&] {
+      test.Config(options).ConfigEp(DefaultWebGpuExecutionProvider()).RunWithConfig();
+    });
+  }
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_Features) {
+  RunWebGpuPrefillFeatures(false);
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_Features_Dispatch) {
+  if (!SupportsCooperativeDispatchObservation()) {
+    GTEST_SKIP() << "Requires Qualcomm cooperative prefill and timestamp queries";
+  }
+  RunWebGpuPrefillFeatures(true);
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_PastAndPadding) {
+  if (!SupportsQualcommCooperativePrefill()) {
+    GTEST_SKIP() << "Requires Qualcomm FP16 and 128-lane subgroup control";
+  }
+  constexpr int sequence_length = 129;
+  constexpr int num_heads = 6;
+  constexpr int kv_num_heads = 2;
+  constexpr int head_size = 128;
+  for (int past_length : {0, 17}) {
+    SCOPED_TRACE(past_length);
+    const int batch_size = past_length == 0 ? 2 : 1;
+    const int capacity = past_length + sequence_length + 11;
+    InlinedVector<MLFloat16> query(batch_size * sequence_length * num_heads * head_size, MLFloat16(0.0f));
+    InlinedVector<MLFloat16> key(batch_size * sequence_length * kv_num_heads * head_size, MLFloat16(0.0f));
+    InlinedVector<MLFloat16> value(key.size());
+    InlinedVector<MLFloat16> past_key(batch_size * kv_num_heads * capacity * head_size, MLFloat16(0.0f));
+    InlinedVector<MLFloat16> past_value(past_key.size(), MLFloat16(0.0f));
+    InlinedVector<MLFloat16> expected(query.size());
+    InlinedVector<int32_t> seqlens_k(batch_size);
+    const auto value_at = [](int batch, int kv_head, int row, int dim) {
+      return batch * 0.5f + kv_head * 0.25f + (row % 7) * 0.125f + (dim % 4) * 0.0625f;
+    };
+    for (int batch = 0; batch < batch_size; ++batch) {
+      const int valid_keys = past_length == 0 && batch == 1 ? 101 : past_length + sequence_length;
+      seqlens_k[batch] = valid_keys - 1;
+      for (int kv_head = 0; kv_head < kv_num_heads; ++kv_head) {
+        for (int row = 0; row < past_length; ++row) {
+          for (int dim = 0; dim < head_size; ++dim) {
+            past_value[((batch * kv_num_heads + kv_head) * capacity + row) * head_size + dim] =
+                MLFloat16(value_at(batch, kv_head, row, dim));
+          }
+        }
+        for (int row = 0; row < sequence_length; ++row) {
+          for (int dim = 0; dim < head_size; ++dim) {
+            value[((batch * sequence_length + row) * kv_num_heads + kv_head) * head_size + dim] =
+                MLFloat16(value_at(batch, kv_head, past_length + row, dim));
+          }
+        }
+      }
+      // Zero queries give uniform probabilities over the visible causal prefix.
+      for (int row = 0; row < sequence_length; ++row) {
+        const int count = std::min(valid_keys, past_length + row + 1);
+        for (int head = 0; head < num_heads; ++head) {
+          for (int dim = 0; dim < head_size; ++dim) {
+            float sum = 0.0f;
+            for (int k = 0; k < count; ++k) {
+              sum += value_at(batch, head / (num_heads / kv_num_heads), k, dim);
+            }
+            expected[((batch * sequence_length + row) * num_heads + head) * head_size + dim] =
+                MLFloat16(sum / count);
+          }
+        }
+      }
+    }
+    OpTester test("GroupQueryAttention", 1, kMSDomain);
+    test.AddAttribute<int64_t>("num_heads", num_heads);
+    test.AddAttribute<int64_t>("kv_num_heads", kv_num_heads);
+    test.AddInput<MLFloat16>("query", {batch_size, sequence_length, num_heads * head_size}, query.data(), query.size());
+    test.AddInput<MLFloat16>("key", {batch_size, sequence_length, kv_num_heads * head_size}, key.data(), key.size());
+    test.AddInput<MLFloat16>("value", {batch_size, sequence_length, kv_num_heads * head_size}, value.data(), value.size());
+    test.AddInput<MLFloat16>("past_key", {batch_size, kv_num_heads, capacity, head_size}, past_key.data(), past_key.size());
+    test.AddInput<MLFloat16>("past_value", {batch_size, kv_num_heads, capacity, head_size}, past_value.data(), past_value.size());
+    test.AddInput<int32_t>("seqlens_k", {batch_size}, seqlens_k.data(), seqlens_k.size());
+    test.AddInput<int32_t>("total_sequence_length", {1}, {past_length + sequence_length}, true);
+    test.AddOutput<MLFloat16>("output", {batch_size, sequence_length, num_heads * head_size}, expected.data(), expected.size());
+    test.SetOutputAbsErr("output", 0.001f);
+    SessionOptions options;
+    ASSERT_STATUS_OK(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1"));
+    test.Config(options).ConfigEp(DefaultWebGpuExecutionProvider()).RunWithConfig();
+  }
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_Dispatch) {
+  RunCooperativePrefillTest(false);
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_GraphCaptureReplay) {
+  RunCooperativePrefillTest(true);
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_SegmentedBuffers) {
+  if (!SupportsCooperativeDispatchObservation()) {
+    GTEST_SKIP() << "Requires Qualcomm cooperative prefill and timestamp queries";
+  }
+  {
+    auto ep = DefaultWebGpuExecutionProvider();
+    const auto& context = webgpu::WebGpuContextFactory::GetContext(ep->GetDeviceId());
+    if (context.DeviceLimits().maxStorageBuffersPerShaderStage < 9) {
+      GTEST_SKIP() << "Segmented Q, K, V and output require nine storage bindings";
+    }
+  }
+  for (bool capture : {false, true}) {
+    for (int window : {-1, 17}) {
+      CooperativePrefillTester test;
+      // Q/output span two bindings; the 512-position caches also cross the
+      // 256-position binding boundary in each run, including graph replays.
+      test.RunPrefill(capture, window, /*binding_limit=*/65536);
+    }
+  }
+}
+
+TEST(GroupQueryAttentionTest, WebGPU_FP16_CooperativePrefill_SmallQueryScale) {
+  if (!SupportsCooperativeDispatchObservation()) {
+    GTEST_SKIP() << "Requires Qualcomm cooperative prefill and timestamp queries";
+  }
+  constexpr int sequence_length = 129;
+  constexpr int head_size = 128;
+  constexpr float query_value = 0x1p-14f;
+  constexpr float key_value = 65504.0f;
+  constexpr float scale = 0x1p-13f;
+  const InlinedVector<MLFloat16> query(sequence_length * head_size, MLFloat16(query_value));
+  InlinedVector<MLFloat16> key(query.size());
+  InlinedVector<MLFloat16> value(query.size());
+  for (int row = 0; row < sequence_length; ++row) {
+    const float sign = row % 2 == 0 ? 1.0f : -1.0f;
+    std::fill_n(key.begin() + row * head_size, head_size, MLFloat16(sign * key_value));
+    std::fill_n(value.begin() + row * head_size, head_size, MLFloat16(sign));
+  }
+  for (float softcap : {0.0f, 0.05f}) {
+    SCOPED_TRACE(softcap);
+    // Scaling Q in FP16 would round 2^-27 to zero and erase these scores.
+    double score = static_cast<double>(query_value) * key_value * head_size * scale;
+    if (softcap > 0.0f) {
+      score = std::tanh(score / softcap) * softcap;
+    }
+    InlinedVector<MLFloat16> expected(query.size());
+    for (int row = 0; row < sequence_length; ++row) {
+      const double positive = (row / 2 + 1) * std::exp(score);
+      const double negative = ((row + 1) / 2) * std::exp(-score);
+      const auto result = MLFloat16(static_cast<float>((positive - negative) / (positive + negative)));
+      std::fill_n(expected.begin() + row * head_size, head_size, result);
+    }
+    OpTester test("GroupQueryAttention", 1, kMSDomain);
+    test.AddAttribute<int64_t>("num_heads", 1);
+    test.AddAttribute<int64_t>("kv_num_heads", 1);
+    test.AddAttribute<float>("scale", scale);
+    test.AddAttribute<float>("softcap", softcap);
+    test.AddInput<MLFloat16>("query", {1, sequence_length, head_size}, query.data(), query.size());
+    test.AddInput<MLFloat16>("key", {1, sequence_length, head_size}, key.data(), key.size());
+    test.AddInput<MLFloat16>("value", {1, sequence_length, head_size}, value.data(), value.size());
+    test.AddOptionalInputEdge<MLFloat16>();  // past_key
+    test.AddOptionalInputEdge<MLFloat16>();  // past_value
+    test.AddInput<int32_t>("seqlens_k", {1}, {sequence_length - 1});
+    test.AddInput<int32_t>("total_sequence_length", {1}, {sequence_length}, true);
+    test.AddOutput<MLFloat16>("output", {1, sequence_length, head_size}, expected.data(), expected.size());
+    test.SetOutputAbsErr("output", 0.001f);
+    SessionOptions options;
+    ASSERT_STATUS_OK(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1"));
+    RunWithCooperativeDispatchCheck(true, [&] {
+      test.Config(options).ConfigEp(DefaultWebGpuExecutionProvider()).RunWithConfig();
+    });
+  }
+}
+
+}  // namespace onnxruntime::test
+
+#endif  // !defined(ORT_MINIMAL_BUILD) && !defined(DISABLE_CONTRIB_OPS)
diff --git a/onnxruntime/test/providers/webgpu/contrib_ops/multihead_attention_test.cc b/onnxruntime/test/providers/webgpu/contrib_ops/multihead_attention_test.cc
new file mode 100644
index 0000000000000..0347e62ce9926
--- /dev/null
+++ b/onnxruntime/test/providers/webgpu/contrib_ops/multihead_attention_test.cc
@@ -0,0 +1,123 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#if !defined(ORT_MINIMAL_BUILD) && !defined(DISABLE_CONTRIB_OPS)
+
+#include <algorithm>
+#include <cmath>
+#include <limits>
+
+#include "gtest/gtest.h"
+
+#include "core/common/inlined_containers.h"
+#include "core/session/onnxruntime_session_options_config_keys.h"
+#include "test/providers/provider_test_utils.h"
+#include "test/providers/webgpu/attention_test_utils.h"
+
+namespace onnxruntime::test {
+
+namespace {
+
+void RunNonCausalPrefill(bool verify_dispatch) {
+  if (!WebGpuSupportsFp16()) {
+    GTEST_SKIP() << "WebGPU shader-f16 support is unavailable";
+  }
+  constexpr int sequence_length = 129;
+  constexpr int num_heads = 2;
+  constexpr int head_size = 128;
+  constexpr int hidden_size = num_heads * head_size;
+  for (int kv_length : {83, 145}) {
+    for (bool bnsh_query : {false, true}) {
+      SCOPED_TRACE(testing::Message() << "kv_length=" << kv_length << " bnsh_query=" << bnsh_query);
+      InlinedVector<MLFloat16> query(sequence_length * hidden_size);
+      InlinedVector<MLFloat16> key(kv_length * hidden_size);
+      InlinedVector<MLFloat16> value(key.size());
+      InlinedVector<MLFloat16> bias(sequence_length * kv_length);
+      InlinedVector<MLFloat16> expected(query.size());
+      for (int row = 0; row < sequence_length; ++row) {
+        for (int head = 0; head < num_heads; ++head) {
+          for (int d = 0; d < head_size; ++d) {
+            query[row * hidden_size + head * head_size + d] =
+                MLFloat16(((row * 7 + head * 11 + d * 3) % 29 - 14) * 0.125f);
+          }
+        }
+      }
+      for (int k = 0; k < kv_length; ++k) {
+        for (int d = 0; d < hidden_size; ++d) {
+          key[k * hidden_size + d] =
+              MLFloat16(((k * 11 + (d / head_size) * 5 + (d % head_size) * 7) % 31 - 15) * 0.125f);
+          value[k * hidden_size + d] = MLFloat16((k % 11) * 0.125f + (d % 7) * 0.0625f);
+        }
+      }
+      for (int row = 0; row < sequence_length; ++row) {
+        for (int k = 0; k < kv_length; ++k) {
+          bias[row * kv_length + k] = MLFloat16(((k + row) % 13) * 0.25f);
+        }
+        for (int head = 0; head < num_heads; ++head) {
+          InlinedVector<double> scores(kv_length);
+          double maximum = -std::numeric_limits<double>::infinity();
+          for (int k = 0; k < kv_length; ++k) {
+            double dot = 0.0;
+            for (int d = 0; d < head_size; ++d) {
+              dot += static_cast<double>(query[row * hidden_size + head * head_size + d].ToFloat()) *
+                     key[k * hidden_size + head * head_size + d].ToFloat();
+            }
+            scores[k] = dot / std::sqrt(static_cast<double>(head_size)) + bias[row * kv_length + k].ToFloat();
+            maximum = std::max(maximum, scores[k]);
+          }
+          double denominator = 0.0;
+          for (int k = 0; k < kv_length; ++k) {
+            scores[k] = std::exp(scores[k] - maximum);
+            denominator += scores[k];
+          }
+          for (int d = 0; d < head_size; ++d) {
+            double sum = 0.0;
+            for (int k = 0; k < kv_length; ++k) {
+              sum += scores[k] * value[k * hidden_size + head * head_size + d].ToFloat();
+            }
+            expected[row * hidden_size + head * head_size + d] = MLFloat16(static_cast<float>(sum / denominator));
+          }
+        }
+      }
+      OpTester test("MultiHeadAttention", 1, kMSDomain);
+      test.AddAttribute<int64_t>("num_heads", num_heads);
+      test.AddAttribute<int64_t>("unidirectional", 0);
+      test.AddInput<MLFloat16>("query", {1, sequence_length, hidden_size}, query.data(), query.size());
+      test.AddInput<MLFloat16>("key", {1, kv_length, hidden_size}, key.data(), key.size());
+      test.AddInput<MLFloat16>("value", {1, kv_length, hidden_size}, value.data(), value.size());
+      if (bnsh_query) {
+        // Projection bias makes the EP transpose Q to BNSH before flash attention.
+        const InlinedVector<MLFloat16> projection_bias(3 * hidden_size, MLFloat16(0.0f));
+        test.AddInput<MLFloat16>("bias", {3 * hidden_size}, projection_bias.data(), projection_bias.size());
+      } else {
+        test.AddOptionalInputEdge<MLFloat16>();
+      }
+      test.AddOptionalInputEdge<int32_t>();  // key_padding_mask
+      test.AddInput<MLFloat16>("attention_bias", {1, 1, sequence_length, kv_length}, bias.data(), bias.size());
+      test.AddOutput<MLFloat16>("output", {1, sequence_length, hidden_size}, expected.data(), expected.size());
+      test.SetOutputAbsErr("output", 0.001f);
+      SessionOptions options;
+      ASSERT_STATUS_OK(options.config_options.AddConfigEntry(kOrtSessionOptionsDisableCPUEPFallback, "1"));
+      RunWithCooperativeDispatchCheck(verify_dispatch, [&] {
+        test.Config(options).ConfigEp(DefaultWebGpuExecutionProvider()).RunWithConfig();
+      });
+    }
+  }
+}
+
+}  // namespace
+
+TEST(MultiHeadAttentionTest, WebGPU_FP16_CooperativePrefill_NonCausal) {
+  RunNonCausalPrefill(false);
+}
+
+TEST(MultiHeadAttentionTest, WebGPU_FP16_CooperativePrefill_NonCausal_Dispatch) {
+  if (!SupportsCooperativeDispatchObservation()) {
+    GTEST_SKIP() << "Requires Qualcomm cooperative prefill and timestamp queries";
+  }
+  RunNonCausalPrefill(true);
+}
+
+}  // namespace onnxruntime::test
+
+#endif  // !defined(ORT_MINIMAL_BUILD) && !defined(DISABLE_CONTRIB_OPS)
diff --git a/onnxruntime/test/providers/webgpu/contrib_ops/paged_attention_test.cc b/onnxruntime/test/providers/webgpu/contrib_ops/paged_attention_test.cc
new file mode 100644
index 0000000000000..dc729bfe60439
--- /dev/null
+++ b/onnxruntime/test/providers/webgpu/contrib_ops/paged_attention_test.cc
@@ -0,0 +1,61 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#if !defined(ORT_MINIMAL_BUILD) && !defined(DISABLE_CONTRIB_OPS)
+
+#include "gtest/gtest.h"
+
+#include "test/contrib_ops/paged_attention_op_test_helper.h"
+#include "test/providers/provider_test_utils.h"
+#include "test/providers/webgpu/attention_test_utils.h"
+
+namespace onnxruntime::test {
+
+namespace {
+
+void RunRaggedPrefill(bool verify_dispatch) {
+  if (!WebGpuSupportsFp16()) {
+    GTEST_SKIP() << "WebGPU shader-f16 support is unavailable";
+  }
+  for (bool causal : {false, true}) {
+    for (int window : {-1, 17}) {
+      SCOPED_TRACE(causal);
+      SCOPED_TRACE(window);
+      PagedAttentionTestCase c{};
+      c.batch_size = 2;
+      c.token_count = 230;
+      c.num_heads = 6;
+      c.kv_num_heads = 2;
+      c.head_size = 128;
+      c.block_size = 16;  // Forces dense fallback instead of fused paged prefill.
+      c.num_blocks = 20;
+      c.max_num_blocks_per_seq = 10;
+      c.is_causal = causal;
+      c.local_window_size = window;
+      c.cumulative_seqlens_q = {0, 129, 230};
+      c.past_seqlens = {17, 7};
+      c.block_table = {18, 2, 16, 4, 14, 6, 12, 8, 10, 0,
+                       19, 1, 17, 3, 15, 5, 13, 7, 11, 9};
+      RunWithCooperativeDispatchCheck(verify_dispatch, [&] {
+        RunPagedAttentionTest(c, DefaultWebGpuExecutionProvider(), /*output_abs_error=*/1e-4f);
+      });
+    }
+  }
+}
+
+}  // namespace
+
+TEST(PagedAttention, WebGpu_CooperativePrefill_RaggedDenseFallback) {
+  RunRaggedPrefill(false);
+}
+
+TEST(PagedAttention, WebGpu_CooperativePrefill_RaggedDenseFallback_Dispatch) {
+  if (!SupportsCooperativeDispatchObservation()) {
+    GTEST_SKIP() << "Requires Qualcomm cooperative prefill and timestamp queries";
+  }
+  RunRaggedPrefill(true);
+}
+
+}  // namespace onnxruntime::test
+
+#endif  // !defined(ORT_MINIMAL_BUILD) && !defined(DISABLE_CONTRIB_OPS)