sync 2e7068faf55e

Browse files

Files changed (11) hide show

README.md +69 -0
build/webgpu/bench.json +106 -0
build/webgpu/manifest.json +273 -0
build/webgpu/mean-variance-normalization-serial-rows.wgsl.jinja +62 -0
build/webgpu/mean-variance-normalization-subgroup.wgsl.jinja +136 -0
build/webgpu/metadata.json +23 -0
build/webgpu/noop.wgsl.jinja +4 -0
build/webgpu/norm-flat-apply.wgsl.jinja +14 -0
build/webgpu/norm-flat-splitk-combine.wgsl.jinja +20 -0
build/webgpu/norm-flat-splitk-partials.wgsl.jinja +42 -0
build/webgpu/test.json +582 -0

README.md CHANGED Viewed

@@ -1,3 +1,72 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.MeanVarianceNormalization
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Normalizes each group as `(X - mean) / sqrt(variance)`, reducing over `axes` (default `[0, 2, 3]`).
+See the [ONNX `MeanVarianceNormalization` spec](https://onnx.ai/onnx/operators/onnx__MeanVarianceNormalization.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Input tensor to normalize. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | same as `X` | Normalized tensor with the same shape as `X`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axes` | `[0,2,3]` | Axes that share a mean and variance; negative values count from the back. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, 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)
+- [`norm-flat-apply.wgsl.jinja`](build/webgpu/norm-flat-apply.wgsl.jinja)
+- [`norm-flat-splitk-combine.wgsl.jinja`](build/webgpu/norm-flat-splitk-combine.wgsl.jinja)
+- [`norm-flat-splitk-partials.wgsl.jinja`](build/webgpu/norm-flat-splitk-partials.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+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] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,106 @@

+{
+  "op": "ai.onnx.MeanVarianceNormalization",
+  "cases": [
+    {
+      "name": "1x32x32x32_default_axes",
+      "preset": "smoke",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [1, 32, 32, 32] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 32, 32] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "8x64x128x128_default_axes",
+      "preset": "smoke",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 64, 128, 128] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 128, 128] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "16x64x64x64_spatial_axes",
+      "preset": "smoke",
+      "attrs": { "axes": [2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [16, 64, 64, 64] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16, 64, 64, 64] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "f16_vec4_8x64x128x128_default_axes",
+      "preset": "smoke",
+      "vars": { "dtype": "float16" },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": { "x": { "dtype": "float16", "shape": [8, 64, 128, 128], "dist": "normal", "seed": 900, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [8, 64, 128, 128] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "3 * dtypeBytes(args.dtype) * numel(shapes.x)" }]
+      }
+    },
+    {
+      "name": "f16_all_axes_flat_split_1x1x256x256",
+      "preset": "smoke",
+      "attrs": { "axes": [0, 1, 2, 3] },
+      "inputs": { "x": { "dtype": "float16", "shape": [1, 1, 256, 256], "dist": "normal", "seed": 906, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 256, 256] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 2 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "vec4_align_cliff_8x64x128x130_unaligned_W",
+      "preset": "smoke",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 64, 128, 130], "dist": "normal", "seed": 901, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 128, 130] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "channel_axis_scalar_strided_8x64x128x128",
+      "preset": "smoke",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 64, 128, 128], "dist": "normal", "seed": 902, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 128, 128] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "channel_axis_wg_pow2_cliff_8x65x128x128",
+      "preset": "stress",
+      "attrs": { "axes": [1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 65, 128, 128], "dist": "normal", "seed": 905, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 65, 128, 128] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "rank3_scalar_8x4096x1024_reduce_hidden",
+      "preset": "stress",
+      "provenance": { "notes": "Stress-only capacity case: input plus output occupy exactly 256 MiB of GPU storage." },
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 4096, 1024], "dist": "normal", "seed": 903, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 4096, 1024] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "rank3_rows1_decode_hidden768",
+      "preset": "smoke",
+      "attrs": { "axes": [2] },
+      "inputs": { "x": { "dtype": "float32", "shape": [1, 1, 768], "dist": "normal", "seed": 904, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 768] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "stress_all_axes_single_workgroup_rank4_3p9m",
+      "preset": "stress",
+      "attrs": { "axes": [0, 1, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [64, 64, 64, 15], "dist": "normal", "seed": 950, "scale": 0.5 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [64, 64, 64, 15] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "3 * 4 * numel(shapes.x)" }] }
+    },
+    {
+      "name": "stress_tiny_reduce_launchbound_wg2_4m_rows",
+      "preset": "stress",
+      "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)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,273 @@

+{
+  "domain": "ai.onnx",
+  "name": "MeanVarianceNormalization",
+  "sinceVersion": 13,
+  "description": "Normalizes each group as `(X - mean) / sqrt(variance)`, reducing over `axes` (default `[0, 2, 3]`).",
+  "inputs": [{ "role": "X", "dtype": "T", "description": "Input tensor to normalize." }],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.X",
+      "description": "Normalized tensor with the same shape as `X`.",
+      "shape": "shapes.X"
+    }
+  ],
+  "attributes": { "axes": [0, 2, 3] },
+  "attributeDescriptions": { "axes": "Axes that share a mean and variance; negative values count from the back." },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "SERIAL_WORKGROUP_SIZE": 256,
+    "SERIAL_TINY_WORKGROUP_SIZE": 64,
+    "SERIAL_MAX_REDUCTION": 128,
+    "SERIAL_MIN_ROWS": 256,
+    "TREE_MEDIUM_WORKGROUP_SIZE": 64,
+    "VEC4_MIN_REDUCTION": 8,
+    "FLAT_SPLIT_MIN_ELEMENTS": 65536,
+    "FLAT_SPLIT_TARGET_ELEMENTS": 4096,
+    "MAX_FLAT_SPLITS": 256
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "foldedDispatchCapacity": "device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension",
+    "shapeContract": "ranks.X >= 3 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)",
+    "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))",
+    "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))))",
+    "maxWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "minSubgroupSize": "device.adapterInfo.subgroupMinSize if has(device.adapterInfo, \"subgroupMinSize\") else 1",
+    "useSubgroups": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and minSubgroupSize > 0",
+    "scalarWorkgroupSize": "min(maxWorkgroupSize, tunables.TREE_MEDIUM_WORKGROUP_SIZE) if not useSubgroups and reduceCount > tunables.TREE_MEDIUM_WORKGROUP_SIZE and reduceCount <= 2 * tunables.TREE_MEDIUM_WORKGROUP_SIZE else min(maxWorkgroupSize, max(1, pow2ceil(reduceCount)))",
+    "vectorWorkgroupSize": "min(maxWorkgroupSize, max(1, pow2ceil(ceilDiv(reduceCount, 4))))",
+    "serialWorkgroupSize": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, (tunables.SERIAL_TINY_WORKGROUP_SIZE if reduceCount <= 4 else tunables.SERIAL_WORKGROUP_SIZE))",
+    "rowDispatchFits": "rowCount <= foldedDispatchCapacity",
+    "serialDispatchFits": "ceilDiv(rowCount, serialWorkgroupSize) <= foldedDispatchCapacity",
+    "applyDispatchFits": "ceilDiv(numel(shapes.Y), maxWorkgroupSize) <= foldedDispatchCapacity",
+    "scalarStorageFits": "scalarWorkgroupSize * 8 <= device.limits.maxComputeWorkgroupStorageSize",
+    "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 <= device.limits.maxComputeWorkgroupsPerDimension and flatScratchBytes <= device.limits.maxStorageBufferBindingSize and flatScratchBytes <= device.limits.maxBufferSize and maxWorkgroupSize * 8 <= device.limits.maxComputeWorkgroupStorageSize and applyDispatchFits"
+  },
+  "bindingSets": {
+    "rows": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$ioElement" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "rows", "type": "u32", "value": "rowCount" }] }
+      }
+    ],
+    "flatPartials": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "vec2<f32>" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
+      }
+    ],
+    "flatCombine": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "partials",
+        "semantic": "partials",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec2<f32>"
+      },
+      { "name": "stats", "semantic": "stats", "buffer": { "type": "storage" }, "elementType": "f32", "length": 2 },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
+      }
+    ],
+    "flatApply": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "stats",
+        "semantic": "stats",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 2
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" }] }
+      }
+    ],
+    "noopParams": [
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "dummy", "type": "u32", "value": 0 }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "empty_noop",
+      "priority": 200,
+      "when": ["shapeContract", "numel(shapes.X) == 0"],
+      "passes": [
+        {
+          "id": "noop",
+          "name": "MeanVarianceNormalization.Empty",
+          "shader": "noop.wgsl.jinja",
+          "bindings": "noopParams",
+          "dispatch": { "x": 0 }
+        }
+      ]
+    },
+    {
+      "id": "all_axes_flat_split",
+      "priority": 120,
+      "when": ["shapeContract", "numel(shapes.X) > 0", "allAxesReduced", "numel(shapes.X) >= tunables.FLAT_SPLIT_MIN_ELEMENTS", "flatPathFits"],
+      "constants": { "scalar": "dtypes.T" },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[flatSplit, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[2]" }
+      ],
+      "passes": [
+        {
+          "id": "partials",
+          "name": "MeanVarianceNormalization.FlatPartials",
+          "source": {
+            "shader": "norm-flat-splitk-partials.wgsl.jinja",
+            "inputs": { "workgroupSize": "maxWorkgroupSize", "split": "flatSplit", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "flatPartials",
+          "dispatch": { "workgroups": "flatSplit" }
+        },
+        {
+          "id": "combine",
+          "name": "MeanVarianceNormalization.FlatCombine",
+          "source": {
+            "shader": "norm-flat-splitk-combine.wgsl.jinja",
+            "inputs": { "split": "flatSplit", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "flatCombine",
+          "dispatch": { "workgroups": 1 }
+        },
+        {
+          "id": "apply",
+          "name": "MeanVarianceNormalization.FlatApply",
+          "source": {
+            "shader": "norm-flat-apply.wgsl.jinja",
+            "inputs": { "workgroupSize": "maxWorkgroupSize", "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" }
+          },
+          "bindings": "flatApply",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "maxWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "serial_rows",
+      "priority": 115,
+      "when": ["shapeContract", "numel(shapes.X) > 0", "reduceCount <= tunables.SERIAL_MAX_REDUCTION", "rowCount >= tunables.SERIAL_MIN_ROWS", "serialDispatchFits"],
+      "constants": { "scalar": "dtypes.T", "ioElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MeanVarianceNormalization.SerialRows",
+          "source": {
+            "shader": "mean-variance-normalization-serial-rows.wgsl.jinja",
+            "inputs": {
+              "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)"],
+              "reduceCount": "reduceCount",
+              "workgroupSize": "serialWorkgroupSize",
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\""
+            }
+          },
+          "bindings": "rows",
+          "dispatch": { "threads": "rowCount", "workgroupSize": "serialWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "cooperative_vec4",
+      "priority": 110,
+      "when": ["shapeContract", "numel(shapes.X) > 0", "reduceCount >= tunables.VEC4_MIN_REDUCTION", "vec4Eligible", "rowDispatchFits", "vectorStorageFits"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "ioElement": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MeanVarianceNormalization.CooperativeVec4",
+          "source": {
+            "shader": "mean-variance-normalization-subgroup.wgsl.jinja",
+            "inputs": {
+              "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)"],
+              "reduceCount": "reduceCount",
+              "wg": "vectorWorkgroupSize",
+              "minSubgroupSize": "minSubgroupSize",
+              "maxSubgroups": "ceilDiv(vectorWorkgroupSize, minSubgroupSize)",
+              "scalar": "dtypes.T",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "usesF16": "dtypes.T == \"f16\"",
+              "useSubgroups": "useSubgroups",
+              "vectorized": true
+            }
+          },
+          "bindings": "rows",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    },
+    {
+      "id": "cooperative_scalar",
+      "priority": 100,
+      "when": ["shapeContract", "numel(shapes.X) > 0", "rowDispatchFits", "scalarStorageFits"],
+      "constants": { "scalar": "dtypes.T", "ioElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "MeanVarianceNormalization.CooperativeScalar",
+          "source": {
+            "shader": "mean-variance-normalization-subgroup.wgsl.jinja",
+            "inputs": {
+              "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)"],
+              "reduceCount": "reduceCount",
+              "wg": "scalarWorkgroupSize",
+              "minSubgroupSize": "minSubgroupSize",
+              "maxSubgroups": "ceilDiv(scalarWorkgroupSize, minSubgroupSize)",
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "useSubgroups": "useSubgroups",
+              "vectorized": false
+            }
+          },
+          "bindings": "rows",
+          "dispatch": { "workgroups": "rowCount" }
+        }
+      ]
+    }
+  ]
+}

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

	@@ -0,0 +1,62 @@

+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// One invocation owns one normalization group. For many short groups, adjacent
+// invocations visit the same reduced coordinate, making their strided scalar
+// loads adjacent across execution lanes.
+//
+// The reduction is shifted by its first value before accumulating moments. This
+// stabilizes the variance calculation while retaining the zero-variance ONNX
+// result as 0/0 = NaN.
+{% macro nd_offset(target, index, reduced, indent) %}
+{% set rank = source.xShape | length %}
+{% for i in range(rank) %}
+{% if source.reduce[i] if reduced else not source.reduce[i] %}
+{% set pa = namespace(v=1) %}
+{% set st = namespace(v=1) %}
+{% for j in range(i + 1, rank) %}
+{% set st.v = st.v * source.xShape[j] %}
+{% if source.reduce[j] if reduced else not source.reduce[j] %}{% set pa.v = pa.v * source.xShape[j] %}{% endif %}
+{% endfor %}
+{{ indent }}{{ target }} = {{ target }} + (({{ index }} / {{ pa.v }}u) % {{ source.xShape[i] }}u) * {{ st.v }}u;
+{% endif %}
+{% endfor %}
+{% endmacro %}
+const WG: u32 = {{ source.workgroupSize }}u;
+const R: u32 = {{ source.reduceCount }}u;
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  // `threads` dispatches fold past the WebGPU x-dimension limit into y.
+  let row = gid.x + gid.y * nwg.x * WG;
+  if (row >= params.rows) {
+    return;
+  }
+  var base_off = 0u;
+{{ nd_offset("base_off", "row", 0, "  ") }}
+  let shift = f32(x[base_off]);
+  var sum_d = 0.0;
+  var sum_d2 = 0.0;
+  for (var r = 0u; r < R; r = r + 1u) {
+    var off = base_off;
+{{ nd_offset("off", "r", 1, "    ") }}    let d = f32(x[off]) - shift;
+    sum_d = sum_d + d;
+    sum_d2 = sum_d2 + d * d;
+  }
+  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, "    ") }}    y[off] = {{ source.scalar }}((f32(x[off]) - mean) / denom);
+  }
+}

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

	@@ -0,0 +1,136 @@

+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{% if source.useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% macro nd_offset(target, index, reduced, indent) %}
+{% set rank = source.xShape | length %}
+{% for i in range(rank) %}
+{% if source.reduce[i] if reduced else not source.reduce[i] %}
+{% set pa = namespace(v=1) %}
+{% set st = namespace(v=1) %}
+{% for j in range(i + 1, rank) %}
+{% set st.v = st.v * source.xShape[j] %}
+{% if source.reduce[j] if reduced else not source.reduce[j] %}{% set pa.v = pa.v * source.xShape[j] %}{% endif %}
+{% endfor %}
+{{ indent }}{{ target }} = {{ target }} + (({{ index }} / {{ pa.v }}u) % {{ source.xShape[i] }}u) * {{ st.v }}u;
+{% endif %}
+{% endfor %}
+{% endmacro %}
+const WG: u32 = {{ source.wg }}u;
+const R: u32 = {{ source.reduceCount }}u;
+{% if source.vectorized %}
+const RV: u32 = R / 4u;
+{% endif %}
+{% if source.useSubgroups %}
+{% if source.wg > source.minSubgroupSize %}
+// The device-reported minimum subgroup width bounds cross-subgroup partials.
+const MAX_SG: u32 = {{ source.maxSubgroups }}u;
+var<workgroup> sg_partials: array<vec2<f32>, MAX_SG>;
+{% endif %}
+{% else %}
+var<workgroup> wg_red: array<vec2<f32>, WG>;
+{% endif %}
+fn reduce_pair(value: vec2<f32>{% if not source.useSubgroups or source.wg > source.minSubgroupSize %}, tid: u32{% endif %}{% if source.useSubgroups and source.wg > source.minSubgroupSize %}, sg_lane: u32, sg_size: u32{% endif %}) -> vec2<f32> {
+{% if source.useSubgroups %}
+  let s = vec2<f32>(subgroupAdd(value.x), subgroupAdd(value.y));
+{% if source.wg > source.minSubgroupSize %}
+  if (sg_lane == 0u) {
+    sg_partials[tid / sg_size] = s;
+  }
+  workgroupBarrier();
+  let num_sg = (WG + sg_size - 1u) / sg_size;
+  var total = vec2<f32>(0.0, 0.0);
+  for (var i = 0u; i < num_sg; i = i + 1u) {
+    total = total + sg_partials[i];
+  }
+  return total;
+{% else %}
+  return s;
+{% endif %}
+{% else %}
+  wg_red[tid] = value;
+  workgroupBarrier();
+  for (var step = WG >> 1u; step > 0u; step = step >> 1u) {
+    if (tid < step) {
+      wg_red[tid] = wg_red[tid] + wg_red[tid + step];
+    }
+    workgroupBarrier();
+  }
+  return wg_red[0];
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg_id: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+{%- if source.useSubgroups and source.wg > source.minSubgroupSize %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_size) sg_size: u32
+{%- endif %}
+) {
+  let row = wg_id.x + wg_id.y * nwg.x;
+  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, "  ") }}
+{% if source.vectorized %}
+  let shift = f32(x[base_off / 4u].x);
+{% else %}
+  let shift = f32(x[base_off]);
+{% endif %}
+  var acc = vec2<f32>(0.0, 0.0);
+{% if source.vectorized %}
+  for (var q = tid; q < RV; q = q + WG) {
+    let r = q * 4u;
+{% else %}
+  for (var r = tid; r < R; r = r + WG) {
+{% endif %}
+    var off = base_off;
+{{ nd_offset("off", "r", 1, "    ") }}{% if source.vectorized %}
+    let d = vec4<f32>(x[off / 4u]) - vec4<f32>(shift);
+    acc.x = acc.x + d.x + d.y + d.z + d.w;
+    acc.y = acc.y + dot(d, d);
+{% else %}
+    let d = f32(x[off]) - shift;
+    acc.x = acc.x + d;
+    acc.y = acc.y + d * d;
+{% endif %}
+  }
+{% if source.useSubgroups %}
+  let totals = reduce_pair(acc{% if source.wg > source.minSubgroupSize %}, tid, sg_lane, sg_size{% endif %});
+{% else %}
+  let totals = reduce_pair(acc, tid);
+{% 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 source.vectorized %}
+  for (var q = tid; q < RV; q = q + WG) {
+    let r = q * 4u;
+{% else %}
+  for (var r = tid; r < R; r = r + WG) {
+{% endif %}
+    var off = base_off;
+{{ nd_offset("off", "r", 1, "    ") }}{% if source.vectorized %}
+    let v = vec4<f32>(x[off / 4u]);
+    y[off / 4u] = {{ source.vecType }}((v - vec4<f32>(mean)) / vec4<f32>(denom));
+{% else %}
+    y[off] = {{ source.scalar }}((f32(x[off]) - mean) / denom);
+{% endif %}
+  }
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,23 @@

+{
+  "name": "ai.onnx.MeanVarianceNormalization",
+  "id": "_ai_onnx_meanvariancenormalization_webgpu_09532e8",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "Z7dtgwBXEEaVuYjwvsi+nzdPnqr+T3GbPX6F4A4qUyQ=",
+      "manifest.json": "8549q1hb8wFCRzXlrsGTXzEyEVO5WtP64OtKI3Hq6Ug=",
+      "mean-variance-normalization-serial-rows.wgsl.jinja": "C/TVvmYhaLeq+yKB4nAK535avgH9kcYaeSZKnRY/Df0=",
+      "mean-variance-normalization-subgroup.wgsl.jinja": "eFY8+N5f0/V/npXvheysb002bNEWtJLxNE//2aGzIr0=",
+      "noop.wgsl.jinja": "k/5BMD6UO81N7XlF+t4iSKyt3dbtcqNMCru5aUKNBKE=",
+      "norm-flat-apply.wgsl.jinja": "O8G9eyv748OmW3GMQjqSFbgsiJ/pW/plJAzAbFGvGNc=",
+      "norm-flat-splitk-combine.wgsl.jinja": "uBFuaqUxbf6Qr2uE5huMcdP2+S2nHF+zM3RH0tzAx9o=",
+      "norm-flat-splitk-partials.wgsl.jinja": "YwepfH5ztse2kwIz4xhl80Gw5GEnWNMue0oGZUD47p8=",
+      "test.json": "OAgHM6FyfQCVSV8h47D92KbimR901cmqk1fTTm4g3q8="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.MeanVarianceNormalization" }
+}

build/webgpu/noop.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,4 @@

+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size(1)
+fn main() {}

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

	@@ -0,0 +1,14 @@

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

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

	@@ -0,0 +1,20 @@

+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const SPLIT: u32 = {{ source.split }}u;
+@compute @workgroup_size(1, 1, 1)
+fn main() {
+  var pair = vec2<f32>(0.0);
+  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);
+}

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

	@@ -0,0 +1,42 @@

+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ source.workgroupSize }}u;
+const SPLIT: u32 = {{ source.split }}u;
+var<workgroup> reduction: array<vec2<f32>, WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let part = wg.x;
+  let tid = lid.x;
+  let chunk = (params.count + SPLIT - 1u) / SPLIT;
+  let start = part * chunk;
+  let end = min(start + chunk, params.count);
+  // Shifting by the first element keeps the two-moment variance stable for
+  // tensors with a large common offset. Variance is invariant to the shift.
+  let shift = f32(x[0]);
+  var pair = vec2<f32>(0.0);
+  for (var i = start + tid; i < end; i = i + WG) {
+    let d = f32(x[i]) - shift;
+    pair = pair + vec2<f32>(d, d * d);
+  }
+  reduction[tid] = pair;
+  workgroupBarrier();
+  for (var stride = WG >> 1u; stride > 0u; stride = stride >> 1u) {
+    if (tid < stride) {
+      reduction[tid] = reduction[tid] + reduction[tid + stride];
+    }
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    partials[part] = reduction[0];
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,582 @@

+{
+  "op": "ai.onnx.MeanVarianceNormalization",
+  "fixtureArrays": {
+    "ort_axes_subset_rank5_0_2_4_input_x": [0.6369617, 0.2697867, 0.0409735, 0.0165276, 0.8132702, 0.9127556, 0.6066358, 0.7294966, 0.543625, 0.9350724, 0.8158536, 0.0027385, 0.8574043, 0.0335856, 0.7296554, 0.1756556, 0.8631789, 0.5414612, 0.2997119, 0.4226872, 0.0283197, 0.1242833, 0.6706244, 0.6471895, 0.6153851, 0.3836776, 0.9972099, 0.9808353, 0.685542, 0.6504593, 0.6884467, 0.3889214]
+  },
+  "cases": [
+    {
+      "name": "rank5_serial_rows_channel_axis_f32",
+      "provenance": {
+        "notes": "Exercises the coalesced serial-row path for a realistic channel-only reduction with many independent spatial groups."
+      },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 0.5 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 32, 8, 8, 4], "tolerance": 0.00002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "vec4_reduced_suffix_w6_hw12_f32",
+      "provenance": {
+        "notes": "The innermost dimension is not vec4-aligned, but the contiguous reduced HxW suffix is; vectors may safely cross an H/W boundary."
+      },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.071, "cosStep": 0.113, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2, 6], "tolerance": 0.00002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "dispatch_cliff_rank3_over_16m_elements",
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8388608, 1, 2],
+          "data": { "kind": "cycle", "values": [0.1, 0.9, 0.3, 0.7] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8388608, 1, 2], "tolerance": 0.002 } }
+    },
+    {
+      "name": "subgroup_default_axes_r768",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8, 16, 24],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8, 16, 24] } }
+    },
+    {
+      "name": "subgroup_keep_inner_axes_0_1",
+      "attrs": { "axes": [0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 8, 6, 10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 8, 6, 10] } }
+    },
+    {
+      "name": "subgroup_rank3_axes_0_1_r300",
+      "attrs": { "axes": [0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [10, 30, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [10, 30, 3] } }
+    },
+    {
+      "name": "ort_gen_rank3_axes_0_2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/gen_mvn_test_data.py",
+        "test": "shape [2, 3, 2] reduced over axes [0, 2]",
+        "notes": "Small deterministic generator case with a non-contiguous reduction axis set."
+      },
+      "attrs": { "axes": [0, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.6369617, 0.2697867, 0.0409735, 0.0165276, 0.8132702, 0.9127556, 0.6066358, 0.7294966, 0.543625, 0.9350724, 0.8158536, 0.0027385]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "tolerance": 0.00001,
+          "data": {
+            "kind": "values",
+            "values": [0.438269436, -1.67241299, -0.899517715, -0.963612974, 0.48143062, 0.751848817, 0.263942957, 0.970200479, 0.418393701, 1.44473696, 0.488452703, -1.72173214]
+          }
+        }
+      }
+    },
+    {
+      "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": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "values", "values": [1e-20, -1e-20, 2e-20, -2e-20] }
+        }
+      },
+      "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": "Rank-4 default/spatial axes companion: the variance is positive but subnormal, so the normalized output should be finite."
+      },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [1e-20, -1e-20, 2e-20, -2e-20] }
+        }
+      },
+      "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": "Rank-5 companion: the reduced variance is positive but subnormal, so normalization should produce finite values."
+      },
+      "attrs": { "axes": [0, 2, 3, 4] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 2, 2],
+          "data": { "kind": "values", "values": [1e-20, -1e-20, 2e-20, -2e-20] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1, 2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "subgroup_large_offset_cancellation",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [4000.25, 4001.5, 3999.75, 4000.875, 4001.125, 3998.5, 4000.0, 4002.25, -2000.5, -2001.25, -1999.875, -2000.125, -2002.0, -1998.75, -2000.625, -2001.5]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 4] } }
+    },
+    {
+      "name": "default_axes_f32",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 1, 2] } },
+      "tolerance": 0.00001
+    },
+    {
+      "name": "spatial_axes_f16",
+      "attrs": { "axes": [2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 2, 2] } },
+      "tolerance": 0.002
+    },
+    {
+      "name": "channel_axis_only_f32",
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 4.0, 6.0, 8.0, 10.0, -1.0, -2.0, -3.0, -4.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "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": {
+          "dtype": "float32",
+          "shape": [1, 65, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 65, 2, 2], "tolerance": 0.00002, "relTolerance": 0.00002 } }
+    },
+    {
+      "name": "negative_spatial_axes_f32",
+      "attrs": { "axes": [-2, -1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 4.0, 8.0, 16.0, 32.0, -1.0, -2.0, -4.0, -8.0, -16.0, -32.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_spatial_axes_2x8x32x32",
+      "provenance": {
+        "notes": "Compact sibling for the spatial-axes MVN benchmark; preserves axes=[2,3] over many channel planes without benchmark-scale tensors."
+      },
+      "attrs": { "axes": [2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.017, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8, 32, 32], "tolerance": 0.00001, "relTolerance": 0.00001 } }
+    },
+    {
+      "name": "all_axes_f32",
+      "attrs": { "axes": [0, 1, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [-4.0, -2.0, -1.0, 0.0, 1.0, 2.0, 4.0, 8.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "all_axes_zero_variance_nan",
+      "attrs": { "axes": [0, 1, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [1, 2, 2, 2], "data": { "kind": "constant", "value": 7.0 } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "ort_all_axes_rank3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
+        "test": "MeanVarianceNormalizationTest.AllAxes",
+        "notes": "Covers ORT's rank-3 all-axes shape using the same deterministic value pattern."
+      },
+      "attrs": { "axes": [0, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [-5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, -5.0, -4.0, -3.0, -2.0, -1.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_default_axes_rank4_two_batches",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
+        "test": "MeanVarianceNormalizationTest.DefaultAxes"
+      },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [3.0, -3.0, -1.0, 1.0, 2.0, -1.0, -2.0, -2.0, -2.0, 4.0, 1.0, 4.0, 0.0, -2.0, -2.0, -4.0, 5.0, 7.0, 5.0, -5.0, -5.0, 3.0, 4.0, 4.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_all_axes_rank4_two_batches",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
+        "test": "MeanVarianceNormalizationTest.AllAxes",
+        "notes": "Uses ORT's rank-4 all-axes shape with deterministic nonconstant values."
+      },
+      "attrs": { "axes": [0, 1, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [-5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, -5.0, -4.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_axes_subset_rank5_0_2_4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
+        "test": "MeanVarianceNormalizationTest.AxesSubsets5D"
+      },
+      "attrs": { "axes": [0, 2, 4] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axes_subset_rank5_0_2_4_input_x" } }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [0.3508345, -0.7870349, -1.4605863, -1.5525494, 0.8972119, 1.2055154, 0.6673803, 1.1295706, -0.168333, 1.3134559, 0.6321192, -1.7208749, 1.0194501, -2.0990413, 0.3826789, -1.220487, 1.0518781, 0.0548801, -0.4872377, -0.0246164, -1.5353374, -1.2379477, 0.9080993, 0.8199395, 0.1033084, -0.7737996, 1.1569287, 1.1095439, 0.3688809, 0.2360785, 0.2634291, -0.6033379]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_axes_subset_rank5_1_2_3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
+        "test": "MeanVarianceNormalizationTest.AxesSubsets5D"
+      },
+      "attrs": { "axes": [1, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axes_subset_rank5_0_2_4_input_x" } }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [0.0260567, -0.3008327, -2.3950341, -0.9652744, 0.7422773, 1.3860379, -0.0971367, 0.905246, -0.3531062, 1.4445876, 0.7527716, -1.001451, 0.9215636, -0.9205217, 0.4026078, -0.5477917, 0.8924309, 0.1015335, -1.0632416, -0.4004898, -2.0051854, -1.6617567, 0.2241158, 0.5484163, 0.0323921, -0.5653723, 1.3576236, 1.9586404, 0.2758914, 0.5622366, 0.2859732, -0.543208]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_axes_subset_rank5_0_1_4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/mean_variance_normalization_test.cc",
+        "test": "MeanVarianceNormalizationTest.AxesSubsets5D"
+      },
+      "attrs": { "axes": [0, 1, 4] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axes_subset_rank5_0_2_4_input_x" } }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2, 2],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [0.1843672, -1.5822912, -1.0098907, -1.0706838, 0.8350127, 1.1118552, 0.1472972, 0.8161319, -0.2647213, 1.6187237, 0.9171133, -1.1049752, 0.9578265, -1.3346525, 0.8169968, -2.1988905, 1.2728089, -0.2751323, -0.3664493, -0.0606291, -1.3493062, -1.0822638, 0.4956413, 0.3680653, 0.0805517, -1.0343067, 1.3681179, 1.3273969, 0.4795772, 0.3819509, 0.5926631, -1.0379052]
+          }
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_mvn",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mvn" },
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 3, 3, 1],
+          "data": {
+            "kind": "values",
+            "values": [0.8439682722091675, 0.5665143728256226, 0.058367349207401276, 0.029163669794797897, 0.12964272499084473, 0.5060197114944458, 0.7953830361366272, 0.9411345720291138, 0.9546573162078857, 0.17730942368507385, 0.46192094683647156, 0.264804482460022, 0.6746842265129089, 0.01665256917476654, 0.6247307658195496, 0.9240844249725342, 0.9722340703010559, 0.1196569874882698, 0.41356155276298523, 0.9129372835159302, 0.5933007597923279, 0.8192993402481079, 0.7862604260444641, 0.11799798905849457, 0.692484438419342, 0.5411941409111023, 0.07513222843408585]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 3, 3, 1], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [0, 8, 16, 24], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 8, 16, 24], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_vec4_reduce_innermost_only",
+      "attrs": { "axes": [3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 16], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "vec4_dispatch_fold_over_65535_rows_axis3",
+      "attrs": { "axes": [3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [70000, 1, 1, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [70000, 1, 1, 8], "tolerance": 0.0002 } }
+    },
+    {
+      "name": "reduce_size1_axis_zero_variance_nan_rank4",
+      "attrs": { "axes": [2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 1, 2],
+          "data": { "kind": "values", "values": [1.0, -3.0, 2.5, 100.0, -7.0, 0.0, 42.0, -0.5] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 1, 2],
+          "tolerance": 0,
+          "allowNaN": true,
+          "data": { "kind": "values", "values": ["NaN", "NaN", "NaN", "NaN", "NaN", "NaN", "NaN", "NaN"] }
+        }
+      }
+    },
+    {
+      "name": "vec4_negative_innermost_axis_axis_minus1",
+      "attrs": { "axes": [-1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.07, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5, 8], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "kept_channel_dim_zero_empty_rank4",
+      "attrs": { "axes": [0, 2, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 0, 16, 24], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 0, 16, 24], "tolerance": 0 } }
+    },
+    {
+      "name": "all_axes_flat_splitk_65536",
+      "attrs": { "axes": [0, 1, 2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 256, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 256, 256], "tolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 1, 256, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 256, 256], "tolerance": 0.004, "relTolerance": 0.004 } }
+    },
+    {
+      "name": "rank6_channel_axis_generic_geometry",
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 2, 1, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 2.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 4, 1, 2, 1, 2], "tolerance": 0.00002, "relTolerance": 0.00002 }
+      }
+    },
+    {
+      "name": "rank8_channel_axis_generic_geometry",
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 2, 1, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.27 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 1, 2, 1, 2, 2, 2], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "rank5_serial_rows_channel_axis_f16",
+      "provenance": {
+        "notes": "float16 on the coalesced serial-row route, which needs at least SERIAL_MIN_ROWS independent rows and a reduction within SERIAL_MAX_REDUCTION. Only f32 cases and a bench had reached it."
+      },
+      "attrs": { "axes": [1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 8, 8, 4], "tolerance": 0.02 } }
+    }
+  ]
+}