sync 91d990483a17

Browse files

Files changed (9) hide show

README.md +16 -11
build/webgpu/batch-normalization-nc-vec4.wgsl.jinja +19 -16
build/webgpu/batch-normalization-nchw-flat-vec4.wgsl.jinja +63 -0
build/webgpu/batch-normalization-nchw-vec4.wgsl.jinja +15 -15
build/webgpu/batch-normalization-nchw.wgsl.jinja +15 -18
build/webgpu/bench.json +28 -1
build/webgpu/manifest.json +114 -200
build/webgpu/metadata.json +18 -9
build/webgpu/test.json +66 -4

README.md CHANGED Viewed

@@ -18,19 +18,19 @@ See the [ONNX `BatchNormalization` spec](https://onnx.ai/onnx/operators/onnx__Ba
 ## 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
@@ -50,20 +50,25 @@ Default values (overridable per request):
 ## 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.

 ## Inputs
+| Name | Upstream name | 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` | — | `T` | `1` | — | Per-channel scale tensor with shape `(C)`. | required |
+| `b` | `B` | `T` | `1` | — | Per-channel bias tensor with shape `(C)`. | required |
+| `inputMean` | `input_mean` | `T` | `1` | — | Precomputed estimated mean tensor with shape `(C)` used for inference. | required |
+| `inputVar` | `input_var` | `T` | `1` | — | Precomputed estimated variance tensor with shape `(C)` used for inference. | required |
 ## Outputs
+| Name | Upstream name | 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
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`batch-normalization-nc-vec4.wgsl.jinja`](build/webgpu/batch-normalization-nc-vec4.wgsl.jinja)
+- [`batch-normalization-nchw-flat-vec4.wgsl.jinja`](build/webgpu/batch-normalization-nchw-flat-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`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

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

@@ -1,30 +1,30 @@
 {% 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;
   }
@@ -36,9 +36,12 @@
 // 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);
 }

 {% 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 dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; 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 dispatch's
+  // per-axis workgroup fold width (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
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 // 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>) {
 {{ flat_index_2d("i", "params.count4") }}
   let channel4 = i % params.channels4;
   let alpha = scale[channel4] * inverseSqrt(input_var[channel4] + vec4<f32>(params.epsilon));
+  // Subtract the mean before scaling. The expanded form
+  // x * alpha + (bias - mean * alpha) can lose the small residual (x - mean)
+  // when |mean| is much larger than |x - mean|.
+  y[i] = (x[i] - input_mean[channel4]) * alpha + bias[channel4];
 }

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

	@@ -0,0 +1,63 @@

+{% 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 dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; 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 dispatch's
+  // per-axis workgroup fold width (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
+  // dispatch's per-axis workgroup fold width.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// Inference-only vec4 specialization for channel planes whose length is not a
+// multiple of four. The vec4 grid covers the flat tensor, so a group can cross a
+// channel boundary. Boundary groups resolve each component's channel separately;
+// all other groups share one channel. Each component computes:
+//   (x - mean) * inverseSqrt(var + epsilon) * scale + bias
+fn normalize_at(value: f32, channel: u32) -> f32 {
+  let normalized = (value - input_mean[channel]) * inverseSqrt(input_var[channel] + params.epsilon);
+  return normalized * scale[channel] + bias[channel];
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("i", "params.count4") }}
+  let base = i * 4u;
+  let first = (base / params.spatial) % params.channels;
+  let last = ((base + 3u) / params.spatial) % params.channels;
+  let value = x[i];
+  if (first == last) {
+    let normalized = (value - vec4<f32>(input_mean[first])) * inverseSqrt(input_var[first] + params.epsilon);
+    y[i] = normalized * vec4<f32>(scale[first]) + vec4<f32>(bias[first]);
+    return;
+  }
+  y[i] = vec4<f32>(
+    normalize_at(value.x, first),
+    normalize_at(value.y, ((base + 1u) / params.spatial) % params.channels),
+    normalize_at(value.z, ((base + 2u) / params.spatial) % params.channels),
+    normalize_at(value.w, last));
+}

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

@@ -1,30 +1,30 @@
 {% 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;
   }
@@ -38,7 +38,7 @@
 // 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);

 {% 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 dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; 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 dispatch's
+  // per-axis workgroup fold width (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
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 // calculation is:
 //   (x - mean) * inverseSqrt(var + epsilon) * scale + bias
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: 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);

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

@@ -1,42 +1,39 @@
 {% 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;

 {% 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 dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; 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 dispatch's
+  // per-axis workgroup fold width (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
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {%- endif %}
 {% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d("index") }}
   let spatial = params.height * params.width;
   let channel = (index / spatial) % params.channels;

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.BatchNormalization",
   "cases": [
     {
       "name": "nchw_4x64x112x112_vec4",
@@ -84,6 +83,34 @@
           }
         ]
       }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "nchw_4x64x112x112_vec4",
           }
         ]
       }
+    },
+    {
+      "name": "batchnorm-nchw-f32-c256-plane49-unaligned",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [128, 256, 7, 7] },
+        "scale": { "dtype": "float32", "shape": [256] },
+        "b": { "dtype": "float32", "shape": [256] },
+        "inputMean": { "dtype": "float32", "shape": [256] },
+        "inputVar": { "dtype": "float32", "shape": [256], "min": 0.5, "max": 2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [128, 256, 7, 7] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.x) * 4 * 2" }] }
+    },
+    {
+      "name": "batchnorm-nchw-f32-c256-plane64-aligned-control",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [98, 256, 8, 8] },
+        "scale": { "dtype": "float32", "shape": [256] },
+        "b": { "dtype": "float32", "shape": [256] },
+        "inputMean": { "dtype": "float32", "shape": [256] },
+        "inputVar": { "dtype": "float32", "shape": [256], "min": 0.5, "max": 2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [98, 256, 8, 8] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.x) * 4 * 2" }] }
     }
   ]
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -2,241 +2,155 @@
   "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" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "BatchNormalization",
   "sinceVersion": 15,
+  "inputs": {
+    "x": { "onnx": "X", "dtype": "T" },
+    "scale": { "dtype": "T", "rank": 1 },
+    "b": { "onnx": "B", "dtype": "T", "rank": 1 },
+    "inputMean": { "onnx": "input_mean", "dtype": "T", "rank": 1 },
+    "inputVar": { "onnx": "input_var", "dtype": "T", "rank": 1 }
   },
+  "outputs": { "y": { "onnx": "Y", "dtype": "T", "rank": "ranks.x", "shape": "shapes.x" } },
+  "attributes": { "epsilon": { "default": 0.00001 }, "momentum": { "default": 0.9 }, "training_mode": { "default": 0 } },
+  "attributeConstraints": { "momentum": { "values": [0.9] }, "training_mode": { "values": [0] } },
   "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
+    "normalizationParamsOk": "ranks.scale == 1 and ranks.b == 1 and ranks.inputMean == 1 and ranks.inputVar == 1 and dim(shapes.scale, 0) == dim(shapes.x, 1) and dim(shapes.b, 0) == dim(shapes.x, 1) and dim(shapes.inputMean, 0) == dim(shapes.x, 1) and dim(shapes.inputVar, 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"
   },
+  "when": ["inferenceContractOk"],
+  "bindings": {
+    "scale": { "buffer": "read-only-storage", "elementType": "$T" },
+    "bias": { "arg": "b", "buffer": "read-only-storage", "elementType": "$T" },
+    "input_mean": { "arg": "inputMean", "buffer": "read-only-storage", "elementType": "$T" },
+    "input_var": { "arg": "inputVar", "buffer": "read-only-storage", "elementType": "$T" },
+    "x_2": { "name": "x", "buffer": "read-only-storage", "elementType": "vec4<f32>" },
+    "scale_2": { "name": "scale", "buffer": "read-only-storage", "elementType": "vec4<f32>" },
+    "bias_2": { "arg": "b", "name": "bias", "buffer": "read-only-storage", "elementType": "vec4<f32>" },
+    "input_mean_2": {
+      "arg": "inputMean",
+      "name": "input_mean",
+      "buffer": "read-only-storage",
+      "elementType": "vec4<f32>"
+    },
+    "input_var_2": { "arg": "inputVar", "name": "input_var", "buffer": "read-only-storage", "elementType": "vec4<f32>" },
+    "y_2": { "name": "y", "buffer": "storage", "elementType": "vec4<f32>" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "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" }
+      ]
+    },
+    "params_3": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count4", "type": "u32", "value": "numel(shapes.y) / 4" },
+        { "name": "spatial", "type": "u32", "value": "inner(shapes.x, 1)" },
+        { "name": "channels", "type": "u32", "value": "dim(shapes.x, 1)" },
+        { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+      ]
+    },
+    "params_4": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "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" }
+      ]
+    }
   },
   "variants": [
     {
       "id": "inference_scalar",
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "BatchNormalization.InferenceScalar",
           "shader": "batch-normalization-nchw.wgsl.jinja",
+          "bindings": [
+            "x",
+            "scale",
+            "bias",
+            "input_mean",
+            "input_var",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "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)" }
+              ]
+            }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "nc_inference_vec4",
       "priority": 110,
+      "when": ["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": ["x_2", "scale_2", "bias_2", "input_mean_2", "input_var_2", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "nchw_inference_flat_vec4",
+      "priority": 90,
+      "when": ["dtypes.T == \"f32\"", "ranks.x >= 3", "inner(shapes.x, 1) % 4 != 0 and numel(shapes.x) % 4 == 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "BatchNormalization.InferenceFlatVec4",
+          "shader": "batch-normalization-nchw-flat-vec4.wgsl.jinja",
+          "bindings": ["x_2", "scale", "bias", "input_mean", "input_var", "y_2", "params_3"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "nchw_inference_vec4",
       "priority": 100,
+      "when": ["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": ["x_2", "scale", "bias", "input_mean", "input_var", "y_2", "params_4"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,29 @@
 {
   "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" }
 }

 {
   "name": "ai.onnx.BatchNormalization",
+  "id": "_ai_onnx_batchnormalization_webgpu_4d2bfed",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "batch-normalization-nc-vec4.wgsl.jinja": "yZ3Hf6SM+D6EW4GvvG5Q0N5tuxQlys5l78zAfqwH0IE=",
+      "batch-normalization-nchw-flat-vec4.wgsl.jinja": "ZU77NvidV7PpSQC2Fhc/pDgtSlG6EypswF7MWBi+qzw=",
+      "batch-normalization-nchw-vec4.wgsl.jinja": "L4z18z+y+0qESVx8FkmHJhTORNeEAOPz1KoV/e7WkCE=",
+      "batch-normalization-nchw.wgsl.jinja": "bv0TT2Ry2691M8bUjXYC1tlsrPZ+vPsGgdcUKSj59pg=",
+      "bench.json": "yPXcRY22XiVLwm5V3NsowDKaM6BKq2z0Batz3+SR+cI=",
+      "manifest.json": "PRmol3fOKZDuk/f21hi3uxRSbxsOhP7RQd+KA2zkW0E=",
+      "test.json": "AAN+Vhz8gU2l1tY8tLPNyRAYMpP+8dGalRvLSrzR/5M="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "inference_scalar": ["batch-normalization-nchw.wgsl.jinja"],
+      "nc_inference_vec4": ["batch-normalization-nc-vec4.wgsl.jinja"],
+      "nchw_inference_flat_vec4": ["batch-normalization-nchw-flat-vec4.wgsl.jinja"],
+      "nchw_inference_vec4": ["batch-normalization-nchw-vec4.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "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]
   },
@@ -83,7 +82,7 @@
       "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": {
@@ -190,7 +189,7 @@
       "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": {
@@ -450,7 +449,7 @@
       "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": {
@@ -615,6 +614,69 @@
         "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 } }
     }
   ]
 }

 {
   "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]
   },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
         "test": "BatchNormTest.PositiveTestCase",
+        "notes": "On the vec4 inference path, a subnormal scale must not collapse an otherwise ordinary normalized channel to zero."
       },
       "attrs": { "epsilon": 0.00001 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
         "test": "BatchNormTest.PositiveTestCase",
+        "notes": "Zero variance with a large epsilon must still produce a finite, correctly scaled output."
       },
       "attrs": { "epsilon": 0.25 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc",
         "test": "BatchNormTest.BatchNorm2d_fp16",
+        "notes": "A compact rank-4 float16 tensor exercises two-dimensional inference normalization."
       },
       "attrs": { "epsilon": 0.00001 },
       "inputs": {
         "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 } }
+    },
+    {
+      "name": "nc_inference_rank2_vec4_large_mean_small_residual",
+      "provenance": {
+        "notes": "Inputs sit within 0.1..0.4 of channel means of magnitude 125..1000 with variance 1e-6 (alpha = 1000, epsilon 0). The expanded form x*alpha + (bias - mean*alpha) rounds two ~1e6 products to a 1/16 ulp before cancelling and loses ~0.04 of a ~100 result; (x - mean)*alpha + bias keeps the residual exact."
+      },
+      "attrs": { "epsilon": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": {
+            "kind": "values",
+            "values": [1000.4, -1000.1999999999999, 500.9, -500.5, 250.4, -249.79999999999998, 126.30000000000001, -125.5, 1000.1999999999999, -1000.4, 500.5, -500.9, 249.79999999999998, -250.4, 125.5, -126.30000000000001]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [8], "data": { "kind": "constant", "value": 1.0 } },
+        "b": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.5, -0.5, 0.25, -0.25, 0.125, -0.125, 0.0625, -0.0625] }
+        },
+        "inputMean": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1000.3, -1000.3, 500.7, -500.7, 250.1, -250.1, 125.9, -125.9] }
+        },
+        "inputVar": { "dtype": "float32", "shape": [8], "data": { "kind": "constant", "value": 0.000001 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.002 } }
+    },
+    {
+      "name": "nchw_inference_plane9_flat_vec4",
+      "provenance": {
+        "source": "ONNX BatchNormalization-15 specification",
+        "notes": "A 3x3 channel plane is not a multiple of four, so four consecutive elements can belong to two channels and each channel's own scale, bias, mean and variance must still reach its own elements. The parameters differ per channel so a group taken from the wrong channel changes the result."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 1.145098, 1.126169, 0.970775, 0.726453, 0.451691, 0.205283, 0.035887, -0.026454, 0.024608, 0.170881, 0.373162, 0.578522, 0.730102, 0.77774, 0.687602, 0.449162, 0.078279, -0.384177, -0.878017, -1.332846, -1.679122, -1.859379, -1.837756, -1.606254, -1.186673, -0.627873, 0.001237, 0.621873, 1.15709, 1.543172, 1.738877, 1.731074, 1.535902, 1.195315, 0.76962]
+          }
+        },
+        "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": [0.25, -0.5] } },
+        "inputVar": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [4.0, 0.25] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.375, 0.447548, 0.438084, 0.360387, 0.238226, 0.100845, -0.022358, -0.107056, -0.138227, 1.524598, 1.670868, 1.873145, 2.0785, 2.230077, 2.277714, 2.187578, 1.949143, 1.578267, -0.317088, -0.564008, -0.791422, -0.96456, -1.054688, -1.043877, -0.928126, -0.718336, -0.438936, 1.501227, 2.121851, 2.657057, 3.043131, 3.238832, 3.231029, 3.035861, 2.695281, 2.269595]
+          },
+          "tolerance": 0.00001
+        }
+      }
     }
   ]
 }