diff --git a/backends/webgpu/runtime/WebGPUBackend.cpp b/backends/webgpu/runtime/WebGPUBackend.cpp
index ba7f6f13e64..e1497bc4ed8 100644
--- a/backends/webgpu/runtime/WebGPUBackend.cpp
+++ b/backends/webgpu/runtime/WebGPUBackend.cpp
@@ -79,8 +79,32 @@ Result<DelegateHandle*> WebGPUBackend::init(
     return Error::DelegateInvalidCompatibility;
   }
 
+  // Load-time backend option (BackendOption / LoadBackendOptionsMap), keyed by
+  // the registered backend name; default false. Mirrors the CoreML/XNNPACK
+  // runtime-spec pattern -- no compile flag and no .pte re-export needed.
+  bool enable_f16_kv_cache = false;
+  {
+    Result<bool> spec = context.get_runtime_spec<bool>("enable_f16_kv_cache");
+    if (spec.ok()) {
+      enable_f16_kv_cache = spec.get();
+    }
+  }
+  bool enable_f16_accumulate_gemm = false;
+  {
+    Result<bool> spec =
+        context.get_runtime_spec<bool>("enable_f16_accumulate_gemm");
+    if (spec.ok()) {
+      enable_f16_accumulate_gemm = spec.get();
+    }
+  }
+
   try {
-    graph->build(flatbuffer_data, constant_data, context.get_named_data_map());
+    graph->build(
+        flatbuffer_data,
+        constant_data,
+        context.get_named_data_map(),
+        enable_f16_kv_cache,
+        enable_f16_accumulate_gemm);
   } catch (const std::exception& e) {
     ET_LOG(Error, "WebGPU graph build failed: %s", e.what());
     graph->~WebGPUGraph();
diff --git a/backends/webgpu/runtime/WebGPUGraph.cpp b/backends/webgpu/runtime/WebGPUGraph.cpp
index eb10fb44e28..d9ab6467882 100644
--- a/backends/webgpu/runtime/WebGPUGraph.cpp
+++ b/backends/webgpu/runtime/WebGPUGraph.cpp
@@ -90,6 +90,49 @@ WGPUBuffer WebGPUGraph::create_scratch_buffer(size_t nbytes) {
   return buffer;
 }
 
+WGPUBuffer WebGPUGraph::acquire_scratch(size_t nbytes) {
+  nbytes = nbytes > 0 ? nbytes : 4;
+  // Best-fit reuse: smallest free slot with size in [nbytes, 2*nbytes] -- the
+  // 2x cap stops a large Cmax-sized buffer from backing a tiny request. Never
+  // reuse an in_use slot (co-live safety).
+  ScratchSlot* best = nullptr;
+  for (auto& s : scratch_pool_) {
+    // s.size - nbytes (safe: s.size >= nbytes) avoids overflowing 2 * nbytes.
+    if (!s.in_use && s.size >= nbytes && s.size - nbytes <= nbytes) {
+      if (best == nullptr || s.size < best->size) {
+        best = &s;
+      }
+    }
+  }
+  if (best != nullptr) {
+    best->in_use = true;
+    return best->buffer;
+  }
+  // None reusable -> create a new slot (freed in the dtor, like
+  // scratch_buffers_).
+  WGPUBufferDescriptor buf_desc = {};
+  buf_desc.size = nbytes;
+  buf_desc.usage = WGPUBufferUsage_Storage | WGPUBufferUsage_CopyDst |
+      WGPUBufferUsage_CopySrc;
+  buf_desc.mappedAtCreation = false;
+  WGPUBuffer buffer = wgpuDeviceCreateBuffer(device_, &buf_desc);
+  scratch_pool_.push_back({buffer, nbytes, true});
+  return buffer;
+}
+
+void WebGPUGraph::release_scratch(WGPUBuffer buffer) {
+  if (!buffer) {
+    return;
+  }
+  for (auto& s : scratch_pool_) {
+    if (s.buffer == buffer) {
+      s.in_use = false;
+      return;
+    }
+  }
+  // Not a pooled buffer -> no-op; the dtor frees it via scratch_buffers_.
+}
+
 WGPUBuffer WebGPUGraph::make_uniform_buffer(const void* data, size_t size) {
   WGPUBufferDescriptor desc = {};
   desc.size = size;
@@ -267,6 +310,11 @@ WebGPUGraph::~WebGPUGraph() {
       wgpuBufferRelease(buf);
     }
   }
+  for (auto& s : scratch_pool_) {
+    if (s.buffer) {
+      wgpuBufferRelease(s.buffer);
+    }
+  }
   for (auto& buf : owned_uniform_buffers_) {
     if (buf) {
       wgpuBufferRelease(buf);
@@ -310,7 +358,9 @@ WebGPUGraph::~WebGPUGraph() {
 void WebGPUGraph::build(
     const void* flatbuffer_data,
     const uint8_t* constant_data,
-    const executorch::runtime::NamedDataMap* named_data_map) {
+    const executorch::runtime::NamedDataMap* named_data_map,
+    bool f16_kv_cache,
+    bool f16_accumulate_gemm) {
   if (!device_) {
     auto* ctx = get_default_webgpu_context();
     if (ctx) {
@@ -331,6 +381,15 @@ void WebGPUGraph::build(
   constant_data_ = constant_data;
   named_data_map_ = named_data_map;
 
+  // f16 KV cache (runtime opt-in): store K/V caches as f16 iff the opt-in is
+  // set AND the device negotiated shader-f16 (fail-closed).
+  const WebGPUContext* kv_ctx = get_default_webgpu_context();
+  kv_f16_ = f16_kv_cache && (kv_ctx != nullptr && kv_ctx->shader_f16_supported);
+
+  // f16-accumulate q4gsw steel prefill GEMM (runtime opt-in). QuantizedLinear
+  // additionally gates the kernel on the negotiated shader-f16 feature.
+  f16_accumulate_gemm_ = f16_accumulate_gemm;
+
   // Phase 1: Create all values
   const auto* values = graph->values();
   const int num_vals = values ? values->size() : 0;
@@ -358,6 +417,13 @@ void WebGPUGraph::build(
       if (!a) {
         continue;
       }
+      // f16 KV: tag sdpa K/V cache values (args[3],[4]) for half-size alloc.
+      // Inert unless kv_f16_ (runtime opt-in) is set.
+      if (kv_f16_ && a->size() > 4 &&
+          oc->name()->str() == "sdpa_with_kv_cache.default") {
+        kv_cache_ids_.insert(static_cast<int>(a->Get(3)));
+        kv_cache_ids_.insert(static_cast<int>(a->Get(4)));
+      }
       for (unsigned j = 0; j < a->size(); j++) {
         int id = static_cast<int>(a->Get(j));
         if (is_prepack && j == 0) {
@@ -378,6 +444,29 @@ void WebGPUGraph::build(
     }
   }
 
+  // f16 KV defensive guard: fail loud if a non-sdpa op reads an f16 cache.
+  // Inert unless kv_f16_ (runtime opt-in) is set.
+  if (kv_f16_ && !kv_cache_ids_.empty() && chain_prescan) {
+    for (unsigned ci = 0; ci < chain_prescan->size(); ci++) {
+      const auto* oc = chain_prescan->Get(ci);
+      const std::string nm = oc->name()->str();
+      if (nm == "sdpa_with_kv_cache.default" || nm == kPrepackOpName) {
+        continue;
+      }
+      const auto* a = oc->args();
+      if (!a) {
+        continue;
+      }
+      for (unsigned j = 0; j < a->size(); j++) {
+        if (kv_cache_ids_.count(static_cast<int>(a->Get(j))) != 0) {
+          throw std::runtime_error(
+              "WebGPU f16 KV: cache tensor consumed by non-sdpa op '" + nm +
+              "' would misread the f16 buffer");
+        }
+      }
+    }
+  }
+
   for (int i = 0; i < num_vals; i++) {
     const auto* val = values->Get(i);
     if (!val || val->value_type() == vkgraph::GraphTypes::NONE) {
@@ -407,6 +496,23 @@ void WebGPUGraph::build(
         tensor.cur_dims = tensor.dims;
         tensor.cur_nbytes = tensor.nbytes;
 
+        // f16 KV cache: dedicated half-size array<f16> buffer. WebGPU
+        // zero-initializes freshly-created buffers, so no explicit clear is
+        // needed. Inert unless kv_f16_ (runtime opt-in) is set.
+        if (kv_f16_ && kv_cache_ids_.count(i) != 0) {
+          tensor.elem_size = 2;
+          tensor.nbytes = numel * 2;
+          tensor.cur_nbytes = tensor.nbytes;
+          tensor_mem_obj_ids_[i] = -1;
+          WGPUBufferDescriptor buf_desc = {};
+          buf_desc.size = std::max(tensor.nbytes, size_t(4));
+          buf_desc.usage = WGPUBufferUsage_Storage | WGPUBufferUsage_CopyDst |
+              WGPUBufferUsage_CopySrc;
+          buf_desc.mappedAtCreation = false;
+          tensor.buffer = wgpuDeviceCreateBuffer(device_, &buf_desc);
+          break;
+        }
+
         int constant_id = vk_tensor->constant_id();
         int mem_obj_id = vk_tensor->mem_obj_id();
 
diff --git a/backends/webgpu/runtime/WebGPUGraph.h b/backends/webgpu/runtime/WebGPUGraph.h
index 0ebcd8071f9..66f0e401de5 100644
--- a/backends/webgpu/runtime/WebGPUGraph.h
+++ b/backends/webgpu/runtime/WebGPUGraph.h
@@ -104,7 +104,9 @@ class WebGPUGraph {
   void build(
       const void* flatbuffer_data,
       const uint8_t* constant_data,
-      const executorch::runtime::NamedDataMap* named_data_map = nullptr);
+      const executorch::runtime::NamedDataMap* named_data_map = nullptr,
+      bool f16_kv_cache = false,
+      bool f16_accumulate_gemm = false);
 
   // Copy input tensor data from host pointers into GPU buffers.
   void copy_inputs(const std::vector<InputData>& inputs);
@@ -267,6 +269,35 @@ class WebGPUGraph {
   // Graph-owned scratch storage buffer for fused-op intermediates (e.g. SDPA).
   WGPUBuffer create_scratch_buffer(size_t nbytes);
 
+  // Reusable scratch pool for SINGLE-OP-LIFETIME fused-op scratch (SDPA
+  // attn_weights/softmax, FlashDecoding partials). acquire_scratch() reuses a
+  // free slot (best-fit, size in [n,2n]) or creates one; the caller RELEASES it
+  // at op-lowering scope exit (use ScopedScratch), so N layers' scratch reuses
+  // a small constant of buffers instead of N x held to graph teardown.
+  // Correctness: WebGPU/Dawn auto-inserts RAW hazard barriers between
+  // dispatches on a shared storage buffer regardless of pass structure -- the
+  // SAME guarantee mem_obj_id aliasing already relies on -- so reuse is
+  // bit-identical. Never hand a still-in_use slot to a co-live requester.
+  WGPUBuffer acquire_scratch(size_t nbytes);
+  void release_scratch(WGPUBuffer buffer);
+  // RAII: releases an acquired scratch slot when the op-lowering scope exits
+  // (leak-safe vs early returns).
+  struct ScopedScratch {
+    WebGPUGraph* g = nullptr;
+    WGPUBuffer buf = nullptr;
+    ScopedScratch(WebGPUGraph* graph, WGPUBuffer b) : g(graph), buf(b) {}
+    ~ScopedScratch() {
+      if (g && buf) {
+        g->release_scratch(buf);
+      }
+    }
+    ScopedScratch(const ScopedScratch&) = delete;
+    ScopedScratch& operator=(const ScopedScratch&) = delete;
+    operator WGPUBuffer() const {
+      return buf;
+    }
+  };
+
   // Create a mapped-at-creation uniform buffer from `size` bytes and track it
   // in the memory stats. Shared helper for ops needing a uniform Params buffer.
   WGPUBuffer make_uniform_buffer(const void* data, size_t size);
@@ -314,6 +345,23 @@ class WebGPUGraph {
     return value_types_[id];
   }
 
+ public:
+  // True when the sdpa K/V cache is stored f16-packed (runtime opt-in).
+  bool kv_f16() const {
+    return kv_f16_;
+  }
+
+  // True when the q4gsw steel prefill GEMM uses the lossy f16-accumulate kernel
+  // (runtime opt-in; perplexity-gated, not bit-exact).
+  bool f16_accumulate_gemm() const {
+    return f16_accumulate_gemm_;
+  }
+
+ private:
+  bool kv_f16_ = false;
+  std::unordered_set<int> kv_cache_ids_;
+  bool f16_accumulate_gemm_ = false;
+
  private:
   WGPUInstance instance_ = nullptr;
   WGPUDevice device_ = nullptr;
@@ -366,6 +414,16 @@ class WebGPUGraph {
   // Long-lived scratch storage buffers for fused ops (e.g. SDPA temporaries).
   std::vector<WGPUBuffer> scratch_buffers_;
 
+  // Reusable scratch pool: single-op-lifetime buffers recycled across ops
+  // (acquire_scratch/release_scratch). Each slot is freed in the dtor. See
+  // acquire_scratch() for the reuse policy.
+  struct ScratchSlot {
+    WGPUBuffer buffer = nullptr;
+    size_t size = 0;
+    bool in_use = false;
+  };
+  std::vector<ScratchSlot> scratch_pool_;
+
   // Uniform buffers owned for the graph's lifetime; released in the dtor.
   std::vector<WGPUBuffer> owned_uniform_buffers_;
 
diff --git a/backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp b/backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp
index 20f7024ca82..b0728764310 100644
--- a/backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp
+++ b/backends/webgpu/runtime/ops/quantized_linear/QuantizedLinear.cpp
@@ -12,6 +12,8 @@
 #include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
 #include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_coop4_bicol_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_shmem_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_f16acc_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_wgsl.h>
@@ -270,7 +272,24 @@ void q4gsw_linear_impl(WebGPUGraph& graph, const std::vector<int>& args) {
   if (use_steel) {
     const WebGPUContext* ctx = get_default_webgpu_context();
     if (ctx != nullptr && ctx->shader_f16_supported) {
-      shader_src = kQ4gswLinearGemmSteelHalfWGSL;
+      // Packed-word dequant: bit-exact to the steel `half` kernel but loads
+      // each u32 weight word once + hoists the per-column scale (half re-reads
+      // them ~8x/~16x). Needs group_size % BK == 0 so the hoisted scale is
+      // constant across the BK tile; else the per-nibble `half` kernel.
+      shader_src = (gs % kQ4gswSteelBK == 0u)
+          ? kQ4gswLinearGemmSteelHalfPwdqWGSL
+          : kQ4gswLinearGemmSteelHalfWGSL;
+    }
+  }
+  // f16-accumulate: pwdq staging with an f16 register accumulator.
+  // Lossy (f16 accumulate over K) -> opt-in via the enable_f16_accumulate_gemm
+  // runtime spec (default off), gated on the negotiated shader-f16 feature and
+  // group_size % BK == 0 (same hoisted-scale requirement as pwdq). Overrides
+  // the f32-accumulate steel kernels.
+  if (use_steel && graph.f16_accumulate_gemm() && (gs % kQ4gswSteelBK == 0u)) {
+    const WebGPUContext* ctx = get_default_webgpu_context();
+    if (ctx != nullptr && ctx->shader_f16_supported) {
+      shader_src = kQ4gswLinearGemmSteelHalfPwdqF16accWGSL;
     }
   }
   const uint32_t workgroup_count = compute_q4gsw_workgroup_count(
diff --git a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl
index 788c9ffd941..4d7ab0b1d1e 100644
--- a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl
+++ b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.wgsl
@@ -21,7 +21,16 @@ struct Params {
 // "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
 // The "steel" name + register-tiled dequant-to-shared GEMM structure are
 // inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
-// mlx/backend/metal/kernels/steel).
+// mlx/backend/metal/kernels/steel). One template, four variants:
+//   DTYPE=float           f32 storage/multiply, per-nibble weight staging.
+//   DTYPE=half            f16 storage/multiply, per-nibble weight staging.
+//   PWDQ (half only)      packed-word dequant: load each u32 weight word ONCE,
+//     unpack all 16 nibbles of a column + hoist the per-column scale to one read
+//     (the per-nibble path re-reads each word ~8x). Requires K%BK==0 (steel
+//     route guarantees it) and group_size%BK==0 (hoisted scale across the tile).
+//   ACC=half (PWDQ only)  f16 accumulate with fma(), cast to f32 in the epilogue
+//     -- LOSSY, perplexity-gated, opt-in via a runtime spec. ACC=float is f32
+//     accumulate -- BIT-EXACT to the per-nibble half kernel.
 const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
 var<workgroup> As: array<${buffer_scalar_type(DTYPE)}, 1024>;   // BM*BK
 var<workgroup> Bs: array<${buffer_scalar_type(DTYPE)}, 1024>;   // BK*BN
@@ -34,16 +43,17 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
   let row0 = by * BM;
   let col0 = bx * BN;
   let tid = lid.y * 16u + lid.x;
-  var acc: array<array<f32, 4>, 4>;
+  var acc: array<array<${buffer_scalar_type(ACC)}, 4>, 4>;
   for (var m: u32 = 0u; m < 4u; m = m + 1u) {
-    for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
+    for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = ${"0.0h" if ACC == "half" else "0.0"}; }
   }
   // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
   let ar = tid / 4u;          // 0..63 (row in tile)
   let ac = (tid % 4u) * 4u;   // 0,4,8,12 (K offset, 4 contiguous)
-  // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
-  let br = tid / 16u;         // 0..15 (K within BK)
-  let bc = (tid % 16u) * 4u;  // 0,4,..60 (N offset, 4 contiguous)
+  $if not PWDQ:
+    // B staging coords: 256 threads load 16x64 = 1024 dequant weights -> 4 cols each.
+    let br = tid / 16u;         // 0..15 (K within BK)
+    let bc = (tid % 16u) * 4u;  // 0,4,..60 (N offset, 4 contiguous)
 
   var k0: u32 = 0u;
   loop {
@@ -57,28 +67,53 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
       As[ar * BK + ac + 2u] = ${buffer_scalar_type(DTYPE)}(t_input[base + 2u]);
       As[ar * BK + ac + 3u] = ${buffer_scalar_type(DTYPE)}(t_input[base + 3u]);
     } else {
-      As[ar * BK + ac + 0u] = 0.0; As[ar * BK + ac + 1u] = 0.0;
-      As[ar * BK + ac + 2u] = 0.0; As[ar * BK + ac + 3u] = 0.0;
+      As[ar * BK + ac + 0u] = ${"0.0h" if PWDQ else "0.0"}; As[ar * BK + ac + 1u] = ${"0.0h" if PWDQ else "0.0"};
+      As[ar * BK + ac + 2u] = ${"0.0h" if PWDQ else "0.0"}; As[ar * BK + ac + 3u] = ${"0.0h" if PWDQ else "0.0"};
     }
-    // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
-    let kk = k0 + br;               // K index for this shmem row
-    let scale_row = (kk / params.group_size) * params.padded_N;
-    for (var j: u32 = 0u; j < 4u; j = j + 1u) {
-      let n = col0 + bc + j;
-      var dqv: ${buffer_scalar_type(DTYPE)} = 0.0;
-      if (n < params.N) {
-        let byte_idx = n * params.K_packed + (kk >> 1u);
-        let word = t_weight[byte_idx >> 2u];
-        let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
-        var nib: u32;
-        if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
-        $if DTYPE == "half":
-          dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
-        $else:
-          dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
+    $if PWDQ:
+      // Packed-word dequant: threads [0,BN) each stage one full BK-column of Bs.
+      if (tid < BN) {
+        let c = tid;                 // Bs column within this tile
+        let n = col0 + c;            // global output column
+        if (n < params.N) {
+          // Scale is constant across the BK tile (group_size % BK == 0 for all real
+          // group sizes; K%BK==0 on the steel route), so hoist it to one read.
+          let scale_row = (k0 / params.group_size) * params.padded_N;
+          let scale = f16(t_scales[scale_row + n]);
+          // Column n's 16-nibble K-slice for this tile = two consecutive words.
+          // K_packed multiple of 8 => base_word stays inside column n's own region.
+          let base_word = n * (params.K_packed >> 2u) + (k0 >> 3u);
+          let w0 = t_weight[base_word];
+          let w1 = t_weight[base_word + 1u];
+          for (var br: u32 = 0u; br < BK; br = br + 1u) {
+            let word = select(w1, w0, br < 8u);        // word0 holds K-slice [0,8)
+            let nib = (word >> ((br & 7u) * 4u)) & 0x0Fu;
+            Bs[br * BN + c] = f16(i32(nib) - 8) * scale;
+          }
+        } else {
+          for (var br: u32 = 0u; br < BK; br = br + 1u) { Bs[br * BN + c] = 0.0h; }
+        }
+      }
+    $else:
+      // stage DEQUANTIZED weights into Bs[k][n]: 4 contiguous N per thread.
+      let kk = k0 + br;               // K index for this shmem row
+      let scale_row = (kk / params.group_size) * params.padded_N;
+      for (var j: u32 = 0u; j < 4u; j = j + 1u) {
+        let n = col0 + bc + j;
+        var dqv: ${buffer_scalar_type(DTYPE)} = 0.0;
+        if (n < params.N) {
+          let byte_idx = n * params.K_packed + (kk >> 1u);
+          let word = t_weight[byte_idx >> 2u];
+          let b = (word >> ((byte_idx & 3u) * 8u)) & 0xFFu;
+          var nib: u32;
+          if ((kk & 1u) == 0u) { nib = b & 0x0Fu; } else { nib = (b >> 4u) & 0x0Fu; }
+          $if DTYPE == "half":
+            dqv = f16(i32(nib) - 8) * f16(t_scales[scale_row + n]);
+          $else:
+            dqv = f32(i32(nib) - 8) * t_scales[scale_row + n];
+        }
+        Bs[br * BN + bc + j] = dqv;
       }
-      Bs[br * BN + bc + j] = dqv;
-    }
     workgroupBarrier();
     for (var k: u32 = 0u; k < BK; k = k + 1u) {
       var a: array<${buffer_scalar_type(DTYPE)}, 4>;
@@ -86,7 +121,9 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
       for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
       for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
       for (var m: u32 = 0u; m < 4u; m = m + 1u) {
-        $if DTYPE == "half":
+        $if ACC == "half":
+          for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = fma(a[m], bvec[n], acc[m][n]); }
+        $elif DTYPE == "half":
           for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
         $else:
           for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + a[m] * bvec[n]; }
@@ -100,7 +137,7 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
       let r = row0 + lid.y * 4u + m;
       let c = col0 + lid.x * 4u + n;
       if (r < params.M && c < params.N) {
-        var v = acc[m][n];
+        var v = ${"f32(acc[m][n])" if ACC == "half" else "acc[m][n]"};
         if (params.has_bias != 0u) { v = v + t_bias[c]; }
         t_out[r * params.N + c] = v;
       }
diff --git a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml
index 1505d8b9924..5a2cae5e499 100644
--- a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml
+++ b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel.yaml
@@ -1,11 +1,22 @@
 q4gsw_linear_gemm_steel:
   parameter_names_with_default_values:
     DTYPE: float
-  generate_variant_forall:
-    DTYPE:
-      - VALUE: float
-        SUFFIX: ""
-      - VALUE: half
-        SUFFIX: half
+    PWDQ: false
+    ACC: float
   shader_variants:
     - NAME: q4gsw_linear_gemm_steel
+      DTYPE: float
+      PWDQ: false
+      ACC: float
+    - NAME: q4gsw_linear_gemm_steel_half
+      DTYPE: half
+      PWDQ: false
+      ACC: float
+    - NAME: q4gsw_linear_gemm_steel_half_pwdq
+      DTYPE: half
+      PWDQ: true
+      ACC: float
+    - NAME: q4gsw_linear_gemm_steel_half_pwdq_f16acc
+      DTYPE: half
+      PWDQ: true
+      ACC: half
diff --git a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_f16acc_wgsl.h b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_f16acc_wgsl.h
new file mode 100644
index 00000000000..efefd7edce1
--- /dev/null
+++ b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_f16acc_wgsl.h
@@ -0,0 +1,141 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#pragma once
+
+#include <cstdint>
+
+namespace executorch::backends::webgpu {
+
+// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
+// wgsl-sha256: 36b3d3f9dd08a529909c13ec7d66cd0cf392c347ca047a4d38453b3c295f72ce
+inline constexpr const char* kQ4gswLinearGemmSteelHalfPwdqF16accWGSL = R"(
+enable f16;
+@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
+@group(0) @binding(1) var<storage, read> t_input: array<f32>;
+@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
+@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
+@group(0) @binding(4) var<storage, read> t_bias: array<f32>;
+
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  K_packed: u32,
+  group_size: u32,
+  padded_N: u32,
+  has_bias: u32,
+  _pad: u32,
+}
+@group(0) @binding(5) var<uniform> params: Params;
+
+// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
+// The "steel" name + register-tiled dequant-to-shared GEMM structure are
+// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
+// mlx/backend/metal/kernels/steel). One template, four variants:
+//   DTYPE=float           f32 storage/multiply, per-nibble weight staging.
+//   DTYPE=half            f16 storage/multiply, per-nibble weight staging.
+//   PWDQ (half only)      packed-word dequant: load each u32 weight word ONCE,
+//     unpack all 16 nibbles of a column + hoist the per-column scale to one read
+//     (the per-nibble path re-reads each word ~8x). Requires K%BK==0 (steel
+//     route guarantees it) and group_size%BK==0 (hoisted scale across the tile).
+//   ACC=half (PWDQ only)  f16 accumulate with fma(), cast to f32 in the epilogue
+//     -- LOSSY, perplexity-gated, opt-in via a runtime spec. ACC=float is f32
+//     accumulate -- BIT-EXACT to the per-nibble half kernel.
+const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
+var<workgroup> As: array<f16, 1024>;   // BM*BK
+var<workgroup> Bs: array<f16, 1024>;   // BK*BN
+@compute @workgroup_size(16, 16)
+fn main(@builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  let nbN = (params.N + BN - 1u) / BN;
+  let bx = wid.x % nbN;      // decode 2D tile id from 1D dispatch
+  let by = wid.x / nbN;
+  let row0 = by * BM;
+  let col0 = bx * BN;
+  let tid = lid.y * 16u + lid.x;
+  var acc: array<array<f16, 4>, 4>;
+  for (var m: u32 = 0u; m < 4u; m = m + 1u) {
+    for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0h; }
+  }
+  // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
+  let ar = tid / 4u;          // 0..63 (row in tile)
+  let ac = (tid % 4u) * 4u;   // 0,4,8,12 (K offset, 4 contiguous)
+
+  var k0: u32 = 0u;
+  loop {
+    if (k0 >= params.K) { break; }
+    // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
+    let arow = row0 + ar;
+    if (arow < params.M) {
+      let base = arow * params.K + k0 + ac;
+      As[ar * BK + ac + 0u] = f16(t_input[base]);
+      As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
+      As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
+      As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
+    } else {
+      As[ar * BK + ac + 0u] = 0.0h; As[ar * BK + ac + 1u] = 0.0h;
+      As[ar * BK + ac + 2u] = 0.0h; As[ar * BK + ac + 3u] = 0.0h;
+    }
+    // Packed-word dequant: threads [0,BN) each stage one full BK-column of Bs.
+    if (tid < BN) {
+      let c = tid;                 // Bs column within this tile
+      let n = col0 + c;            // global output column
+      if (n < params.N) {
+        // Scale is constant across the BK tile (group_size % BK == 0 for all real
+        // group sizes; K%BK==0 on the steel route), so hoist it to one read.
+        let scale_row = (k0 / params.group_size) * params.padded_N;
+        let scale = f16(t_scales[scale_row + n]);
+        // Column n's 16-nibble K-slice for this tile = two consecutive words.
+        // K_packed multiple of 8 => base_word stays inside column n's own region.
+        let base_word = n * (params.K_packed >> 2u) + (k0 >> 3u);
+        let w0 = t_weight[base_word];
+        let w1 = t_weight[base_word + 1u];
+        for (var br: u32 = 0u; br < BK; br = br + 1u) {
+          let word = select(w1, w0, br < 8u);        // word0 holds K-slice [0,8)
+          let nib = (word >> ((br & 7u) * 4u)) & 0x0Fu;
+          Bs[br * BN + c] = f16(i32(nib) - 8) * scale;
+        }
+      } else {
+        for (var br: u32 = 0u; br < BK; br = br + 1u) { Bs[br * BN + c] = 0.0h; }
+      }
+    }
+    workgroupBarrier();
+    for (var k: u32 = 0u; k < BK; k = k + 1u) {
+      var a: array<f16, 4>;
+      var bvec: array<f16, 4>;
+      for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
+      for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
+      for (var m: u32 = 0u; m < 4u; m = m + 1u) {
+        for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = fma(a[m], bvec[n], acc[m][n]); }
+      }
+    }
+    workgroupBarrier();
+    k0 = k0 + BK;
+  }
+  for (var m: u32 = 0u; m < 4u; m = m + 1u) {
+    for (var n: u32 = 0u; n < 4u; n = n + 1u) {
+      let r = row0 + lid.y * 4u + m;
+      let c = col0 + lid.x * 4u + n;
+      if (r < params.M && c < params.N) {
+        var v = f32(acc[m][n]);
+        if (params.has_bias != 0u) { v = v + t_bias[c]; }
+        t_out[r * params.N + c] = v;
+      }
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kQ4gswLinearGemmSteelHalfPwdqF16accWorkgroupSizeX =
+    16;
+inline constexpr uint32_t kQ4gswLinearGemmSteelHalfPwdqF16accWorkgroupSizeY =
+    16;
+inline constexpr uint32_t kQ4gswLinearGemmSteelHalfPwdqF16accWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_wgsl.h b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_wgsl.h
new file mode 100644
index 00000000000..46057de2340
--- /dev/null
+++ b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_pwdq_wgsl.h
@@ -0,0 +1,139 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#pragma once
+
+#include <cstdint>
+
+namespace executorch::backends::webgpu {
+
+// @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
+// wgsl-sha256: 1f916bcd30dbbbcc7eca37e795ecc26e3c72e645ccd2c361fa0ac4e66f1a174a
+inline constexpr const char* kQ4gswLinearGemmSteelHalfPwdqWGSL = R"(
+enable f16;
+@group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
+@group(0) @binding(1) var<storage, read> t_input: array<f32>;
+@group(0) @binding(2) var<storage, read> t_weight: array<u32>;
+@group(0) @binding(3) var<storage, read> t_scales: array<f32>;
+@group(0) @binding(4) var<storage, read> t_bias: array<f32>;
+
+struct Params {
+  M: u32,
+  N: u32,
+  K: u32,
+  K_packed: u32,
+  group_size: u32,
+  padded_N: u32,
+  has_bias: u32,
+  _pad: u32,
+}
+@group(0) @binding(5) var<uniform> params: Params;
+
+// "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
+// The "steel" name + register-tiled dequant-to-shared GEMM structure are
+// inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
+// mlx/backend/metal/kernels/steel). One template, four variants:
+//   DTYPE=float           f32 storage/multiply, per-nibble weight staging.
+//   DTYPE=half            f16 storage/multiply, per-nibble weight staging.
+//   PWDQ (half only)      packed-word dequant: load each u32 weight word ONCE,
+//     unpack all 16 nibbles of a column + hoist the per-column scale to one read
+//     (the per-nibble path re-reads each word ~8x). Requires K%BK==0 (steel
+//     route guarantees it) and group_size%BK==0 (hoisted scale across the tile).
+//   ACC=half (PWDQ only)  f16 accumulate with fma(), cast to f32 in the epilogue
+//     -- LOSSY, perplexity-gated, opt-in via a runtime spec. ACC=float is f32
+//     accumulate -- BIT-EXACT to the per-nibble half kernel.
+const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
+var<workgroup> As: array<f16, 1024>;   // BM*BK
+var<workgroup> Bs: array<f16, 1024>;   // BK*BN
+@compute @workgroup_size(16, 16)
+fn main(@builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  let nbN = (params.N + BN - 1u) / BN;
+  let bx = wid.x % nbN;      // decode 2D tile id from 1D dispatch
+  let by = wid.x / nbN;
+  let row0 = by * BM;
+  let col0 = bx * BN;
+  let tid = lid.y * 16u + lid.x;
+  var acc: array<array<f32, 4>, 4>;
+  for (var m: u32 = 0u; m < 4u; m = m + 1u) {
+    for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = 0.0; }
+  }
+  // A staging coords: 256 threads load 64x16 = 1024 f32 -> 4 rows each (4 contiguous K).
+  let ar = tid / 4u;          // 0..63 (row in tile)
+  let ac = (tid % 4u) * 4u;   // 0,4,8,12 (K offset, 4 contiguous)
+
+  var k0: u32 = 0u;
+  loop {
+    if (k0 >= params.K) { break; }
+    // stage activations (edge-masked on M; K is a multiple of BK for our shapes)
+    let arow = row0 + ar;
+    if (arow < params.M) {
+      let base = arow * params.K + k0 + ac;
+      As[ar * BK + ac + 0u] = f16(t_input[base]);
+      As[ar * BK + ac + 1u] = f16(t_input[base + 1u]);
+      As[ar * BK + ac + 2u] = f16(t_input[base + 2u]);
+      As[ar * BK + ac + 3u] = f16(t_input[base + 3u]);
+    } else {
+      As[ar * BK + ac + 0u] = 0.0h; As[ar * BK + ac + 1u] = 0.0h;
+      As[ar * BK + ac + 2u] = 0.0h; As[ar * BK + ac + 3u] = 0.0h;
+    }
+    // Packed-word dequant: threads [0,BN) each stage one full BK-column of Bs.
+    if (tid < BN) {
+      let c = tid;                 // Bs column within this tile
+      let n = col0 + c;            // global output column
+      if (n < params.N) {
+        // Scale is constant across the BK tile (group_size % BK == 0 for all real
+        // group sizes; K%BK==0 on the steel route), so hoist it to one read.
+        let scale_row = (k0 / params.group_size) * params.padded_N;
+        let scale = f16(t_scales[scale_row + n]);
+        // Column n's 16-nibble K-slice for this tile = two consecutive words.
+        // K_packed multiple of 8 => base_word stays inside column n's own region.
+        let base_word = n * (params.K_packed >> 2u) + (k0 >> 3u);
+        let w0 = t_weight[base_word];
+        let w1 = t_weight[base_word + 1u];
+        for (var br: u32 = 0u; br < BK; br = br + 1u) {
+          let word = select(w1, w0, br < 8u);        // word0 holds K-slice [0,8)
+          let nib = (word >> ((br & 7u) * 4u)) & 0x0Fu;
+          Bs[br * BN + c] = f16(i32(nib) - 8) * scale;
+        }
+      } else {
+        for (var br: u32 = 0u; br < BK; br = br + 1u) { Bs[br * BN + c] = 0.0h; }
+      }
+    }
+    workgroupBarrier();
+    for (var k: u32 = 0u; k < BK; k = k + 1u) {
+      var a: array<f16, 4>;
+      var bvec: array<f16, 4>;
+      for (var m: u32 = 0u; m < 4u; m = m + 1u) { a[m] = As[(lid.y * 4u + m) * BK + k]; }
+      for (var n: u32 = 0u; n < 4u; n = n + 1u) { bvec[n] = Bs[k * BN + lid.x * 4u + n]; }
+      for (var m: u32 = 0u; m < 4u; m = m + 1u) {
+        for (var n: u32 = 0u; n < 4u; n = n + 1u) { acc[m][n] = acc[m][n] + f32(a[m] * bvec[n]); }
+      }
+    }
+    workgroupBarrier();
+    k0 = k0 + BK;
+  }
+  for (var m: u32 = 0u; m < 4u; m = m + 1u) {
+    for (var n: u32 = 0u; n < 4u; n = n + 1u) {
+      let r = row0 + lid.y * 4u + m;
+      let c = col0 + lid.x * 4u + n;
+      if (r < params.M && c < params.N) {
+        var v = acc[m][n];
+        if (params.has_bias != 0u) { v = v + t_bias[c]; }
+        t_out[r * params.N + c] = v;
+      }
+    }
+  }
+}
+)";
+
+inline constexpr uint32_t kQ4gswLinearGemmSteelHalfPwdqWorkgroupSizeX = 16;
+inline constexpr uint32_t kQ4gswLinearGemmSteelHalfPwdqWorkgroupSizeY = 16;
+inline constexpr uint32_t kQ4gswLinearGemmSteelHalfPwdqWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h
index 7e9363e3b36..ae03f63fb5c 100644
--- a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h
+++ b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_half_wgsl.h
@@ -13,7 +13,7 @@
 namespace executorch::backends::webgpu {
 
 // @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
-// wgsl-sha256: e3c21e7db7c18f6e085de71e283988f0bd3b2543807ddc17774a1c607e69c766
+// wgsl-sha256: 00cdd2f2fb98a5c7343d16fdf7e59f1b840e180cec3f82bf9b569513c0a45396
 inline constexpr const char* kQ4gswLinearGemmSteelHalfWGSL = R"(
 enable f16;
 @group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
@@ -37,7 +37,16 @@ struct Params {
 // "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
 // The "steel" name + register-tiled dequant-to-shared GEMM structure are
 // inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
-// mlx/backend/metal/kernels/steel).
+// mlx/backend/metal/kernels/steel). One template, four variants:
+//   DTYPE=float           f32 storage/multiply, per-nibble weight staging.
+//   DTYPE=half            f16 storage/multiply, per-nibble weight staging.
+//   PWDQ (half only)      packed-word dequant: load each u32 weight word ONCE,
+//     unpack all 16 nibbles of a column + hoist the per-column scale to one read
+//     (the per-nibble path re-reads each word ~8x). Requires K%BK==0 (steel
+//     route guarantees it) and group_size%BK==0 (hoisted scale across the tile).
+//   ACC=half (PWDQ only)  f16 accumulate with fma(), cast to f32 in the epilogue
+//     -- LOSSY, perplexity-gated, opt-in via a runtime spec. ACC=float is f32
+//     accumulate -- BIT-EXACT to the per-nibble half kernel.
 const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
 var<workgroup> As: array<f16, 1024>;   // BM*BK
 var<workgroup> Bs: array<f16, 1024>;   // BK*BN
diff --git a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h
index 9e73a0c9a66..71b0f45bdf7 100644
--- a/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h
+++ b/backends/webgpu/runtime/ops/quantized_linear/q4gsw_linear_gemm_steel_wgsl.h
@@ -13,7 +13,7 @@
 namespace executorch::backends::webgpu {
 
 // @generated from q4gsw_linear_gemm_steel.wgsl - DO NOT EDIT.
-// wgsl-sha256: 43536b16026d3b62d77087b86289885606c04834d9783c3c871512d6789ee6f6
+// wgsl-sha256: dd771b9ab096410f3ad0d9259bef7816e41330a434325ea28baa5abbfb2841d2
 inline constexpr const char* kQ4gswLinearGemmSteelWGSL = R"(
 @group(0) @binding(0) var<storage, read_write> t_out: array<f32>;
 @group(0) @binding(1) var<storage, read> t_input: array<f32>;
@@ -36,7 +36,16 @@ struct Params {
 // "steel" prefill GEMM (M>1): 64x64 tile, 256 threads; K%16==0 host-guarded.
 // The "steel" name + register-tiled dequant-to-shared GEMM structure are
 // inspired by MLX's steel GEMM kernels (github.com/ml-explore/mlx,
-// mlx/backend/metal/kernels/steel).
+// mlx/backend/metal/kernels/steel). One template, four variants:
+//   DTYPE=float           f32 storage/multiply, per-nibble weight staging.
+//   DTYPE=half            f16 storage/multiply, per-nibble weight staging.
+//   PWDQ (half only)      packed-word dequant: load each u32 weight word ONCE,
+//     unpack all 16 nibbles of a column + hoist the per-column scale to one read
+//     (the per-nibble path re-reads each word ~8x). Requires K%BK==0 (steel
+//     route guarantees it) and group_size%BK==0 (hoisted scale across the tile).
+//   ACC=half (PWDQ only)  f16 accumulate with fma(), cast to f32 in the epilogue
+//     -- LOSSY, perplexity-gated, opt-in via a runtime spec. ACC=float is f32
+//     accumulate -- BIT-EXACT to the per-nibble half kernel.
 const BM: u32 = 64u; const BN: u32 = 64u; const BK: u32 = 16u;
 var<workgroup> As: array<f32, 1024>;   // BM*BK
 var<workgroup> Bs: array<f32, 1024>;   // BK*BN
diff --git a/backends/webgpu/runtime/ops/sdpa/Sdpa.cpp b/backends/webgpu/runtime/ops/sdpa/Sdpa.cpp
index 17918863a6e..50321ba4bdf 100644
--- a/backends/webgpu/runtime/ops/sdpa/Sdpa.cpp
+++ b/backends/webgpu/runtime/ops/sdpa/Sdpa.cpp
@@ -9,10 +9,13 @@
 #include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
 #include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
 #include <executorch/backends/webgpu/runtime/ops/OperatorRegistry.h>
+#include <executorch/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights_half_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out_half_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa/sdpa_softmax_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa_fd_decode/SdpaFdDecode.h>
+#include <executorch/backends/webgpu/runtime/ops/update_cache/update_cache_half_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/update_cache/update_cache_wgsl.h>
 
 #include <webgpu/webgpu.h>
@@ -255,9 +258,13 @@ static WGPUBuffer record_update_cache_dispatch(
   WGPUBuffer ubuf = graph.make_uniform_buffer(&uc, sizeof(uc));
   BufferBinding bindings[2] = {
       {cache.buffer, cache.nbytes}, {src.buffer, src.nbytes}};
+  const char* uc_src = kUpdateCacheWGSL;
+  if (graph.kv_f16()) {
+    uc_src = kUpdateCacheHalfWGSL;
+  }
   build_dispatch(
       graph,
-      kUpdateCacheWGSL,
+      uc_src,
       bindings,
       2,
       ubuf,
@@ -474,8 +481,11 @@ void sdpa_with_kv_cache_impl(WebGPUGraph& graph, const std::vector<int>& args) {
   }
 
   // QK/softmax scratch — allocated only on the non-FD path (Hq*S*Cmax prefill).
-  WGPUBuffer attn_weights = graph.create_scratch_buffer(aw_bytes);
-  WGPUBuffer attn_weights_softmax = graph.create_scratch_buffer(aw_bytes);
+  WGPUBuffer attn_weights = graph.acquire_scratch(aw_bytes);
+  WebGPUGraph::ScopedScratch attn_weights_guard(&graph, attn_weights);
+  WGPUBuffer attn_weights_softmax = graph.acquire_scratch(aw_bytes);
+  WebGPUGraph::ScopedScratch attn_weights_softmax_guard(
+      &graph, attn_weights_softmax);
 
   // --- Dispatch 3: QK -> attn_weights. One thread per TM x TN tile.
   {
@@ -494,9 +504,13 @@ void sdpa_with_kv_cache_impl(WebGPUGraph& graph, const std::vector<int>& args) {
         {attn_weights, aw_bytes},
         {q.buffer, q.nbytes},
         {k_cache.buffer, k_cache.nbytes}};
+    const char* qk_src = kSdpaComputeAttnWeightsWGSL;
+    if (graph.kv_f16()) {
+      qk_src = kSdpaComputeAttnWeightsHalfWGSL;
+    }
     build_dispatch(
         graph,
-        kSdpaComputeAttnWeightsWGSL,
+        qk_src,
         bindings,
         3,
         ubuf,
@@ -547,9 +561,13 @@ void sdpa_with_kv_cache_impl(WebGPUGraph& graph, const std::vector<int>& args) {
         {out.buffer, out.nbytes},
         {attn_weights_softmax, aw_bytes},
         {v_cache.buffer, v_cache.nbytes}};
+    const char* av_src = kSdpaComputeOutWGSL;
+    if (graph.kv_f16()) {
+      av_src = kSdpaComputeOutHalfWGSL;
+    }
     build_dispatch(
         graph,
-        kSdpaComputeOutWGSL,
+        av_src,
         bindings,
         3,
         ubuf,
diff --git a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.wgsl b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.wgsl
index 014f0039048..097d87ecce0 100644
--- a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.wgsl
+++ b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.wgsl
@@ -1,6 +1,8 @@
+$if DTYPE == "half":
+  enable f16;
 @group(0) @binding(0) var<storage, read_write> t_attn_weights: array<f32>;
 @group(0) @binding(1) var<storage, read> t_q: array<vec4<f32>>;
-@group(0) @binding(2) var<storage, read> t_k_cache: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read> t_k_cache: array<${buffer_gvec_type(DTYPE, 4)}>;
 
 struct Params {
   S: u32,
@@ -36,7 +38,10 @@ fn load_k_vec4(c: u32, kvh: u32, d4: u32) -> vec4<f32> {
     return vec4<f32>(0.0, 0.0, 0.0, 0.0);
   }
   let base = c * params.Hkv * params.D + kvh * params.D + d4;
-  return t_k_cache[base / 4u];
+  $if DTYPE == "half":
+    return vec4<f32>(t_k_cache[base / 4u]);
+  $else:
+    return t_k_cache[base / 4u];
 }
 
 fn store_qk(s: u32, c: u32, h: u32, raw: f32) {
diff --git a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.yaml b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.yaml
new file mode 100644
index 00000000000..d031e2864f4
--- /dev/null
+++ b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights.yaml
@@ -0,0 +1,11 @@
+sdpa_compute_attn_weights:
+  parameter_names_with_default_values:
+    DTYPE: float
+  generate_variant_forall:
+    DTYPE:
+      - VALUE: float
+        SUFFIX: ""
+      - VALUE: half
+        SUFFIX: half
+  shader_variants:
+    - NAME: sdpa_compute_attn_weights
diff --git a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights_half_wgsl.h b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights_half_wgsl.h
new file mode 100644
index 00000000000..dc6f5858c7d
--- /dev/null
+++ b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_attn_weights_half_wgsl.h
@@ -0,0 +1,145 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#pragma once
+
+#include <cstdint>
+
+namespace executorch::backends::webgpu {
+
+// @generated from sdpa_compute_attn_weights.wgsl - DO NOT EDIT.
+// wgsl-sha256: c8795d66b9b51516795fb0113b21fd55086b4a54d9a4b81c7f394ffd96d117b3
+inline constexpr const char* kSdpaComputeAttnWeightsHalfWGSL = R"(
+enable f16;
+@group(0) @binding(0) var<storage, read_write> t_attn_weights: array<f32>;
+@group(0) @binding(1) var<storage, read> t_q: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read> t_k_cache: array<vec4<f16>>;
+
+struct Params {
+  S: u32,
+  Hq: u32,
+  Hkv: u32,
+  D: u32,
+  context_len: u32,
+  input_pos: u32,
+  g: u32,
+  scale: f32,
+}
+@group(0) @binding(3) var<uniform> params: Params;
+
+// WGSL forbids literal -inf; large finite negative is a WGSL-safe stand-in.
+const NEG_INF: f32 = -1.0e30;
+
+override wg_size: u32 = 64;
+
+const TM: u32 = 4u;
+const TN: u32 = 4u;
+
+// D is a multiple of 4 (host-guarded), so a d4 chunk is fully in-bounds — no per-lane check.
+fn load_q_vec4(s: u32, h: u32, d4: u32) -> vec4<f32> {
+  if (s >= params.S) {
+    return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  }
+  let base = s * params.Hq * params.D + h * params.D + d4;
+  return t_q[base / 4u];
+}
+
+fn load_k_vec4(c: u32, kvh: u32, d4: u32) -> vec4<f32> {
+  if (c >= params.context_len) {
+    return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  }
+  let base = c * params.Hkv * params.D + kvh * params.D + d4;
+  return vec4<f32>(t_k_cache[base / 4u]);
+}
+
+fn store_qk(s: u32, c: u32, h: u32, raw: f32) {
+  if (s >= params.S || c >= params.context_len) {
+    return;
+  }
+  var val = raw * params.scale;
+  // Causal mask: position c may not attend beyond s + input_pos.
+  if (c > s + params.input_pos) {
+    val = NEG_INF;
+  }
+  let idx = h * params.S * params.context_len + s * params.context_len + c;
+  t_attn_weights[idx] = val;
+}
+
+@compute @workgroup_size(wg_size, 1, 1)
+fn main(
+    @builtin(global_invocation_id) gid: vec3<u32>,
+    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
+  let nrt = (params.S + TM - 1u) / TM;
+  let nct = (params.context_len + TN - 1u) / TN;
+  let tiles = nrt * nct;
+  let total = tiles * params.Hq;
+  // 2D dispatch fold: recover the linear tile index across x/y.
+  let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
+  if (idx >= total) {
+    return;
+  }
+
+  let h = idx / tiles;
+  let rem = idx % tiles;
+  let row_tile = rem / nct;
+  let col_tile = rem % nct;
+  let kvh = h / params.g;
+  let s0 = row_tile * TM;
+  let c0 = col_tile * TN;
+
+  var acc: array<vec4<f32>, 4>;
+  acc[0] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  acc[1] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  acc[2] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  acc[3] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+
+  // Skip fully-masked causal tiles; mirrors Vulkan attn_weights_tiled.glsl.
+  let skip_tile = c0 > s0 + (TM - 1u) + params.input_pos;
+  var d4: u32 = 0u;
+  loop {
+    if (d4 >= params.D || skip_tile) {
+      break;
+    }
+    var q: array<vec4<f32>, TM>;
+    var k: array<vec4<f32>, TN>;
+    for (var i: u32 = 0u; i < TM; i = i + 1u) {
+      q[i] = load_q_vec4(s0 + i, h, d4);
+    }
+    for (var j: u32 = 0u; j < TN; j = j + 1u) {
+      k[j] = load_k_vec4(c0 + j, kvh, d4);
+    }
+    for (var i: u32 = 0u; i < TM; i = i + 1u) {
+      acc[i] += vec4<f32>(
+          dot(q[i], k[0]),
+          dot(q[i], k[1]),
+          dot(q[i], k[2]),
+          dot(q[i], k[3]));
+    }
+    d4 = d4 + 4u;
+  }
+
+  var m: u32 = 0u;
+  loop {
+    if (m >= TM) {
+      break;
+    }
+    let av = acc[m];
+    store_qk(s0 + m, c0 + 0u, h, av.x);
+    store_qk(s0 + m, c0 + 1u, h, av.y);
+    store_qk(s0 + m, c0 + 2u, h, av.z);
+    store_qk(s0 + m, c0 + 3u, h, av.w);
+    m = m + 1u;
+  }
+}
+)";
+
+inline constexpr uint32_t kSdpaComputeAttnWeightsHalfWorkgroupSizeX = 64;
+inline constexpr uint32_t kSdpaComputeAttnWeightsHalfWorkgroupSizeY = 1;
+inline constexpr uint32_t kSdpaComputeAttnWeightsHalfWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.wgsl b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.wgsl
index 713345c0afa..53067447832 100644
--- a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.wgsl
+++ b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.wgsl
@@ -1,6 +1,8 @@
+$if DTYPE == "half":
+  enable f16;
 @group(0) @binding(0) var<storage, read_write> t_out: array<vec4<f32>>;
 @group(0) @binding(1) var<storage, read> t_attn_weights_softmax: array<f32>;
-@group(0) @binding(2) var<storage, read> t_v_cache: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read> t_v_cache: array<${buffer_gvec_type(DTYPE, 4)}>;
 
 struct Params {
   S: u32,
@@ -38,7 +40,10 @@ fn load_v_d4(c: u32, kvh: u32, d0: u32) -> vec4<f32> {
     return vec4<f32>(0.0, 0.0, 0.0, 0.0);
   }
   let base = c * params.Hkv * params.D + kvh * params.D + d0;
-  return t_v_cache[base / 4u];
+  $if DTYPE == "half":
+    return vec4<f32>(t_v_cache[base / 4u]);
+  $else:
+    return t_v_cache[base / 4u];
 }
 
 // Branch-free loaders for the aligned body: caller guarantees c4..c4+3 < context_len.
@@ -52,7 +57,10 @@ fn load_a_vec4_nc(s: u32, h: u32, c4: u32) -> vec4<f32> {
 
 fn load_v_d4_nc(c: u32, kvh: u32, d0: u32) -> vec4<f32> {
   let base = c * params.Hkv * params.D + kvh * params.D + d0;
-  return t_v_cache[base / 4u];
+  $if DTYPE == "half":
+    return vec4<f32>(t_v_cache[base / 4u]);
+  $else:
+    return t_v_cache[base / 4u];
 }
 
 fn store_out_vec4(s: u32, d0: u32, h: u32, val: vec4<f32>) {
diff --git a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.yaml b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.yaml
new file mode 100644
index 00000000000..14b418f8a60
--- /dev/null
+++ b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out.yaml
@@ -0,0 +1,11 @@
+sdpa_compute_out:
+  parameter_names_with_default_values:
+    DTYPE: float
+  generate_variant_forall:
+    DTYPE:
+      - VALUE: float
+        SUFFIX: ""
+      - VALUE: half
+        SUFFIX: half
+  shader_variants:
+    - NAME: sdpa_compute_out
diff --git a/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out_half_wgsl.h b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out_half_wgsl.h
new file mode 100644
index 00000000000..228921af58b
--- /dev/null
+++ b/backends/webgpu/runtime/ops/sdpa/sdpa_compute_out_half_wgsl.h
@@ -0,0 +1,163 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#pragma once
+
+#include <cstdint>
+
+namespace executorch::backends::webgpu {
+
+// @generated from sdpa_compute_out.wgsl - DO NOT EDIT.
+// wgsl-sha256: ed9709c966538edf2cbc6be97c284b89a9d921b6a4dbf115c6cbd76af301a1be
+inline constexpr const char* kSdpaComputeOutHalfWGSL = R"(
+enable f16;
+@group(0) @binding(0) var<storage, read_write> t_out: array<vec4<f32>>;
+@group(0) @binding(1) var<storage, read> t_attn_weights_softmax: array<f32>;
+@group(0) @binding(2) var<storage, read> t_v_cache: array<vec4<f16>>;
+
+struct Params {
+  S: u32,
+  Hq: u32,
+  Hkv: u32,
+  D: u32,
+  context_len: u32,
+  g: u32,
+  _pad0: u32,
+  _pad1: u32,
+}
+@group(0) @binding(3) var<uniform> params: Params;
+
+override wg_size: u32 = 64;
+
+const TM: u32 = 4u;
+const TN: u32 = 4u;
+
+// Checked loaders mask context lanes past context_len (D%4==0, host-guarded).
+fn load_a_vec4(s: u32, h: u32, c4: u32) -> vec4<f32> {
+  var r = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  if (s >= params.S) {
+    return r;
+  }
+  let base = h * params.S * params.context_len + s * params.context_len;
+  if (c4 + 0u < params.context_len) { r.x = t_attn_weights_softmax[base + c4 + 0u]; }
+  if (c4 + 1u < params.context_len) { r.y = t_attn_weights_softmax[base + c4 + 1u]; }
+  if (c4 + 2u < params.context_len) { r.z = t_attn_weights_softmax[base + c4 + 2u]; }
+  if (c4 + 3u < params.context_len) { r.w = t_attn_weights_softmax[base + c4 + 3u]; }
+  return r;
+}
+
+fn load_v_d4(c: u32, kvh: u32, d0: u32) -> vec4<f32> {
+  if (c >= params.context_len) {
+    return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  }
+  let base = c * params.Hkv * params.D + kvh * params.D + d0;
+  return vec4<f32>(t_v_cache[base / 4u]);
+}
+
+// Branch-free loaders for the aligned body: caller guarantees c4..c4+3 < context_len.
+fn load_a_vec4_nc(s: u32, h: u32, c4: u32) -> vec4<f32> {
+  if (s >= params.S) {
+    return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  }
+  let base = h * params.S * params.context_len + s * params.context_len + c4;
+  return vec4<f32>(t_attn_weights_softmax[base], t_attn_weights_softmax[base + 1u], t_attn_weights_softmax[base + 2u], t_attn_weights_softmax[base + 3u]);
+}
+
+fn load_v_d4_nc(c: u32, kvh: u32, d0: u32) -> vec4<f32> {
+  let base = c * params.Hkv * params.D + kvh * params.D + d0;
+  return vec4<f32>(t_v_cache[base / 4u]);
+}
+
+fn store_out_vec4(s: u32, d0: u32, h: u32, val: vec4<f32>) {
+  if (s >= params.S) {
+    return;
+  }
+  let idx = s * params.Hq * params.D + h * params.D + d0;
+  t_out[idx / 4u] = val;
+}
+
+@compute @workgroup_size(wg_size, 1, 1)
+fn main(
+    @builtin(global_invocation_id) gid: vec3<u32>,
+    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
+  let nrt = (params.S + TM - 1u) / TM;
+  let nct = (params.D + TN - 1u) / TN;
+  let tiles = nrt * nct;
+  let total = tiles * params.Hq;
+  // 2D dispatch fold: recover the linear tile index across x/y.
+  let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
+  if (idx >= total) {
+    return;
+  }
+
+  let h = idx / tiles;
+  let rem = idx % tiles;
+  let row_tile = rem / nct;
+  let col_tile = rem % nct;
+  let kvh = h / params.g;
+  let s0 = row_tile * TM;
+  let d0 = col_tile * TN;
+
+  var acc: array<vec4<f32>, 4>;
+  acc[0] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  acc[1] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  acc[2] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+  acc[3] = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+
+  // Branch-free aligned body + checked tail; mirrors Vulkan out_tiled.glsl.
+  let ctx_aligned = params.context_len - (params.context_len & 3u);
+  var c4: u32 = 0u;
+  loop {
+    if (c4 >= ctx_aligned) {
+      break;
+    }
+    let a0 = load_a_vec4_nc(s0 + 0u, h, c4);
+    let a1 = load_a_vec4_nc(s0 + 1u, h, c4);
+    let a2 = load_a_vec4_nc(s0 + 2u, h, c4);
+    let a3 = load_a_vec4_nc(s0 + 3u, h, c4);
+    let v0 = load_v_d4_nc(c4 + 0u, kvh, d0);
+    let v1 = load_v_d4_nc(c4 + 1u, kvh, d0);
+    let v2 = load_v_d4_nc(c4 + 2u, kvh, d0);
+    let v3 = load_v_d4_nc(c4 + 3u, kvh, d0);
+    acc[0] += a0.x * v0 + a0.y * v1 + a0.z * v2 + a0.w * v3;
+    acc[1] += a1.x * v0 + a1.y * v1 + a1.z * v2 + a1.w * v3;
+    acc[2] += a2.x * v0 + a2.y * v1 + a2.z * v2 + a2.w * v3;
+    acc[3] += a3.x * v0 + a3.y * v1 + a3.z * v2 + a3.w * v3;
+    c4 = c4 + 4u;
+  }
+  if (c4 < params.context_len) {
+    let a0 = load_a_vec4(s0 + 0u, h, c4);
+    let a1 = load_a_vec4(s0 + 1u, h, c4);
+    let a2 = load_a_vec4(s0 + 2u, h, c4);
+    let a3 = load_a_vec4(s0 + 3u, h, c4);
+    let v0 = load_v_d4(c4 + 0u, kvh, d0);
+    let v1 = load_v_d4(c4 + 1u, kvh, d0);
+    let v2 = load_v_d4(c4 + 2u, kvh, d0);
+    let v3 = load_v_d4(c4 + 3u, kvh, d0);
+    acc[0] += a0.x * v0 + a0.y * v1 + a0.z * v2 + a0.w * v3;
+    acc[1] += a1.x * v0 + a1.y * v1 + a1.z * v2 + a1.w * v3;
+    acc[2] += a2.x * v0 + a2.y * v1 + a2.z * v2 + a2.w * v3;
+    acc[3] += a3.x * v0 + a3.y * v1 + a3.z * v2 + a3.w * v3;
+  }
+
+  var m: u32 = 0u;
+  loop {
+    if (m >= TM) {
+      break;
+    }
+    store_out_vec4(s0 + m, d0, h, acc[m]);
+    m = m + 1u;
+  }
+}
+)";
+
+inline constexpr uint32_t kSdpaComputeOutHalfWorkgroupSizeX = 64;
+inline constexpr uint32_t kSdpaComputeOutHalfWorkgroupSizeY = 1;
+inline constexpr uint32_t kSdpaComputeOutHalfWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/sdpa_fd_decode/SdpaFdDecode.cpp b/backends/webgpu/runtime/ops/sdpa_fd_decode/SdpaFdDecode.cpp
index 70108beb892..ffd3b24dce3 100644
--- a/backends/webgpu/runtime/ops/sdpa_fd_decode/SdpaFdDecode.cpp
+++ b/backends/webgpu/runtime/ops/sdpa_fd_decode/SdpaFdDecode.cpp
@@ -12,6 +12,7 @@
 #include <executorch/backends/webgpu/runtime/WebGPUUtils.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa_fd_decode/SdpaFdDecode.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_reduce_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split_half_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split_wgsl.h>
 
 #include <webgpu/webgpu.h>
@@ -185,8 +186,10 @@ void sdpa_fd_decode_dispatch(
       static_cast<uint64_t>(kSdpaFdMaxSplits) * static_cast<uint64_t>(D);
   const uint64_t pml_floats = static_cast<uint64_t>(Hq) *
       static_cast<uint64_t>(kSdpaFdMaxSplits) * 2ull;
-  WGPUBuffer part_o = graph.create_scratch_buffer(po_floats * sizeof(float));
-  WGPUBuffer part_ml = graph.create_scratch_buffer(pml_floats * sizeof(float));
+  WGPUBuffer part_o = graph.acquire_scratch(po_floats * sizeof(float));
+  WebGPUGraph::ScopedScratch part_o_guard(&graph, part_o);
+  WGPUBuffer part_ml = graph.acquire_scratch(pml_floats * sizeof(float));
+  WebGPUGraph::ScopedScratch part_ml_guard(&graph, part_ml);
 
   // Pass 1: split (Hq*num_splits WGs) -> writes part_o, part_ml.
   FdSplitParams sp = {};
@@ -218,9 +221,13 @@ void sdpa_fd_decode_dispatch(
       static_cast<uint32_t>(split_threads),
       kSdpaFdSplitWorkgroupSizeX,
       "fd_split");
+  const char* split_shader = kSdpaFdSplitWGSL;
+  if (graph.kv_f16()) {
+    split_shader = kSdpaFdSplitHalfWGSL;
+  }
   build_dispatch(
       graph,
-      kSdpaFdSplitWGSL,
+      split_shader,
       split_bindings,
       5,
       2,
diff --git a/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.wgsl b/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.wgsl
index c14c6bd07bd..da67489cc4d 100644
--- a/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.wgsl
+++ b/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.wgsl
@@ -1,8 +1,10 @@
+$if DTYPE == "half":
+  enable f16;
 @group(0) @binding(0) var<storage, read_write> t_part_o: array<f32>;
 @group(0) @binding(1) var<storage, read_write> t_part_ml: array<f32>;
 @group(0) @binding(2) var<storage, read> t_q: array<f32>;
-@group(0) @binding(3) var<storage, read> t_k_cache: array<f32>;
-@group(0) @binding(4) var<storage, read> t_v_cache: array<f32>;
+@group(0) @binding(3) var<storage, read> t_k_cache: array<${buffer_scalar_type(DTYPE)}>;
+@group(0) @binding(4) var<storage, read> t_v_cache: array<${buffer_scalar_type(DTYPE)}>;
 
 struct Params {
   _pad0: u32,
@@ -66,9 +68,14 @@ fn main(
         let qi = q_base + i4 * 4u;
         let ki = kvbase + i4 * 4u;
         let qv = vec4<f32>(t_q[qi], t_q[qi + 1u], t_q[qi + 2u], t_q[qi + 3u]);
-        let kvv = vec4<f32>(
-            t_k_cache[ki], t_k_cache[ki + 1u],
-            t_k_cache[ki + 2u], t_k_cache[ki + 3u]);
+        $if DTYPE == "half":
+          let kvv = vec4<f32>(
+              f32(t_k_cache[ki]), f32(t_k_cache[ki + 1u]),
+              f32(t_k_cache[ki + 2u]), f32(t_k_cache[ki + 3u]));
+        $else:
+          let kvv = vec4<f32>(
+              t_k_cache[ki], t_k_cache[ki + 1u],
+              t_k_cache[ki + 2u], t_k_cache[ki + 3u]);
         acc4 = acc4 + qv * kvv;
       }
       s = (acc4.x + acc4.y + acc4.z + acc4.w) * params.scale;
@@ -107,7 +114,10 @@ fn main(
         var acc: f32 = rescale * o_acc[nd];
         for (var j: u32 = 0u; j < n; j = j + 1u) {
           let vbase = (block + j) * kv_row_stride + kv * D;
-          acc = acc + sh_s[j] * t_v_cache[vbase + d];
+          $if DTYPE == "half":
+            acc = acc + sh_s[j] * f32(t_v_cache[vbase + d]);
+          $else:
+            acc = acc + sh_s[j] * t_v_cache[vbase + d];
         }
         o_acc[nd] = acc;
       }
diff --git a/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.yaml b/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.yaml
new file mode 100644
index 00000000000..283e6516996
--- /dev/null
+++ b/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split.yaml
@@ -0,0 +1,11 @@
+sdpa_fd_split:
+  parameter_names_with_default_values:
+    DTYPE: float
+  generate_variant_forall:
+    DTYPE:
+      - VALUE: float
+        SUFFIX: ""
+      - VALUE: half
+        SUFFIX: half
+  shader_variants:
+    - NAME: sdpa_fd_split
diff --git a/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split_half_wgsl.h b/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split_half_wgsl.h
new file mode 100644
index 00000000000..bc69a444edb
--- /dev/null
+++ b/backends/webgpu/runtime/ops/sdpa_fd_decode/sdpa_fd_split_half_wgsl.h
@@ -0,0 +1,156 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#pragma once
+
+#include <cstdint>
+
+namespace executorch::backends::webgpu {
+
+// @generated from sdpa_fd_split.wgsl - DO NOT EDIT.
+// wgsl-sha256: 147fc5775b76f3626dc934b2df5e34148bef12c345789332f098d13143bb646e
+inline constexpr const char* kSdpaFdSplitHalfWGSL = R"(
+enable f16;
+@group(0) @binding(0) var<storage, read_write> t_part_o: array<f32>;
+@group(0) @binding(1) var<storage, read_write> t_part_ml: array<f32>;
+@group(0) @binding(2) var<storage, read> t_q: array<f32>;
+@group(0) @binding(3) var<storage, read> t_k_cache: array<f16>;
+@group(0) @binding(4) var<storage, read> t_v_cache: array<f16>;
+
+struct Params {
+  _pad0: u32,
+  Hkv: u32,
+  D: u32,
+  context_len: u32,
+  g: u32,
+  num_splits: u32,
+  split_len: u32,
+  scale: f32,
+}
+@group(0) @binding(5) var<uniform> params: Params;
+
+const WG_SIZE: u32 = 64u;
+const MAX_SPLITS: u32 = 128u;
+const MAX_D_PER_LANE: u32 = 2u;
+const NEG_INF: f32 = -1.0e30;
+
+// sh_s: block scores then softmax weights; sh_red: max/sum reduction scratch.
+var<workgroup> sh_s: array<f32, WG_SIZE>;
+var<workgroup> sh_red: array<f32, WG_SIZE>;
+
+// FlashDecoding pass 1: per-(head,split) unnormalized softmax partial.
+@compute @workgroup_size(64, 1, 1)
+fn main(
+    @builtin(workgroup_id) wid: vec3<u32>,
+    @builtin(local_invocation_id) lid: vec3<u32>) {
+  let h = wid.x / params.num_splits;
+  let split_i = wid.x % params.num_splits;
+  let t = lid.x;
+  let D = params.D;
+  let D4 = D / 4u; // D is a multiple of 4 (guarded host-side); vec4 QK dot
+  let ctx = params.context_len;
+  let kv = h / params.g;
+  let q_base = h * D;
+  let kv_row_stride = params.Hkv * D;
+
+  let c0 = split_i * params.split_len;
+  var c1 = c0 + params.split_len;
+  if (c1 > ctx) { c1 = ctx; }
+
+  var m: f32 = NEG_INF;
+  var l: f32 = 0.0;
+  var o_acc: array<f32, MAX_D_PER_LANE>;
+  for (var nd: u32 = 0u; nd < MAX_D_PER_LANE; nd = nd + 1u) { o_acc[nd] = 0.0; }
+
+  // Stream the split in blocks of WG_SIZE KV positions.
+  var block: u32 = c0;
+  loop {
+    if (block >= c1) { break; }
+    var n: u32 = c1 - block;
+    if (n > WG_SIZE) { n = WG_SIZE; }
+
+    // Phase 1: lane t computes the full QK dot for position block+t (vec4), one
+    // K row read once. Out-of-block lanes hold NEG_INF (safe for the max).
+    var s: f32 = NEG_INF;
+    if (t < n) {
+      let kvbase = (block + t) * kv_row_stride + kv * D;
+      var acc4 = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+      for (var i4: u32 = 0u; i4 < D4; i4 = i4 + 1u) {
+        let qi = q_base + i4 * 4u;
+        let ki = kvbase + i4 * 4u;
+        let qv = vec4<f32>(t_q[qi], t_q[qi + 1u], t_q[qi + 2u], t_q[qi + 3u]);
+        let kvv = vec4<f32>(
+            f32(t_k_cache[ki]), f32(t_k_cache[ki + 1u]),
+            f32(t_k_cache[ki + 2u]), f32(t_k_cache[ki + 3u]));
+        acc4 = acc4 + qv * kvv;
+      }
+      s = (acc4.x + acc4.y + acc4.z + acc4.w) * params.scale;
+    }
+    sh_s[t] = s;
+
+    // Phase 2a: block max via tree reduction (sh_red written from register s).
+    sh_red[t] = s;
+    workgroupBarrier();
+    for (var stride: u32 = WG_SIZE / 2u; stride > 0u; stride = stride >> 1u) {
+      if (t < stride) { sh_red[t] = max(sh_red[t], sh_red[t + stride]); }
+      workgroupBarrier();
+    }
+    let m_new = max(m, sh_red[0]);
+    let rescale = exp(m - m_new);
+
+    // Phase 2b: each lane exponentiates ITS position once -> p (reuse sh_s),
+    // and reduce the block sum of p.
+    var p_t: f32 = 0.0;
+    if (t < n) { p_t = exp(sh_s[t] - m_new); }
+    workgroupBarrier(); // all reads of sh_s (the scores) done before overwrite
+    sh_s[t] = p_t;
+    sh_red[t] = p_t;
+    workgroupBarrier();
+    for (var stride: u32 = WG_SIZE / 2u; stride > 0u; stride = stride >> 1u) {
+      if (t < stride) { sh_red[t] = sh_red[t] + sh_red[t + stride]; }
+      workgroupBarrier();
+    }
+    l = rescale * l + sh_red[0];
+
+    // Phase 2c: each lane accumulates V for its own output dims over the block,
+    // reading the shared per-position weights (no exp in this loop).
+    for (var nd: u32 = 0u; nd < MAX_D_PER_LANE; nd = nd + 1u) {
+      let d = t + nd * WG_SIZE;
+      if (d < D) {
+        var acc: f32 = rescale * o_acc[nd];
+        for (var j: u32 = 0u; j < n; j = j + 1u) {
+          let vbase = (block + j) * kv_row_stride + kv * D;
+          acc = acc + sh_s[j] * f32(t_v_cache[vbase + d]);
+        }
+        o_acc[nd] = acc;
+      }
+    }
+    m = m_new;
+    workgroupBarrier(); // before the next block overwrites sh_s / sh_red
+    block = block + WG_SIZE;
+  }
+
+  let part = h * MAX_SPLITS + split_i;
+  for (var nd: u32 = 0u; nd < MAX_D_PER_LANE; nd = nd + 1u) {
+    let d = t + nd * WG_SIZE;
+    if (d < D) {
+      t_part_o[part * D + d] = o_acc[nd];
+    }
+  }
+  if (t == 0u) {
+    t_part_ml[part * 2u + 0u] = m;
+    t_part_ml[part * 2u + 1u] = l;
+  }
+}
+)";
+
+inline constexpr uint32_t kSdpaFdSplitHalfWorkgroupSizeX = 64;
+inline constexpr uint32_t kSdpaFdSplitHalfWorkgroupSizeY = 1;
+inline constexpr uint32_t kSdpaFdSplitHalfWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/runtime/ops/update_cache/update_cache.wgsl b/backends/webgpu/runtime/ops/update_cache/update_cache.wgsl
index 62f882ad547..cdda4aac9cd 100644
--- a/backends/webgpu/runtime/ops/update_cache/update_cache.wgsl
+++ b/backends/webgpu/runtime/ops/update_cache/update_cache.wgsl
@@ -1,4 +1,6 @@
-@group(0) @binding(0) var<storage, read_write> t_cache: array<f32>;
+$if DTYPE == "half":
+  enable f16;
+@group(0) @binding(0) var<storage, read_write> t_cache: array<${buffer_scalar_type(DTYPE)}>;
 @group(0) @binding(1) var<storage, read> t_value: array<f32>;
 
 struct Params {
@@ -20,5 +22,8 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   if (params.dst_offset + i >= params.cache_numel) {
     return;
   }
-  t_cache[params.dst_offset + i] = t_value[i];
+  $if DTYPE == "half":
+    t_cache[params.dst_offset + i] = f16(t_value[i]);
+  $else:
+    t_cache[params.dst_offset + i] = t_value[i];
 }
diff --git a/backends/webgpu/runtime/ops/update_cache/update_cache.yaml b/backends/webgpu/runtime/ops/update_cache/update_cache.yaml
new file mode 100644
index 00000000000..4df425c763e
--- /dev/null
+++ b/backends/webgpu/runtime/ops/update_cache/update_cache.yaml
@@ -0,0 +1,11 @@
+update_cache:
+  parameter_names_with_default_values:
+    DTYPE: float
+  generate_variant_forall:
+    DTYPE:
+      - VALUE: float
+        SUFFIX: ""
+      - VALUE: half
+        SUFFIX: half
+  shader_variants:
+    - NAME: update_cache
diff --git a/backends/webgpu/runtime/ops/update_cache/update_cache_half_wgsl.h b/backends/webgpu/runtime/ops/update_cache/update_cache_half_wgsl.h
new file mode 100644
index 00000000000..d31728178a8
--- /dev/null
+++ b/backends/webgpu/runtime/ops/update_cache/update_cache_half_wgsl.h
@@ -0,0 +1,49 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ * All rights reserved.
+ *
+ * This source code is licensed under the BSD-style license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#pragma once
+
+#include <cstdint>
+
+namespace executorch::backends::webgpu {
+
+// @generated from update_cache.wgsl - DO NOT EDIT.
+// wgsl-sha256: 390daabe0d4545311dd5c6768d427fc9d133125bb1143869184c7d7631a88954
+inline constexpr const char* kUpdateCacheHalfWGSL = R"(
+enable f16;
+@group(0) @binding(0) var<storage, read_write> t_cache: array<f16>;
+@group(0) @binding(1) var<storage, read> t_value: array<f32>;
+
+struct Params {
+  numel: u32,
+  dst_offset: u32,
+  cache_numel: u32,
+  _pad0: u32,
+}
+@group(0) @binding(2) var<uniform> params: Params;
+
+override wg_size: u32 = 256;
+
+@compute @workgroup_size(wg_size, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= params.numel) {
+    return;
+  }
+  if (params.dst_offset + i >= params.cache_numel) {
+    return;
+  }
+  t_cache[params.dst_offset + i] = f16(t_value[i]);
+}
+)";
+
+inline constexpr uint32_t kUpdateCacheHalfWorkgroupSizeX = 256;
+inline constexpr uint32_t kUpdateCacheHalfWorkgroupSizeY = 1;
+inline constexpr uint32_t kUpdateCacheHalfWorkgroupSizeZ = 1;
+
+} // namespace executorch::backends::webgpu
diff --git a/backends/webgpu/test/native/test_scratch_buffer.cpp b/backends/webgpu/test/native/test_scratch_buffer.cpp
index 98cf3648c6b..1a8f4fcd96f 100644
--- a/backends/webgpu/test/native/test_scratch_buffer.cpp
+++ b/backends/webgpu/test/native/test_scratch_buffer.cpp
@@ -6,7 +6,8 @@
  * LICENSE file in the root directory of this source tree.
  */
 
-// White-box unit tests for WebGPUGraph::create_scratch_buffer.
+// White-box unit tests for WebGPUGraph scratch buffers:
+// create_scratch_buffer and the acquire_scratch/release_scratch reuse pool.
 
 #include <executorch/backends/webgpu/runtime/WebGPUCompat.h>
 #include <executorch/backends/webgpu/runtime/WebGPUDevice.h>
@@ -19,6 +20,7 @@
 #include <atomic>
 #include <cmath>
 #include <cstdio>
+#include <cstdlib>
 #include <cstring>
 #include <exception>
 #include <vector>
@@ -222,6 +224,59 @@ TEST(ScratchBuffer, Tier3Lifecycle) {
   } // each graph's dtor releases its 256 buffers here
 }
 
+// Tier 4: reuse-pool semantics (acquire_scratch / release_scratch /
+// ScopedScratch). The pool recycles single-op-lifetime scratch across ops so N
+// layers reuse a small constant of buffers instead of N x.
+
+// A released slot is handed back on the next same-size acquire (the reuse win).
+TEST(ScratchPool, ReuseAfterRelease) {
+  WebGPUGraph g;
+  g.set_device(g_device);
+  WGPUBuffer a = g.acquire_scratch(64 * sizeof(float));
+  g.release_scratch(a);
+  WGPUBuffer b = g.acquire_scratch(64 * sizeof(float));
+  EXPECT_EQ(a, b) << "released slot should be reused for a same-size request";
+}
+
+// A still-in_use slot is never handed to a co-live requester (RAW-safety).
+TEST(ScratchPool, NoReuseWhileInUse) {
+  WebGPUGraph g;
+  g.set_device(g_device);
+  WGPUBuffer a = g.acquire_scratch(64 * sizeof(float));
+  WGPUBuffer b = g.acquire_scratch(64 * sizeof(float)); // a not released
+  EXPECT_TRUE(a && b && a != b) << "co-live acquires must be distinct buffers";
+}
+
+// Best-fit 2x cap: a large free slot must not back a much smaller request, but
+// a request it does fit (size in [n, 2n]) reuses it.
+TEST(ScratchPool, BestFitSizeCap) {
+  WebGPUGraph g;
+  g.set_device(g_device);
+  WGPUBuffer big = g.acquire_scratch(1024 * sizeof(float));
+  g.release_scratch(big);
+  // 1024*4 bytes is outside [4, 8], so the big slot is ineligible for 4 bytes.
+  WGPUBuffer tiny = g.acquire_scratch(4);
+  EXPECT_NE(big, tiny)
+      << "oversized slot must not back a tiny request (2x cap)";
+  g.release_scratch(tiny);
+  WGPUBuffer same = g.acquire_scratch(1024 * sizeof(float));
+  EXPECT_EQ(big, same) << "an in-range request should reuse the big slot";
+}
+
+// ScopedScratch releases its slot at scope exit, so the next acquire reuses it.
+TEST(ScratchPool, ScopedScratchReleasesOnScopeExit) {
+  WebGPUGraph g;
+  g.set_device(g_device);
+  WGPUBuffer first = nullptr;
+  {
+    WebGPUGraph::ScopedScratch s(&g, g.acquire_scratch(64 * sizeof(float)));
+    first = s; // operator WGPUBuffer
+    EXPECT_NE(first, nullptr);
+  } // s releases the slot here
+  WGPUBuffer second = g.acquire_scratch(64 * sizeof(float));
+  EXPECT_EQ(first, second) << "slot freed by ScopedScratch should be reused";
+}
+
 int main(int argc, char** argv) {
   ::testing::InitGoogleTest(&argc, argv);
 
diff --git a/backends/webgpu/test/ops/test_quantized_linear.py b/backends/webgpu/test/ops/test_quantized_linear.py
index e9a785dcd6d..72945c37d6d 100644
--- a/backends/webgpu/test/ops/test_quantized_linear.py
+++ b/backends/webgpu/test/ops/test_quantized_linear.py
@@ -70,6 +70,19 @@ class Q4gswConfig:
     # Partial M and N steel tiles under the f16 kernel; exercises f16 boundary
     # masking (the exact-N "steel_f16" shape does not). N%8==0, steel-isolating.
     Q4gswConfig("steel_f16_edge", 70, 1024, 136),  # f16 partial-tile
+    # pwdq (packed-word dequant) backs the f16 steel path at group_size % BK(16)
+    # == 0 (bit-exact to steel_half; steel_f16 above runs it at gs=32). These lock
+    # the gs gate at group sizes those omit: gs=64 stays on pwdq; gs=8 (< BK) falls
+    # back to the per-nibble steel_half kernel (its hoisted-per-BK scale is invalid
+    # there). Same fp64 golden regardless of which kernel runs.
+    Q4gswConfig("pwdq_gs64", 96, 2048, 256, group_size=64),  # pwdq, non-32 group
+    Q4gswConfig("pwdq_gs8", 96, 2048, 256, group_size=8),  # steel_half fallback
+    # pwdqf16acc (f16-accumulate) runs when the enable_f16_accumulate_gemm runtime
+    # spec is set and gs % BK == 0 (perplexity-gated; see the kernel diff). Same
+    # .pte as the f32 configs -- only the accumulator dtype differs -- goldened at a
+    # looser f16-accumulate tol in the native test; deep-K stresses the worst case.
+    Q4gswConfig("pwdqf16acc", 96, 2048, 256),  # f16-accumulate steel (runtime)
+    Q4gswConfig("pwdqf16acc_down", 128, 8192, 2048),  # deep-K f16-accum worst case
     Q4gswConfig("gate_proj_pf", 128, 2048, 8192),  # gate/up prefill (shmem via N)
     Q4gswConfig("down_proj_pf", 128, 8192, 2048),  # down prefill (shmem via K)
     Q4gswConfig("shmem_edge", 130, 4096, 2056),  # partial 32-tile bounds
diff --git a/backends/webgpu/test/test_webgpu_native.cpp b/backends/webgpu/test/test_webgpu_native.cpp
index 5f3d6d788ec..fbdfbd09076 100644
--- a/backends/webgpu/test/test_webgpu_native.cpp
+++ b/backends/webgpu/test/test_webgpu_native.cpp
@@ -11,6 +11,8 @@
 #include <executorch/backends/webgpu/runtime/WebGPUGraph.h>
 #include <executorch/extension/module/module.h>
 #include <executorch/extension/tensor/tensor.h>
+#include <executorch/runtime/backend/backend_options_map.h>
+#include <executorch/runtime/backend/options.h>
 
 #include <gtest/gtest.h>
 
@@ -229,6 +231,15 @@ bool sdpa_within_tol(
     int n,
     float* ma,
     float* mr) {
+  float atol = 1e-4f, rtol = 1e-3f;
+  // f16 KV (runtime opt-in) reads K/V at reduced precision; loosen the tol on a
+  // shader-f16 device to cover that rounding. Harmless for f32 KV (looser
+  // gate).
+  const WebGPUContext* kv_ctx = get_default_webgpu_context();
+  if (kv_ctx != nullptr && kv_ctx->shader_f16_supported) {
+    atol = 2e-3f;
+    rtol = 1e-2f;
+  }
   float max_abs = 0.0f, max_rel = 0.0f;
   bool ok = true;
   for (int i = 0; i < n; i++) {
@@ -236,7 +247,7 @@ bool sdpa_within_tol(
     const float re = ae / std::max(std::abs(golden[i]), 1e-6f);
     max_abs = std::max(max_abs, ae);
     max_rel = std::max(max_rel, re);
-    if (ae > 1e-4f && re > 1e-3f) {
+    if (ae > atol && re > rtol) {
       ok = false;
     }
   }
@@ -284,6 +295,21 @@ const Q4gswConfig kQ4gswConfigs[] = {
     {"steel_f16", 96, 2048, 256, 2.3e-4f, 1e-3f, true, false},
     // Partial M and N steel tiles under the f16 kernel (f16 boundary masking).
     {"steel_f16_edge", 70, 1024, 136, 2.3e-4f, 1e-3f, true, false},
+    // pwdq (packed-word dequant) backs the f16 steel path at group_size % BK ==
+    // 0
+    // (bit-exact to steel_half; the steel_f16 configs above run it at gs=32).
+    // These lock the gs gate at group sizes those omit: gs=64 stays on pwdq;
+    // gs=8 (< BK=16) falls back to the per-nibble steel_half kernel.
+    {"pwdq_gs64", 96, 2048, 256, 2.3e-4f, 1e-3f, true, false},
+    {"pwdq_gs8", 96, 2048, 256, 2.3e-4f, 1e-3f, true, false},
+    // f16-ACCUMULATE steel (pwdqf16acc): lossy, so a wider gate than the
+    // f16-multiply steel_f16 (2.3e-4). f16 accumulation error grows with K, so
+    // the deep-K down shape (K=8192) gets the loosest tol. Perplexity is the
+    // primary quality gate (see the kernel diff); this catches gross bit/index
+    // bugs. gs=32 (% BK == 0) selects pwdqf16acc; the sweep loads these rows
+    // with the enable_f16_accumulate_gemm runtime spec set.
+    {"pwdqf16acc", 96, 2048, 256, 2e-2f, 3e-2f, true, false},
+    {"pwdqf16acc_down", 128, 8192, 2048, 5e-2f, 8e-2f, true, false},
     {"gate_proj_pf", 128, 2048, 8192, 1e-4f, 1e-3f, true, false}, // shmem via N
     {"down_proj_pf", 128, 8192, 2048, 1e-3f, 1e-2f, true, false}, // shmem via K
     {"shmem_edge", 130, 4096, 2056, 1e-4f, 1e-3f, true, false}, // partial tiles
@@ -547,7 +573,18 @@ void test_q4gsw_config(
       cfg.n);
 
   Module module(pte);
-  ASSERT_EQ(module.load_forward(), Error::Ok) << "could not load " << pte;
+  // pwdqf16acc rows exercise the lossy f16-accumulate kernel, a runtime opt-in
+  // (default off); enable it via the backend option keyed by the registered id.
+  if (std::string(cfg.name).rfind("pwdqf16acc", 0) == 0) {
+    BackendOptions<1> opts;
+    opts.set_option("enable_f16_accumulate_gemm", true);
+    LoadBackendOptionsMap map;
+    ASSERT_EQ(map.set_options("VulkanBackend", opts.view()), Error::Ok);
+    ASSERT_EQ(module.load_forward(nullptr, nullptr, &map), Error::Ok)
+        << "could not load " << pte;
+  } else {
+    ASSERT_EQ(module.load_forward(), Error::Ok) << "could not load " << pte;
+  }
 
   const int in_numel = cfg.m * cfg.k;
   const int out_numel = cfg.m * cfg.n;