sync 2e7068faf55e

Browse files

Files changed (10) hide show

README.md +78 -0
build/webgpu/bench.json +257 -0
build/webgpu/instance-normalization-apply.wgsl.jinja +43 -0
build/webgpu/instance-normalization-batched-planes-vec4.wgsl.jinja +75 -0
build/webgpu/instance-normalization-splitk-combine.wgsl.jinja +27 -0
build/webgpu/instance-normalization-splitk-partials.wgsl.jinja +141 -0
build/webgpu/manifest.json +490 -0
build/webgpu/metadata.json +22 -0
build/webgpu/norm-row-stats.wgsl.jinja +170 -0
build/webgpu/test.json +705 -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.InstanceNormalization
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 6
+## Description
+Applies instance normalization to the input: `y = scale * (x - mean) / sqrt(variance + epsilon) + B`, where `mean` and `variance` are computed per instance per channel over the spatial dimensions. Equivalent to batch normalization with a batch size of one per channel.
+See the [ONNX `InstanceNormalization` spec](https://onnx.ai/onnx/operators/onnx__InstanceNormalization.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | — | — | Input tensor of shape `(N x C x D1 x ... x Dn)`; at least 3-D. | required |
+| `scale` | `scale` | `T` | `1` | — | 1-D scale tensor of size C, one scale factor per channel. | required |
+| `B` | `b` | `T` | `1` | — | 1-D bias tensor of size C, one bias value per channel. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | same as `input` | same as `input` | Normalized output tensor; same shape as the input. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `0.00001` | Small constant added to the variance before taking the square root to avoid division by zero. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Device requirements
+Some implementation variants require `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
+## 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
+- [`instance-normalization-apply.wgsl.jinja`](build/webgpu/instance-normalization-apply.wgsl.jinja)
+- [`instance-normalization-batched-planes-vec4.wgsl.jinja`](build/webgpu/instance-normalization-batched-planes-vec4.wgsl.jinja)
+- [`instance-normalization-splitk-combine.wgsl.jinja`](build/webgpu/instance-normalization-splitk-combine.wgsl.jinja)
+- [`instance-normalization-splitk-partials.wgsl.jinja`](build/webgpu/instance-normalization-splitk-partials.wgsl.jinja)
+- [`norm-row-stats.wgsl.jinja`](build/webgpu/norm-row-stats.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.InstanceNormalization", { version: 1 });
+const { output } = await kernel({
+  input: { data: inputData, shape: [1, 2, 1, 3] },
+  scale: { data: scaleData, shape: [2] },
+  b: { data: bData, shape: [2] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,257 @@

+{
+  "op": "ai.onnx.InstanceNormalization",
+  "cases": [
+    {
+      "name": "nchw_4x64x128x128",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 4, "channels": 64, "spatial": 16384 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4, 64, 128, 128], "dist": "normal", "seed": 750, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 751, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 752, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 64, 128, 128] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nchw_f16_4x64x128x128",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "batch": 4, "channels": 64, "spatial": 16384 },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [4, 64, 128, 128], "dist": "normal", "seed": 750, "scale": 0.5 },
+        "scale": { "dtype": "float16", "shape": [64], "dist": "uniform", "seed": 751, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float16", "shape": [64], "dist": "normal", "seed": 752, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 64, 128, 128] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nchw_1x64x56x56",
+      "preset": "smoke",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 64, 56, 56] },
+        "scale": { "dtype": "float32", "shape": [64] },
+        "b": { "dtype": "float32", "shape": [64] }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 64, 56, 56] } }
+    },
+    {
+      "name": "alignment_cliff_rank3_2x64x4095_unaligned",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 2, "channels": 64, "spatial": 4095 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 64, 4095], "dist": "normal", "seed": 810, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 811, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 812, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 64, 4095] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "alignment_healthy_rank3_2x64x4096_aligned",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 2, "channels": 64, "spatial": 4096 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 64, 4096], "dist": "normal", "seed": 810, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 811, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 812, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 64, 4096] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "dispatch_cliff_rows70000_yfold",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 70000, "spatial": 64 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 70000, 64], "dist": "normal", "seed": 820, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [70000], "dist": "uniform", "seed": 821, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [70000], "dist": "normal", "seed": 822, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 70000, 64] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "dispatch_healthy_rows60000_under_cap",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 60000, "spatial": 64 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 60000, 64], "dist": "normal", "seed": 820, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [60000], "dist": "uniform", "seed": 821, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [60000], "dist": "normal", "seed": 822, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 60000, 64] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "single_huge_instance_1x1x4194304_onewg",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 1, "spatial": 4194304 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 1, 4194304], "dist": "normal", "seed": 830, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [1], "dist": "uniform", "seed": 831, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [1], "dist": "normal", "seed": 832, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 4194304] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "onewg_per_plane_tiny_spatial_launchbound_1x65280x3x1",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 65280, "spatial": 3 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 65280, 3, 1], "dist": "normal", "seed": 901, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [65280], "dist": "uniform", "seed": 902, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [65280], "dist": "normal", "seed": 903, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 65280, 3, 1] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "onewg_per_plane_tiny_spatial_launchbound_rank3_1x49152x6",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 49152, "spatial": 6 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 49152, 6], "dist": "normal", "seed": 911, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [49152], "dist": "uniform", "seed": 912, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [49152], "dist": "normal", "seed": 913, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 49152, 6] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "rank7_spatial_flatten_gpu_gap",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 2, "spatial": 32 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 2, 2, 2, 2, 2, 2], "dist": "normal", "seed": 931, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [2], "dist": "uniform", "seed": 932, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [2], "dist": "normal", "seed": 933, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2, 2, 2, 2] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "splitk-priority-cliff-c256-256x256",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for variant coverage",
+        "notes": "Realistic 16.8M-element feature map that pins the selector boundary between plane_subgroup_vec4 and plane_splitk."
+      },
+      "vars": { "dtype": "float32", "batch": 1, "channels": 256, "spatial": 65536 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 256, 256, 256], "dist": "normal", "seed": 991, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [256], "dist": "uniform", "seed": 992, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [256], "dist": "normal", "seed": 993, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 256, 256, 256], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "splitk-priority-cliff-c32-512x512",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for variant coverage",
+        "notes": "Realistic 8.4M-element high-resolution feature map that pins the selector boundary between plane_subgroup_vec4 and plane_splitk."
+      },
+      "vars": { "dtype": "float32", "batch": 1, "channels": 32, "spatial": 262144 },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 32, 512, 512], "dist": "normal", "seed": 994, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [32], "dist": "uniform", "seed": 995, "scale": 0.25, "offset": 1 },
+        "b": { "dtype": "float32", "shape": [32], "dist": "normal", "seed": 996, "scale": 0.1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 32, 512, 512], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/instance-normalization-apply.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,43 @@

+// Apply per-plane statistics: y = (x - mean) * inverseStddev * scale + bias.
+// The vectorized route packs four adjacent spatial values per invocation; each
+// packed load and store remains within one plane.
+{% set vectorized = vectorized if vectorized is defined else false %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% set LOAD_OPEN = "vec4<f32>(" if usesF16 else "" %}
+{% set LOAD_CLOSE = ")" if usesF16 else "" %}
+{% set STORE_OPEN = "vec4<f16>(" if usesF16 else "" %}
+{% set STORE_CLOSE = ")" if usesF16 else "" %}
+{% set CHAN_OPEN = "f32(" if usesF16 else "" %}
+{% set CHAN_CLOSE = ")" if usesF16 else "" %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ applyWorkgroupSize }}u;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{% if not vectorized %}
+  // 2D-folded flat index: gid.y carries the high bits after dispatch folding.
+{% endif %}
+  let index = gid.x + gid.y * nwg.x * WG;
+  if (index >= params.count) {
+    return;
+  }
+{% if vectorized %}
+  // The vectorized path requires each plane to contain a multiple of four
+  // values, so a packed load/store never crosses an instance boundary.
+  let plane = (index * 4u) / params.spatial;
+{% else %}
+  let plane = index / params.spatial;
+{% endif %}
+  let channel = plane % params.channels;
+  let mean = stats[plane * 2u];
+  let inv_std = stats[plane * 2u + 1u];
+{% if vectorized %}
+  let value = {{ LOAD_OPEN }}input[index]{{ LOAD_CLOSE }};
+  output[index] = {{ STORE_OPEN }}(value - vec4<f32>(mean)) * vec4<f32>(inv_std * {{ CHAN_OPEN }}scale[channel]{{ CHAN_CLOSE }}) + vec4<f32>({{ CHAN_OPEN }}bias[channel]{{ CHAN_CLOSE }}){{ STORE_CLOSE }};
+{% else %}
+  output[index] = {{ scalar }}((f32(input[index]) - mean) * inv_std * f32(scale[channel]) + f32(bias[channel]));
+{% endif %}
+}

build/webgpu/instance-normalization-batched-planes-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,75 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// A power-of-two lane cohort reduces one plane while several cohorts share a
+// workgroup.
+const HIDDEN: u32 = {{ hidden }}u;
+const HIDDEN_V4: u32 = {{ hiddenVec }}u;
+const CHANNELS: u32 = {{ channels }}u;
+const EPSILON: f32 = {{ epsilon }};
+const WG: u32 = {{ workgroupSize }}u;
+const LANES: u32 = {{ lanesPerPlane }}u;
+const PLANES_PER_WG: u32 = {{ planesPerWorkgroup }}u;
+var<workgroup> shifted_moments: array<vec2<f32>, WG>;
+var<workgroup> plane_shift: array<f32, PLANES_PER_WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) workgroup: vec3<u32>,
+  @builtin(num_workgroups) workgroup_count: vec3<u32>,
+  @builtin(local_invocation_id) local: vec3<u32>
+) {
+  let tid = local.x;
+  let plane_in_workgroup = tid / LANES;
+  let lane = tid % LANES;
+  let group = workgroup.x + workgroup.y * workgroup_count.x;
+  let row = group * PLANES_PER_WG + plane_in_workgroup;
+  let is_active = row < params.rows;
+  let base = row * HIDDEN_V4;
+  if (is_active && lane == 0u) {
+    plane_shift[plane_in_workgroup] = f32(x[base].x);
+  }
+  workgroupBarrier();
+  let shift = plane_shift[plane_in_workgroup];
+  var moments = vec2<f32>(0.0);
+  if (is_active) {
+    for (var i = lane; i < HIDDEN_V4; i = i + LANES) {
+      let value = vec4<f32>(x[base + i]);
+      let delta = value - vec4<f32>(shift);
+      moments.x = moments.x + delta.x + delta.y + delta.z + delta.w;
+      moments.y = moments.y + dot(delta, delta);
+    }
+  }
+  shifted_moments[tid] = moments;
+  workgroupBarrier();
+  var stride = LANES / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (lane < stride) {
+      shifted_moments[tid] = shifted_moments[tid] + shifted_moments[tid + stride];
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  if (is_active) {
+    let total = shifted_moments[plane_in_workgroup * LANES];
+    let mean_delta = total.x / f32(HIDDEN);
+    let variance = max(total.y / f32(HIDDEN) - mean_delta * mean_delta, 0.0);
+    let mean = shift + mean_delta;
+    let inv_std = inverseSqrt(variance + EPSILON);
+    let channel = row % CHANNELS;
+    let affine_scale = inv_std * f32(scale[channel]);
+    let affine_bias = f32(bias[channel]);
+    for (var i = lane; i < HIDDEN_V4; i = i + LANES) {
+      let value = vec4<f32>(x[base + i]);
+      y[base + i] = {{ vectorScalar }}((value - vec4<f32>(mean)) * vec4<f32>(affine_scale) + vec4<f32>(affine_bias));
+    }
+  }
+}

build/webgpu/instance-normalization-splitk-combine.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,27 @@

+// Fold SPLIT per-plane (sum, sum-of-squares) partials into mean and inverse
+// standard deviation. One thread handles each plane. Variance uses
+// E[x^2] - E[x]^2; max(value, 0) guards against negative rounding residue.
+{{ env.wgsl.resourceDeclarations }}
+const SPLIT: u32 = {{ split }}u;
+const COMBINE_WG: u32 = {{ combineWorkgroupSize }}u;
+@compute @workgroup_size(COMBINE_WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  let plane = gid.x + gid.y * nwg.x * COMBINE_WG;
+  if (plane >= params.planes) {
+    return;
+  }
+  var total = 0.0;
+  var total_sq = 0.0;
+  let b = plane * SPLIT;
+  for (var k = 0u; k < SPLIT; k = k + 1u) {
+    total = total + partials[(b + k) * 2u];
+    total_sq = total_sq + partials[(b + k) * 2u + 1u];
+  }
+  let n = f32(params.spatial);
+  let mean = total / n;
+  let variance = max(total_sq / n - mean * mean, 0.0);
+  stats[plane * 2u] = mean;
+  stats[plane * 2u + 1u] = inverseSqrt(variance + params.epsilon);
+}

build/webgpu/instance-normalization-splitk-partials.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,141 @@

+{% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" %}
+{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
+{%- else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
+{%- endif %}
+{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
+  var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
+  loop {
+{% if form == "head" %}
+{% if breakInline %}
+    if ({{ svar }} == 0u) { break; }
+{% else %}
+    if ({{ svar }} == 0u) {
+      break;
+    }
+{% endif %}
+{% endif %}
+{% if bodyInline %}
+    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
+{% else %}
+    if ({{ idx }} < {{ svar }}) {
+{% for a in arrays %}
+      {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
+{% endfor %}
+    }
+{% endif %}
+{% if form == "head" %}
+{% if barrierFirst %}
+    workgroupBarrier();
+    {{ svar }} = {{ svar }} / 2u;
+{% else %}
+    {{ svar }} = {{ svar }} / 2u;
+    workgroupBarrier();
+{% endif %}
+{% else %}
+    workgroupBarrier();
+    if ({{ svar }} == 1u) {
+      break;
+    }
+    {{ svar }} = {{ svar }} / 2u;
+{% endif %}
+  }
+{%- endmacro %}
+/* Split-K partial sums for tensors with few planes and a large spatial extent.
+   A workgroup-per-plane kernel exposes too little parallelism, so this pass
+   splits each plane across SPLIT workgroups. Each accumulates a raw sum and
+   sum-of-squares over its slice. The combine pass produces mean and inverse
+   standard deviation, and the apply pass normalizes. */
+{% set vectorized = vectorized if vectorized is defined else false %}
+{% set useSubgroups = useSubgroups if useSubgroups is defined else false %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% if useSubgroups %}
+enable subgroups;
+{% endif %}
+{% set LOAD_OPEN = "vec4<f32>(" if usesF16 else "" %}
+{% set LOAD_CLOSE = ")" if usesF16 else "" %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+{% if useSubgroups %}
+// One slot per possible subgroup avoids assuming any mapping from local
+// invocation IDs to subgroup membership.
+var<workgroup> subgroup_partials: array<vec2<f32>, WG>;
+{% else %}
+var<workgroup> red_sum: array<f32, WG>;
+var<workgroup> red_sq: array<f32, WG>;
+{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>{% if useSubgroups %},
+        @builtin(subgroup_invocation_id) subgroup_lane: u32,
+        @builtin(subgroup_id) subgroup_id: u32,
+        @builtin(num_subgroups) num_subgroups: u32{% endif %}) {
+  let plane = wg.x + wg.y * nwg.x;
+  if (plane >= params.planes) {
+    return;
+  }
+  let k = wg.z;
+  let tid = lid.x;
+{% if vectorized %}
+  let spatial = params.spatial / 4u;
+{% else %}
+  let spatial = params.spatial;
+{% endif %}
+  let chunk = (spatial + SPLIT - 1u) / SPLIT;
+  let start = k * chunk;
+  var end = start + chunk;
+  if (end > spatial) { end = spatial; }
+  let base = plane * spatial;
+  var s = 0.0;
+  var sq = 0.0;
+  var i = start + tid;
+  loop {
+    if (i >= end) { break; }
+{% if vectorized %}
+    let v = {{ LOAD_OPEN }}input[base + i]{{ LOAD_CLOSE }};
+    s = s + v.x + v.y + v.z + v.w;
+    sq = sq + dot(v, v);
+{% else %}
+    let v = f32(input[base + i]);
+    s = s + v;
+    sq = sq + v * v;
+{% endif %}
+    i = i + WG;
+  }
+{% if useSubgroups %}
+  let subgroup_total = vec2<f32>(subgroupAdd(s), subgroupAdd(sq));
+  if (subgroup_lane == 0u) {
+    subgroup_partials[subgroup_id] = subgroup_total;
+  }
+  workgroupBarrier();
+  if (tid == 0u) {
+    var total = vec2<f32>(0.0);
+    for (var subgroup = 0u; subgroup < num_subgroups; subgroup = subgroup + 1u) {
+      total = total + subgroup_partials[subgroup];
+    }
+    let idx = (plane * SPLIT + k) * 2u;
+    partials[idx] = total.x;
+    partials[idx + 1u] = total.y;
+  }
+{% else %}
+  red_sum[tid] = s;
+  red_sq[tid] = sq;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["red_sum", "red_sq"], idx="tid", wg="WG", typed=true, form="head", breakInline=true) }}
+  if (tid == 0u) {
+    let idx = (plane * SPLIT + k) * 2u;
+    partials[idx] = red_sum[0];
+    partials[idx + 1u] = red_sq[0];
+  }
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,490 @@

+{
+  "domain": "ai.onnx",
+  "name": "InstanceNormalization",
+  "sinceVersion": 6,
+  "description": "Applies instance normalization to the input: `y = scale * (x - mean) / sqrt(variance + epsilon) + B`, where `mean` and `variance` are computed per instance per channel over the spatial dimensions. Equivalent to batch normalization with a batch size of one per channel.",
+  "inputs": [
+    { "role": "input", "dtype": "T", "description": "Input tensor of shape `(N x C x D1 x ... x Dn)`; at least 3-D." },
+    {
+      "role": "scale",
+      "dtype": "T",
+      "rank": 1,
+      "description": "1-D scale tensor of size C, one scale factor per channel."
+    },
+    { "role": "B", "dtype": "T", "rank": 1, "description": "1-D bias tensor of size C, one bias value per channel." }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.input",
+      "description": "Normalized output tensor; same shape as the input.",
+      "shape": "shapes.input"
+    }
+  ],
+  "attributes": { "epsilon": 0.00001 },
+  "attributeDescriptions": {
+    "epsilon": "Small constant added to the variance before taking the square root to avoid division by zero."
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "scale": { "kind": "tensor", "semantic": "scale", "role": "input" },
+    "b": { "kind": "tensor", "semantic": "B", "role": "input" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "MAX_STATS_SPLITS": 256,
+    "STATS_VALUES_PER_SPLIT": 2048,
+    "SPLIT_STATS_MIN_SPATIAL": 65536,
+    "SPLIT_STATS_MAX_PLANES": 256,
+    "COMBINE_WORKGROUP_SIZE": 64,
+    "BATCHED_MIN_PLANES_PER_WORKGROUP": 8
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
+    "reportedNonWave32Adapter": "not wave32Adapter and (has(device.adapterInfo, \"subgroupMinSize\") or has(device.adapterInfo, \"subgroupMaxSize\"))",
+    "instanceContractOk": "f16Ok(dtypes.T) and ranks.input >= 3 and ranks.output == ranks.input and sameShape(shapes.output, shapes.input) and ranks.scale == 1 and ranks.B == 1 and dim(shapes.scale, 0) == dim(shapes.input, 1) and dim(shapes.B, 0) == dim(shapes.input, 1)",
+    "instancePlanes": "dim(shapes.input, 0) * dim(shapes.input, 1)",
+    "instanceSpatial": "inner(shapes.input, 1)",
+    "normDeviceWorkgroupCap": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "normWorkgroupCap": "max(1, pow2ceil(normDeviceWorkgroupCap + 1) / 2)",
+    "normSubgroupMin": "device.adapterInfo.subgroupMinSize if has(device.adapterInfo, \"subgroupMinSize\") else 1",
+    "normSubgroupMax": "device.adapterInfo.subgroupMaxSize if has(device.adapterInfo, \"subgroupMaxSize\") else 32",
+    "hasSubgroupId": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "instanceScalarWorkgroup": "min(normWorkgroupCap, pow2ceil(instanceSpatial))",
+    "instanceVec4Workgroup": "min(normWorkgroupCap, pow2ceil(instanceSpatial / 4))",
+    "instanceVec4SubgroupEfficient": "not hasSubgroupId or instanceVec4Workgroup >= normSubgroupMin",
+    "instanceBatchedVec4Lanes": "instanceVec4Workgroup",
+    "instanceBatchedVec4PlanesPerWorkgroup": "max(1, floor(normWorkgroupCap / instanceBatchedVec4Lanes))",
+    "instanceBatchedVec4Workgroups": "ceilDiv(instancePlanes, instanceBatchedVec4PlanesPerWorkgroup)",
+    "instanceBatchedVec4StorageBytes": "normWorkgroupCap * 2 * 4 + instanceBatchedVec4PlanesPerWorkgroup * 4",
+    "instanceRowWorkgroupBytes": "normWorkgroupCap * 2 * 4",
+    "instanceStatsBytes": "instancePlanes * 2 * 4",
+    "instanceStatsFits": "instanceStatsBytes <= device.limits.maxStorageBufferBindingSize and instanceStatsBytes <= device.limits.maxBufferSize",
+    "instanceRowCovered": "instanceContractOk and instanceRowWorkgroupBytes <= device.limits.maxComputeWorkgroupStorageSize",
+    "instanceSplitCount": "min(tunables.MAX_STATS_SPLITS, device.limits.maxComputeWorkgroupsPerDimension, pow2ceil(ceilDiv(instanceSpatial, tunables.STATS_VALUES_PER_SPLIT)))",
+    "instancePartialBytes": "instancePlanes * instanceSplitCount * 2 * 4",
+    "splitStatsCovered": "instanceRowCovered and instanceStatsFits and instancePlanes <= tunables.SPLIT_STATS_MAX_PLANES and instancePlanes <= device.limits.maxComputeWorkgroupsPerDimension and instanceSpatial >= tunables.SPLIT_STATS_MIN_SPATIAL and instancePartialBytes <= device.limits.maxStorageBufferBindingSize and instancePartialBytes <= device.limits.maxBufferSize",
+    "splitStatsPreferred": "splitStatsCovered and instancePlanes < normSubgroupMax"
+  },
+  "bindingSets": {
+    "planeIo": [
+      {
+        "name": "x",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      {
+        "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": "y",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$ioElement"
+      }
+    ],
+    "plane": [
+      {
+        "name": "x",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      {
+        "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": "y",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$ioElement"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "dim(shapes.input, 0) * dim(shapes.input, 1)" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(dim(shapes.input, 0) * dim(shapes.input, 1), device.limits.maxComputeWorkgroupsPerDimension))"
+            }
+          ]
+        }
+      }
+    ],
+    "planeBatched": [
+      {
+        "name": "x",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      {
+        "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": "y",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$ioElement"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [{ "name": "rows", "type": "u32", "value": "dim(shapes.input, 0) * dim(shapes.input, 1)" }]
+        }
+      }
+    ],
+    "applyScalar": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$T"
+      },
+      { "name": "stats", "semantic": "stats", "buffer": { "type": "read-only-storage" }, "elementType": "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": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.input, 1)" },
+            { "name": "spatial", "type": "u32", "value": "instanceSpatial" }
+          ]
+        }
+      }
+    ],
+    "applyVec4": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "stats", "semantic": "stats", "buffer": { "type": "read-only-storage" }, "elementType": "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": "output",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.input, 1)" },
+            { "name": "spatial", "type": "u32", "value": "instanceSpatial" }
+          ]
+        }
+      }
+    ],
+    "splitPartials": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$splitInputElement"
+      },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "planes", "type": "u32", "value": "instancePlanes" },
+            { "name": "spatial", "type": "u32", "value": "instanceSpatial" }
+          ]
+        }
+      }
+    ],
+    "splitCombine": [
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      { "name": "stats", "semantic": "stats", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "planes", "type": "u32", "value": "instancePlanes" },
+            { "name": "spatial", "type": "u32", "value": "instanceSpatial" },
+            { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "plane_batched_vec4",
+      "priority": 115,
+      "when": ["instanceRowCovered", "instanceSpatial % 4 == 0", "instanceSpatial >= 4", "instancePlanes >= normWorkgroupCap", "instanceBatchedVec4PlanesPerWorkgroup >= tunables.BATCHED_MIN_PLANES_PER_WORKGROUP", "instanceBatchedVec4StorageBytes <= device.limits.maxComputeWorkgroupStorageSize"],
+      "demoteWhen": ["reportedNonWave32Adapter and instancePlanes <= device.limits.maxComputeWorkgroupsPerDimension"],
+      "constants": {
+        "usesF16": "dtypes.T == \"f16\"",
+        "ioElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "hidden": "instanceSpatial",
+        "hiddenVec": "instanceSpatial / 4",
+        "channels": "dim(shapes.input, 1) if dim(shapes.input, 1) > 0 else 1",
+        "epsilon": "attrs.epsilon",
+        "workgroupSize": "normWorkgroupCap",
+        "lanesPerPlane": "instanceBatchedVec4Lanes",
+        "planesPerWorkgroup": "instanceBatchedVec4PlanesPerWorkgroup"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "InstanceNormalization.PlaneBatchedVec4",
+          "shader": "instance-normalization-batched-planes-vec4.wgsl.jinja",
+          "bindings": "planeBatched",
+          "dispatch": { "workgroups": "instanceBatchedVec4Workgroups" }
+        }
+      ]
+    },
+    {
+      "id": "plane_subgroup_vec4",
+      "priority": 110,
+      "requires": { "features": [] },
+      "when": ["instanceRowCovered", "inner(shapes.input, 1) % 4 == 0", "instanceVec4SubgroupEfficient"],
+      "constants": { "ioElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "InstanceNormalization.plane_subgroup_vec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"instance\"",
+              "vec4": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "instanceSpatial",
+              "wg": "instanceVec4Workgroup",
+              "epsilon": "attrs.epsilon",
+              "channels": "dim(shapes.input, 1) if dim(shapes.input, 1) > 0 else 1",
+              "hiddenVec": "instanceSpatial / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "plane",
+          "dispatch": { "workgroups": "instancePlanes" }
+        }
+      ]
+    },
+    {
+      "id": "plane_subgroup_vec4_scalar_io",
+      "priority": 111,
+      "requires": { "features": ["subgroups"] },
+      "when": ["dtypes.T == \"f32\"", "wave32Adapter", "device.wgslLanguageFeatures.has(\"subgroup_id\")", "instanceRowCovered", "instanceSpatial % 4 == 0", "instanceVec4SubgroupEfficient"],
+      "constants": { "ioElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "InstanceNormalization.plane_subgroup_vec4_scalar_io",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"instance\"",
+              "vec4": true,
+              "scalarIo": true,
+              "scalar": "dtypes.T",
+              "usesF16": false,
+              "hidden": "instanceSpatial",
+              "wg": "instanceVec4Workgroup",
+              "epsilon": "attrs.epsilon",
+              "channels": "dim(shapes.input, 1) if dim(shapes.input, 1) > 0 else 1",
+              "hiddenVec": "instanceSpatial / 4",
+              "vecType": "\"vec4<f32>\"",
+              "combineSubgroups": true
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "plane",
+          "dispatch": { "workgroups": "instancePlanes" }
+        }
+      ]
+    },
+    {
+      "id": "plane_subgroup",
+      "priority": 100,
+      "requires": { "features": [] },
+      "when": ["instanceRowCovered"],
+      "constants": { "ioElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "InstanceNormalization.plane_subgroup",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"instance\"",
+              "vec4": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "instanceSpatial",
+              "wg": "instanceScalarWorkgroup",
+              "epsilon": "attrs.epsilon",
+              "channels": "dim(shapes.input, 1) if dim(shapes.input, 1) > 0 else 1",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "plane",
+          "dispatch": { "workgroups": "instancePlanes" }
+        }
+      ]
+    },
+    {
+      "id": "plane_splitk_vec4",
+      "priority": 121,
+      "when": ["splitStatsPreferred", "instanceSpatial % 4 == 0"],
+      "constants": {
+        "vectorized": true,
+        "usesF16": "dtypes.T == \"f16\"",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "useSubgroups": "hasSubgroupId",
+        "splitInputElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "workgroupSize": "normWorkgroupCap",
+        "split": "instanceSplitCount",
+        "combineWorkgroupSize": "min(tunables.COMBINE_WORKGROUP_SIZE, normWorkgroupCap)",
+        "applyWorkgroupSize": "normWorkgroupCap"
+      },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[instancePlanes * instanceSplitCount, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[instancePlanes, 2]" }
+      ],
+      "passes": [
+        {
+          "id": "partials",
+          "name": "InstanceNormalization.SplitKPartialsVec4",
+          "shader": "instance-normalization-splitk-partials.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "splitPartials",
+          "dispatch": { "workgroups": "instancePlanes", "z": "instanceSplitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "InstanceNormalization.SplitKCombine",
+          "shader": "instance-normalization-splitk-combine.wgsl.jinja",
+          "bindings": "splitCombine",
+          "dispatch": { "threads": "instancePlanes", "workgroupSize": "constants.combineWorkgroupSize" }
+        },
+        {
+          "id": "apply",
+          "name": "InstanceNormalization.ApplyVec4",
+          "shader": "instance-normalization-apply.wgsl.jinja",
+          "bindings": "applyVec4",
+          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "constants.applyWorkgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "plane_splitk",
+      "priority": 120,
+      "when": ["splitStatsPreferred"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "splitInputElement": "dtypes.T",
+        "workgroupSize": "normWorkgroupCap",
+        "split": "instanceSplitCount",
+        "combineWorkgroupSize": "min(tunables.COMBINE_WORKGROUP_SIZE, normWorkgroupCap)",
+        "applyWorkgroupSize": "normWorkgroupCap"
+      },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[instancePlanes * instanceSplitCount, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[instancePlanes, 2]" }
+      ],
+      "passes": [
+        {
+          "id": "partials",
+          "name": "InstanceNormalization.SplitKPartials",
+          "shader": "instance-normalization-splitk-partials.wgsl.jinja",
+          "bindings": "splitPartials",
+          "dispatch": { "workgroups": "instancePlanes", "z": "instanceSplitCount" }
+        },
+        {
+          "id": "combine",
+          "name": "InstanceNormalization.SplitKCombine",
+          "shader": "instance-normalization-splitk-combine.wgsl.jinja",
+          "bindings": "splitCombine",
+          "dispatch": { "threads": "instancePlanes", "workgroupSize": "constants.combineWorkgroupSize" }
+        },
+        {
+          "id": "apply",
+          "name": "InstanceNormalization.Apply",
+          "shader": "instance-normalization-apply.wgsl.jinja",
+          "bindings": "applyScalar",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "constants.applyWorkgroupSize" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,22 @@

+{
+  "name": "ai.onnx.InstanceNormalization",
+  "id": "_ai_onnx_instancenormalization_webgpu_1aea375",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "GgyPTi2dmoXyw7KRjHv05H93qU5QjSYSvl7Sq6uxIQ8=",
+      "instance-normalization-apply.wgsl.jinja": "xg8Pb+LQ5uMHQJNjPgxQAC9te2uiz1hfhLiKSxqD6WQ=",
+      "instance-normalization-batched-planes-vec4.wgsl.jinja": "fj7ftnvxAyb54O/TFqt5sezNAX0zVCyQNZfWZ34EiLE=",
+      "instance-normalization-splitk-combine.wgsl.jinja": "faaEHRq8I1cCglDAVTTpxsfaUtjDKjQOgDF6cobZWhY=",
+      "instance-normalization-splitk-partials.wgsl.jinja": "EPM/k6ud99fto6j+fbpwVBzJbg65sv3aiO4Cg6X9lLw=",
+      "manifest.json": "bM48AIoFxoG9TXkEwvxB57O6GrOiJMrH/MTEU39Fplk=",
+      "norm-row-stats.wgsl.jinja": "ZK1Wy+fDkHUODnGtjaF3FoacpqZidO/r2VxqrHtpBCg=",
+      "test.json": "TTDhrrElGsd6piLKBRdf7VIHBRzb1DwvWSgc5xB3iyk="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.InstanceNormalization" }
+}

build/webgpu/norm-row-stats.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,170 @@

+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{% set combineSubgroups = source.combineSubgroups %}
+{% set scalarIo = source.scalarIo if source.scalarIo is defined else false %}
+{% set reduceThreadParameters = ", sg_lane: u32, sg_id: u32, num_sg: u32"
+    if combineSubgroups else ", tid: u32" %}
+{% set reduceThreadArguments = ", sg_lane, sg_id, num_sg"
+    if combineSubgroups else ", tid" %}
+{% if combineSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Workgroup-parallel single-pass row statistics + fused normalize/affine.
+//
+// One workgroup owns one contiguous normalization span ("row": a last-axis
+// row, an instance plane, or a channel group). Threads stride the row once,
+// accumulating (sum, sum_sq) simultaneously. Partials are reduced either with
+// subgroupAdd plus a shared-memory combine or with a portable shared-memory
+// tree, then every thread applies the fused normalize + affine write.
+//
+// Shifted moments avoid cancellation from a large common offset; clamp the
+// variance to zero before adding EPSILON and taking inverseSqrt.
+const HIDDEN: u32 = {{ source.hidden }}u;
+{% if source.vec4 %}
+const HIDDEN_V: u32 = {{ source.hiddenVec }}u;
+{% endif %}
+const WG: u32 = {{ source.wg }}u;
+const EPSILON: f32 = {{ source.epsilon }};
+const CHANNELS: u32 = {{ source.channels }}u;
+{% if source.vec4 and scalarIo %}
+fn load_vec4(index: u32) -> vec4<f32> {
+  return vec4<f32>(x[index], x[index + 1u], x[index + 2u], x[index + 3u]);
+}
+{% endif %}
+{% if combineSubgroups %}
+var<workgroup> sg_partials: array<vec2<f32>, WG>;
+fn reduce_pair(value: vec2<f32>{{ reduceThreadParameters }}) -> vec2<f32> {
+  let s = vec2<f32>(subgroupAdd(value.x), subgroupAdd(value.y));
+  if (num_sg == 1u) {
+    return s;
+  }
+  if (sg_lane == 0u) {
+    sg_partials[sg_id] = s;
+  }
+  workgroupBarrier();
+  var total = vec2<f32>(0.0);
+  for (var i = 0u; i < num_sg; i++) {
+    total += sg_partials[i];
+  }
+  return total;
+}
+{% else %}
+// Each shared-memory tree reduction deliberately ends with a barrier. It keeps
+// lanes that have read the result from starting a later reduction and
+// overwriting scratch while slower lanes are still reading it.
+var<workgroup> tr0: array<f32, WG>;
+var<workgroup> tr1: array<f32, WG>;
+fn reduce_pair(value: vec2<f32>, tid: u32) -> vec2<f32> {
+  tr0[tid] = value.x;
+  tr1[tid] = value.y;
+  workgroupBarrier();
+  var stride: u32 = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) {
+      tr0[tid] = tr0[tid] + tr0[tid + stride];
+      tr1[tid] = tr1[tid] + tr1[tid + stride];
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  let reduced = vec2<f32>(tr0[0], tr1[0]);
+  workgroupBarrier();
+  return reduced;
+}
+{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg_id: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>{% if combineSubgroups %},
+  @builtin(subgroup_invocation_id) sg_lane: u32,
+  @builtin(subgroup_id) sg_id: u32,
+  @builtin(num_subgroups) num_sg: u32{% endif %}
+) {
+  let row = wg_id.x + wg_id.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+{% if source.vec4 and not scalarIo %}
+  let base = row * HIDDEN_V;
+{% else %}
+  let base = row * HIDDEN;
+{% endif %}
+{% if source.vec4 %}
+{% if scalarIo %}
+  let shift = f32(x[base]);
+{% else %}
+  let shift = f32(x[base].x);
+{% endif %}
+{% else %}
+  let shift = f32(x[base]);
+{% endif %}
+  var acc = vec2<f32>(0.0, 0.0);
+{% if source.vec4 %}
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+{% if scalarIo %}
+    let v = load_vec4(base + i * 4u);
+{% else %}
+    let v = vec4<f32>(x[base + i]);
+{% endif %}
+    let d = v - vec4<f32>(shift);
+    acc.x = acc.x + d.x + d.y + d.z + d.w;
+    acc.y = acc.y + dot(d, d);
+  }
+{% else %}
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let v = f32(x[base + i]);
+    let d = v - shift;
+    acc.x = acc.x + d;
+    acc.y = acc.y + d * d;
+  }
+{% endif %}
+  let totals = reduce_pair(acc{{ reduceThreadArguments }});
+  let mean_d = totals.x / f32(HIDDEN);
+  let variance = max(totals.y / f32(HIDDEN) - mean_d * mean_d, 0.0);
+  let inv = inverseSqrt(variance + EPSILON);
+  let row_mean = shift + mean_d;
+  let c = row % CHANNELS;
+  let ch_scale = f32(scale[c]);
+  let ch_bias = f32(bias[c]);
+{% if source.vec4 %}
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+{% if scalarIo %}
+    let idx = base + i * 4u;
+    let v = load_vec4(idx);
+{% else %}
+    let idx = base + i;
+    let v = vec4<f32>(x[idx]);
+{% endif %}
+{% if scalarIo %}
+    let value = (v - vec4<f32>(row_mean)) * inv * vec4<f32>(ch_scale) + vec4<f32>(ch_bias);
+    y[idx] = value.x;
+    y[idx + 1u] = value.y;
+    y[idx + 2u] = value.z;
+    y[idx + 3u] = value.w;
+{% else %}
+    y[idx] = {{ source.vecType }}((v - vec4<f32>(row_mean)) * inv * vec4<f32>(ch_scale) + vec4<f32>(ch_bias));
+{% endif %}
+  }
+{% else %}
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let idx = base + i;
+    let v = f32(x[idx]);
+    y[idx] = {{ source.scalar }}((v - row_mean) * inv * ch_scale + ch_bias);
+  }
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,705 @@

+{
+  "op": "ai.onnx.InstanceNormalization",
+  "fixtureArrays": {
+    "ort_rank3_batch2_repeated_epsilon_point3_input_input": [3.1513367, 9.283596, 1.4546119, 5.4617004, 8.519701, 1.2382338, 1.7930176, 5.1099434, 7.9195533, 7.638727, 8.065445, 3.8082376, 3.1513367, 9.283596, 1.4546119, 5.4617004, 8.519701, 1.2382338, 1.7930176, 5.1099434, 7.9195533, 7.638727, 8.065445, 3.8082376]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_ncl_16777216",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 16777217],
+          "data": { "kind": "cycle", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 16777217], "tolerance": 0.002 } }
+    },
+    {
+      "name": "subgroup_vec4_2x4x32x32",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 4, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19, "scale": 0.25 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 4, 32, 32], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "subgroup_scalar_1x2x13x5",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 13, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.13, "scale": 0.2 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 13, 5], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "f32_yfold_many_channels_1x70000x4",
+      "provenance": {
+        "notes": "Compact sibling for the y-fold InstanceNormalization benchmark; preserves C>65535 with much smaller spatial work."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 70000, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.019, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [70000],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.023, "scale": 0.25, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [70000],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.029, "scale": 0.1 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 70000, 4], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "f32_batched_planes_1x257x64",
+      "provenance": {
+        "notes": "Compact correctness lock for the feature-independent lane-cohort plane batching used by the 70,000-row dispatch-cliff benchmark."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 257, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [257],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.25, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [257],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.023, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 257, 64], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "nchw",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, -1.0, -2.0, -3.0, -4.0, -5.0, -6.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] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_tiny_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 per-instance variance collapses to zero so normalization is non-finite."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormNCHW",
+        "notes": "Valid epsilon=0 edge: normal inputs produce subnormal per-instance variance but finite order-one normalized outputs."
+      },
+      "attrs": { "epsilon": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [1e-20, -1e-20, 2e-20, -2e-20] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 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/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNorm",
+        "notes": "Subnormal scale is a valid affine parameter; the normalized output should preserve tiny values."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 1.0, -2.0, 2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, -2e-40] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2], "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/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormBatch1",
+        "notes": "Vec4 spatial companion for valid subnormal scale outputs."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [-3.0, -1.0, 1.0, 3.0, 4.0, 2.0, 0.0, -2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, -2e-40] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "ncdhw_rank5",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "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] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "large_mean_small_variance_one_pass_cancellation",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [8192.0, 8193.0, 8194.0, 8195.0, 8196.0, 8197.0, -8192.0, -8193.0, -8194.0, -8195.0, -8196.0, -8197.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] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "zero_variance_bias_only",
+      "attrs": { "epsilon": 0.001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "values", "values": [7.0, 7.0, 7.0, 7.0, -3.0, -3.0, -3.0, -3.0, 0.0, 0.0, 0.0, 0.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [10.0, -2.0, 4.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -1.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "large_epsilon_damps_output",
+      "attrs": { "epsilon": 0.5 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 4.0, 8.0, -1.0, -2.0, -4.0, -8.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.5, -0.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_webgpu_nchw_zero_variance_channel",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormNCHW_webgpu"
+      },
+      "attrs": { "epsilon": 0.009999999776482582 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 2.0, 2.0, 2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_webgpu_nchw_mixed_zero_variance",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormNCHW_webgpu_2"
+      },
+      "attrs": { "epsilon": 0.009999999776482582 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 2.0, 2.0, 2.0, 2.0, 2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_rank3_batch1_epsilon_point3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormBatch1"
+      },
+      "attrs": { "epsilon": 0.3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [3.1513367, 9.283596, 1.4546119, 5.4617004, 8.519701, 1.2382338, 1.7930176, 5.1099434, 7.9195533, 7.638727, 8.065445, 3.8082376]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_rank3_batch2_distinct_epsilon_point3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNorm"
+      },
+      "attrs": { "epsilon": 0.3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [3.1513367, 9.283596, 1.4546119, 5.4617004, 8.519701, 1.2382338, 1.7930176, 5.1099434, 7.9195533, 7.638727, 8.065445, 3.8082376, 2.3667817, 2.8248506, 3.7754705, 5.861325, 5.058735, 3.2787242, 3.6843839, 9.755121, 2.7902672, 7.3974323, 8.283609, 8.488337]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_rank3_batch2_repeated_epsilon_point3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormBatch2"
+      },
+      "attrs": { "epsilon": 0.3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_rank3_batch2_repeated_epsilon_point3_input_input" }
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_rank3_batch1_fp16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormBatch1_fp16",
+        "notes": "Valid fp16 InstanceNormalization coverage for a rank-3 single-batch input."
+      },
+      "attrs": { "epsilon": 0.3 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [3.1513367, 9.283596, 1.4546119, 5.4617004, 8.519701, 1.2382338, 1.7930176, 5.1099434, 7.9195533, 7.638727, 8.065445, 3.8082376]
+          }
+        },
+        "scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 3, 4], "tolerance": 0.005, "relTolerance": 0.005 } }
+    },
+    {
+      "name": "ort_rank3_batch2_fp16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNormBatch2_fp16",
+        "notes": "Valid fp16 InstanceNormalization coverage for a rank-3 two-batch input."
+      },
+      "attrs": { "epsilon": 0.3 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_rank3_batch2_repeated_epsilon_point3_input_input" }
+          }
+        },
+        "scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 3, 4], "tolerance": 0.005, "relTolerance": 0.005 } }
+    },
+    {
+      "name": "ort_rank5_instance_norm_2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/instance_norm_op_test.cc",
+        "test": "InstanceNormalizationOpTest.InstanceNorm_2"
+      },
+      "attrs": { "epsilon": 0.3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 1, 3],
+          "data": {
+            "kind": "values",
+            "values": [2.676342, 4.1100464, 4.570907, 5.8493505, 4.772751, 7.1669755, 2.8400702, 8.903057, 1.2464883, 7.034208, 4.755743, 6.0282083, 2.2634823, 2.7829134, 8.206701, 9.7143545, 3.8208177, 7.2309036, 8.887503, 9.05146, 1.7653979, 1.351493, 2.5284739, 8.903282, 1.8851215, 4.7899685, 9.621006, 5.7984877, 7.226894, 3.8396406, 7.1785083, 8.511631, 1.1645945, 7.751299, 9.89975, 7.733491]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [4.753198, 7.4829206, 1.0010294] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [3.720993, 2.320803, 1.8310473] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 2, 1, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_example_1x2x1x3",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_instancenorm_example" },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, 2.0, 3.0, 4.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_instancenorm_epsilon",
+      "attrs": { "epsilon": 0.009999999776482582 },
+      "inputs": {
+        "input": {
+          "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] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 4, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_instancenorm_epsilon" }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [0, 4, 32, 32], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 4, 32, 32], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_vec4_cnn_1x32x64x64",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 32, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.19, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.13, "scale": 0.4, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float16", "shape": [1, 32, 64, 64], "tolerance": 0.005, "relTolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f16_scalar_oddspatial_1x64x15x17",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 64, 15, 17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11, "scale": 0.4, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.23, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float16", "shape": [1, 64, 15, 17], "tolerance": 0.006, "relTolerance": 0.006 }
+      }
+    },
+    {
+      "name": "f16_vec4_rank3_2x32x1024_eps1e3",
+      "attrs": { "epsilon": 0.001 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [2, 32, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.11, "scale": 0.4, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19, "scale": 0.25 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 32, 1024], "tolerance": 0.005, "relTolerance": 0.005 } }
+    },
+    {
+      "name": "f32_vec4_negscale_largeeps_1x8x16x16",
+      "attrs": { "epsilon": 0.2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 8, 16, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21, "scale": 0.7 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.5, -0.5, 2.0, -1.0, 0.25, -2.5, 1.0, -0.75] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.5, -0.25, 2.0, -2.0, 0.75] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 16, 16], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_splitk_2plane_huge_1x2x131072",
+      "attrs": { "epsilon": 0.000001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 131072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.023, "scale": 0.5 }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.5, -0.75] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 131072], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "empty_zero_spatial_dim_1x4x3x0",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 4, 3, 0], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 3, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_channels",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [2, 0, 4, 4], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "b": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 0, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank6_vec4_general_spatial_flatten_1x2x2x2x2x2",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.5, -0.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2, 2, 2], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "rank7_spatial_flatten",
+      "provenance": {
+        "source": "onnx spec: InstanceNormalization accepts N x C x D1 x ... x Dn",
+        "notes": "Exercises instance normalization over a rank-7 tensor using the flattened spatial extent for each batch/channel row."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.5, -0.5] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2, 2, 2, 2], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "rank5_scalar_nosg_oddspatial_2x3x3x3x3",
+      "attrs": { "epsilon": 0.001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 3, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, -2.0, 0.5] } },
+        "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, -1.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 3, 3, 3], "tolerance": 0.000004 } }
+    },
+    {
+      "name": "rank8_spatial_flatten",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2, 2, 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] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2, 2, 2, 2, 2], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "split_stats_spatial_65536_f16",
+      "provenance": {
+        "notes": "float16 on the split-statistics rungs, which were gated to float32 until the partials pass learned to widen its packed load and the apply pass to narrow its packed store. A spatial extent of SPLIT_STATS_MIN_SPATIAL that is also a multiple of four leaves both the vec4 and the scalar rung eligible, so one case renders both."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 1, 65536],
+          "data": { "kind": "cycle", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "scale": { "dtype": "float16", "shape": [1], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float16", "shape": [1], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 1, 65536], "tolerance": 0.03 } }
+    },
+    {
+      "name": "batched_planes_vec4_p256_s64_f16",
+      "provenance": {
+        "notes": "float16 on the batched-planes path uses 2x128 planes of 64 elements, giving 16 lanes per plane and 16 planes per workgroup. This checks half-precision storage across the batched normalization and apply stages."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [2, 128, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 0.5 }
+        },
+        "scale": { "dtype": "float16", "shape": [128], "data": { "kind": "constant", "value": 1.0 } },
+        "b": { "dtype": "float16", "shape": [128], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 128, 64], "tolerance": 0.03 } }
+    }
+  ]
+}