sync 6fdf6301e2bb

Browse files

Files changed (11) hide show

README.md +10 -3
build/webgpu/bench.json +276 -1
build/webgpu/manifest.json +44 -27
build/webgpu/mean-variance-normalization-packed-rows.wgsl.jinja +90 -0
build/webgpu/mean-variance-normalization-serial-rows.wgsl.jinja +58 -9
build/webgpu/mean-variance-normalization-subgroup.wgsl.jinja +73 -20
build/webgpu/metadata.json +13 -11
build/webgpu/norm-flat-apply.wgsl.jinja +21 -8
build/webgpu/norm-flat-splitk-combine.wgsl.jinja +1 -5
build/webgpu/norm-flat-splitk-partials.wgsl.jinja +0 -3
build/webgpu/test.json +1785 -23

README.md CHANGED Viewed

@@ -42,12 +42,19 @@ Default values (overridable per request):
 | --- | --- |
 | `T` | `float32`, `float16` |
 ## Files
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
-- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`mean-variance-normalization-serial-rows.wgsl.jinja`](build/webgpu/mean-variance-normalization-serial-rows.wgsl.jinja)
 - [`mean-variance-normalization-subgroup.wgsl.jinja`](build/webgpu/mean-variance-normalization-subgroup.wgsl.jinja)
 - [`noop.wgsl.jinja`](build/webgpu/noop.wgsl.jinja)
@@ -58,7 +65,7 @@ Default values (overridable per request):
 ## Use with `@huggingface/kernels`
 ```sh
-npm install --save-exact @huggingface/kernels@0.0.1-preview.2
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
@@ -72,5 +79,5 @@ Replace each `*Data` placeholder with a typed array containing the corresponding
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/ai.onnx.MeanVarianceNormalization", { version: 1 });
-const { y } = await kernel({ x: { data: xData, shape: [2, 2, 1, 2] } });
 ```

 | --- | --- |
 | `T` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `packed_short_groups` — Caches contiguous short groups in vector words, uses scaled shifted moments, and guards scalar access for partial words.
 ## Files
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
+- [`mean-variance-normalization-packed-rows.wgsl.jinja`](build/webgpu/mean-variance-normalization-packed-rows.wgsl.jinja)
 - [`mean-variance-normalization-serial-rows.wgsl.jinja`](build/webgpu/mean-variance-normalization-serial-rows.wgsl.jinja)
 - [`mean-variance-normalization-subgroup.wgsl.jinja`](build/webgpu/mean-variance-normalization-subgroup.wgsl.jinja)
 - [`noop.wgsl.jinja`](build/webgpu/noop.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/ai.onnx.MeanVarianceNormalization", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [1, 1, 2, 2] } });
 ```

build/webgpu/bench.json CHANGED Viewed

@@ -99,7 +99,282 @@
       "attrs": { "axes": [2] },
       "inputs": { "x": { "dtype": "float32", "shape": [4194304, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
       "outputs": { "y": { "dtype": "float32", "shape": [4194304, 1, 2] } },
-      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
     }
   ]
 }

       "attrs": { "axes": [2] },
       "inputs": { "x": { "dtype": "float32", "shape": [4194304, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
       "outputs": { "y": { "dtype": "float32", "shape": [4194304, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] },
+      "provenance": { "notes": "The bandwidth metric counts each input and output element once." }
+    },
+    {
+      "name": "packed_float32_rows1048576_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [1048576, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048576, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] },
+      "tunableSpace": { "PACKED_WORKGROUP_SIZE": [64, 512] }
+    },
+    {
+      "name": "packed_float32_rows4194239_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [4194239, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4194239, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows4194240_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [4194240, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4194240, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows4194241_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [4194241, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4194241, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows4194304_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [4194304, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4194304, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows2097152_r4",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2097152, 1, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2097152, 1, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows255_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [255, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [255, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows256_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [256, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows257_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [257, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows513_r4",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [513, 1, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [513, 1, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float32_rows65537_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [65537, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [65537, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows1048576_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [1048576, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [1048576, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] },
+      "tunableSpace": { "PACKED_WORKGROUP_SIZE": [64, 512] }
+    },
+    {
+      "name": "packed_float16_rows4194239_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [4194239, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [4194239, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows4194240_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [4194240, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [4194240, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows4194241_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [4194241, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [4194241, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows4194304_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [4194304, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [4194304, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows2097152_r4",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [2097152, 1, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [2097152, 1, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows255_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [255, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [255, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows256_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [256, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows257_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [257, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows513_r4",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [513, 1, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [513, 1, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "packed_float16_rows65537_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [65537, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [65537, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_small_float32_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_strided_float32_r2",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [262144, 2, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [262144, 2, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_small_float32_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_strided_float32_r3",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [262144, 3, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [262144, 3, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_small_float32_r4",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 1, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_strided_float32_r4",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [262144, 4, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [262144, 4, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_small_float16_r2",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [2, 1, 2], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 2] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_strided_float16_r2",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float16", "shape": [262144, 2, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [262144, 2, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_small_float16_r3",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [2, 1, 3], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_strided_float16_r3",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float16", "shape": [262144, 3, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [262144, 3, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_small_float16_r4",
+      "preset": "stress",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [2, 1, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "cached_strided_float16_r4",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float16", "shape": [262144, 4, 4], "dist": "normal", "seed": 951, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [262144, 4, 4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 2 * numel(shapes.x)" }] }
     }
   ]
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -16,13 +16,15 @@
     "VEC4_MIN_REDUCTION": { "default": 8 },
     "FLAT_SPLIT_MIN_ELEMENTS": { "default": 65536 },
     "FLAT_SPLIT_TARGET_ELEMENTS": { "default": 4096 },
-    "MAX_FLAT_SPLITS": { "default": 256 }
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "shapeContract": "ranks.x >= 3 and ranks.x <= 8 and ranks.y == ranks.x and sameShape(shapes.y, shapes.x) and f16Ok(dtypes.T)",
     "reduceCount": "(dim(shapes.x, 0) if hasAxis(attrs.axes, 0, ranks.x) else 1) * (dim(shapes.x, 1) if hasAxis(attrs.axes, 1, ranks.x) else 1) * (dim(shapes.x, 2) if hasAxis(attrs.axes, 2, ranks.x) else 1) * (dim(shapes.x, 3) if ranks.x >= 4 and hasAxis(attrs.axes, 3, ranks.x) else 1) * (dim(shapes.x, 4) if ranks.x >= 5 and hasAxis(attrs.axes, 4, ranks.x) else 1) * (dim(shapes.x, 5) if ranks.x >= 6 and hasAxis(attrs.axes, 5, ranks.x) else 1) * (dim(shapes.x, 6) if ranks.x >= 7 and hasAxis(attrs.axes, 6, ranks.x) else 1) * (dim(shapes.x, 7) if ranks.x >= 8 and hasAxis(attrs.axes, 7, ranks.x) else 1)",
     "rowCount": "numel(shapes.x) / max(1, reduceCount)",
     "allAxesReduced": "hasAxis(attrs.axes, 0, ranks.x) and hasAxis(attrs.axes, 1, ranks.x) and hasAxis(attrs.axes, 2, ranks.x) and (ranks.x < 4 or hasAxis(attrs.axes, 3, ranks.x)) and (ranks.x < 5 or hasAxis(attrs.axes, 4, ranks.x)) and (ranks.x < 6 or hasAxis(attrs.axes, 5, ranks.x)) and (ranks.x < 7 or hasAxis(attrs.axes, 6, ranks.x)) and (ranks.x < 8 or hasAxis(attrs.axes, 7, ranks.x))",
     "vec4Eligible": "((ranks.x == 3 and hasAxis(attrs.axes, 2, 3) and (dim(shapes.x, 2) % 4 == 0 or (hasAxis(attrs.axes, 1, 3) and dim(shapes.x, 1) * dim(shapes.x, 2) % 4 == 0) or (hasAxis(attrs.axes, 0, 3) and hasAxis(attrs.axes, 1, 3) and numel(shapes.x) % 4 == 0))) or (ranks.x == 4 and hasAxis(attrs.axes, 3, 4) and (dim(shapes.x, 3) % 4 == 0 or (hasAxis(attrs.axes, 2, 4) and dim(shapes.x, 2) * dim(shapes.x, 3) % 4 == 0) or (hasAxis(attrs.axes, 1, 4) and hasAxis(attrs.axes, 2, 4) and dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) % 4 == 0) or (hasAxis(attrs.axes, 0, 4) and hasAxis(attrs.axes, 1, 4) and hasAxis(attrs.axes, 2, 4) and numel(shapes.x) % 4 == 0))) or (ranks.x == 5 and hasAxis(attrs.axes, 4, 5) and (dim(shapes.x, 4) % 4 == 0 or (hasAxis(attrs.axes, 3, 5) and dim(shapes.x, 3) * dim(shapes.x, 4) % 4 == 0) or (hasAxis(attrs.axes, 2, 5) and hasAxis(attrs.axes, 3, 5) and dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) % 4 == 0) or (hasAxis(attrs.axes, 1, 5) and hasAxis(attrs.axes, 2, 5) and hasAxis(attrs.axes, 3, 5) and dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) % 4 == 0) or (allAxesReduced and numel(shapes.x) % 4 == 0))) or (ranks.x == 6 and hasAxis(attrs.axes, 5, 6) and (dim(shapes.x, 5) % 4 == 0 or (hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (hasAxis(attrs.axes, 3, 6) and hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 3) * dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (hasAxis(attrs.axes, 2, 6) and hasAxis(attrs.axes, 3, 6) and hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (hasAxis(attrs.axes, 1, 6) and hasAxis(attrs.axes, 2, 6) and hasAxis(attrs.axes, 3, 6) and hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (allAxesReduced and numel(shapes.x) % 4 == 0))))",
@@ -39,19 +41,19 @@
     "vectorStorageFits": "vectorWorkgroupSize * 8 <= device.limits.maxComputeWorkgroupStorageSize",
     "flatSplit": "min(tunables.MAX_FLAT_SPLITS, pow2ceil(ceilDiv(numel(shapes.x), tunables.FLAT_SPLIT_TARGET_ELEMENTS)))",
     "flatScratchBytes": "flatSplit * 8",
-    "flatPathFits": "flatSplit <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and flatScratchBytes <= device.limits.maxStorageBufferBindingSize and flatScratchBytes <= device.limits.maxBufferSize and maxWorkgroupSize * 8 <= device.limits.maxComputeWorkgroupStorageSize and applyDispatchFits"
   },
   "when": ["shapeContract"],
   "bindings": {
-    "x": { "buffer": "read-only-storage", "elementType": "$scalar" },
-    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.x)" }] },
-    "x_2": { "name": "x", "buffer": "read-only-storage", "elementType": "$ioElement" },
-    "y_2": { "name": "y", "buffer": "storage", "elementType": "$ioElement" },
-    "params_4": {
-      "name": "params",
-      "buffer": "uniform",
-      "struct": [{ "name": "rows", "type": "u32", "value": "rowCount" }]
-    }
   },
   "variants": [
     {
@@ -87,18 +89,17 @@
             "splitSpec": "flatSplit",
             "usesF16Spec": "dtypes.T == \"f16\""
           },
-          "bindings": ["x", { "name": "partials", "buffer": "storage", "elementType": "vec2<f32>" }, "params"],
-          "dispatch": { "x": "min(flatSplit, 65535)", "y": "ceilDiv(flatSplit, 65535)", "z": 1 }
         },
         {
           "id": "combine",
           "name": "MeanVarianceNormalization.FlatCombine",
           "shader": "norm-flat-splitk-combine.wgsl.jinja",
-          "derive": { "splitSpec": "flatSplit", "usesF16Spec": "dtypes.T == \"f16\"" },
           "bindings": [
-            "x",
             { "name": "partials", "buffer": "read-only-storage", "elementType": "vec2<f32>" },
-            { "name": "stats", "buffer": "storage", "elementType": "f32", "length": 2 },
             "params"
           ],
           "dispatch": { "x": "min(1, 65535)", "y": "ceilDiv(1, 65535)", "z": 1 }
@@ -107,11 +108,7 @@
           "id": "apply",
           "name": "MeanVarianceNormalization.FlatApply",
           "shader": "norm-flat-apply.wgsl.jinja",
-          "derive": {
-            "workgroupSizeSpec": "maxWorkgroupSize",
-            "scalar": "dtypes.T",
-            "usesF16Spec": "dtypes.T == \"f16\""
-          },
           "bindings": [
             "x",
             { "name": "stats", "buffer": "read-only-storage", "elementType": "f32", "length": 2 },
@@ -126,6 +123,29 @@
         }
       ]
     },
     {
       "id": "serial_rows",
       "priority": 115,
@@ -140,10 +160,9 @@
             "xShape": "shapes.x",
             "reduce": ["hasAxis(attrs.axes, 0, ranks.x)", "hasAxis(attrs.axes, 1, ranks.x)", "hasAxis(attrs.axes, 2, ranks.x)", "hasAxis(attrs.axes, 3, ranks.x)", "hasAxis(attrs.axes, 4, ranks.x)", "hasAxis(attrs.axes, 5, ranks.x)", "hasAxis(attrs.axes, 6, ranks.x)", "hasAxis(attrs.axes, 7, ranks.x)"],
             "workgroupSizeSpec": "serialWorkgroupSize",
-            "scalar": "dtypes.T",
             "usesF16Spec": "dtypes.T == \"f16\""
           },
-          "bindings": ["x_2", "y_2", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((rowCount), (serialWorkgroupSize)), 65535)",
             "y": "ceilDiv(ceilDiv((rowCount), (serialWorkgroupSize)), 65535)",
@@ -162,7 +181,6 @@
           "id": "main",
           "name": "MeanVarianceNormalization.CooperativeVec4",
           "shader": "mean-variance-normalization-subgroup.wgsl.jinja",
-          "subgroupCollectivesWidth": "portable",
           "derive": {
             "xShape": "shapes.x",
             "reduce": ["hasAxis(attrs.axes, 0, ranks.x)", "hasAxis(attrs.axes, 1, ranks.x)", "hasAxis(attrs.axes, 2, ranks.x)", "hasAxis(attrs.axes, 3, ranks.x)", "hasAxis(attrs.axes, 4, ranks.x)", "hasAxis(attrs.axes, 5, ranks.x)", "hasAxis(attrs.axes, 6, ranks.x)", "hasAxis(attrs.axes, 7, ranks.x)"],
@@ -172,7 +190,7 @@
             "usesF16Spec": "dtypes.T == \"f16\"",
             "vectorizedSpec": true
           },
-          "bindings": ["x_2", "y_2", "params_4"],
           "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
@@ -187,7 +205,6 @@
           "id": "main",
           "name": "MeanVarianceNormalization.CooperativeScalar",
           "shader": "mean-variance-normalization-subgroup.wgsl.jinja",
-          "subgroupCollectivesWidth": "portable",
           "derive": {
             "xShape": "shapes.x",
             "reduce": ["hasAxis(attrs.axes, 0, ranks.x)", "hasAxis(attrs.axes, 1, ranks.x)", "hasAxis(attrs.axes, 2, ranks.x)", "hasAxis(attrs.axes, 3, ranks.x)", "hasAxis(attrs.axes, 4, ranks.x)", "hasAxis(attrs.axes, 5, ranks.x)", "hasAxis(attrs.axes, 6, ranks.x)", "hasAxis(attrs.axes, 7, ranks.x)"],
@@ -196,7 +213,7 @@
             "usesF16Spec": "dtypes.T == \"f16\"",
             "vectorizedSpec": false
           },
-          "bindings": ["x_2", "y_2", "params_4"],
           "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]

     "VEC4_MIN_REDUCTION": { "default": 8 },
     "FLAT_SPLIT_MIN_ELEMENTS": { "default": 65536 },
     "FLAT_SPLIT_TARGET_ELEMENTS": { "default": 4096 },
+    "MAX_FLAT_SPLITS": { "default": 256 },
+    "PACKED_WORKGROUP_SIZE": { "default": 512 }
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "shapeContract": "ranks.x >= 3 and ranks.x <= 8 and ranks.y == ranks.x and sameShape(shapes.y, shapes.x) and f16Ok(dtypes.T)",
     "reduceCount": "(dim(shapes.x, 0) if hasAxis(attrs.axes, 0, ranks.x) else 1) * (dim(shapes.x, 1) if hasAxis(attrs.axes, 1, ranks.x) else 1) * (dim(shapes.x, 2) if hasAxis(attrs.axes, 2, ranks.x) else 1) * (dim(shapes.x, 3) if ranks.x >= 4 and hasAxis(attrs.axes, 3, ranks.x) else 1) * (dim(shapes.x, 4) if ranks.x >= 5 and hasAxis(attrs.axes, 4, ranks.x) else 1) * (dim(shapes.x, 5) if ranks.x >= 6 and hasAxis(attrs.axes, 5, ranks.x) else 1) * (dim(shapes.x, 6) if ranks.x >= 7 and hasAxis(attrs.axes, 6, ranks.x) else 1) * (dim(shapes.x, 7) if ranks.x >= 8 and hasAxis(attrs.axes, 7, ranks.x) else 1)",
+    "shortGroup": "reduceCount >= 2 and reduceCount <= 4",
     "rowCount": "numel(shapes.x) / max(1, reduceCount)",
     "allAxesReduced": "hasAxis(attrs.axes, 0, ranks.x) and hasAxis(attrs.axes, 1, ranks.x) and hasAxis(attrs.axes, 2, ranks.x) and (ranks.x < 4 or hasAxis(attrs.axes, 3, ranks.x)) and (ranks.x < 5 or hasAxis(attrs.axes, 4, ranks.x)) and (ranks.x < 6 or hasAxis(attrs.axes, 5, ranks.x)) and (ranks.x < 7 or hasAxis(attrs.axes, 6, ranks.x)) and (ranks.x < 8 or hasAxis(attrs.axes, 7, ranks.x))",
     "vec4Eligible": "((ranks.x == 3 and hasAxis(attrs.axes, 2, 3) and (dim(shapes.x, 2) % 4 == 0 or (hasAxis(attrs.axes, 1, 3) and dim(shapes.x, 1) * dim(shapes.x, 2) % 4 == 0) or (hasAxis(attrs.axes, 0, 3) and hasAxis(attrs.axes, 1, 3) and numel(shapes.x) % 4 == 0))) or (ranks.x == 4 and hasAxis(attrs.axes, 3, 4) and (dim(shapes.x, 3) % 4 == 0 or (hasAxis(attrs.axes, 2, 4) and dim(shapes.x, 2) * dim(shapes.x, 3) % 4 == 0) or (hasAxis(attrs.axes, 1, 4) and hasAxis(attrs.axes, 2, 4) and dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) % 4 == 0) or (hasAxis(attrs.axes, 0, 4) and hasAxis(attrs.axes, 1, 4) and hasAxis(attrs.axes, 2, 4) and numel(shapes.x) % 4 == 0))) or (ranks.x == 5 and hasAxis(attrs.axes, 4, 5) and (dim(shapes.x, 4) % 4 == 0 or (hasAxis(attrs.axes, 3, 5) and dim(shapes.x, 3) * dim(shapes.x, 4) % 4 == 0) or (hasAxis(attrs.axes, 2, 5) and hasAxis(attrs.axes, 3, 5) and dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) % 4 == 0) or (hasAxis(attrs.axes, 1, 5) and hasAxis(attrs.axes, 2, 5) and hasAxis(attrs.axes, 3, 5) and dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) % 4 == 0) or (allAxesReduced and numel(shapes.x) % 4 == 0))) or (ranks.x == 6 and hasAxis(attrs.axes, 5, 6) and (dim(shapes.x, 5) % 4 == 0 or (hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (hasAxis(attrs.axes, 3, 6) and hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 3) * dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (hasAxis(attrs.axes, 2, 6) and hasAxis(attrs.axes, 3, 6) and hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (hasAxis(attrs.axes, 1, 6) and hasAxis(attrs.axes, 2, 6) and hasAxis(attrs.axes, 3, 6) and hasAxis(attrs.axes, 4, 6) and dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.x, 4) * dim(shapes.x, 5) % 4 == 0) or (allAxesReduced and numel(shapes.x) % 4 == 0))))",
     "vectorStorageFits": "vectorWorkgroupSize * 8 <= device.limits.maxComputeWorkgroupStorageSize",
     "flatSplit": "min(tunables.MAX_FLAT_SPLITS, pow2ceil(ceilDiv(numel(shapes.x), tunables.FLAT_SPLIT_TARGET_ELEMENTS)))",
     "flatScratchBytes": "flatSplit * 8",
+    "flatPathFits": "flatSplit <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and flatScratchBytes <= device.limits.maxStorageBufferBindingSize and flatScratchBytes <= device.limits.maxBufferSize and maxWorkgroupSize * 8 <= device.limits.maxComputeWorkgroupStorageSize and applyDispatchFits",
+    "reducedSpan": "1 + ((dim(shapes.x, 0) - 1) * inner(shapes.x, 0) if ranks.x > 0 and hasAxis(attrs.axes, 0, ranks.x) else 0) + ((dim(shapes.x, 1) - 1) * inner(shapes.x, 1) if ranks.x > 1 and hasAxis(attrs.axes, 1, ranks.x) else 0) + ((dim(shapes.x, 2) - 1) * inner(shapes.x, 2) if ranks.x > 2 and hasAxis(attrs.axes, 2, ranks.x) else 0) + ((dim(shapes.x, 3) - 1) * inner(shapes.x, 3) if ranks.x > 3 and hasAxis(attrs.axes, 3, ranks.x) else 0) + ((dim(shapes.x, 4) - 1) * inner(shapes.x, 4) if ranks.x > 4 and hasAxis(attrs.axes, 4, ranks.x) else 0) + ((dim(shapes.x, 5) - 1) * inner(shapes.x, 5) if ranks.x > 5 and hasAxis(attrs.axes, 5, ranks.x) else 0) + ((dim(shapes.x, 6) - 1) * inner(shapes.x, 6) if ranks.x > 6 and hasAxis(attrs.axes, 6, ranks.x) else 0) + ((dim(shapes.x, 7) - 1) * inner(shapes.x, 7) if ranks.x > 7 and hasAxis(attrs.axes, 7, ranks.x) else 0)",
+    "packedRows": "4 if reduceCount % 2 != 0 else (2 if reduceCount % 4 != 0 else 1)",
+    "packedVectors": "reduceCount * packedRows / 4",
+    "packedWorkgroupSize": "tunables.PACKED_WORKGROUP_SIZE if tunables.PACKED_WORKGROUP_SIZE <= deviceWorkgroupCap else min(tunables.SERIAL_TINY_WORKGROUP_SIZE, deviceWorkgroupCap)"
   },
   "when": ["shapeContract"],
   "bindings": {
+    "x": { "elementType": "$scalar" },
+    "params": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.x)" }] },
+    "x_main": { "name": "x", "elementType": "$ioElement" },
+    "y_main": { "name": "y", "elementType": "$ioElement" },
+    "params_rows": { "name": "params", "struct": [{ "name": "rows", "type": "u32", "value": "rowCount" }] }
   },
   "variants": [
     {
             "splitSpec": "flatSplit",
             "usesF16Spec": "dtypes.T == \"f16\""
           },
+          "bindings": ["x", { "name": "partials", "elementType": "vec2<f32>" }, "params"],
+          "dispatch": { "x": "flatSplit" }
         },
         {
           "id": "combine",
           "name": "MeanVarianceNormalization.FlatCombine",
           "shader": "norm-flat-splitk-combine.wgsl.jinja",
+          "derive": { "splitSpec": "flatSplit" },
           "bindings": [
             { "name": "partials", "buffer": "read-only-storage", "elementType": "vec2<f32>" },
+            { "name": "stats", "elementType": "f32", "length": 2 },
             "params"
           ],
           "dispatch": { "x": "min(1, 65535)", "y": "ceilDiv(1, 65535)", "z": 1 }
           "id": "apply",
           "name": "MeanVarianceNormalization.FlatApply",
           "shader": "norm-flat-apply.wgsl.jinja",
+          "derive": { "workgroupSizeSpec": "maxWorkgroupSize", "usesF16Spec": "dtypes.T == \"f16\"" },
           "bindings": [
             "x",
             { "name": "stats", "buffer": "read-only-storage", "elementType": "f32", "length": 2 },
         }
       ]
     },
+    {
+      "id": "packed_short_groups",
+      "priority": 116,
+      "when": ["numel(shapes.x) > 0", "shortGroup", "rowCount >= tunables.SERIAL_MIN_ROWS", "reducedSpan == reduceCount", "ceilDiv(rowCount / packedRows, packedWorkgroupSize) <= foldedDispatchCapacity"],
+      "derive": {
+        "packedTail": "rowCount % packedRows != 0",
+        "ioElement": "dtypes.T if packedTail else \"vec4<\" ~ dtypes.T ~ \">\"",
+        "usesF16Spec": "dtypes.T == \"f16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MeanVarianceNormalization.PackedShortGroups",
+          "shader": "mean-variance-normalization-packed-rows.wgsl.jinja",
+          "bindings": ["x_main", "y_main", "params_rows"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(rowCount, packedRows)), (packedWorkgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(rowCount, packedRows)), (packedWorkgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
     {
       "id": "serial_rows",
       "priority": 115,
             "xShape": "shapes.x",
             "reduce": ["hasAxis(attrs.axes, 0, ranks.x)", "hasAxis(attrs.axes, 1, ranks.x)", "hasAxis(attrs.axes, 2, ranks.x)", "hasAxis(attrs.axes, 3, ranks.x)", "hasAxis(attrs.axes, 4, ranks.x)", "hasAxis(attrs.axes, 5, ranks.x)", "hasAxis(attrs.axes, 6, ranks.x)", "hasAxis(attrs.axes, 7, ranks.x)"],
             "workgroupSizeSpec": "serialWorkgroupSize",
             "usesF16Spec": "dtypes.T == \"f16\""
           },
+          "bindings": ["x_main", "y_main", "params_rows"],
           "dispatch": {
             "x": "min(ceilDiv((rowCount), (serialWorkgroupSize)), 65535)",
             "y": "ceilDiv(ceilDiv((rowCount), (serialWorkgroupSize)), 65535)",
           "id": "main",
           "name": "MeanVarianceNormalization.CooperativeVec4",
           "shader": "mean-variance-normalization-subgroup.wgsl.jinja",
           "derive": {
             "xShape": "shapes.x",
             "reduce": ["hasAxis(attrs.axes, 0, ranks.x)", "hasAxis(attrs.axes, 1, ranks.x)", "hasAxis(attrs.axes, 2, ranks.x)", "hasAxis(attrs.axes, 3, ranks.x)", "hasAxis(attrs.axes, 4, ranks.x)", "hasAxis(attrs.axes, 5, ranks.x)", "hasAxis(attrs.axes, 6, ranks.x)", "hasAxis(attrs.axes, 7, ranks.x)"],
             "usesF16Spec": "dtypes.T == \"f16\"",
             "vectorizedSpec": true
           },
+          "bindings": ["x_main", "y_main", "params_rows"],
           "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
           "id": "main",
           "name": "MeanVarianceNormalization.CooperativeScalar",
           "shader": "mean-variance-normalization-subgroup.wgsl.jinja",
           "derive": {
             "xShape": "shapes.x",
             "reduce": ["hasAxis(attrs.axes, 0, ranks.x)", "hasAxis(attrs.axes, 1, ranks.x)", "hasAxis(attrs.axes, 2, ranks.x)", "hasAxis(attrs.axes, 3, ranks.x)", "hasAxis(attrs.axes, 4, ranks.x)", "hasAxis(attrs.axes, 5, ranks.x)", "hasAxis(attrs.axes, 6, ranks.x)", "hasAxis(attrs.axes, 7, ranks.x)"],
             "usesF16Spec": "dtypes.T == \"f16\"",
             "vectorizedSpec": false
           },
+          "bindings": ["x_main", "y_main", "params_rows"],
           "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]

build/webgpu/mean-variance-normalization-packed-rows.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,90 @@

+{% macro normalize_small_group(tag) %}
+{% set count = reduceCount | int %}
+{% if not usesF16Spec %}
+  var magnitude{{ tag }} = abs(value{{ tag }}_0);
+{% for c in range(1, count) %}
+  magnitude{{ tag }} = max(magnitude{{ tag }}, abs(value{{ tag }}_{{ c }}));
+{% endfor %}
+  let exponent{{ tag }} = (bitcast<u32>(magnitude{{ tag }}) >> 23u) & 255u;
+  let scale{{ tag }} = bitcast<f32>((254u - clamp(exponent{{ tag }}, 1u, 253u)) << 23u);
+{% for c in range(count) %}
+  let scaled{{ tag }}_{{ c }} = value{{ tag }}_{{ c }} * scale{{ tag }};
+{% endfor %}
+{% endif %}
+{% set prefix = "value" if usesF16Spec else "scaled" %}
+  var sum{{ tag }} = 0.0;
+{% for c in range(count) %}
+  let d{{ tag }}_{{ c }} = {{ subtract(prefix ~ tag ~ "_" ~ c, prefix ~ tag ~ "_0") }};
+  sum{{ tag }} = sum{{ tag }} + d{{ tag }}_{{ c }};
+{% endfor %}
+  let mean{{ tag }} = sum{{ tag }} / {{ count }}.0;
+  var square{{ tag }} = 0.0;
+{% for c in range(count) %}
+  let centered{{ tag }}_{{ c }} = {{ subtract("d" ~ tag ~ "_" ~ c, "mean" ~ tag) }};
+  square{{ tag }} = square{{ tag }} + centered{{ tag }}_{{ c }} * centered{{ tag }}_{{ c }};
+{% endfor %}
+  let denom{{ tag }} = sqrt(square{{ tag }} / {{ count }}.0);
+{% for c in range(count) %}
+  let n{{ tag }}_{{ c }} = centered{{ tag }}_{{ c }} / denom{{ tag }};
+{% endfor %}
+{% endmacro %}
+{% macro subtract(value, shift, vector=false) %}
+fma(-1.0, {{ shift }}, {{ value }}){% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+{% set tail = packedTail %}
+{% set R = reduceCount | int %}
+{% set rows = packedRows | int %}
+{% set vectors = packedVectors | int %}
+{% macro element(index) %}v{{ (index / 4) | int }}.{{ "xyzw"[index % 4] }}{% endmacro %}
+{% if tail %}
+fn load_element(index: u32) -> f32 {
+  if (index < params.rows * {{ R }}u) {
+    return f32(x[index]);
+  }
+  return 0.0;
+}
+{% endif %}
+// An invocation retains a whole number of short groups in complete vec4 words.
+// Keeping differences in shifted coordinates avoids rounding the mean back
+// onto the input's larger offset. Explicit fma preserves these intermediate
+// differences through compiler arithmetic reassociation.
+@compute @workgroup_size({{ packedWorkgroupSize }}, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let block = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ packedWorkgroupSize }}u;
+  if (block >= {% if tail %}(params.rows + {{ rows - 1 }}u){% else %}params.rows{% endif %}{% if rows != 1 %} / {{ rows }}u{% endif %}) {
+    return;
+  }
+  let base = block{% if vectors != 1 %} * {{ vectors }}u{% endif %};
+{% for v in range(vectors) %}
+{% if tail %}
+  let v{{ v }} = vec4<f32>(
+{% for c in range(4) %}
+    load_element(base * 4u + {{ v * 4 + c }}u){% if not loop.last %},{% endif %}
+{% endfor %}
+  );
+{% else %}
+  let v{{ v }} = vec4<f32>(x[base{% if v != 0 %} + {{ v }}u{% endif %}]);
+{% endif %}
+{% endfor %}
+{% for row in range(rows) %}
+{% for c in range(R) %}
+  let value{{ row }}_{{ c }} = {{ element(row * R + c) }};
+{% endfor %}
+{{ normalize_small_group(row) }}
+{% endfor %}
+{% for v in range(vectors) %}
+  {% if tail %}let out{{ v }}{% else %}y[base{% if v != 0 %} + {{ v }}u{% endif %}]{% endif %} = {% if usesF16Spec %}vec4<f16>({% endif %}vec4<f32>(
+{% for c in range(4) %}
+    n{{ ((v * 4 + c) / R) | int }}_{{ (v * 4 + c) % R }}{% if not loop.last %},{% endif %}
+{% endfor %}
+  ){% if usesF16Spec %}){% endif %};
+{% if tail %}
+{% for c in range(4) %}
+  if (base * 4u + {{ v * 4 + c }}u < params.rows * {{ R }}u) {
+    y[base * 4u + {{ v * 4 + c }}u] = out{{ v }}.{{ "xyzw"[c] }};
+  }
+{% endfor %}
+{% endif %}
+{% endfor %}
+}

build/webgpu/mean-variance-normalization-serial-rows.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,54 @@
-{% if usesF16Spec %}
-enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // One invocation owns one normalization group. For many short groups, adjacent
@@ -25,20 +73,21 @@ enable f16;
 {% endfor %}
 {% endmacro %}
 const WG: u32 = {{ workgroupSizeSpec }}u;
 const R: u32 = {{ reduceCount }}u;
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(global_invocation_id) gid: vec3<u32>
 ) {
-  // `threads` dispatches fold past the WebGPU x-dimension limit into y.
-  let row = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
-  if (row >= params.rows) {
-    return;
-  }
   var base_off = 0u;
 {{ nd_offset("base_off", "row", 0) | indent(2, true) }}
   let shift = f32(x[base_off]);
   var sum_d = 0.0;
   var sum_d2 = 0.0;
@@ -51,11 +100,11 @@ fn main(
   let mean_d = sum_d / f32(R);
   let variance = max(sum_d2 / f32(R) - mean_d * mean_d, 0.0);
-  let mean = shift + mean_d;
   let denom = sqrt(variance);
   for (var r = 0u; r < R; r = r + 1u) {
     var off = base_off;
-{{ nd_offset("off", "r", 1) | indent(4, true) }}    y[off] = {{ scalar }}((f32(x[off]) - mean) / denom);
   }
 }

+{% macro normalize_small_group(tag) %}
+{% set count = reduceCount | int %}
+{% if not usesF16Spec %}
+  var magnitude{{ tag }} = abs(value{{ tag }}_0);
+{% for c in range(1, count) %}
+  magnitude{{ tag }} = max(magnitude{{ tag }}, abs(value{{ tag }}_{{ c }}));
+{% endfor %}
+  let exponent{{ tag }} = (bitcast<u32>(magnitude{{ tag }}) >> 23u) & 255u;
+  let scale{{ tag }} = bitcast<f32>((254u - clamp(exponent{{ tag }}, 1u, 253u)) << 23u);
+{% for c in range(count) %}
+  let scaled{{ tag }}_{{ c }} = value{{ tag }}_{{ c }} * scale{{ tag }};
+{% endfor %}
 {% endif %}
+{% set prefix = "value" if usesF16Spec else "scaled" %}
+  var sum{{ tag }} = 0.0;
+{% for c in range(count) %}
+  let d{{ tag }}_{{ c }} = {{ subtract(prefix ~ tag ~ "_" ~ c, prefix ~ tag ~ "_0") }};
+  sum{{ tag }} = sum{{ tag }} + d{{ tag }}_{{ c }};
+{% endfor %}
+  let mean{{ tag }} = sum{{ tag }} / {{ count }}.0;
+  var square{{ tag }} = 0.0;
+{% for c in range(count) %}
+  let centered{{ tag }}_{{ c }} = {{ subtract("d" ~ tag ~ "_" ~ c, "mean" ~ tag) }};
+  square{{ tag }} = square{{ tag }} + centered{{ tag }}_{{ c }} * centered{{ tag }}_{{ c }};
+{% endfor %}
+  let denom{{ tag }} = sqrt(square{{ tag }} / {{ count }}.0);
+{% for c in range(count) %}
+  let n{{ tag }}_{{ c }} = centered{{ tag }}_{{ c }} / denom{{ tag }};
+{% endfor %}
+{% endmacro %}
+{% macro indexed_small_group(index_offset, vectorized=false) %}
+{% for c in range(reduceCount | int) %}
+  var off{{ c }} = base_off;
+{{ index_offset("off" ~ c, c ~ "u", 1) | indent(2, true) }}
+  let valueg_{{ c }} = f32(x[off{{ c }}]);
+{% endfor %}
+{{ normalize_small_group("g") }}
+{% for c in range(reduceCount | int) %}
+  y[off{{ c }}] = {{ scalar }}(ng_{{ c }});
+{% endfor %}
+{% endmacro %}
+{% macro subtract(value, shift, vector=false) %}
+fma(-1.0, {{ shift }}, {{ value }}){% endmacro %}
+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 // One invocation owns one normalization group. For many short groups, adjacent
 {% endfor %}
 {% endmacro %}
 const WG: u32 = {{ workgroupSizeSpec }}u;
+{% if not shortGroup %}
 const R: u32 = {{ reduceCount }}u;
+{% endif %}
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(global_invocation_id) gid: vec3<u32>
 ) {
+{{ flat_index_2d("WG", "row", "params.rows") }}
   var base_off = 0u;
 {{ nd_offset("base_off", "row", 0) | indent(2, true) }}
+{% if shortGroup %}
+{{ indexed_small_group(nd_offset) }}
+{% else %}
   let shift = f32(x[base_off]);
   var sum_d = 0.0;
   var sum_d2 = 0.0;
   let mean_d = sum_d / f32(R);
   let variance = max(sum_d2 / f32(R) - mean_d * mean_d, 0.0);
   let denom = sqrt(variance);
   for (var r = 0u; r < R; r = r + 1u) {
     var off = base_off;
+{{ nd_offset("off", "r", 1) | indent(4, true) }}    y[off] = {{ scalar }}({{ subtract(subtract("f32(x[off])", "shift"), "mean_d") }} / denom);
   }
+{% endif %}
 }

build/webgpu/mean-variance-normalization-subgroup.wgsl.jinja CHANGED Viewed

@@ -1,7 +1,47 @@
-{% if usesF16Spec %}
-enable f16;
 {% endif %}
-{% if useSubgroups %}
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
@@ -21,6 +61,7 @@ enable subgroups;
 {% endfor %}
 {% endmacro %}
 const WG: u32 = {{ wg }}u;
 const R: u32 = {{ reduceCount }}u;
 {% if vectorizedSpec %}
 const RV: u32 = R / 4u;
@@ -35,7 +76,7 @@ var<workgroup> sg_partials: array<vec2<f32>, WG>;
 var<workgroup> wg_red: array<vec2<f32>, WG>;
 {% endif %}
-fn reduce_pair(value: vec2<f32>, tid: u32{% if useSubgroups %}, sg_size: u32{% endif %}) -> vec2<f32> {
 {% if useSubgroups %}
   let s = vec2<f32>(subgroupAdd(value.x), subgroupAdd(value.y));
   // The whole workgroup is one subgroup: the collective above already covers it
@@ -46,15 +87,20 @@ fn reduce_pair(value: vec2<f32>, tid: u32{% if useSubgroups %}, sg_size: u32{% e
     return s;
   }
   // Cross-subgroup fold that assumes nothing about which invocations share a
-  // subgroup, how many subgroups there are, or which of a subgroup's lanes are
-  // active: every invocation owns the slot at its own index, the elected lane
   // publishes its subgroup pair there and every other lane publishes the sum
-  // identity. Each subgroup then folds all WG slots — lane `rank`, its dense
-  // position among the active lanes, walks slots rank, rank + count, ... — and
-  // one more collective merges the lane partials, so every slot is added
-  // exactly once at any legal width and partition.
-  let rank = subgroupExclusiveAdd(1u);
-  let count = subgroupAdd(1u);
   sg_partials[tid] = select(vec2<f32>(0.0, 0.0), s, rank == 0u);
   workgroupBarrier();
   var total = vec2<f32>(0.0, 0.0);
@@ -75,23 +121,30 @@ fn reduce_pair(value: vec2<f32>, tid: u32{% if useSubgroups %}, sg_size: u32{% e
 {% endif %}
 }
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg_id: vec3<u32>,
-  @builtin(local_invocation_id) lid: vec3<u32>
-{%- if useSubgroups %},
-  @builtin(subgroup_size) sg_size: u32
-{%- endif %}
 ) {
   let row = wg_id.x + wg_id.y * {{ DISPATCH_FOLD_WIDTH }}u;
   if (row >= params.rows) {
     return;
   }
   let tid = lid.x;
   // Base offset from the kept-axis coordinates.
   var base_off = 0u;
 {{ nd_offset("base_off", "row", 0) | indent(2, true) }}
 {% if vectorizedSpec %}
   let shift = f32(x[base_off / 4u].x);
 {% else %}
@@ -116,10 +169,9 @@ fn main(
 {% endif %}
   }
-  let totals = reduce_pair(acc, tid{% if useSubgroups %}, sg_size{% endif %});
   let mean_d = totals.x / f32(R);
   let variance = max(totals.y / f32(R) - mean_d * mean_d, 0.0);
-  let mean = shift + mean_d;
   let denom = sqrt(variance);
 {% if vectorizedSpec %}
@@ -131,9 +183,10 @@ fn main(
     var off = base_off;
 {{ nd_offset("off", "r", 1) | indent(4, true) }}{% if vectorizedSpec %}
     let v = vec4<f32>(x[off / 4u]);
-    y[off / 4u] = {{ vecType }}((v - vec4<f32>(mean)) / vec4<f32>(denom));
 {% else %}
-    y[off] = {{ scalar }}((f32(x[off]) - mean) / denom);
 {% endif %}
   }
 }

+{% macro normalize_small_group(tag) %}
+{% set count = reduceCount | int %}
+{% if not usesF16Spec %}
+  var magnitude{{ tag }} = abs(value{{ tag }}_0);
+{% for c in range(1, count) %}
+  magnitude{{ tag }} = max(magnitude{{ tag }}, abs(value{{ tag }}_{{ c }}));
+{% endfor %}
+  let exponent{{ tag }} = (bitcast<u32>(magnitude{{ tag }}) >> 23u) & 255u;
+  let scale{{ tag }} = bitcast<f32>((254u - clamp(exponent{{ tag }}, 1u, 253u)) << 23u);
+{% for c in range(count) %}
+  let scaled{{ tag }}_{{ c }} = value{{ tag }}_{{ c }} * scale{{ tag }};
+{% endfor %}
 {% endif %}
+{% set prefix = "value" if usesF16Spec else "scaled" %}
+  var sum{{ tag }} = 0.0;
+{% for c in range(count) %}
+  let d{{ tag }}_{{ c }} = {{ subtract(prefix ~ tag ~ "_" ~ c, prefix ~ tag ~ "_0") }};
+  sum{{ tag }} = sum{{ tag }} + d{{ tag }}_{{ c }};
+{% endfor %}
+  let mean{{ tag }} = sum{{ tag }} / {{ count }}.0;
+  var square{{ tag }} = 0.0;
+{% for c in range(count) %}
+  let centered{{ tag }}_{{ c }} = {{ subtract("d" ~ tag ~ "_" ~ c, "mean" ~ tag) }};
+  square{{ tag }} = square{{ tag }} + centered{{ tag }}_{{ c }} * centered{{ tag }}_{{ c }};
+{% endfor %}
+  let denom{{ tag }} = sqrt(square{{ tag }} / {{ count }}.0);
+{% for c in range(count) %}
+  let n{{ tag }}_{{ c }} = centered{{ tag }}_{{ c }} / denom{{ tag }};
+{% endfor %}
+{% endmacro %}
+{% macro indexed_small_group(index_offset, vectorized=false) %}
+{% for c in range(reduceCount | int) %}
+  var off{{ c }} = base_off;
+{{ index_offset("off" ~ c, c ~ "u", 1) | indent(2, true) }}
+  let valueg_{{ c }} = f32(x[off{{ c }}{% if vectorized %} / 4u][off{{ c }} % 4u{% endif %}]);
+{% endfor %}
+{{ normalize_small_group("g") }}
+{% for c in range(reduceCount | int) %}
+  y[off{{ c }}{% if vectorized %} / 4u][off{{ c }} % 4u{% endif %}] = {{ scalar }}(ng_{{ c }});
+{% endfor %}
+{% endmacro %}
+{% macro subtract(value, shift, vector=false) %}
+fma({% if vector %}vec4<f32>(-1.0), vec4<f32>({{ shift }}){% else %}-1.0, {{ shift }}{% endif %}, {{ value }}){% endmacro %}
+{% if useSubgroups and not shortGroup %}
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% endfor %}
 {% endmacro %}
 const WG: u32 = {{ wg }}u;
+{% if not shortGroup %}
 const R: u32 = {{ reduceCount }}u;
 {% if vectorizedSpec %}
 const RV: u32 = R / 4u;
 var<workgroup> wg_red: array<vec2<f32>, WG>;
 {% endif %}
+fn reduce_pair(value: vec2<f32>, tid: u32{% if useSubgroups %}, sg_lane: u32, sg_size: u32{% endif %}) -> vec2<f32> {
 {% if useSubgroups %}
   let s = vec2<f32>(subgroupAdd(value.x), subgroupAdd(value.y));
   // The whole workgroup is one subgroup: the collective above already covers it
     return s;
   }
   // Cross-subgroup fold that assumes nothing about which invocations share a
+  // subgroup or how many subgroups there are (it does require the uniform
+  // control flow this entry point already has): every invocation owns the slot
+  // at its own index, the elected lane
   // publishes its subgroup pair there and every other lane publishes the sum
+  // identity. Each subgroup then folds all WG slots — lane `sg_lane` walks slots
+  // sg_lane, sg_lane + sg_size, ... — and one more collective merges the lane
+  // partials, so every slot is added exactly once at any legal width.
+  // The coordinates are the BUILTINS, never `subgroupExclusiveAdd(1u)` /
+  // `subgroupAdd(1u)`: those agree with them under this contract, but a driver
+  // in the wild answers a claim derived from them with zero and leaves most of
+  // the workgroup's slots unclaimed. Every lane of this entry point is active
+  // here, so the full subgroup width is the right stride.
+  let rank = sg_lane;
+  let count = sg_size;
   sg_partials[tid] = select(vec2<f32>(0.0, 0.0), s, rank == 0u);
   workgroupBarrier();
   var total = vec2<f32>(0.0, 0.0);
 {% endif %}
 }
+{% endif %}
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wg_id: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>{% if useSubgroups and not shortGroup %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_size) sg_size: u32{% endif %}
 ) {
   let row = wg_id.x + wg_id.y * {{ DISPATCH_FOLD_WIDTH }}u;
   if (row >= params.rows) {
     return;
   }
+{% if shortGroup %}
+  if (lid.x != 0u) { return; }
+{% else %}
   let tid = lid.x;
+{% endif %}
   // Base offset from the kept-axis coordinates.
   var base_off = 0u;
 {{ nd_offset("base_off", "row", 0) | indent(2, true) }}
+{% if shortGroup %}
+{{ indexed_small_group(nd_offset, vectorizedSpec) }}
+{% else %}
 {% if vectorizedSpec %}
   let shift = f32(x[base_off / 4u].x);
 {% else %}
 {% endif %}
   }
+  let totals = reduce_pair(acc, tid{% if useSubgroups %}, sg_lane, sg_size{% endif %});
   let mean_d = totals.x / f32(R);
   let variance = max(totals.y / f32(R) - mean_d * mean_d, 0.0);
   let denom = sqrt(variance);
 {% if vectorizedSpec %}
     var off = base_off;
 {{ nd_offset("off", "r", 1) | indent(4, true) }}{% if vectorizedSpec %}
     let v = vec4<f32>(x[off / 4u]);
+    y[off / 4u] = {{ vecType }}({{ subtract(subtract("v", "shift", true), "mean_d", true) }} / vec4<f32>(denom));
 {% else %}
+    y[off] = {{ scalar }}({{ subtract(subtract("f32(x[off])", "shift"), "mean_d") }} / denom);
 {% endif %}
   }
+{% endif %}
 }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,29 +1,31 @@
 {
   "name": "ai.onnx.MeanVarianceNormalization",
-  "id": "_ai_onnx_meanvariancenormalization_webgpu_b716a78",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "zH0PcLLzlKTcTGM2WzPi8keFUjvDcCnttgcQLxJ9WuU=",
-      "manifest.json": "Snfz0B5Kkd7enbnP+ETHCh/vyME+C3vu4puxVB1xF3c=",
-      "mean-variance-normalization-serial-rows.wgsl.jinja": "qcoGxwmc0jShVeozos6da6awjOJbbIIKLfoMCTyzNuk=",
-      "mean-variance-normalization-subgroup.wgsl.jinja": "GTbYa5JP9Z186OoP6/BZ6w9CyZD2iihXHHl2Cil91pE=",
       "noop.wgsl.jinja": "k/5BMD6UO81N7XlF+t4iSKyt3dbtcqNMCru5aUKNBKE=",
-      "norm-flat-apply.wgsl.jinja": "fJkwUp+SOc5+EV2M+skzaoOyf4cCAj6RE/l9EFpRzlw=",
-      "norm-flat-splitk-combine.wgsl.jinja": "qrWEZGCTfaNnkDSBX7p6j+pnJHKCaAwFmHIXe0DJQC0=",
-      "norm-flat-splitk-partials.wgsl.jinja": "mQFvRpYjvQsoFwdFiShP0nMYFrWbaGmWxYIMXrSi4Sc=",
-      "test.json": "bGHJmclwK4mv21Rm4JYKnWsrwe5GeHdDnBXy+q8fIj8="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "empty_noop": ["noop.wgsl.jinja"],
       "all_axes_flat_split": ["norm-flat-apply.wgsl.jinja", "norm-flat-splitk-combine.wgsl.jinja", "norm-flat-splitk-partials.wgsl.jinja"],
       "serial_rows": ["mean-variance-normalization-serial-rows.wgsl.jinja"],
       "cooperative_vec4": ["mean-variance-normalization-subgroup.wgsl.jinja"],
       "cooperative_scalar": ["mean-variance-normalization-subgroup.wgsl.jinja"]

 {
   "name": "ai.onnx.MeanVarianceNormalization",
+  "id": "_ai_onnx_meanvariancenormalization_webgpu_85e599c",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "evedCasBZFZVduU8ildFMV2Akf76pploR2KtoAy9EIw=",
+      "manifest.json": "iBS4rpeWaBACd7GbnoAYLUB96cJVBEV90PYewcY97vc=",
+      "mean-variance-normalization-packed-rows.wgsl.jinja": "oD40PGfvnT00FNOwz1Kt3uLBo//T6s8ItMXu6AG2Daw=",
+      "mean-variance-normalization-serial-rows.wgsl.jinja": "FzHqHCKU5zFvX4wSxRU+taXe5o4rjkfLckYYx92LztI=",
+      "mean-variance-normalization-subgroup.wgsl.jinja": "Odpm4LwlNRrB7HrK4UI3tRvervrlOfWHcW06frw/s1Y=",
       "noop.wgsl.jinja": "k/5BMD6UO81N7XlF+t4iSKyt3dbtcqNMCru5aUKNBKE=",
+      "norm-flat-apply.wgsl.jinja": "F+5aySjp34jig9aiKiCFcg9fEpGNRNlcenEqMbpzZBY=",
+      "norm-flat-splitk-combine.wgsl.jinja": "hwWrqyz1DrwsOhxMR6OySSAFGWG8aiYb0MsiLdpp1kk=",
+      "norm-flat-splitk-partials.wgsl.jinja": "JYD6TeaKsOyk0ZL+yHUhqh65kHSqL4h2NlC261PqtKY=",
+      "test.json": "PI5ZucTE7woRMpLWN33z6rsoTAnlIsm+udDXiaOzDm4="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
       "empty_noop": ["noop.wgsl.jinja"],
       "all_axes_flat_split": ["norm-flat-apply.wgsl.jinja", "norm-flat-splitk-combine.wgsl.jinja", "norm-flat-splitk-partials.wgsl.jinja"],
+      "packed_short_groups": ["mean-variance-normalization-packed-rows.wgsl.jinja"],
       "serial_rows": ["mean-variance-normalization-serial-rows.wgsl.jinja"],
       "cooperative_vec4": ["mean-variance-normalization-subgroup.wgsl.jinja"],
       "cooperative_scalar": ["mean-variance-normalization-subgroup.wgsl.jinja"]

build/webgpu/norm-flat-apply.wgsl.jinja CHANGED Viewed

@@ -1,13 +1,26 @@
-{% if usesF16Spec %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ workgroupSizeSpec }}, 1, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ workgroupSizeSpec }}u;
-  if (i >= params.count) {
-    return;
-  }
-  y[i] = {{ scalar }}((f32(x[i]) - stats[0]) / stats[1]);
 }

+{% macro subtract(value, shift, vector=false) %}
+fma(-1.0, {{ shift }}, {{ value }}){% endmacro %}
+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ workgroupSizeSpec }}, 1, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(workgroupSizeSpec) }}
+  // Divide the shifted difference BEFORE the mean is subtracted, so no sum ever
+  // mixes a value of order 1 with one of order 1e-8. Subtracting both shifts
+  // first and dividing once is exact as written, but a driver that reassociates
+  // it into x - (x[0] + mean) loses the mean entirely: for inputs one ULP apart
+  // the mean is far below an ULP of x[0], so the result would be {0, 2.12} where
+  // {-0.71, 1.41} is owed. After the division both terms are order 1, so no
+  // reassociation of them can drop one, at the cost of one extra division per
+  // element.
+  let d = {{ subtract("f32(x[i])", "f32(x[0])") }};
+  y[i] = {{ scalar }}(d / stats[1] - stats[0] / stats[1]);
 }

build/webgpu/norm-flat-splitk-combine.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,3 @@
-{% if usesF16Spec %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const SPLIT: u32 = {{ splitSpec }}u;
@@ -11,10 +8,9 @@ fn main() {
   for (var part = 0u; part < SPLIT; part = part + 1u) {
     pair = pair + partials[part];
   }
-  let shift = f32(x[0]);
   let n = f32(params.count);
   let mean_d = pair.x / n;
   let variance = max(pair.y / n - mean_d * mean_d, 0.0);
-  stats[0] = shift + mean_d;
   stats[1] = sqrt(variance);
 }

 {{ env.wgsl.resourceDeclarations }}
 const SPLIT: u32 = {{ splitSpec }}u;
   for (var part = 0u; part < SPLIT; part = part + 1u) {
     pair = pair + partials[part];
   }
   let n = f32(params.count);
   let mean_d = pair.x / n;
   let variance = max(pair.y / n - mean_d * mean_d, 0.0);
+  stats[0] = mean_d;
   stats[1] = sqrt(variance);
 }

build/webgpu/norm-flat-splitk-partials.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,3 @@
-{% if usesF16Spec %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ workgroupSizeSpec }}u;





1	{{ env.wgsl.resourceDeclarations }}
2
3	const WG: u32 = {{ workgroupSizeSpec }}u;

build/webgpu/test.json CHANGED Viewed

@@ -114,15 +114,11 @@
       }
     },
     {
-      "name": "f32_tiny_variance_axis2_gpu_gap",
-      "skipGpu": {
-        "category": "permanent",
-        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal reduced-axis variance collapses to zero so normalization yields Infinity."
-      },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
         "test": "MeanVarianceNormalizationTest.DefaultAxes",
-        "notes": "Valid finite inputs with a positive subnormal variance along the reduced axis; output should be finite instead of +/-Infinity."
       },
       "attrs": { "axes": [2] },
       "inputs": {
@@ -135,15 +131,11 @@
       "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001 } }
     },
     {
-      "name": "f32_tiny_variance_default_axes_rank4_gpu_gap",
-      "skipGpu": {
-        "category": "permanent",
-        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal reduced variance collapses to zero so normalization is non-finite (rank-4)."
-      },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
         "test": "MeanVarianceNormalizationTest.DefaultAxes",
-        "notes": "With rank-4 default spatial axes, a positive subnormal variance must produce a finite normalized output."
       },
       "attrs": { "axes": [0, 2, 3] },
       "inputs": {
@@ -156,15 +148,11 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 2], "tolerance": 0.00001 } }
     },
     {
-      "name": "f32_tiny_variance_default_axes_rank5_gpu_gap",
-      "skipGpu": {
-        "category": "permanent",
-        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal reduced variance collapses to zero so normalization is non-finite (rank-5)."
-      },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
         "test": "MeanVarianceNormalizationTest.DefaultAxes",
-        "notes": "With a rank-5 input, a positive subnormal reduced variance must produce finite normalized values."
       },
       "attrs": { "axes": [0, 2, 3, 4] },
       "inputs": {
@@ -230,9 +218,7 @@
     },
     {
       "name": "channel_axis_no_subgroup_wg64_tail65_f32",
-      "provenance": {
-        "notes": "Locks the 65-value channel reduction where the portable reduction uses 64 fully occupied lanes plus one tail value instead of a half-empty 128-lane tree."
-      },
       "attrs": { "axes": [1] },
       "inputs": {
         "x": {
@@ -524,7 +510,7 @@
     {
       "name": "all_axes_flat_split_f16_65536",
       "provenance": {
-        "notes": "Locks the widened-f32 split-reduction path for large f16 tensors; scratch stays f32 while input/output storage remains f16."
       },
       "attrs": { "axes": [0, 1, 2, 3] },
       "inputs": {
@@ -565,7 +551,7 @@
     {
       "name": "rank5_serial_rows_channel_axis_f16",
       "provenance": {
-        "notes": "Float16 data exercises the coalesced serial-row path with at least SERIAL_MIN_ROWS independent rows and a reduction within SERIAL_MAX_REDUCTION."
       },
       "attrs": { "axes": [1] },
       "inputs": {
@@ -593,6 +579,1782 @@
       "outputs": {
         "y": { "dtype": "float32", "shape": [1, 4, 2, 2, 2, 2, 64], "tolerance": 0.00002, "relTolerance": 0.00002 }
       }
     }
   ]
 }

       }
     },
     {
+      "name": "f32_tiny_variance_axis2",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
         "test": "MeanVarianceNormalizationTest.DefaultAxes",
+        "notes": "The input values are normal f32 values, but their unscaled variance is subnormal. Normalization must remain finite."
       },
       "attrs": { "axes": [2] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001 } }
     },
     {
+      "name": "f32_tiny_variance_default_axes_rank4",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
         "test": "MeanVarianceNormalizationTest.DefaultAxes",
+        "notes": "The input values are normal f32 values, but their unscaled variance is subnormal. Normalization must remain finite."
       },
       "attrs": { "axes": [0, 2, 3] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 2], "tolerance": 0.00001 } }
     },
     {
+      "name": "f32_tiny_variance_default_axes_rank5",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
         "test": "MeanVarianceNormalizationTest.DefaultAxes",
+        "notes": "The input values are normal f32 values, but their unscaled variance is subnormal. Normalization must remain finite."
       },
       "attrs": { "axes": [0, 2, 3, 4] },
       "inputs": {
     },
     {
       "name": "channel_axis_no_subgroup_wg64_tail65_f32",
+      "provenance": { "notes": "A 65-value channel reduction checks the final value beyond a 64-element boundary." },
       "attrs": { "axes": [1] },
       "inputs": {
         "x": {
     {
       "name": "all_axes_flat_split_f16_65536",
       "provenance": {
+        "notes": "A large float16 tensor reduced across all axes checks accurate accumulation and float16 output conversion."
       },
       "attrs": { "axes": [0, 1, 2, 3] },
       "inputs": {
     {
       "name": "rank5_serial_rows_channel_axis_f16",
       "provenance": {
+        "notes": "A float16, rank-5 input reduces over axis 1 (32 channels) across 512 independent groups; verifies normalization accuracy at this row count and reduction size in float16."
       },
       "attrs": { "axes": [1] },
       "inputs": {
       "outputs": {
         "y": { "dtype": "float32", "shape": [1, 4, 2, 2, 2, 2, 64], "tolerance": 0.00002, "relTolerance": 0.00002 }
       }
+    },
+    {
+      "name": "packed_boundary_float32_rows255_r2",
+      "provenance": {
+        "notes": "255 independent 2-element groups (axis 2) sit one row below a 256-row count; verifies the float32 output, including the last row, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [255, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [255, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows256_r2",
+      "provenance": {
+        "notes": "256 independent 2-element groups (axis 2) exactly reach a 256-row count; verifies every float32 row, with none left over, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows257_r2",
+      "provenance": {
+        "notes": "257 independent 2-element groups (axis 2) sit one row above a 256-row count; verifies the extra float32 row, beyond the 256, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows513_r2",
+      "provenance": {
+        "notes": "513 independent 2-element groups (axis 2) sit one row above twice a 256-row count; verifies the extra float32 row at this larger scale matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [513, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [513, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows255_r3",
+      "provenance": {
+        "notes": "255 independent 3-element groups (axis 2) sit one row below a 256-row count; verifies the float32 output, including the last row, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [255, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [255, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows256_r3",
+      "provenance": {
+        "notes": "256 independent 3-element groups (axis 2) exactly reach a 256-row count; verifies every float32 row, with none left over, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows257_r3",
+      "provenance": {
+        "notes": "257 independent 3-element groups (axis 2) sit one row above a 256-row count; verifies the extra float32 row, beyond the 256, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows513_r3",
+      "provenance": {
+        "notes": "513 independent 3-element groups (axis 2) sit one row above twice a 256-row count; verifies the extra float32 row at this larger scale matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [513, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [513, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows255_r4",
+      "provenance": {
+        "notes": "255 independent 4-element groups (axis 2) sit one row below a 256-row count; verifies the float32 output, including the last row, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [255, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [255, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows256_r4",
+      "provenance": {
+        "notes": "256 independent 4-element groups (axis 2) exactly reach a 256-row count; verifies every float32 row, with none left over, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows257_r4",
+      "provenance": {
+        "notes": "257 independent 4-element groups (axis 2) sit one row above a 256-row count; verifies the extra float32 row, beyond the 256, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_boundary_float32_rows513_r4",
+      "provenance": {
+        "notes": "513 independent 4-element groups (axis 2) sit one row above twice a 256-row count; verifies the extra float32 row at this larger scale matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [513, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [513, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows2_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [1.0, 1.0000001192092896] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows2_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [1.0, 1.0000001192092896, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows2_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 8],
+          "data": {
+            "kind": "cycle",
+            "values": [1.0, 1.0000001192092896, 1.0, 1.0, 1.0000001192092896, 1.0, 1.0, 1.0000001192092896]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 8], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows256_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 2],
+          "data": { "kind": "cycle", "values": [1.0, 1.0000001192092896] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows257_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [1.0, 1.0000001192092896, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows256_r4",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 4],
+          "data": { "kind": "cycle", "values": [1.0, 1.0000001192092896, 1.0, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_unit_rows256_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 8],
+          "data": {
+            "kind": "cycle",
+            "values": [1.0, 1.0000001192092896, 1.0, 1.0, 1.0000001192092896, 1.0, 1.0, 1.0000001192092896]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 8], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows2_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [100000000.0, 100000008.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows2_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [100000000.0, 100000008.0, 100000000.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows2_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 8],
+          "data": {
+            "kind": "cycle",
+            "values": [100000000.0, 100000008.0, 100000000.0, 100000000.0, 100000008.0, 100000000.0, 100000000.0, 100000008.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 8], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows256_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 2],
+          "data": { "kind": "cycle", "values": [100000000.0, 100000008.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows257_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [100000000.0, 100000008.0, 100000000.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows256_r4",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 4],
+          "data": { "kind": "cycle", "values": [100000000.0, 100000008.0, 100000000.0, 100000000.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_float32_large_rows256_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [256, 1, 8],
+          "data": {
+            "kind": "cycle",
+            "values": [100000000.0, 100000008.0, 100000000.0, 100000000.0, 100000008.0, 100000000.0, 100000000.0, 100000008.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [256, 1, 8], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "shifted_mean_flat_float32",
+      "provenance": {
+        "notes": "The split combine stores a shifted mean and the apply pass preserves both subtractions."
+      },
+      "attrs": { "axes": [0, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 65536],
+          "data": { "kind": "cycle", "values": [1.0, 1.0000001192092896, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 65536], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_layout_float32_contiguous_axes",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 2, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 2, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_layout_float32_singleton_axes",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 257, 3, 1],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 257, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_layout_float32_strided_axis",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 2, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 2, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "packed_layout_float32_rank8",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [-1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 1, 1, 1, 1, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [257, 1, 1, 1, 1, 1, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "packed_zero_variance_float32",
+      "provenance": { "notes": "Zero variance preserves the normalization contract, including the final partial word." },
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [257, 1, 3], "data": { "kind": "constant", "value": 7.0 } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "tolerance": 0.00001,
+          "relTolerance": 0.00001,
+          "allowNaN": true
+        }
+      }
+    },
+    {
+      "name": "packed_limited_workgroup_float32",
+      "provenance": {
+        "notes": "An oversized requested workgroup must still produce correct normalization within the device workgroup limit."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [257, 1, 3], "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } },
+      "tunables": { "PACKED_WORKGROUP_SIZE": 2048 }
+    },
+    {
+      "name": "packed_boundary_float16_rows255_r2",
+      "provenance": {
+        "notes": "255 independent 2-element groups (axis 2) sit one row below a 256-row count; verifies the float16 output, including the last row, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [255, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [255, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows256_r2",
+      "provenance": {
+        "notes": "256 independent 2-element groups (axis 2) exactly reach a 256-row count; verifies every float16 row, with none left over, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows257_r2",
+      "provenance": {
+        "notes": "257 independent 2-element groups (axis 2) sit one row above a 256-row count; verifies the extra float16 row, beyond the 256, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows513_r2",
+      "provenance": {
+        "notes": "513 independent 2-element groups (axis 2) sit one row above twice a 256-row count; verifies the extra float16 row at this larger scale matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [513, 1, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [513, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows255_r3",
+      "provenance": {
+        "notes": "255 independent 3-element groups (axis 2) sit one row below a 256-row count; verifies the float16 output, including the last row, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [255, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [255, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows256_r3",
+      "provenance": {
+        "notes": "256 independent 3-element groups (axis 2) exactly reach a 256-row count; verifies every float16 row, with none left over, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows257_r3",
+      "provenance": {
+        "notes": "257 independent 3-element groups (axis 2) sit one row above a 256-row count; verifies the extra float16 row, beyond the 256, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows513_r3",
+      "provenance": {
+        "notes": "513 independent 3-element groups (axis 2) sit one row above twice a 256-row count; verifies the extra float16 row at this larger scale matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [513, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [513, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows255_r4",
+      "provenance": {
+        "notes": "255 independent 4-element groups (axis 2) sit one row below a 256-row count; verifies the float16 output, including the last row, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [255, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [255, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows256_r4",
+      "provenance": {
+        "notes": "256 independent 4-element groups (axis 2) exactly reach a 256-row count; verifies every float16 row, with none left over, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows257_r4",
+      "provenance": {
+        "notes": "257 independent 4-element groups (axis 2) sit one row above a 256-row count; verifies the extra float16 row, beyond the 256, matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_boundary_float16_rows513_r4",
+      "provenance": {
+        "notes": "513 independent 4-element groups (axis 2) sit one row above twice a 256-row count; verifies the extra float16 row at this larger scale matches the reference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [513, 1, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [513, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows2_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [2, 1, 2], "data": { "kind": "cycle", "values": [1.0, 1.0009765625] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows2_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [2, 1, 3], "data": { "kind": "cycle", "values": [1.0, 1.0009765625, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows2_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 8],
+          "data": { "kind": "cycle", "values": [1.0, 1.0009765625, 1.0, 1.0, 1.0009765625, 1.0, 1.0, 1.0009765625] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 8], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows256_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [256, 1, 2], "data": { "kind": "cycle", "values": [1.0, 1.0009765625] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows257_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [1.0, 1.0009765625, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows256_r4",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 4],
+          "data": { "kind": "cycle", "values": [1.0, 1.0009765625, 1.0, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_unit_rows256_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 8],
+          "data": { "kind": "cycle", "values": [1.0, 1.0009765625, 1.0, 1.0, 1.0009765625, 1.0, 1.0, 1.0009765625] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 8], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows2_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [2, 1, 2], "data": { "kind": "cycle", "values": [1024.0, 1025.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows2_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [2, 1, 3], "data": { "kind": "cycle", "values": [1024.0, 1025.0, 1024.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows2_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 8],
+          "data": { "kind": "cycle", "values": [1024.0, 1025.0, 1024.0, 1024.0, 1025.0, 1024.0, 1024.0, 1025.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 8], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows256_r2",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [256, 1, 2], "data": { "kind": "cycle", "values": [1024.0, 1025.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows257_r3",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [1024.0, 1025.0, 1024.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows256_r4",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 4],
+          "data": { "kind": "cycle", "values": [1024.0, 1025.0, 1024.0, 1024.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_float16_large_rows256_r8",
+      "provenance": {
+        "notes": "Adjacent representable inputs retain their centered differences without rounding an absolute mean."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [256, 1, 8],
+          "data": { "kind": "cycle", "values": [1024.0, 1025.0, 1024.0, 1024.0, 1025.0, 1024.0, 1024.0, 1025.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [256, 1, 8], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "shifted_mean_flat_float16",
+      "provenance": {
+        "notes": "The split combine stores a shifted mean and the apply pass preserves both subtractions."
+      },
+      "attrs": { "axes": [0, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 1, 65536],
+          "data": { "kind": "cycle", "values": [1.0, 1.0009765625, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 65536], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_layout_float16_contiguous_axes",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 2, 2],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 2, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_layout_float16_singleton_axes",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 257, 3, 1],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 257, 3, 1], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_layout_float16_strided_axis",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 2, 4],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 2, 4], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "packed_layout_float16_rank8",
+      "provenance": {
+        "notes": "Checks reduced-span contiguity and rank-general indexing, including a strided fallback."
+      },
+      "attrs": { "axes": [-1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [257, 1, 1, 1, 1, 1, 1, 3],
+          "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5, 1.25, -0.25, 0.875] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [257, 1, 1, 1, 1, 1, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 }
+      }
+    },
+    {
+      "name": "packed_zero_variance_float16",
+      "provenance": { "notes": "Zero variance preserves the normalization contract, including the final partial word." },
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float16", "shape": [257, 1, 3], "data": { "kind": "constant", "value": 7.0 } } },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [257, 1, 3], "tolerance": 0.002, "relTolerance": 0.001, "allowNaN": true }
+      }
+    },
+    {
+      "name": "packed_limited_workgroup_float16",
+      "provenance": {
+        "notes": "An oversized requested workgroup must still produce correct normalization within the device workgroup limit."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [257, 1, 3], "data": { "kind": "cycle", "values": [0.125, -0.75, 0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [257, 1, 3], "tolerance": 0.002, "relTolerance": 0.001 } },
+      "tunables": { "PACKED_WORKGROUP_SIZE": 2048 }
+    },
+    {
+      "name": "shifted_mean_original_close_pair",
+      "provenance": {
+        "notes": "Two nearly identical values within a row isolate centered-output precision at very small variance."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [-0.35695281624794006, -0.35694074630737305] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r2_e-100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [-7.888609052210118e-31, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r2_e-60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [-8.673617379884035e-19, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r2_e60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [-1152921504606847000.0, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r2_e100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [-1.2676506002282294e+30, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_minimum_normal_rows2_r2",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [1.1754943508222875e-38, 1.175494490952134e-38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_maximum_finite_rows2_r2",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "cycle", "values": [3.4028234663852886e+38, 3.4028232635611926e+38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r2_e-100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [-7.888609052210118e-31, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r2_e-60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [-8.673617379884035e-19, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r2_e60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [-1152921504606847000.0, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r2_e100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [-1.2676506002282294e+30, 0.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_minimum_normal_rows257_r2",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [1.1754943508222875e-38, 1.175494490952134e-38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_maximum_finite_rows257_r2",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [3.4028234663852886e+38, 3.4028232635611926e+38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_strided_f32_r2_e-100",
+      "provenance": { "notes": "Strided groups use the same cached exponent-safe moments as contiguous groups." },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 2, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-7.888609052210118e-31, -7.888609052210118e-31, -7.888609052210118e-31, -7.888609052210118e-31, 0.0, 0.0, 0.0, 0.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 2, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_strided_f32_r2_e100",
+      "provenance": { "notes": "Strided groups use the same cached exponent-safe moments as contiguous groups." },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 2, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.2676506002282294e+30, -1.2676506002282294e+30, -1.2676506002282294e+30, -1.2676506002282294e+30, 0.0, 0.0, 0.0, 0.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 2, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r3_e-100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [-7.888609052210118e-31, 0.0, 7.888609052210118e-31] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r3_e-60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [-8.673617379884035e-19, 0.0, 8.673617379884035e-19] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r3_e60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [-1152921504606847000.0, 0.0, 1152921504606847000.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r3_e100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [-1.2676506002282294e+30, 0.0, 1.2676506002282294e+30] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_minimum_normal_rows2_r3",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "cycle", "values": [1.1754943508222875e-38, 1.175494490952134e-38, 1.1754946310819804e-38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_maximum_finite_rows2_r3",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": {
+            "kind": "cycle",
+            "values": [3.4028234663852886e+38, 3.4028232635611926e+38, 3.4028230607370965e+38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r3_e-100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [-7.888609052210118e-31, 0.0, 7.888609052210118e-31] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r3_e-60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [-8.673617379884035e-19, 0.0, 8.673617379884035e-19] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r3_e60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [-1152921504606847000.0, 0.0, 1152921504606847000.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r3_e100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [-1.2676506002282294e+30, 0.0, 1.2676506002282294e+30] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_minimum_normal_rows257_r3",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": { "kind": "cycle", "values": [1.1754943508222875e-38, 1.175494490952134e-38, 1.1754946310819804e-38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_maximum_finite_rows257_r3",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": {
+            "kind": "cycle",
+            "values": [3.4028234663852886e+38, 3.4028232635611926e+38, 3.4028230607370965e+38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_strided_f32_r3_e-100",
+      "provenance": { "notes": "Strided groups use the same cached exponent-safe moments as contiguous groups." },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 3, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-7.888609052210118e-31, -7.888609052210118e-31, -7.888609052210118e-31, -7.888609052210118e-31, 0.0, 0.0, 0.0, 0.0, 7.888609052210118e-31, 7.888609052210118e-31, 7.888609052210118e-31, 7.888609052210118e-31]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 3, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_strided_f32_r3_e100",
+      "provenance": { "notes": "Strided groups use the same cached exponent-safe moments as contiguous groups." },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 3, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.2676506002282294e+30, -1.2676506002282294e+30, -1.2676506002282294e+30, -1.2676506002282294e+30, 0.0, 0.0, 0.0, 0.0, 1.2676506002282294e+30, 1.2676506002282294e+30, 1.2676506002282294e+30, 1.2676506002282294e+30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 3, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r4_e-100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-7.888609052210118e-31, 0.0, 7.888609052210118e-31, 1.5777218104420236e-30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r4_e-60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-8.673617379884035e-19, 0.0, 8.673617379884035e-19, 1.734723475976807e-18]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r4_e60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1152921504606847000.0, 0.0, 1152921504606847000.0, 2305843009213694000.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows2_r4_e100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.2676506002282294e+30, 0.0, 1.2676506002282294e+30, 2.535301200456459e+30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_minimum_normal_rows2_r4",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [1.1754943508222875e-38, 1.175494490952134e-38, 1.1754946310819804e-38, 1.1754947712118268e-38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_maximum_finite_rows2_r4",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [3.4028234663852886e+38, 3.4028232635611926e+38, 3.4028230607370965e+38, 3.4028228579130005e+38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r4_e-100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-7.888609052210118e-31, 0.0, 7.888609052210118e-31, 1.5777218104420236e-30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r4_e-60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-8.673617379884035e-19, 0.0, 8.673617379884035e-19, 1.734723475976807e-18]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r4_e60",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1152921504606847000.0, 0.0, 1152921504606847000.0, 2305843009213694000.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_f32_rows257_r4_e100",
+      "provenance": {
+        "notes": "Finite normal inputs cover squared-deviation underflow/overflow and neighboring safe exponent controls."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.2676506002282294e+30, 0.0, 1.2676506002282294e+30, 2.535301200456459e+30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_minimum_normal_rows257_r4",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [1.1754943508222875e-38, 1.175494490952134e-38, 1.1754946310819804e-38, 1.1754947712118268e-38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "adjacent_maximum_finite_rows257_r4",
+      "provenance": {
+        "notes": "Adjacent finite f32 values at an exponent endpoint retain their normalized differences."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [3.4028234663852886e+38, 3.4028232635611926e+38, 3.4028230607370965e+38, 3.4028228579130005e+38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_strided_f32_r4_e-100",
+      "provenance": { "notes": "Strided groups use the same cached exponent-safe moments as contiguous groups." },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 4, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-7.888609052210118e-31, -7.888609052210118e-31, -7.888609052210118e-31, -7.888609052210118e-31, 0.0, 0.0, 0.0, 0.0, 7.888609052210118e-31, 7.888609052210118e-31, 7.888609052210118e-31, 7.888609052210118e-31, 1.5777218104420236e-30, 1.5777218104420236e-30, 1.5777218104420236e-30, 1.5777218104420236e-30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 4, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_strided_f32_r4_e100",
+      "provenance": { "notes": "Strided groups use the same cached exponent-safe moments as contiguous groups." },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 4, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.2676506002282294e+30, -1.2676506002282294e+30, -1.2676506002282294e+30, -1.2676506002282294e+30, 0.0, 0.0, 0.0, 0.0, 1.2676506002282294e+30, 1.2676506002282294e+30, 1.2676506002282294e+30, 1.2676506002282294e+30, 2.535301200456459e+30, 2.535301200456459e+30, 2.535301200456459e+30, 2.535301200456459e+30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 4, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "mixed_normal_exponents_r2",
+      "provenance": {
+        "notes": "Widely separated finite exponents normalize without overflowing a centered difference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 2],
+          "data": { "kind": "cycle", "values": [1.1754943508222875e-38, 3.4028234663852886e+38] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 2], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "mixed_normal_exponents_r3",
+      "provenance": {
+        "notes": "Widely separated finite exponents normalize without overflowing a centered difference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 3],
+          "data": {
+            "kind": "cycle",
+            "values": [1.1754943508222875e-38, 3.4028234663852886e+38, -3.4028234663852886e+38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 3], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "mixed_normal_exponents_r4",
+      "provenance": {
+        "notes": "Widely separated finite exponents normalize without overflowing a centered difference."
+      },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [257, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [1.1754943508222875e-38, 3.4028234663852886e+38, -3.4028234663852886e+38, 2.350988701644575e-38]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [257, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "exponent_forced_vec4_float32_e-100",
+      "provenance": { "notes": "A four-value vector group exercises cached moments with vector storage bindings." },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-7.888609052210118e-31, 0.0, 7.888609052210118e-31, 1.5777218104420236e-30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } },
+      "tunables": { "VEC4_MIN_REDUCTION": 4 }
+    },
+    {
+      "name": "exponent_forced_vec4_float32_e100",
+      "provenance": { "notes": "A four-value vector group exercises cached moments with vector storage bindings." },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.2676506002282294e+30, 0.0, 1.2676506002282294e+30, 2.535301200456459e+30]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 4], "tolerance": 0.00001, "relTolerance": 0.00001 } },
+      "tunables": { "VEC4_MIN_REDUCTION": 4 }
+    },
+    {
+      "name": "exponent_forced_vec4_float16_e-14",
+      "provenance": { "notes": "A four-value vector group exercises cached moments with vector storage bindings." },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 4],
+          "data": { "kind": "cycle", "values": [-0.00006103515625, 0.0, 0.00006103515625, 0.0001220703125] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } },
+      "tunables": { "VEC4_MIN_REDUCTION": 4 }
+    },
+    {
+      "name": "exponent_forced_vec4_float16_e14",
+      "provenance": { "notes": "A four-value vector group exercises cached moments with vector storage bindings." },
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 1, 4],
+          "data": { "kind": "cycle", "values": [-16384.0, 0.0, 16384.0, 32768.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 1, 4], "tolerance": 0.002, "relTolerance": 0.001 } },
+      "tunables": { "VEC4_MIN_REDUCTION": 4 }
     }
   ]
 }