sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +78 -0
build/webgpu/batch-normalization-nc-vec4.wgsl.jinja +44 -0
build/webgpu/batch-normalization-nchw-vec4.wgsl.jinja +46 -0
build/webgpu/batch-normalization-nchw.wgsl.jinja +52 -0
build/webgpu/bench.json +89 -0
build/webgpu/manifest.json +244 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +620 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.BatchNormalization
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 15
+## Description
+Applies inference-mode batch normalization: `Y = (X - input_mean) / sqrt(input_var + epsilon) * scale + B`. This package supports `training_mode=0`, rank-2-or-higher inputs, and a common float16 or float32 dtype for every tensor. ONNX training mode is intentionally not implemented because this inference-only release does not expose its required running-mean and running-variance outputs.
+See the [ONNX `BatchNormalization` spec](https://onnx.ai/onnx/operators/onnx__BatchNormalization.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Input data tensor with shape `(N, C, D1, ..., Dn)`, normalized independently per channel using the supplied estimated statistics. | required |
+| `scale` | `scale` | `T` | `1` | — | Per-channel scale tensor with shape `(C)`. | required |
+| `B` | `b` | `T` | `1` | — | Per-channel bias tensor with shape `(C)`. | required |
+| `input_mean` | `inputMean` | `T` | `1` | — | Precomputed estimated mean tensor with shape `(C)` used for inference. | required |
+| `input_var` | `inputVar` | `T` | `1` | — | Precomputed estimated variance tensor with shape `(C)` used for inference. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | same as `X` | Batch-normalized output tensor with the same shape as `X`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `0.00001` | Small value added to the variance before taking the square root to avoid division by zero. |
+| `momentum` | `0.9` | Standard ONNX running-statistics momentum. This inference-only package accepts the default `0.9`; non-default values are reserved for the unsupported training-state update. |
+| `training_mode` | `0` | Execution mode. This inference-only package supports the default value 0; value 1 is rejected because the ONNX training outputs are not exposed. |
+## 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
+- [`batch-normalization-nc-vec4.wgsl.jinja`](build/webgpu/batch-normalization-nc-vec4.wgsl.jinja)
+- [`batch-normalization-nchw-vec4.wgsl.jinja`](build/webgpu/batch-normalization-nchw-vec4.wgsl.jinja)
+- [`batch-normalization-nchw.wgsl.jinja`](build/webgpu/batch-normalization-nchw.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.BatchNormalization", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [2, 3] },
+  scale: { data: scaleData, shape: [3] },
+  b: { data: bData, shape: [3] },
+  inputMean: { data: inputMeanData, shape: [3] },
+  inputVar: { data: inputVarData, shape: [3] },
+});
+```

build/webgpu/batch-normalization-nc-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,44 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// Rank-2 [N, C] inference vec4 specialization. Vectors run across adjacent channels, so
+// scale/bias/mean/var are also bound as vec4<f32> and C must be divisible by 4.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("i", "params.count4") }}
+  let channel4 = i % params.channels4;
+  let alpha = scale[channel4] * inverseSqrt(input_var[channel4] + vec4<f32>(params.epsilon));
+  y[i] = x[i] * alpha + (bias[channel4] - input_mean[channel4] * alpha);
+}

build/webgpu/batch-normalization-nchw-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,46 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// Inference-only vec4 specialization: 128-bit loads/stores over x/y. Gated to
+// spatial % 4 == 0 so each vec4 stays inside one channel; the per-component
+// calculation is:
+//   (x - mean) * inverseSqrt(var + epsilon) * scale + bias
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("i", "params.count4") }}
+  let channel = (i / params.spatial4) % params.channels;
+  let normalized = (x[i] - vec4<f32>(input_mean[channel])) * inverseSqrt(input_var[channel] + params.epsilon);
+  y[i] = normalized * vec4<f32>(scale[channel]) + vec4<f32>(bias[channel]);
+}

build/webgpu/batch-normalization-nchw.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,52 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index") }}
+  let spatial = params.height * params.width;
+  let channel = (index / spatial) % params.channels;
+{% if usesF16 %}
+  // f16 tensors use an f32 epsilon, so widen the normalization arithmetic to
+  // f32 and narrow only on store. WGSL does not allow the mixed f16 + f32 add.
+  let normalized = (f32(x[index]) - f32(input_mean[channel])) * inverseSqrt(f32(input_var[channel]) + params.epsilon);
+  y[index] = {{ scalar }}(normalized * f32(scale[channel]) + f32(bias[channel]));
+{% else %}
+  let normalized = (x[index] - input_mean[channel]) * inverseSqrt(input_var[channel] + params.epsilon);
+  y[index] = normalized * scale[channel] + bias[channel];
+{% endif %}
+}

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,89 @@

+{
+  "op": "ai.onnx.BatchNormalization",
+  "cases": [
+    {
+      "name": "nchw_4x64x112x112_vec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 4, "channels": 64, "height": 112, "width": 112 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4, 64, 112, 112], "dist": "normal", "seed": 770, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 771, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 772, "scale": 0.1 },
+        "inputMean": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 773, "scale": 0.1 },
+        "inputVar": { "dtype": "float32", "shape": [64], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 64, 112, 112] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.batch * args.channels * args.height * args.width * dtypeBytes(args.dtype) * 2 + args.channels * dtypeBytes(args.dtype) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nchw_1x64x112x112",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 64, "height": 112, "width": 112 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 112, 112] },
+        "scale": { "dtype": "float32", "shape": [64] },
+        "b": { "dtype": "float32", "shape": [64] },
+        "inputMean": { "dtype": "float32", "shape": [64] },
+        "inputVar": { "dtype": "float32", "shape": [64], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 112, 112] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.batch * args.channels * args.height * args.width * dtypeBytes(args.dtype) * 2 + args.channels * dtypeBytes(args.dtype) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nc_4096x768",
+      "vars": { "dtype": "float32", "batch": 4096, "channels": 768, "height": 1, "width": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4096, 768] },
+        "scale": { "dtype": "float32", "shape": [768] },
+        "b": { "dtype": "float32", "shape": [768] },
+        "inputMean": { "dtype": "float32", "shape": [768] },
+        "inputVar": { "dtype": "float32", "shape": [768], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4096, 768] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.batch * args.channels * dtypeBytes(args.dtype) * 2 + args.channels * dtypeBytes(args.dtype) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nchw_inference_scalar_no_vec4_1x64x63x63",
+      "preset": "edge",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 64, "height": 63, "width": 63 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 63, 63], "dist": "normal", "seed": 6301, "scale": 1 },
+        "scale": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 6302, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 6303, "scale": 0.1 },
+        "inputMean": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 6304, "scale": 0.1 },
+        "inputVar": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 6305, "scale": 0.5, "offset": 1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 63, 63] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "args.batch * args.channels * args.height * args.width * dtypeBytes(args.dtype) * 2 + args.channels * dtypeBytes(args.dtype) * 4"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,244 @@

+{
+  "domain": "ai.onnx",
+  "name": "BatchNormalization",
+  "sinceVersion": 15,
+  "description": "Applies inference-mode batch normalization: `Y = (X - input_mean) / sqrt(input_var + epsilon) * scale + B`. This package supports `training_mode=0`, rank-2-or-higher inputs, and a common float16 or float32 dtype for every tensor. ONNX training mode is intentionally not implemented because this inference-only release does not expose its required running-mean and running-variance outputs.",
+  "inputs": [
+    {
+      "role": "X",
+      "dtype": "T",
+      "description": "Input data tensor with shape `(N, C, D1, ..., Dn)`, normalized independently per channel using the supplied estimated statistics."
+    },
+    { "role": "scale", "dtype": "T", "rank": 1, "description": "Per-channel scale tensor with shape `(C)`." },
+    { "role": "B", "dtype": "T", "rank": 1, "description": "Per-channel bias tensor with shape `(C)`." },
+    {
+      "role": "input_mean",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Precomputed estimated mean tensor with shape `(C)` used for inference."
+    },
+    {
+      "role": "input_var",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Precomputed estimated variance tensor with shape `(C)` used for inference."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.X",
+      "description": "Batch-normalized output tensor with the same shape as `X`.",
+      "shape": "shapes.X"
+    }
+  ],
+  "attributes": { "epsilon": 0.00001, "momentum": 0.9, "training_mode": 0 },
+  "attributeDescriptions": {
+    "epsilon": "Small value added to the variance before taking the square root to avoid division by zero.",
+    "momentum": "Standard ONNX running-statistics momentum. This inference-only package accepts the default `0.9`; non-default values are reserved for the unsupported training-state update.",
+    "training_mode": "Execution mode. This inference-only package supports the default value 0; value 1 is rejected because the ONNX training outputs are not exposed."
+  },
+  "attributeConstraints": {
+    "momentum": { "values": [0.9], "comparison": "float32" },
+    "training_mode": { "values": [0] }
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "scale": { "kind": "tensor", "semantic": "scale", "role": "input" },
+    "b": { "kind": "tensor", "semantic": "B", "role": "input" },
+    "inputMean": { "kind": "tensor", "semantic": "input_mean", "role": "input" },
+    "inputVar": { "kind": "tensor", "semantic": "input_var", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "normalizationParamsOk": "ranks.scale == 1 and ranks.B == 1 and ranks.input_mean == 1 and ranks.input_var == 1 and dim(shapes.scale, 0) == dim(shapes.X, 1) and dim(shapes.B, 0) == dim(shapes.X, 1) and dim(shapes.input_mean, 0) == dim(shapes.X, 1) and dim(shapes.input_var, 0) == dim(shapes.X, 1)",
+    "inferenceContractOk": "f16Ok(dtypes.T) and ranks.X >= 2 and ranks.Y == ranks.X and sameShape(shapes.Y, shapes.X) and normalizationParamsOk"
+  },
+  "bindingSets": {
+    "ncInferenceVec4": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec4<f32>"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec4<f32>"
+      },
+      {
+        "name": "bias",
+        "arg": "b",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec4<f32>"
+      },
+      {
+        "name": "input_mean",
+        "arg": "inputMean",
+        "semantic": "input_mean",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec4<f32>"
+      },
+      {
+        "name": "input_var",
+        "arg": "inputVar",
+        "semantic": "input_var",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec4<f32>"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "vec4<f32>" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count4", "type": "u32", "value": "numel(shapes.Y) / 4" },
+            { "name": "channels4", "type": "u32", "value": "dim(shapes.X, 1) / 4" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "spatialInferenceVec4": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec4<f32>"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      { "name": "bias", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      {
+        "name": "input_mean",
+        "arg": "inputMean",
+        "semantic": "input_mean",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      {
+        "name": "input_var",
+        "arg": "inputVar",
+        "semantic": "input_var",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "vec4<f32>" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count4", "type": "u32", "value": "numel(shapes.Y) / 4" },
+            { "name": "spatial4", "type": "u32", "value": "inner(shapes.X, 1) / 4" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.X, 1)" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ],
+    "inferenceScalar": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      { "name": "bias", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      {
+        "name": "input_mean",
+        "arg": "inputMean",
+        "semantic": "input_mean",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      {
+        "name": "input_var",
+        "arg": "inputVar",
+        "semantic": "input_var",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "channels", "type": "u32", "value": "dim(shapes.X, 1)" },
+            { "name": "height", "type": "u32", "value": "1 if ranks.X == 2 else dim(shapes.X, 2)" },
+            { "name": "width", "type": "u32", "value": "1 if ranks.X == 2 else inner(shapes.X, 2)" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "inference_scalar",
+      "when": ["inferenceContractOk"],
+      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "BatchNormalization.InferenceScalar",
+          "shader": "batch-normalization-nchw.wgsl.jinja",
+          "bindings": "inferenceScalar",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nc_inference_vec4",
+      "priority": 110,
+      "when": ["inferenceContractOk", "dtypes.T == \"f32\"", "ranks.X == 2", "dim(shapes.X, 1) % 4 == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "BatchNormalization.NcInferenceVec4",
+          "shader": "batch-normalization-nc-vec4.wgsl.jinja",
+          "bindings": "ncInferenceVec4",
+          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw_inference_vec4",
+      "priority": 100,
+      "when": ["inferenceContractOk", "dtypes.T == \"f32\"", "ranks.X >= 3", "inner(shapes.X, 1) % 4 == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "BatchNormalization.InferenceVec4",
+          "shader": "batch-normalization-nchw-vec4.wgsl.jinja",
+          "bindings": "spatialInferenceVec4",
+          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.BatchNormalization",
+  "id": "_ai_onnx_batchnormalization_webgpu_36b417d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "batch-normalization-nc-vec4.wgsl.jinja": "Lrc9lNARoZohL7Uv2G+JcdoTcEs6dqsaiP1+NPW1rsk=",
+      "batch-normalization-nchw-vec4.wgsl.jinja": "Rsd2siANpuQMlk5/ueUS07ZKscDp/jF8EQKsOyQAXN8=",
+      "batch-normalization-nchw.wgsl.jinja": "TfDSthFwA8AdlFykWHTOroAgZeC/TS7i8Wt7OKHyJbc=",
+      "bench.json": "GitYvHBWRsJ0jrz/w4vwR6ryXBDINuJ0gS+Fcrml9Gc=",
+      "manifest.json": "YZlooDSV8vRLHoVdW2ELx4Bhy0CGOB8w5jIIBn7j5ig=",
+      "test.json": "uHgc/mVjiAXJHTJLWPa9ZswwhfBjhvH0QmiRAzqFncE="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BatchNormalization" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,620 @@

+{
+  "op": "ai.onnx.BatchNormalization",
+  "fixtureArrays": {
+    "ort_positive_single_channel_7x7_exact_default_epsilon_input_x": [0.329876, -0.287158, -0.411425, 0.473621, 0.18156, -0.170596, -0.329516, -0.170733, -0.121664, 0.4372, -0.485668, 0.218049, -0.360263, 0.107016, 0.45358, 0.325056, 0.15995, 0.098852, -0.283453, -0.373051, 0.257542, 0.0614853, -0.0592363, 0.434488, -0.0179583, 0.398374, -0.451602, -0.132009, -0.174468, -0.0247169, 0.418897, -0.47159, -0.131925, 0.470943, 0.118357, 0.155664, 0.370062, -0.279229, 0.240311, -0.451034, 0.249178, -0.294496, 0.13683, -0.0806475, -0.309849, -0.450604, -0.28048, -0.420197, -0.433369]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_nc_inference_rank2",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [16776961, 1], "data": { "kind": "linspace", "start": -2.0, "end": 2.0 } },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 1.5 } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.25 } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.1 } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [16776961, 1], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "nchw_inference",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, -1.0, -2.0, -3.0, -4.0, 2.0, 4.0, 6.0, 8.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 0.5, 2.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 1.0, -1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.5, -2.5, 5.0] } },
+        "inputVar": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.25, 1.25, 5.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_subnormal_variance_epsilon_zero_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 variance collapses to zero so inverseSqrt yields Infinity instead of a finite value."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.SpatialNoBatch_2",
+        "notes": "Valid epsilon=0 edge: subnormal input variance and tiny normal centered values should produce finite normalized outputs, not infinities."
+      },
+      "attrs": { "epsilon": 0 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [1e-20, -1e-20] } },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_subnormal_scale_rank3_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 affine scale collapses to zero, losing the tiny output it should preserve."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.PositiveTestCase",
+        "notes": "Subnormal scale is a valid affine parameter; inference output should preserve the tiny normalized values."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [-1.0, 0.0, 2.0] } },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 3], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "f32_subnormal_scale_rank4_vec4_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 affine scale collapses to zero (vec4 path)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.PositiveTestCase",
+        "notes": "Vec4 inference companion: subnormal scale should not collapse an otherwise ordinary normalized channel to zero."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, 2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-2e-40] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 2], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "nc_inference_rank2",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 0.5, 2.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 1.0, -1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.5, 3.5, 4.5] } },
+        "inputVar": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.25, 2.25, 2.25] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "nc_inference_rank2_vec4",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 0.5, 2.0, -1.0, 0.25, 1.5, -0.75, 0.8] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.25, 0.5, -0.5, 0.75, -0.25] }
+        },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.2, -0.3, 0.4, -0.5, 0.1, -0.2, 0.3, -0.4] }
+        },
+        "inputVar": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 0.75, 1.25, 2.0, 0.5, 1.5, 0.9, 1.1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ncl_inference_rank3",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, -1.0, -2.0, -3.0, -4.0, 2.0, 4.0, 6.0, 8.0, 0.5, 1.5, 2.5, 3.5, -0.5, -1.5, -2.5, -3.5, 3.0, 5.0, 7.0, 9.0]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 0.5, 2.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 1.0, -1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, -2.5, 5.0] } },
+        "inputVar": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.25, 1.25, 5.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ncdhw_inference_rank5",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 0.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [4.5, -4.5] } },
+        "inputVar": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [5.25, 5.25] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_variance_large_epsilon",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.PositiveTestCase",
+        "notes": "Project regression adapted from ORT BatchNormalization positive cases to stress zero variance with a large epsilon."
+      },
+      "attrs": { "epsilon": 0.25 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 1.25, 0.75, 2.0, -3.0, -2.5, -4.0, -3.5] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -2.0] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, -1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -3.0] } },
+        "inputVar": { "dtype": "float32", "shape": [2], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_positive_single_channel_7x7_generated",
+      "attrs": { "epsilon": 0.00001 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.PositiveTestCase",
+        "notes": "Uses ORT's single-channel rank-4 shape and parameter values; X is generated deterministically."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19, "scale": 0.75 }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.589433] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.384622] } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-2.45673] } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.37998] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 7, 7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_pytorch_rank3_4x5x3_generated",
+      "attrs": { "epsilon": 0.00001 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.BatchNorm1d_3d_Pytorch",
+        "notes": "Uses ORT's rank-3 Pytorch-style shape and parameter values; X is generated deterministically."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 5, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.05, "scale": 1.25 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [0.36102, 0.592982, 0.808513, 0.0531484, 0.0960613] }
+        },
+        "b": { "dtype": "float32", "shape": [5], "data": { "kind": "constant", "value": 0.0 } },
+        "inputMean": { "dtype": "float32", "shape": [5], "data": { "kind": "constant", "value": 0.0 } },
+        "inputVar": { "dtype": "float32", "shape": [5], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 5, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_positive_single_channel_7x7_exact_default_epsilon",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.PositiveTestCaseDefaultEpsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 7, 7],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_single_channel_7x7_exact_default_epsilon_input_x" }
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.589433] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.384622] } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-2.45673] } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.37998] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 7, 7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_pytorch_rank3_4x5x3_exact",
+      "attrs": { "epsilon": 0.00001 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.BatchNorm1d_3d_Pytorch"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 5, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.948241, 1.23591, -0.39321, 1.4254, -0.730771, 0.439872, 0.0265089, 0.8748, -0.197505, 0.962646, 0.421469, 1.94512, 0.234179, -0.931897, -0.214905, -0.982965, -0.495436, 0.81949, -0.796605, -0.758605, 0.665557, 0.0909539, 1.10448, 1.91214, -1.97433, -2.26429, -0.384419, -0.226564, 0.230568, 0.533968, -1.31382, -0.156257, 0.532323, -0.16714, 0.971087, 0.600249, 0.858778, 0.423108, -0.414433, -1.17608, 0.673753, 0.278517, -2.19044, -0.161453, 1.17092, -0.155138, -0.094729, 0.19479, -1.17344, -0.213813, 0.118659, -2.39525, 0.257687, 0.784609, 0.297942, 1.10277, -1.58026, 0.197625, 0.0432784, 1.12924]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [0.36102, 0.592982, 0.808513, 0.0531484, 0.0960613] }
+        },
+        "b": { "dtype": "float32", "shape": [5], "data": { "kind": "constant", "value": 0.0 } },
+        "inputMean": { "dtype": "float32", "shape": [5], "data": { "kind": "constant", "value": 0.0 } },
+        "inputVar": { "dtype": "float32", "shape": [5], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 5, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_positive_single_channel_5d_exact",
+      "attrs": { "epsilon": 0.00001 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.PositiveTestCase_5D"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 7, 7, 1],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_single_channel_7x7_exact_default_epsilon_input_x" }
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.589433] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.384622] } },
+        "inputMean": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-2.45673] } },
+        "inputVar": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.37998] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 7, 7, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_pytorch_rank4_2x3x6x6_generated",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.BatchNorm2d_Pytorch",
+        "notes": "Uses the ORT PyTorch rank-4 shape and normalization parameters with deterministic fixture data to avoid committing the large literal tensor."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 6, 6],
+          "data": { "kind": "fillFloat32", "scale": 1.4, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [0.736494, 0.580251, 0.374834] }
+        },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } },
+        "inputMean": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } },
+        "inputVar": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 6, 6], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_pytorch_rank5_2x3x4x4x4_generated",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.BatchNorm3d_Pytorch",
+        "notes": "Uses the ORT PyTorch rank-5 shape and normalization parameters with deterministic fixture data to avoid committing the large literal tensor."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 4, 4],
+          "data": { "kind": "fillFloat32", "scale": 1.1, "sinStep": 0.013, "cosStep": 0.019 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [0.241661, 0.960798, 0.474727] }
+        },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } },
+        "inputMean": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } },
+        "inputVar": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 4, 4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_batchnorm_epsilon",
+      "attrs": { "epsilon": 0.009999999776482582 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.40746182203292847, 1.3439544439315796, -0.8182209730148315, 0.08270993828773499, -1.2910584211349487, -0.6611042022705078, -1.1801910400390625, 0.19764263927936554, 0.4138999879360199, 1.1973220109939575, 1.88335382938385, 0.7142238020896912, 2.2843334674835205, 1.5641025304794312, 0.6111037135124207, -0.8773633241653442, -1.6210875511169434, -0.5816730260848999, -0.5378339290618896, -1.5560237169265747, -0.05446484312415123, -1.8112788200378418, -0.6311752200126648, -0.9281591773033142, 1.4907219409942627, 0.19549933075904846, -0.4716043472290039, 1.8123546838760376, -2.294137477874756, 0.6512093544006348, -1.1304965019226074, -0.777346670627594, 1.115938425064087, 1.3394529819488525, -1.7674336433410645, 0.42441245913505554, 1.0893090963363647, -0.38418567180633545, 0.6322014331817627, -0.5496559143066406, 0.5211257338523865, 0.10834954679012299, 0.26166847348213196, -0.9147553443908691, 0.8582378029823303, 0.09433428943157196, -1.4859038591384888, -1.9005842208862305, -1.1375792026519775, -1.7620388269424438, -0.28862321376800537, 1.0479822158813477, 0.24995754659175873, 0.046904463320970535, -1.032243013381958, 0.4031856954097748, -0.6840592622756958, 1.262322187423706, -2.005556583404541, -0.33203038573265076, -0.2961004078388214, -2.218360662460327, -0.18350288271903992, 0.3923080563545227, 0.241634801030159, 0.10393591225147247, -0.8295711874961853, 0.4927593767642975, 0.09011279046535492, -0.9975675344467163, -0.80003821849823, 0.2070755809545517, 0.5234630107879639, -0.6993948221206665, 0.9137058258056641, -0.6727848052978516, 0.1333245038986206, 0.42689600586891174, -0.01284939143806696, -0.35224831104278564, 0.8194665908813477, 0.5219877362251282, 1.1972599029541016, -0.3824862241744995, 0.6916618943214417, 0.3538850247859955, 1.047585368156433, -0.4238962233066559, -3.514768123626709, -1.3431566953659058, 1.4255061149597168, 0.2285820096731186, -0.257663756608963, 0.0503707155585289, -1.3802108764648438, -0.26167207956314087, -0.1793796867132187, -0.6927706003189087, 1.137826919555664, -0.1691572517156601, -0.7639136910438538, -0.49807310104370117, -0.3628911077976227, 0.2639603018760681, -0.6296418905258179, -0.47225841879844666, -1.5133610963821411, 1.1076246500015259, 0.1762387454509735, -0.940353512763977, 0.9295943379402161, -1.0627949237823486, -0.8864062428474426, 1.92134690284729, -0.4597805142402649, -1.0890344381332397, 0.9841172695159912, -1.1592062711715698, -0.43653708696365356, 1.0092445611953735]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [0.7133895754814148, -0.7280577421188354, 0.8395164608955383] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [1.2390209436416626, -1.784803867340088, -0.7961858510971069] }
+        },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [-1.4005413055419922, -0.18435057997703552, -1.391193151473999] }
+        },
+        "inputVar": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [0.044612299650907516, 0.7997958660125732, 0.07695644348859787] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_batchnorm_epsilon" }
+    },
+    {
+      "name": "onnx_backend_batchnorm_example",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902, -0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834, -0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235, 1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314, -0.8707971572875977, -0.5788496732711792, -0.3115525245666504, 0.056165341287851334, -1.1651498079299927, 0.9008265137672424, 0.4656624495983124, -1.5362436771392822, 1.4882521629333496, 1.895889163017273, 1.1787796020507812, -0.1799248307943344, -1.0707526206970215, 1.0544517040252686, -0.4031769335269928, 1.222445011138916, 0.2082749754190445, 0.9766390323638916, 0.3563663959503174, 0.7065731883049011, 0.01050002034753561, 1.7858705520629883, 0.12691208720207214, 0.4019893705844879, 1.8831506967544556, -1.3477590084075928, -1.2704850435256958, 0.969396710395813, -1.1731233596801758, 1.9436211585998535, -0.4136189818382263, -0.747454822063446, 1.922942042350769, 1.4805147647857666, 1.8675589561462402, 0.9060446619987488, -0.8612256646156311, 1.910064935684204, -0.26800337433815, 0.8024563789367676, 0.9472519755363464, -0.15501008927822113, 0.6140793561935425, 0.922206699848175]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [0.37642553448677063, -1.0994007587432861, 0.29823818802833557] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [1.3263858556747437, -0.694567859172821, -0.14963454008102417] }
+        },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [-0.4351535439491272, 1.8492637872695923, 0.6722947359085083] }
+        },
+        "inputVar": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [0.9755215048789978, 0.855803370475769, 0.011714084073901176] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_batchnorm_example" }
+    },
+    {
+      "name": "ort_fp16_inference_rank4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
+        "test": "BatchNormTest.BatchNorm2d_fp16",
+        "notes": "Compact rank-4 fp16 inference projection of ORT's BatchNorm2d_fp16 coverage."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, -1.0, 0.0, 1.0, 2.0] }
+        },
+        "scale": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [1.5, -2.0] } },
+        "b": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, 1.0] } },
+        "inputMean": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [2.5, 0.5] } },
+        "inputVar": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [1.25, 1.25] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 2, 2], "tolerance": 0.002 } }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 3, 2, 2], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "constant", "value": 0.0 } },
+        "inputMean": { "dtype": "float32", "shape": [3], "data": { "kind": "constant", "value": 0.0 } },
+        "inputVar": { "dtype": "float32", "shape": [3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 3, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_dim_nc_vec4_rank2",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 8], "data": { "kind": "values", "values": [] } },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 0.5, 2.0, -1.0, 0.25, 1.5, -0.75, 0.8] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.25, 0.5, -0.5, 0.75, -0.25] }
+        },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.2, -0.3, 0.4, -0.5, 0.1, -0.2, 0.3, -0.4] }
+        },
+        "inputVar": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 0.75, 1.25, 2.0, 0.5, 1.5, 0.9, 1.1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "nc_vec4_rank2_multigroup_3x12",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.23, "scale": 1.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "values", "values": [1.0, 0.5, 2.0, -1.0, 0.25, 1.5, -0.75, 0.8, 1.2, -0.6, 0.9, -1.3] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.25, 0.5, -0.5, 0.75, -0.25, 0.1, -0.1, 0.3, -0.3] }
+        },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "values", "values": [0.2, -0.3, 0.4, -0.5, 0.1, -0.2, 0.3, -0.4, 0.05, -0.05, 0.15, -0.15] }
+        },
+        "inputVar": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "values", "values": [1.0, 0.75, 1.25, 2.0, 0.5, 1.5, 0.9, 1.1, 0.6, 1.4, 0.8, 1.2] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 12], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "nchw_inference_vec4_channel_boundary_1x4x1x4",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, -1.0, -2.0, -3.0, -4.0, 10.0, 20.0, 30.0, 40.0, 0.5, 1.5, 2.5, 3.5]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.5, 2.0, -1.5] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.25] } },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [2.5, -2.5, 25.0, 2.0] }
+        },
+        "inputVar": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1.25, 1.25, 125.0, 1.25] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 1, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_f16_inference_near_zero_variance_multichannel",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 3, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [8.0, 9.0, 10.0, 11.0, -4.0, -3.0, -2.0, -1.0, 100.0, 120.0, 140.0, 160.0]
+          }
+        },
+        "scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [2.0, 1.0, 0.5] } },
+        "b": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.5, -1.0, 2.0] } },
+        "inputMean": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [9.5, -2.5, 130.0] } },
+        "inputVar": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.01, 1.25, 400.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 2, 2], "tolerance": 0.01 } }
+    },
+    {
+      "name": "inference_rank6_spatial_flatten",
+      "attrs": { "epsilon": 0.00001, "training_mode": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 1, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, -4.0, -2.0, 0.0, 2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 0.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, -1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.5, -1.0] } },
+        "inputVar": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.25, 5.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 1, 2, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "inference_rank8_spatial_flatten",
+      "attrs": { "epsilon": 0.00001, "training_mode": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 1, 2, 1, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.27 }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.5, -0.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, -1.0] } },
+        "inputMean": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.1, -0.2] } },
+        "inputVar": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 1, 2, 1, 2, 2], "tolerance": 0.000001 } }
+    }
+  ]
+}