diff --git a/backends/webgpu/runtime/WebGPUShaderRegistry.cpp b/backends/webgpu/runtime/WebGPUShaderRegistry.cpp
index 1cf4b952a5f..b475a6cd999 100644
--- a/backends/webgpu/runtime/WebGPUShaderRegistry.cpp
+++ b/backends/webgpu/runtime/WebGPUShaderRegistry.cpp
@@ -13,8 +13,6 @@
 #include <executorch/backends/webgpu/runtime/ops/adamw/adamw_step_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/add/binary_add_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/addmm/addmm_tiled_wgsl.h>
-#include <executorch/backends/webgpu/runtime/ops/amax/amax_wgsl.h>
-#include <executorch/backends/webgpu/runtime/ops/amin/amin_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/argmax/arg_reduce_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/avg_pool2d/avg_pool2d_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/batch_norm/batch_norm_wgsl.h>
@@ -48,6 +46,8 @@
 #include <executorch/backends/webgpu/runtime/ops/et_vk_sdpa/et_vk_sdpa_qk_entry_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/et_vk_sdpa/et_vk_sdpa_qk_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/expand_copy/expand_copy_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/extrema/amax_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/extrema/amin_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/fill/fill_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/flip/flip_wgsl.h>
 #include <executorch/backends/webgpu/runtime/ops/floor_divide/binary_floor_divide_wgsl.h>
diff --git a/backends/webgpu/runtime/ops/amax/Reduce.cpp b/backends/webgpu/runtime/ops/amax/Reduce.cpp
index 14675d97e96..c1dff62e035 100644
--- a/backends/webgpu/runtime/ops/amax/Reduce.cpp
+++ b/backends/webgpu/runtime/ops/amax/Reduce.cpp
@@ -9,7 +9,7 @@
 #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/amax/amax_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/extrema/amax_wgsl.h>
 
 #include <webgpu/webgpu.h>
 
diff --git a/backends/webgpu/runtime/ops/amin/Reduce.cpp b/backends/webgpu/runtime/ops/amin/Reduce.cpp
index fbe574fdf0b..24a5ebeb826 100644
--- a/backends/webgpu/runtime/ops/amin/Reduce.cpp
+++ b/backends/webgpu/runtime/ops/amin/Reduce.cpp
@@ -9,7 +9,7 @@
 #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/amin/amin_wgsl.h>
+#include <executorch/backends/webgpu/runtime/ops/extrema/amin_wgsl.h>
 
 #include <webgpu/webgpu.h>
 
diff --git a/backends/webgpu/runtime/ops/amin/amin.wgsl b/backends/webgpu/runtime/ops/amin/amin.wgsl
deleted file mode 100644
index 4778800ab3d..00000000000
--- a/backends/webgpu/runtime/ops/amin/amin.wgsl
+++ /dev/null
@@ -1,49 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_rows: u32,
-  reduce_size: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-// Cooperative shared-memory reduction; mirrors Vulkan reduce.glsl (a group of
-// threads co-operates per reduction row, partials aggregated in shared memory).
-// Fixed upper bound (>= any clamped wg_size); only [0, wg_size) is used.
-var<workgroup> partials: array<f32, 256>;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(workgroup_id) wid: vec3<u32>,
-    @builtin(local_invocation_id) lid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    // One workgroup per reduction row; 2D-fold lifts the 65535 grid cap.
-    let row = wid.x + wid.y * num_workgroups.x;
-    if (row >= params.num_rows) {
-        return;
-    }
-    let base = row * params.reduce_size;
-
-    // Each thread reduces a strided slice of the row into a partial. Seed with
-    // the row's first element (always valid; reduce_size >= 1) so threads that
-    // own no element contribute a real value, not an out-of-range identity.
-    var acc = input[base];
-    var i = lid.x;
-    while (i < params.reduce_size) {
-        acc = min(acc, input[base + i]);
-        i = i + wg_size;
-    }
-    partials[lid.x] = acc;
-    workgroupBarrier();
-
-    // Thread 0 aggregates the wg_size partials (mirrors Vulkan's group aggregate).
-    if (lid.x == 0u) {
-        var m = partials[0];
-        for (var t = 1u; t < wg_size; t = t + 1u) {
-            m = min(m, partials[t]);
-        }
-        output[row] = m;
-    }
-}
diff --git a/backends/webgpu/runtime/ops/amax/amax_wgsl.h b/backends/webgpu/runtime/ops/extrema/amax_wgsl.h
similarity index 97%
rename from backends/webgpu/runtime/ops/amax/amax_wgsl.h
rename to backends/webgpu/runtime/ops/extrema/amax_wgsl.h
index 48ec8f20e27..0b7a4d2238e 100644
--- a/backends/webgpu/runtime/ops/amax/amax_wgsl.h
+++ b/backends/webgpu/runtime/ops/extrema/amax_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from amax.wgsl - DO NOT EDIT.
+// @generated from extrema.wgsl - DO NOT EDIT.
 // wgsl-sha256: 35fc059d7c72caa17f9cb1128823ecfd8f75be4ce24b6cd4f9629a97b52f64c0
 inline constexpr const char* kAmaxWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/amin/amin_wgsl.h b/backends/webgpu/runtime/ops/extrema/amin_wgsl.h
similarity index 97%
rename from backends/webgpu/runtime/ops/amin/amin_wgsl.h
rename to backends/webgpu/runtime/ops/extrema/amin_wgsl.h
index 40a97c67a63..8b8bf5456b1 100644
--- a/backends/webgpu/runtime/ops/amin/amin_wgsl.h
+++ b/backends/webgpu/runtime/ops/extrema/amin_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from amin.wgsl - DO NOT EDIT.
+// @generated from extrema.wgsl - DO NOT EDIT.
 // wgsl-sha256: 8cb6035ae4d34eb2a6cc973d93d9847905722e967239c96033fccfe3a1943cb2
 inline constexpr const char* kAminWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/amax/amax.wgsl b/backends/webgpu/runtime/ops/extrema/extrema.wgsl
similarity index 94%
rename from backends/webgpu/runtime/ops/amax/amax.wgsl
rename to backends/webgpu/runtime/ops/extrema/extrema.wgsl
index 2f23b9c38fa..6c55d60d9b7 100644
--- a/backends/webgpu/runtime/ops/amax/amax.wgsl
+++ b/backends/webgpu/runtime/ops/extrema/extrema.wgsl
@@ -32,7 +32,7 @@ fn main(
     var acc = input[base];
     var i = lid.x;
     while (i < params.reduce_size) {
-        acc = max(acc, input[base + i]);
+        acc = ${REDUCE_FN}(acc, input[base + i]);
         i = i + wg_size;
     }
     partials[lid.x] = acc;
@@ -42,7 +42,7 @@ fn main(
     if (lid.x == 0u) {
         var m = partials[0];
         for (var t = 1u; t < wg_size; t = t + 1u) {
-            m = max(m, partials[t]);
+            m = ${REDUCE_FN}(m, partials[t]);
         }
         output[row] = m;
     }
diff --git a/backends/webgpu/runtime/ops/extrema/extrema.yaml b/backends/webgpu/runtime/ops/extrema/extrema.yaml
new file mode 100644
index 00000000000..a85da849c60
--- /dev/null
+++ b/backends/webgpu/runtime/ops/extrema/extrema.yaml
@@ -0,0 +1,13 @@
+# 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.
+
+extrema:
+  parameter_names_with_default_values:
+    REDUCE_FN: max
+  shader_variants:
+    - NAME: amax
+    - NAME: amin
+      REDUCE_FN: min
diff --git a/backends/webgpu/runtime/ops/unary/abs.wgsl b/backends/webgpu/runtime/ops/unary/abs.wgsl
deleted file mode 100644
index e3e10c75dd9..00000000000
--- a/backends/webgpu/runtime/ops/unary/abs.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = abs(x);
-}
diff --git a/backends/webgpu/runtime/ops/unary/abs_wgsl.h b/backends/webgpu/runtime/ops/unary/abs_wgsl.h
index 3d2873e69c4..0efb72d5189 100644
--- a/backends/webgpu/runtime/ops/unary/abs_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/abs_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from abs.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 39d3c163fdf6a92286828f4b3217e00294e3ca5634a878ed5fd34e3b1cdf0a27
 inline constexpr const char* kAbsWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/cos.wgsl b/backends/webgpu/runtime/ops/unary/cos.wgsl
deleted file mode 100644
index c2bafe7b248..00000000000
--- a/backends/webgpu/runtime/ops/unary/cos.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = cos(x);
-}
diff --git a/backends/webgpu/runtime/ops/unary/cos_wgsl.h b/backends/webgpu/runtime/ops/unary/cos_wgsl.h
index 422b5618146..4ca99df88f3 100644
--- a/backends/webgpu/runtime/ops/unary/cos_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/cos_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from cos.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 9df78873e5fae98d347c26db2a02b047ea3d5d2c93f0761cb9ac6995f9a71ab2
 inline constexpr const char* kCosWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/exp_wgsl.h b/backends/webgpu/runtime/ops/unary/exp_wgsl.h
index 28b6c67cb18..cbf85fd415a 100644
--- a/backends/webgpu/runtime/ops/unary/exp_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/exp_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from exp.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 3171399bc36acf9c1cb2a03c2a31038318203c4c63ab03c4881df7a660346020
 inline constexpr const char* kExpWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/hardswish.wgsl b/backends/webgpu/runtime/ops/unary/hardswish.wgsl
deleted file mode 100644
index 2278caf3664..00000000000
--- a/backends/webgpu/runtime/ops/unary/hardswish.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = select(select(x * (x + 3.0) / 6.0, x, x >= 3.0), 0.0, x <= -3.0);
-}
diff --git a/backends/webgpu/runtime/ops/unary/hardswish_wgsl.h b/backends/webgpu/runtime/ops/unary/hardswish_wgsl.h
index 0c991547b9a..43f104c7f8f 100644
--- a/backends/webgpu/runtime/ops/unary/hardswish_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/hardswish_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from hardswish.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: c874a15ef6cdaec71187296016cc2a1515f5e7c889b97dfa8fd4b278e6e2c3d5
 inline constexpr const char* kHardswishWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/neg.wgsl b/backends/webgpu/runtime/ops/unary/neg.wgsl
deleted file mode 100644
index c977957957d..00000000000
--- a/backends/webgpu/runtime/ops/unary/neg.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = -x;
-}
diff --git a/backends/webgpu/runtime/ops/unary/neg_wgsl.h b/backends/webgpu/runtime/ops/unary/neg_wgsl.h
index d528c45fea0..c4c7bb989ac 100644
--- a/backends/webgpu/runtime/ops/unary/neg_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/neg_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from neg.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 8851b9f42d14153f6f04484fee2f8bf67bda26dea892ff48768e09e6ad49cee1
 inline constexpr const char* kNegWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/round.wgsl b/backends/webgpu/runtime/ops/unary/round.wgsl
deleted file mode 100644
index 2269ca59988..00000000000
--- a/backends/webgpu/runtime/ops/unary/round.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = round(x);
-}
diff --git a/backends/webgpu/runtime/ops/unary/round_wgsl.h b/backends/webgpu/runtime/ops/unary/round_wgsl.h
index 209305bc855..8c805ba4f80 100644
--- a/backends/webgpu/runtime/ops/unary/round_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/round_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from round.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 8f3e0edbeb81aa50f35e691c78554e8057fa8d78fe8a86454f4f42e5e8871452
 inline constexpr const char* kRoundWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/rsqrt.wgsl b/backends/webgpu/runtime/ops/unary/rsqrt.wgsl
deleted file mode 100644
index 1f50c4f66ac..00000000000
--- a/backends/webgpu/runtime/ops/unary/rsqrt.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = 1.0 / sqrt(x);
-}
diff --git a/backends/webgpu/runtime/ops/unary/rsqrt_wgsl.h b/backends/webgpu/runtime/ops/unary/rsqrt_wgsl.h
index 83e74a1b8cf..58bd05d010e 100644
--- a/backends/webgpu/runtime/ops/unary/rsqrt_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/rsqrt_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from rsqrt.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 108765d5a23b87473f34651875d08abf2a5fa8980bd92fc8cbe3617295097747
 inline constexpr const char* kRsqrtWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/sin.wgsl b/backends/webgpu/runtime/ops/unary/sin.wgsl
deleted file mode 100644
index ffd2a07ea8a..00000000000
--- a/backends/webgpu/runtime/ops/unary/sin.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = sin(x);
-}
diff --git a/backends/webgpu/runtime/ops/unary/sin_wgsl.h b/backends/webgpu/runtime/ops/unary/sin_wgsl.h
index f22229b2342..54184a1ccb8 100644
--- a/backends/webgpu/runtime/ops/unary/sin_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/sin_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from sin.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: e5762804773659d348fddddcef4935807ae6fe7d92c92eb17a2f44aae8f2c5b9
 inline constexpr const char* kSinWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/sqrt.wgsl b/backends/webgpu/runtime/ops/unary/sqrt.wgsl
deleted file mode 100644
index e34ff440007..00000000000
--- a/backends/webgpu/runtime/ops/unary/sqrt.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = sqrt(x);
-}
diff --git a/backends/webgpu/runtime/ops/unary/sqrt_wgsl.h b/backends/webgpu/runtime/ops/unary/sqrt_wgsl.h
index 260a4ce4266..42dcdb838c0 100644
--- a/backends/webgpu/runtime/ops/unary/sqrt_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/sqrt_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from sqrt.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 008534ae365969f5c180b42e8d6d0b131df78f181e5435abbcafc3ffb8be8aac
 inline constexpr const char* kSqrtWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/tanh.wgsl b/backends/webgpu/runtime/ops/unary/tanh.wgsl
deleted file mode 100644
index 5a541e21699..00000000000
--- a/backends/webgpu/runtime/ops/unary/tanh.wgsl
+++ /dev/null
@@ -1,21 +0,0 @@
-@group(0) @binding(0) var<storage, read> input: array<f32>;
-@group(0) @binding(1) var<storage, read_write> output: array<f32>;
-
-struct Params {
-  num_elements: u32,
-}
-@group(0) @binding(2) var<uniform> params: Params;
-
-override wg_size: u32 = 256u;
-
-@compute @workgroup_size(wg_size)
-fn main(
-    @builtin(global_invocation_id) gid: vec3<u32>,
-    @builtin(num_workgroups) num_workgroups: vec3<u32>) {
-    let idx = gid.x + gid.y * (num_workgroups.x * wg_size);
-    if (idx >= params.num_elements) {
-        return;
-    }
-    let x = input[idx];
-    output[idx] = tanh(clamp(x, -15.0, 15.0));
-}
diff --git a/backends/webgpu/runtime/ops/unary/tanh_wgsl.h b/backends/webgpu/runtime/ops/unary/tanh_wgsl.h
index eef1ecd91af..51ba4d3919c 100644
--- a/backends/webgpu/runtime/ops/unary/tanh_wgsl.h
+++ b/backends/webgpu/runtime/ops/unary/tanh_wgsl.h
@@ -12,7 +12,7 @@
 
 namespace executorch::backends::webgpu {
 
-// @generated from tanh.wgsl - DO NOT EDIT.
+// @generated from unary.wgsl - DO NOT EDIT.
 // wgsl-sha256: 5bd7eb1c6411940d84a9b311884f35b39f15b82103b14bab02902290ed6b0339
 inline constexpr const char* kTanhWGSL = R"(
 @group(0) @binding(0) var<storage, read> input: array<f32>;
diff --git a/backends/webgpu/runtime/ops/unary/exp.wgsl b/backends/webgpu/runtime/ops/unary/unary.wgsl
similarity index 94%
rename from backends/webgpu/runtime/ops/unary/exp.wgsl
rename to backends/webgpu/runtime/ops/unary/unary.wgsl
index b69aa509d8e..d974a2f3319 100644
--- a/backends/webgpu/runtime/ops/unary/exp.wgsl
+++ b/backends/webgpu/runtime/ops/unary/unary.wgsl
@@ -17,5 +17,5 @@ fn main(
         return;
     }
     let x = input[idx];
-    output[idx] = exp(x);
+    output[idx] = ${OPERATOR};
 }
diff --git a/backends/webgpu/runtime/ops/unary/unary.yaml b/backends/webgpu/runtime/ops/unary/unary.yaml
new file mode 100644
index 00000000000..84c0dd660d7
--- /dev/null
+++ b/backends/webgpu/runtime/ops/unary/unary.yaml
@@ -0,0 +1,29 @@
+# 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.
+
+unary:
+  parameter_names_with_default_values:
+    OPERATOR: "abs(x)"
+  shader_variants:
+    - NAME: abs
+    - NAME: cos
+      OPERATOR: "cos(x)"
+    - NAME: exp
+      OPERATOR: "exp(x)"
+    - NAME: hardswish
+      OPERATOR: "select(select(x * (x + 3.0) / 6.0, x, x >= 3.0), 0.0, x <= -3.0)"
+    - NAME: neg
+      OPERATOR: "-x"
+    - NAME: round
+      OPERATOR: "round(x)"
+    - NAME: rsqrt
+      OPERATOR: "1.0 / sqrt(x)"
+    - NAME: sin
+      OPERATOR: "sin(x)"
+    - NAME: sqrt
+      OPERATOR: "sqrt(x)"
+    - NAME: tanh
+      OPERATOR: "tanh(clamp(x, -15.0, 15.0))"
diff --git a/backends/webgpu/test/native/test_compute_dispatch.cpp b/backends/webgpu/test/native/test_compute_dispatch.cpp
index b2b0d028567..6ed7229604b 100644
--- a/backends/webgpu/test/native/test_compute_dispatch.cpp
+++ b/backends/webgpu/test/native/test_compute_dispatch.cpp
@@ -440,6 +440,18 @@ void record_resize_probe(WebGPUGraph&, const ResizeProbeContext& context) {
   ++*context.calls;
 }
 
+struct ThrowingResizeContext {
+  bool* fail;
+  int* calls;
+};
+
+void maybe_throw_resize(WebGPUGraph&, const ThrowingResizeContext& context) {
+  ++*context.calls;
+  if (*context.fail) {
+    throw std::runtime_error("resize hook failure");
+  }
+}
+
 struct GridPickerContext {
   int tensor_id;
   uint32_t x_bias;
@@ -487,6 +499,25 @@ TEST(WebGPUResizeHooks, TypedRegistrationOwnsContextCopies) {
   int tensor_calls = 0;
   int symint_observed = 0;
   int symint_calls = 0;
+  using ResizeProbeFn = void (*)(WebGPUGraph&, const ResizeProbeContext&);
+  EXPECT_THROW(
+      graph.add_tensor_resize_hook(
+          kResizeQ,
+          static_cast<ResizeProbeFn>(nullptr),
+          ResizeProbeContext{0, &tensor_observed, &tensor_calls}),
+      std::runtime_error);
+  EXPECT_THROW(
+      graph.add_tensor_resize_hook(
+          kResizeSymInt,
+          record_resize_probe,
+          ResizeProbeContext{0, &tensor_observed, &tensor_calls}),
+      std::runtime_error);
+  EXPECT_THROW(
+      graph.add_resize_hook(
+          kResizeQ,
+          record_resize_probe,
+          ResizeProbeContext{0, &symint_observed, &symint_calls}),
+      std::runtime_error);
   {
     ResizeProbeContext tensor_context = {17, &tensor_observed, &tensor_calls};
     ResizeProbeContext symint_context = {23, &symint_observed, &symint_calls};
@@ -510,6 +541,36 @@ TEST(WebGPUResizeHooks, TypedRegistrationOwnsContextCopies) {
   EXPECT_EQ(symint_calls, 1);
 }
 
+TEST(WebGPUResizeHooks, RestoresDirtyTriggerWhenHookThrows) {
+  WebGPUGraph graph;
+  build_resize_test_graph(graph);
+  bool hook_fails = true;
+  int hook_calls = 0;
+  int picker_calls = 0;
+  graph.add_tensor_resize_hook(
+      kResizeQ,
+      maybe_throw_resize,
+      ThrowingResizeContext{&hook_fails, &hook_calls});
+  const size_t dispatch = graph.add_dynamic_compute_dispatch(
+      make_dynamic_test_descriptor(graph, "dynamic_hook_retry"),
+      kResizeQ,
+      pick_tensor_grid,
+      GridPickerContext{kResizeQ, 1u, 2u, &picker_calls, nullptr});
+  picker_calls = 0;
+
+  graph.resize_input(kResizeQ, {4, 3});
+  EXPECT_THROW(graph.propagate_resize(), std::runtime_error);
+  EXPECT_EQ(hook_calls, 1);
+  EXPECT_EQ(picker_calls, 0);
+  expect_dispatch_grid(graph, dispatch, 9u, 10u);
+
+  hook_fails = false;
+  EXPECT_NO_THROW(graph.propagate_resize());
+  EXPECT_EQ(hook_calls, 2);
+  EXPECT_EQ(picker_calls, 1);
+  expect_dispatch_grid(graph, dispatch, 5u, 5u);
+}
+
 TEST(WebGPUDynamicDispatch, InitializesAndIsolatesTriggeredGrids) {
   WebGPUGraph graph;
   build_resize_test_graph(graph);
@@ -615,12 +676,60 @@ TEST(WebGPUDynamicDispatch, HandlesCascadesAndStagesPickerFailures) {
   EXPECT_THROW(graph.propagate_resize(), std::runtime_error);
   expect_dispatch_grid(graph, first, 9u, 10u);
   expect_dispatch_grid(graph, second, 11u, 12u);
+  second_fails = false;
+  EXPECT_NO_THROW(graph.propagate_resize());
+  expect_dispatch_grid(graph, first, 5u, 5u);
+  expect_dispatch_grid(graph, second, 7u, 7u);
 }
 
 TEST(WebGPUDynamicDispatch, RejectsRouteOverlapWithoutPoisoningRegistry) {
   WebGPUGraph graph;
   build_resize_test_graph(graph);
   int calls = 0;
+  const size_t dispatches_before_invalid_trigger = graph.num_dispatches();
+  EXPECT_THROW(
+      graph.add_dynamic_compute_dispatch(
+          make_dynamic_test_descriptor(graph, "dynamic_negative_trigger"),
+          -1,
+          pick_tensor_grid,
+          GridPickerContext{kResizeQ, 0u, 0u, &calls, nullptr}),
+      std::runtime_error);
+  EXPECT_THROW(
+      graph.add_dynamic_compute_dispatch(
+          make_dynamic_test_descriptor(graph, "dynamic_oob_trigger"),
+          graph.num_values(),
+          pick_tensor_grid,
+          GridPickerContext{kResizeQ, 0u, 0u, &calls, nullptr}),
+      std::runtime_error);
+  EXPECT_THROW(
+      graph.add_dynamic_compute_dispatch(
+          make_dynamic_test_descriptor(graph, "dynamic_symint_trigger"),
+          kResizeSymInt,
+          pick_tensor_grid,
+          GridPickerContext{kResizeQ, 0u, 0u, &calls, nullptr}),
+      std::runtime_error);
+  EXPECT_EQ(calls, 0);
+  EXPECT_EQ(graph.num_dispatches(), dispatches_before_invalid_trigger);
+
+  auto pick_zero_grid = [](const WebGPUGraph&, const WebGPUDispatchGrid& grid) {
+    return grid;
+  };
+  EXPECT_THROW(
+      graph.add_dynamic_compute_dispatch(
+          make_dynamic_test_descriptor(graph, "dynamic_zero_x"),
+          kResizeQ,
+          +pick_zero_grid,
+          WebGPUDispatchGrid{0u, 1u}),
+      std::runtime_error);
+  EXPECT_THROW(
+      graph.add_dynamic_compute_dispatch(
+          make_dynamic_test_descriptor(graph, "dynamic_zero_y"),
+          kResizeQ,
+          +pick_zero_grid,
+          WebGPUDispatchGrid{1u, 0u}),
+      std::runtime_error);
+  EXPECT_EQ(graph.num_dispatches(), dispatches_before_invalid_trigger);
+
   bool initial_fails = true;
   const size_t dispatches_before_failure = graph.num_dispatches();
   EXPECT_THROW(
diff --git a/backends/webgpu/test/op_tests/cases.py b/backends/webgpu/test/op_tests/cases.py
index be46d35b520..aa8f000bceb 100644
--- a/backends/webgpu/test/op_tests/cases.py
+++ b/backends/webgpu/test/op_tests/cases.py
@@ -112,7 +112,15 @@
     DequantizeConstModule,
     QuantizeModule,
 )
-from executorch.backends.webgpu.test.ops.test_reduce import AmaxModule, AminModule
+from executorch.backends.webgpu.test.ops.test_reduce import (
+    amax_sign_trap_input,
+    amax_tie_input,
+    AmaxModule,
+    amin_sign_trap_input,
+    amin_tie_input,
+    AminModule,
+    EXTREMA_CONFIGS,
+)
 from executorch.backends.webgpu.test.ops.test_repeat import RepeatModule
 from executorch.backends.webgpu.test.ops.test_rms_norm import (
     _CASES,
@@ -489,35 +497,42 @@ def _floor_divide_suite() -> WebGPUTestSuite:
     )
 
 
-def _reduce_suite(module_cls) -> WebGPUTestSuite:
-    # Last-dim reduction; both keepdim variants over a 2d and a 3d shape.
-    return WebGPUTestSuite(
-        module_factory=lambda keepdim: module_cls(keepdim),
-        cases=[
-            Case(name="keepdim_2d", construct={"keepdim": True}, inputs=((M1, M2),)),
-            Case(name="nodim_2d", construct={"keepdim": False}, inputs=((M1, M2),)),
-            Case(
-                name="keepdim_3d",
-                construct={"keepdim": True},
-                inputs=((S, S1, S2),),
-            ),
+def _reduce_suite(module_cls, op: str) -> WebGPUTestSuite:
+    generators = {
+        ("amax", "sign_trap"): amax_sign_trap_input,
+        ("amax", "tie"): amax_tie_input,
+        ("amin", "sign_trap"): amin_sign_trap_input,
+        ("amin", "tie"): amin_tie_input,
+    }
+    cases = []
+    for name, shape, dim, keepdim, input_class in EXTREMA_CONFIGS:
+        inputs = (shape,)
+        kwargs = {}
+        if input_class != "default":
+            inputs = (InputSpec(shape=shape, gen=generators[(op, input_class)]),)
+            kwargs = {"atol": 0.0, "rtol": 0.0}
+        cases.append(
             Case(
-                name="nodim_3d",
-                construct={"keepdim": False},
-                inputs=((S, S1, S2),),
-            ),
-        ],
+                name=name,
+                construct={"keepdim": keepdim, "dim": dim},
+                inputs=inputs,
+                **kwargs,
+            )
+        )
+    return WebGPUTestSuite(
+        module_factory=module_cls,
+        cases=cases,
     )
 
 
 @register_op_test("amax")
 def _amax_suite() -> WebGPUTestSuite:
-    return _reduce_suite(AmaxModule)
+    return _reduce_suite(AmaxModule, "amax")
 
 
 @register_op_test("amin")
 def _amin_suite() -> WebGPUTestSuite:
-    return _reduce_suite(AminModule)
+    return _reduce_suite(AminModule, "amin")
 
 
 @register_op_test("flip")
diff --git a/backends/webgpu/test/ops/test_reduce.py b/backends/webgpu/test/ops/test_reduce.py
index c7b91d8156c..cdfd7eed96c 100644
--- a/backends/webgpu/test/ops/test_reduce.py
+++ b/backends/webgpu/test/ops/test_reduce.py
@@ -19,12 +19,14 @@
 
 from __future__ import annotations
 
+import math
 import unittest
 
 import torch
 
 from executorch.backends.vulkan.partitioner.vulkan_partitioner import VulkanPartitioner
 from executorch.exir import to_edge_transform_and_lower
+from executorch.exir.backend.utils import get_delegates, get_non_lowered_nodes
 
 
 class ReduceModule(torch.nn.Module):
@@ -61,6 +63,70 @@ def _det_input(shape) -> torch.Tensor:
     return ((flat % 17) - 8).div(16.0).reshape(shape)
 
 
+# Shared structural and manifest-driven extrema case authority.
+EXTREMA_CONFIGS = (
+    ("keepdim_2d", (37, 41), -1, True, "default"),
+    ("nodim_2d", (37, 41), -1, False, "default"),
+    ("keepdim_3d", (5, 7, 11), -1, True, "default"),
+    ("nodim_3d", (5, 7, 11), -1, False, "default"),
+    ("sign_trap_63_drop", (3, 63), -1, False, "sign_trap"),
+    ("tie_64_keep", (2, 64), -1, True, "tie"),
+    ("tie_65_posdim_drop", (2, 65), 1, False, "tie"),
+    ("sign_trap_255_keep", (2, 255), -1, True, "sign_trap"),
+    ("tie_256_drop", (2, 256), -1, False, "tie"),
+    ("tie_257_posdim_keep", (2, 257), 1, True, "tie"),
+)
+
+
+def _sign_trap_input(shape, *, for_max: bool) -> torch.Tensor:
+    magnitude = ((torch.arange(math.prod(shape), dtype=torch.float32) % 31) + 1).div(
+        8.0
+    )
+    rows = magnitude.reshape(-1, shape[-1])
+    for row_index, row in enumerate(rows):
+        row.add_(float(row_index))
+    values = -magnitude if for_max else magnitude
+    return values.reshape(shape)
+
+
+def amax_sign_trap_input(shape) -> torch.Tensor:
+    return _sign_trap_input(shape, for_max=True)
+
+
+def amin_sign_trap_input(shape) -> torch.Tensor:
+    return _sign_trap_input(shape, for_max=False)
+
+
+def _tie_input(shape, *, for_max: bool) -> torch.Tensor:
+    values = ((torch.arange(math.prod(shape), dtype=torch.float32) % 29) - 14).div(8.0)
+    rows = values.reshape(-1, shape[-1])
+    for row_index, row in enumerate(rows):
+        extreme = 16.0 + float(row_index)
+        if not for_max:
+            extreme = -extreme
+        row[-2] = extreme
+        row[-1] = extreme
+    return values.reshape(shape)
+
+
+def amax_tie_input(shape) -> torch.Tensor:
+    return _tie_input(shape, for_max=True)
+
+
+def amin_tie_input(shape) -> torch.Tensor:
+    return _tie_input(shape, for_max=False)
+
+
+def _extrema_input(op: str, input_class: str, shape) -> torch.Tensor:
+    if input_class == "default":
+        return _det_input(shape)
+    if input_class == "sign_trap":
+        return (
+            amax_sign_trap_input(shape) if op == "amax" else amin_sign_trap_input(shape)
+        )
+    return amax_tie_input(shape) if op == "amax" else amin_tie_input(shape)
+
+
 def _export(m: torch.nn.Module, x: torch.Tensor):
     ep = torch.export.export(m, (x,))
     return to_edge_transform_and_lower(
@@ -128,21 +194,55 @@ def export_reduce_model(
 
 
 class AmaxModule(torch.nn.Module):
-    def __init__(self, keepdim: bool) -> None:
+    def __init__(self, keepdim: bool, dim: int = -1) -> None:
         super().__init__()
         self.keepdim = keepdim
+        self.dim = dim
 
     def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return torch.amax(x, dim=-1, keepdim=self.keepdim)
+        return torch.amax(x, dim=self.dim, keepdim=self.keepdim)
 
 
 class AminModule(torch.nn.Module):
-    def __init__(self, keepdim: bool) -> None:
+    def __init__(self, keepdim: bool, dim: int = -1) -> None:
         super().__init__()
         self.keepdim = keepdim
+        self.dim = dim
 
     def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return torch.amin(x, dim=-1, keepdim=self.keepdim)
+        return torch.amin(x, dim=self.dim, keepdim=self.keepdim)
+
+
+class TestExtrema(unittest.TestCase):
+    def test_config_contract(self) -> None:
+        self.assertEqual(
+            EXTREMA_CONFIGS,
+            (
+                ("keepdim_2d", (37, 41), -1, True, "default"),
+                ("nodim_2d", (37, 41), -1, False, "default"),
+                ("keepdim_3d", (5, 7, 11), -1, True, "default"),
+                ("nodim_3d", (5, 7, 11), -1, False, "default"),
+                ("sign_trap_63_drop", (3, 63), -1, False, "sign_trap"),
+                ("tie_64_keep", (2, 64), -1, True, "tie"),
+                ("tie_65_posdim_drop", (2, 65), 1, False, "tie"),
+                ("sign_trap_255_keep", (2, 255), -1, True, "sign_trap"),
+                ("tie_256_drop", (2, 256), -1, False, "tie"),
+                ("tie_257_posdim_keep", (2, 257), 1, True, "tie"),
+            ),
+        )
+
+    def test_exports_fully_delegated(self) -> None:
+        for op, module_cls in (("amax", AmaxModule), ("amin", AminModule)):
+            for name, shape, dim, keepdim, input_class in EXTREMA_CONFIGS:
+                with self.subTest(op=op, config=name):
+                    x = _extrema_input(op, input_class, shape)
+                    ep = torch.export.export(module_cls(keepdim, dim).eval(), (x,))
+                    edge = to_edge_transform_and_lower(
+                        ep, partitioner=[VulkanPartitioner()]
+                    )
+                    graph = edge.exported_program().graph_module.graph
+                    self.assertEqual(len(get_delegates(graph)), 1)
+                    self.assertEqual(get_non_lowered_nodes(graph), [])
 
 
 if __name__ == "__main__":
diff --git a/backends/webgpu/test/test_wgsl_codegen.py b/backends/webgpu/test/test_wgsl_codegen.py
index b8c239483f2..d437a0a42e8 100644
--- a/backends/webgpu/test/test_wgsl_codegen.py
+++ b/backends/webgpu/test/test_wgsl_codegen.py
@@ -213,7 +213,7 @@ def test_generated_output_manifest_digest(self) -> None:
         self.assertEqual(len(outputs), 134)
         self.assertEqual(
             digest.hexdigest(),
-            "ab15be30e7cfa2cb2f6fa7743d3b9f03535f5cc0b88d64d9b6bad8f777efda25",
+            "ef97dca2336315ee2c8b0f9e896c6aa082834ae94948bda3ccb42b1145f2bd27",
         )
 
     def test_rope_hf_reconstructs_full_2d_grid_stride(self) -> None:
@@ -952,9 +952,95 @@ def test_to_copy_convert_template_roundtrip_byte_identical(self) -> None:
         )
         self.assertEqual(
             hashlib.sha256(g.registry_path().read_bytes()).hexdigest(),
-            "ce1777820bffe77e7cdda312f86a3fd41a090f8f282a6add66666446d06b1608",
+            "74f972fce4077f12a52dfcf67a0d20ebeea47748283ced9e5c0bcffd659fef74",
         )
 
+    def test_extrema_template_roundtrip_byte_identical(self) -> None:
+        extrema_dir = g.BACKEND_ROOT / "runtime/ops/extrema"
+        template_path = extrema_dir / "extrema.wgsl"
+        spec = g.parse_template_spec(template_path.with_suffix(".yaml"))
+        variants = {params["NAME"]: params for params in spec[template_path.stem]}
+        expected = {
+            "amax": (
+                "max",
+                "35fc059d7c72caa17f9cb1128823ecfd8f75be4ce24b6cd4f9629a97b52f64c0",
+            ),
+            "amin": (
+                "min",
+                "8cb6035ae4d34eb2a6cc973d93d9847905722e967239c96033fccfe3a1943cb2",
+            ),
+        }
+        self.assertEqual(set(variants), set(expected))
+        template = template_path.read_text()
+
+        for name, (reduce_fn, expected_hash) in expected.items():
+            params = variants[name]
+            self.assertEqual(params["REDUCE_FN"], reduce_fn)
+            expanded = g.preprocess(template, {**g.WGSL_HELPERS, **params})
+            self.assertEqual(g.wgsl_sha256(expanded), expected_hash)
+
+            header_path = extrema_dir / f"{name}_wgsl.h"
+            header = header_path.read_text()
+            body = header.split('R"(', 1)[1].split(')";', 1)[0][1:]
+            self.assertEqual(body, expanded)
+            self.assertEqual(g.embedded_sha256(header), expected_hash)
+            self.assertEqual(g.parse_workgroup_size(body), (256, 1, 1))
+
+        entries = {entry.name: entry for entry in g.registry_entries()}
+        for name in expected:
+            self.assertEqual(
+                entries[name].include,
+                f"runtime/ops/extrema/{name}_wgsl.h",
+            )
+            self.assertEqual(entries[name].symbol, g.symbol_base(name))
+
+        handler_hashes = {
+            "amax": "57f929b9f3087dc32403c3587884ce2ed4be2d03c4e80ff7035428b52e7e0e51",
+            "amin": "5dc947d4781a67df953b9c5970c4ab2119317b29657f983a9d308dfdf123dede",
+        }
+        for name, expected_hash in handler_hashes.items():
+            handler = g.BACKEND_ROOT / f"runtime/ops/{name}/Reduce.cpp"
+            self.assertEqual(
+                hashlib.sha256(handler.read_bytes()).hexdigest(), expected_hash
+            )
+
+        self.assertEqual(
+            hashlib.sha256(g.registry_path().read_bytes()).hexdigest(),
+            "74f972fce4077f12a52dfcf67a0d20ebeea47748283ced9e5c0bcffd659fef74",
+        )
+
+    def test_unary_template_roundtrip_byte_identical(self) -> None:
+        unary_dir = g.BACKEND_ROOT / "runtime/ops/unary"
+        template_path = unary_dir / "unary.wgsl"
+        spec = g.parse_template_spec(template_path.with_suffix(".yaml"))
+        variants = {params["NAME"]: params for params in spec[template_path.stem]}
+        expected = {
+            "abs": "39d3c163fdf6a92286828f4b3217e00294e3ca5634a878ed5fd34e3b1cdf0a27",
+            "cos": "9df78873e5fae98d347c26db2a02b047ea3d5d2c93f0761cb9ac6995f9a71ab2",
+            "exp": "3171399bc36acf9c1cb2a03c2a31038318203c4c63ab03c4881df7a660346020",
+            "hardswish": "c874a15ef6cdaec71187296016cc2a1515f5e7c889b97dfa8fd4b278e6e2c3d5",
+            "neg": "8851b9f42d14153f6f04484fee2f8bf67bda26dea892ff48768e09e6ad49cee1",
+            "round": "8f3e0edbeb81aa50f35e691c78554e8057fa8d78fe8a86454f4f42e5e8871452",
+            "rsqrt": "108765d5a23b87473f34651875d08abf2a5fa8980bd92fc8cbe3617295097747",
+            "sin": "e5762804773659d348fddddcef4935807ae6fe7d92c92eb17a2f44aae8f2c5b9",
+            "sqrt": "008534ae365969f5c180b42e8d6d0b131df78f181e5435abbcafc3ffb8be8aac",
+            "tanh": "5bd7eb1c6411940d84a9b311884f35b39f15b82103b14bab02902290ed6b0339",
+        }
+        self.assertEqual(set(variants), set(expected))
+        template = template_path.read_text()
+        entries = {entry.name: entry for entry in g.registry_entries()}
+        for name, expected_hash in expected.items():
+            expanded = g.preprocess(template, {**g.WGSL_HELPERS, **variants[name]})
+            self.assertEqual(g.wgsl_sha256(expanded), expected_hash)
+            header = (unary_dir / f"{name}_wgsl.h").read_text()
+            body = header.split('R"(', 1)[1].split(')";', 1)[0][1:]
+            self.assertEqual(body, expanded)
+            self.assertEqual(g.embedded_sha256(header), expected_hash)
+            self.assertEqual(g.parse_workgroup_size(body), (256, 1, 1))
+            self.assertEqual(entries[name].include, f"runtime/ops/unary/{name}_wgsl.h")
+            self.assertEqual(entries[name].symbol, g.symbol_base(name))
+        self.assertTrue({"clamp", "pow_scalar"}.isdisjoint(variants))
+
     def test_rms_norm_half_variant_is_type_correct(self) -> None:
         # A DTYPE=half expansion must emit compilable WGSL: `enable f16;`, an f32
         # accumulator, loads widened to f32 for the reduction, and the store