sync 2e7068faf55e

Browse files

Files changed (9) hide show

README.md +77 -0
build/webgpu/bench.json +156 -0
build/webgpu/group-normalization-splitk-apply.wgsl.jinja +36 -0
build/webgpu/group-normalization-splitk-partials.wgsl.jinja +34 -0
build/webgpu/group-normalization-stash-f16-serial.wgsl.jinja +114 -0
build/webgpu/manifest.json +275 -0
build/webgpu/metadata.json +21 -0
build/webgpu/norm-row-stats.wgsl.jinja +152 -0
build/webgpu/test.json +711 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.GroupNormalization
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 21
+## Description
+Applies group normalization to the input: `y = scale * (x - mean) / sqrt(variance + epsilon) + bias`, where mean and variance are computed per instance per group of channels. The number of groups `num_groups` must divide the channel count `C` evenly; when `num_groups == C` this is equivalent to InstanceNormalization, and when `num_groups == 1` it is equivalent to LayerNormalization. The normalization stage supports TensorProto `stash_type` values `1` (float32) and `10` (float16).
+See the [ONNX `GroupNormalization` spec](https://onnx.ai/onnx/operators/onnx__GroupNormalization.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Input data tensor of shape `(N x C x D1 x ... x Dn)` where `N` is batch size and `C` is the number of channels. | required |
+| `scale` | `scale` | `T` | `1` | — | Scale tensor of shape `(C)`, one value per channel. | required |
+| `bias` | `bias` | `T` | `1` | — | Bias tensor of shape `(C)`, one value per channel. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | same as `X` | Normalized output tensor of the same shape as `X`. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `0.00001` | Small value added to the variance denominator to avoid division by zero. |
+| `stash_type` | `1` | TensorProto element type used for the normalization stage: `1` computes in float32, while `10` computes in float16. Normalized values are cast back to the input type before scale and bias are applied. |
+| `num_groups` | — | Required number of groups to divide the channels into; must be a divisor of `C`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`group-normalization-splitk-apply.wgsl.jinja`](build/webgpu/group-normalization-splitk-apply.wgsl.jinja)
+- [`group-normalization-splitk-partials.wgsl.jinja`](build/webgpu/group-normalization-splitk-partials.wgsl.jinja)
+- [`group-normalization-stash-f16-serial.wgsl.jinja`](build/webgpu/group-normalization-stash-f16-serial.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.GroupNormalization", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 2, 3] },
+  scale: { data: scaleData, shape: [2] },
+  bias: { data: biasData, shape: [2] },
+}, {
+  attrs: { num_groups: 1 },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,156 @@

+{
+  "op": "ai.onnx.GroupNormalization",
+  "cases": [
+    {
+      "name": "4x64x128x128_g32",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 4, "channels": 64, "spatial": 16384 },
+      "attrs": { "num_groups": 32, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4, 64, 128, 128], "dist": "normal", "seed": 740, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 741, "scale": 0.25, "offset": 1 },
+        "bias": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 742, "scale": 0.1 }
+      },
+      "outputs": { "y": { "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": "1x32x32x32_g8",
+      "attrs": { "num_groups": 8, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 32, 32] },
+        "scale": { "dtype": "float32", "shape": [32], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float32", "shape": [32], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 32, 32] } }
+    },
+    {
+      "name": "f16_vec4_4x64x128x128_g32",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "batch": 4, "channels": 64, "spatial": 16384 },
+      "attrs": { "num_groups": 32, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "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 },
+        "bias": { "dtype": "float16", "shape": [64], "dist": "normal", "seed": 752, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4, 64, 128, 128] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "vec4_spatial_aligned_2x128x56x56_g32",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 2, "channels": 128, "spatial": 3136 },
+      "attrs": { "num_groups": 32, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 128, 56, 56], "dist": "normal", "seed": 760, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [128], "dist": "uniform", "seed": 761, "scale": 0.25, "offset": 1 },
+        "bias": { "dtype": "float32", "shape": [128], "dist": "normal", "seed": 762, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 56, 56] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "scalar_spatial_unaligned_2x128x53x53_g32",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 2, "channels": 128, "spatial": 2809 },
+      "attrs": { "num_groups": 32, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 128, 53, 53], "dist": "normal", "seed": 770, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [128], "dist": "uniform", "seed": 771, "scale": 0.25, "offset": 1 },
+        "bias": { "dtype": "float32", "shape": [128], "dist": "normal", "seed": 772, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 53, 53] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "tiny_spatial_stem_1x32x7x7_g8",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "batch": 1, "channels": 32, "spatial": 49 },
+      "attrs": { "num_groups": 8, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 7, 7], "dist": "normal", "seed": 780, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [32], "dist": "uniform", "seed": 781, "scale": 0.25, "offset": 1 },
+        "bias": { "dtype": "float32", "shape": [32], "dist": "normal", "seed": 782, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 7, 7] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "scalar_lowocc_layernorm_equiv_2x320x65x65_g1",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "batch": 2, "channels": 320, "spatial": 4225 },
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 320, 65, 65], "dist": "normal", "seed": 901, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [320], "dist": "uniform", "seed": 902, "scale": 0.25, "offset": 1 },
+        "bias": { "dtype": "float32", "shape": [320], "dist": "normal", "seed": 903, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 320, 65, 65], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "scalar_cliff_unaligned_spatial_8x256x63x63_g32",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "batch": 8, "channels": 256, "spatial": 3969 },
+      "attrs": { "num_groups": 32, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [8, 256, 63, 63], "dist": "normal", "seed": 911, "scale": 0.5 },
+        "scale": { "dtype": "float32", "shape": [256], "dist": "uniform", "seed": 912, "scale": 0.25, "offset": 1 },
+        "bias": { "dtype": "float32", "shape": [256], "dist": "normal", "seed": 913, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 256, 63, 63], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.batch * args.channels * args.spatial * 2 + args.channels * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/group-normalization-splitk-apply.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,36 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const SPATIAL: u32 = {{ spatial }}u;
+const CPG: u32 = {{ channelsPerGroup }}u;
+const GROUPS: u32 = {{ numGroups }}u;
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+const EPSILON: f32 = {{ epsilon }};
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row = wg.x;
+  let part = wg.z;
+  if (row >= params.rows) { return; }
+  var pair = vec2<f32>(0.0);
+  for (var p = 0u; p < SPLIT; p = p + 1u) { pair = pair + partials[row * SPLIT + p]; }
+  let base = row * HIDDEN;
+  let shift = f32(x[base]);
+  let mean_d = pair.x / f32(HIDDEN);
+  let variance = max(pair.y / f32(HIDDEN) - mean_d * mean_d, 0.0);
+  let mean = shift + mean_d;
+  let inv_std = inverseSqrt(variance + EPSILON);
+  let group = row % GROUPS;
+  let chunk = (HIDDEN + SPLIT - 1u) / SPLIT;
+  let start = part * chunk;
+  let end = min(start + chunk, HIDDEN);
+  for (var d = start + lid.x; d < end; d = d + WG) {
+    let channel = group * CPG + d / SPATIAL;
+    let index = base + d;
+    y[index] = {{ scalar }}((f32(x[index]) - mean) * inv_std * f32(scale[channel]) + f32(bias[channel]));
+  }
+}

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

	@@ -0,0 +1,34 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+var<workgroup> reduction: array<vec2<f32>, WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row = wg.x;
+  let part = wg.z;
+  if (row >= params.rows) { return; }
+  let tid = lid.x;
+  let chunk = (HIDDEN + SPLIT - 1u) / SPLIT;
+  let start = part * chunk;
+  let end = min(start + chunk, HIDDEN);
+  let base = row * HIDDEN;
+  let shift = f32(x[base]);
+  var pair = vec2<f32>(0.0);
+  for (var d = start + tid; d < end; d = d + WG) {
+    let value = f32(x[base + d]) - shift;
+    pair = pair + vec2<f32>(value, value * value);
+  }
+  reduction[tid] = pair;
+  workgroupBarrier();
+  for (var stride = WG >> 1u; stride > 0u; stride = stride >> 1u) {
+    if (tid < stride) { reduction[tid] = reduction[tid] + reduction[tid + stride]; }
+    workgroupBarrier();
+  }
+  if (tid == 0u) { partials[row * SPLIT + part] = reduction[0]; }
+}

build/webgpu/group-normalization-stash-f16-serial.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,114 @@

+{% if usesF16 %}enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const SPATIAL: u32 = {{ spatial }}u;
+const CPG: u32 = {{ channelsPerGroup }}u;
+const GROUPS: u32 = {{ numGroups }}u;
+const EPSILON: f32 = {{ epsilon }};
+fn round_f16_bits_rte(value: f32) -> u32 {
+  let bits = bitcast<u32>(value);
+  let sign = (bits >> 16u) & 0x8000u;
+  let exponent_f32 = (bits >> 23u) & 0xffu;
+  let mantissa_f32 = bits & 0x7fffffu;
+  if (exponent_f32 == 0xffu) {
+    if (mantissa_f32 != 0u) {
+      return 0x7e00u;
+    }
+    return sign | 0x7c00u;
+  }
+  var exponent_f16 = i32(exponent_f32) - 127 + 15;
+  if (exponent_f16 >= 0x1f) {
+    return sign | 0x7c00u;
+  }
+  if (exponent_f16 <= 0) {
+    if (exponent_f16 < -10) {
+      return sign;
+    }
+    let significand = mantissa_f32 | 0x800000u;
+    let shift = u32(14 - exponent_f16);
+    let halfway = 1u << (shift - 1u);
+    let discarded = significand & ((1u << shift) - 1u);
+    var fraction = significand >> shift;
+    if (discarded > halfway || (discarded == halfway && (fraction & 1u) == 1u)) {
+      fraction = fraction + 1u;
+    }
+    return sign | fraction;
+  }
+  let halfway = 1u << 12u;
+  let discarded = mantissa_f32 & 0x1fffu;
+  var mantissa_f16 = mantissa_f32 >> 13u;
+  if (discarded > halfway || (discarded == halfway && (mantissa_f16 & 1u) == 1u)) {
+    mantissa_f16 = mantissa_f16 + 1u;
+    if (mantissa_f16 == 0x400u) {
+      mantissa_f16 = 0u;
+      exponent_f16 = exponent_f16 + 1;
+    }
+  }
+  if (exponent_f16 >= 0x1f) {
+    return sign | 0x7c00u;
+  }
+  return sign | (u32(exponent_f16) << 10u) | mantissa_f16;
+}
+fn widen_f16_bits(value: u32) -> f32 {
+  return unpack2x16float(value & 0xffffu).x;
+}
+// ONNX GroupNormalization-21 expresses the stash_type=FLOAT16 stage as a
+// graph of f16 tensor operators. One invocation owns a complete group so each
+// intermediate addition and arithmetic stage remains f16.
+@compute @workgroup_size(1, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>) {
+  let row = wg.x + wg.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let base = row * HIDDEN;
+  // Keep loop-carried stage values as binary16 bits. An f32 carrier could be
+  // reassociated through pack/unpack and retain wider accumulation precision.
+  var sum = round_f16_bits_rte(0.0);
+  var square_sum = round_f16_bits_rte(0.0);
+  for (var d = 0u; d < HIDDEN; d = d + 1u) {
+    let value = round_f16_bits_rte(f32(x[base + d]));
+    sum = round_f16_bits_rte(widen_f16_bits(sum) + widen_f16_bits(value));
+    let squared = round_f16_bits_rte(widen_f16_bits(value) * widen_f16_bits(value));
+    square_sum = round_f16_bits_rte(widen_f16_bits(square_sum) + widen_f16_bits(squared));
+  }
+  let mean = round_f16_bits_rte(widen_f16_bits(sum) / f32(HIDDEN));
+  let mean_square = round_f16_bits_rte(widen_f16_bits(square_sum) / f32(HIDDEN));
+  let mean_squared = round_f16_bits_rte(widen_f16_bits(mean) * widen_f16_bits(mean));
+  let variance = round_f16_bits_rte(widen_f16_bits(mean_square) - widen_f16_bits(mean_squared));
+  let epsilon_f16 = round_f16_bits_rte(EPSILON);
+  let variance_epsilon = round_f16_bits_rte(widen_f16_bits(variance) + widen_f16_bits(epsilon_f16));
+  let stddev = round_f16_bits_rte(sqrt(widen_f16_bits(variance_epsilon)));
+  let group = row % GROUPS;
+  for (var d = 0u; d < HIDDEN; d = d + 1u) {
+    let index = base + d;
+    let channel = group * CPG + d / SPATIAL;
+    let value = round_f16_bits_rte(f32(x[index]));
+    let centered = round_f16_bits_rte(widen_f16_bits(value) - widen_f16_bits(mean));
+    let normalized_f16 = round_f16_bits_rte(widen_f16_bits(centered) / widen_f16_bits(stddev));
+    {% if scalar == "f16" %}
+    let scaled = round_f16_bits_rte(widen_f16_bits(normalized_f16) * f32(scale[channel]));
+    let biased = round_f16_bits_rte(widen_f16_bits(scaled) + f32(bias[channel]));
+    y[index] = f16(widen_f16_bits(biased));
+    {% else %}
+    // The function casts Normalized back to T before its T-typed affine stage.
+    // For float32 X, scale and bias must therefore remain float32 operations.
+    let normalized = widen_f16_bits(normalized_f16);
+    let scaled = fma(normalized, scale[channel], 0.0);
+    y[index] = fma(1.0, scaled, bias[channel]);
+    {% endif %}
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,275 @@

+{
+  "domain": "ai.onnx",
+  "name": "GroupNormalization",
+  "sinceVersion": 21,
+  "description": "Applies group normalization to the input: `y = scale * (x - mean) / sqrt(variance + epsilon) + bias`, where mean and variance are computed per instance per group of channels. The number of groups `num_groups` must divide the channel count `C` evenly; when `num_groups == C` this is equivalent to InstanceNormalization, and when `num_groups == 1` it is equivalent to LayerNormalization. The normalization stage supports TensorProto `stash_type` values `1` (float32) and `10` (float16).",
+  "inputs": [
+    {
+      "role": "X",
+      "dtype": "T",
+      "description": "Input data tensor of shape `(N x C x D1 x ... x Dn)` where `N` is batch size and `C` is the number of channels."
+    },
+    { "role": "scale", "dtype": "T", "rank": 1, "description": "Scale tensor of shape `(C)`, one value per channel." },
+    { "role": "bias", "dtype": "T", "rank": 1, "description": "Bias tensor of shape `(C)`, one value per channel." }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.X",
+      "description": "Normalized output tensor of the same shape as `X`.",
+      "shape": "shapes.X"
+    }
+  ],
+  "attributes": { "epsilon": 0.00001, "stash_type": 1 },
+  "attributeDescriptions": {
+    "epsilon": "Small value added to the variance denominator to avoid division by zero.",
+    "num_groups": "Required number of groups to divide the channels into; must be a divisor of `C`.",
+    "stash_type": "TensorProto element type used for the normalization stage: `1` computes in float32, while `10` computes in float16. Normalized values are cast back to the input type before scale and bias are applied."
+  },
+  "attributeConstraints": { "num_groups": { "required": true }, "stash_type": { "values": [1, 10] } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "scale": { "kind": "tensor", "semantic": "scale", "role": "input" },
+    "bias": { "kind": "tensor", "semantic": "bias", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "MAX_STATS_SPLITS": 64,
+    "STATS_VALUES_PER_SPLIT": 4096,
+    "SPLIT_STATS_MIN_HIDDEN": 65536,
+    "SPLIT_STATS_MAX_ROWS": 256
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "groupAttributesOk": "attrs.num_groups >= 1",
+    "groupShapeOk": "groupAttributesOk and f16Ok(dtypes.T) and ranks.X >= 3 and ranks.scale == 1 and ranks.bias == 1 and ranks.Y == ranks.X and sameShape(shapes.Y, shapes.X) and dim(shapes.scale, 0) == dim(shapes.X, 1) and dim(shapes.bias, 0) == dim(shapes.X, 1) and dim(shapes.X, 1) % attrs.num_groups == 0",
+    "groupContractOk": "groupShapeOk and attrs.stash_type == onnxDtypeCode(\"float32\")",
+    "groupStashF16Ok": "groupShapeOk and attrs.stash_type == onnxDtypeCode(\"float16\")",
+    "groupRows": "dim(shapes.X, 0) * attrs.num_groups if groupAttributesOk else 0",
+    "groupSpatial": "inner(shapes.X, 1)",
+    "groupChannelsPerGroup": "dim(shapes.X, 1) / attrs.num_groups if groupAttributesOk else 0",
+    "groupHidden": "groupChannelsPerGroup * groupSpatial",
+    "normDeviceWorkgroupCap": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "normWorkgroupCap": "max(1, pow2ceil(normDeviceWorkgroupCap + 1) / 2)",
+    "groupScalarWorkgroup": "min(normWorkgroupCap, pow2ceil(groupHidden))",
+    "groupVec4Workgroup": "min(normWorkgroupCap, pow2ceil(groupHidden / 4))",
+    "hasSubgroupId": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "groupRowWorkgroupBytes": "normWorkgroupCap * 2 * 4",
+    "groupRowCovered": "groupContractOk and groupRowWorkgroupBytes <= device.limits.maxComputeWorkgroupStorageSize",
+    "groupSplitCount": "min(tunables.MAX_STATS_SPLITS, device.limits.maxComputeWorkgroupsPerDimension, pow2ceil(ceilDiv(groupHidden, tunables.STATS_VALUES_PER_SPLIT)))",
+    "groupPartialBytes": "groupRows * groupSplitCount * 2 * 4",
+    "groupSplitCovered": "groupRowCovered and groupRows <= tunables.SPLIT_STATS_MAX_ROWS and groupRows <= device.limits.maxComputeWorkgroupsPerDimension and groupHidden >= tunables.SPLIT_STATS_MIN_HIDDEN and groupPartialBytes <= device.limits.maxStorageBufferBindingSize and groupPartialBytes <= device.limits.maxBufferSize"
+  },
+  "bindingSets": {
+    "norm": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "bias",
+        "arg": "bias",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$ioElement" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "groupRows" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(groupRows, device.limits.maxComputeWorkgroupsPerDimension))"
+            }
+          ]
+        }
+      }
+    ],
+    "splitPartials": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "vec2<f32>" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "rows", "type": "u32", "value": "groupRows" }] }
+      }
+    ],
+    "splitApply": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "bias",
+        "arg": "bias",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "partials",
+        "semantic": "partials",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "vec2<f32>"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "rows", "type": "u32", "value": "groupRows" }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "group_stash_f16_serial",
+      "priority": 1000,
+      "when": "groupStashF16Ok",
+      "constants": {
+        "scalar": "dtypes.T",
+        "ioElement": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "groupHidden",
+        "spatial": "groupSpatial",
+        "channelsPerGroup": "groupChannelsPerGroup",
+        "numGroups": "attrs.num_groups",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNormalization.StashF16Serial",
+          "shader": "group-normalization-stash-f16-serial.wgsl.jinja",
+          "bindings": "norm",
+          "dispatch": { "workgroups": "groupRows" }
+        }
+      ]
+    },
+    {
+      "id": "group_splitk",
+      "priority": 120,
+      "when": ["groupSplitCovered"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "groupHidden",
+        "spatial": "groupSpatial",
+        "channelsPerGroup": "groupChannelsPerGroup",
+        "numGroups": "attrs.num_groups",
+        "workgroupSize": "normWorkgroupCap",
+        "split": "groupSplitCount",
+        "epsilon": "attrs.epsilon"
+      },
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[groupRows * groupSplitCount, 2]" }],
+      "passes": [
+        {
+          "id": "partials",
+          "name": "GroupNormalization.SplitKPartials",
+          "shader": "group-normalization-splitk-partials.wgsl.jinja",
+          "bindings": "splitPartials",
+          "dispatch": { "workgroups": "groupRows", "z": "groupSplitCount" }
+        },
+        {
+          "id": "apply",
+          "name": "GroupNormalization.SplitKApply",
+          "shader": "group-normalization-splitk-apply.wgsl.jinja",
+          "bindings": "splitApply",
+          "dispatch": { "workgroups": "groupRows", "z": "groupSplitCount" }
+        }
+      ]
+    },
+    {
+      "id": "group_subgroup_vec4",
+      "priority": 110,
+      "when": ["groupRowCovered", "groupSpatial % 4 == 0"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "ioElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNormalization.group_subgroup_vec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"group\"",
+              "vec4": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "groupHidden",
+              "wg": "groupVec4Workgroup",
+              "epsilon": "attrs.epsilon",
+              "numGroups": "attrs.num_groups",
+              "cpg": "groupChannelsPerGroup",
+              "hiddenVec": "groupHidden / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "spatialVec": "groupSpatial / 4",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "norm",
+          "dispatch": { "workgroups": "groupRows" }
+        }
+      ]
+    },
+    {
+      "id": "group_subgroup",
+      "priority": 100,
+      "when": ["groupRowCovered"],
+      "constants": { "scalar": "dtypes.T", "ioElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNormalization.group_subgroup",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"group\"",
+              "vec4": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "groupHidden",
+              "wg": "groupScalarWorkgroup",
+              "epsilon": "attrs.epsilon",
+              "numGroups": "attrs.num_groups",
+              "cpg": "groupChannelsPerGroup",
+              "spatial": "groupSpatial",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "norm",
+          "dispatch": { "workgroups": "groupRows" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,21 @@

+{
+  "name": "ai.onnx.GroupNormalization",
+  "id": "_ai_onnx_groupnormalization_webgpu_da50174",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "R60LEBSM13qdMqFSIdwpekWs9JTpBlgud7FNGLzPDdg=",
+      "group-normalization-splitk-apply.wgsl.jinja": "7Q93kvTqmCixZN4Z7GpxYdlxwaxiAUMjH7AMKHq8pfc=",
+      "group-normalization-splitk-partials.wgsl.jinja": "6FNPr81N83OW3DWYMjJln/B5zUsKTdinV610BRzzRvs=",
+      "group-normalization-stash-f16-serial.wgsl.jinja": "Wez9kqS+lzZASbm2BpWZSsHuNTT4rqsus0NiyhmGmvY=",
+      "manifest.json": "nWQpLNZpKQYFo8Pkd/yNT3xboHmv9HZJ5S28OKU86TE=",
+      "norm-row-stats.wgsl.jinja": "RCAuBcGKN1mg+npi5UTIqD4zDk0+B8DQj6u2TH+4IXU=",
+      "test.json": "KbF1vAUrTyrgdXUaq67pwqn41JjEAcHxnO2dImDHUiM="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.GroupNormalization" }
+}

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

	@@ -0,0 +1,152 @@

+{% 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; normalize as
+// (x - mean) / sqrt(variance + EPSILON).
+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 NUM_GROUPS: u32 = {{ source.numGroups }}u;
+const CPG: u32 = {{ source.cpg }}u;
+{% if source.vec4 %}
+const SPATIAL_V: u32 = {{ source.spatialVec }}u;
+{% else %}
+const SPATIAL: u32 = {{ source.spatial }}u;
+{% 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 %}
+  let shift = f32(x[base].x);
+{% 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) {
+    let v = vec4<f32>(x[base + i]);
+    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 denom = sqrt(variance + EPSILON);
+  let row_mean = shift + mean_d;
+  let g_ch_base = (row % NUM_GROUPS) * CPG;
+{% if source.vec4 %}
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+    let idx = base + i;
+    let v = vec4<f32>(x[idx]);
+    let ch = g_ch_base + i / SPATIAL_V;
+    let normed = (v - vec4<f32>(row_mean)) / vec4<f32>(denom);
+    y[idx] = {{ source.vecType }}(normed * vec4<f32>(f32(scale[ch])) + vec4<f32>(f32(bias[ch])));
+  }
+{% else %}
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let idx = base + i;
+    let v = f32(x[idx]);
+    let ch = g_ch_base + i / SPATIAL;
+    let normed = (v - row_mean) / denom;
+    y[idx] = {{ source.scalar }}(normed * f32(scale[ch]) + f32(bias[ch]));
+  }
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,711 @@

+{
+  "op": "ai.onnx.GroupNormalization",
+  "fixtureArrays": {
+    "batch2_groups2_f32_input_x": [1, 2, 3, 4, 5, 6, 7, 8, -1, -2, -3, -4, 8, 7, 6, 5],
+    "ort_rank3_equivalent_instance_norm_groups_equal_channels_input_x": [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],
+    "ort_rank3_group_size_n_input_x": [-1.1258, -1.1524, -0.2506, -0.4339, 0.8487, 0.692, -0.316, -2.1152, 0.3223, -1.2633, 0.35, 0.3081, 0.1198, 1.2377, 1.1168, -0.2473, -1.3527, -1.6959, 0.5667, 0.7935, 0.5988, -1.5551, -0.3414, 1.853, 0.7502, -0.5855, -0.1734, 0.1835, 1.3894, 1.5863, 0.9463, -0.8437, -0.6136, 0.0316, -0.4927, 0.2484, 0.4397, 0.1124, 0.6408, 0.4412, -0.1023, 0.7924, -0.2897, 0.0525, 0.5229, 2.3022, -1.4689, -1.5867],
+    "ort_rank3_group_size_n_output_y": [-0.759, -0.7848, 0.0914, -0.0867, 1.1595, 1.0073, 0.0278, -1.7203, 0.648, -0.8926, 0.6749, 0.6343, 0.0232, 0.9274, 0.8296, -0.2738, -1.1679, -1.4456, 0.3846, 0.5681, 0.4107, -1.3317, -0.3499, 1.4252, 0.5772, -0.8298, -0.3957, -0.0198, 1.2505, 1.458, 0.7838, -1.1017, -0.8594, -0.1798, -0.732, 0.0486, 0.2541, -0.0377, 0.4334, 0.2554, -0.2291, 0.5686, -0.3962, -0.0911, 0.3282, 1.9145, -1.4475, -1.5525],
+    "onnx_backend_group_normalization_example_input_x": [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]
+  },
+  "cases": [
+    {
+      "name": "subgroup_vec4_2x8x16x16_g2",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8, 16, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.11, "scale": 0.5 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19, "scale": 0.25 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8, 16, 16], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "subgroup_scalar_1x4x7x7_g2",
+      "attrs": { "num_groups": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 7, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.13, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 7, 7], "tolerance": 0.000002 } }
+    },
+    {
+      "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; each group's subnormal variance collapses to zero so normalization is non-finite."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.Equivalent_InstanceNorm_G_C",
+        "notes": "Valid epsilon=0 edge: each group has tiny normal centered values with positive subnormal variance, so normalization should be finite."
+      },
+      "attrs": { "num_groups": 2, "epsilon": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 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] } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 1, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_subnormal_scale_instance_groups_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-channel scale collapses to zero, losing the tiny affine output."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.Equivalent_InstanceNorm_G_C",
+        "notes": "Subnormal per-channel scale is valid and should preserve tiny affine outputs after otherwise ordinary normalization."
+      },
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 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] } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 1, 2], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "f32_subnormal_scale_grouped_spatial_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 per-channel scale collapses to zero (grouped spatial vec4 path)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.GroupSize_N",
+        "notes": "Larger grouped companion: subnormal scales should survive across channels and spatial positions."
+      },
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [-3.0, -1.0, 1.0, 3.0, -2.0, 0.0, 2.0, 4.0, 4.0, 2.0, 0.0, -2.0, 3.0, 1.0, -1.0, -3.0]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, -2e-40, 3e-40, -4e-40] }
+        },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 2], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "nchw_groups2_f32",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.5, 0.5, 2.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.25, -0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 1, 2] } },
+      "tolerance": 0.00001
+    },
+    {
+      "name": "instance_equivalent_f16",
+      "attrs": { "num_groups": 4, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "scale": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } },
+        "bias": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 4, 1, 2] } },
+      "tolerance": 0.002
+    },
+    {
+      "name": "layernorm_equivalent_one_group_f32",
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 1],
+          "data": { "kind": "values", "values": [-3.0, -1.0, 0.0, 2.0, 4.0, 6.0, 8.0, 10.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.5, -1.0, 2.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 1.0, -2.0, 3.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 1], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "zero_variance_bias_only",
+      "attrs": { "num_groups": 2, "epsilon": 0.001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 2],
+          "data": { "kind": "values", "values": [7.0, 7.0, 7.0, 7.0, -3.0, -3.0, -3.0, -3.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [10.0, -2.0, 3.0, 4.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -1.0, 2.0, -3.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "batch2_groups2_f32",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 1, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/batch2_groups2_f32_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.5, 0.5, 2.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.25, -0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 1, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "rank5_ncdhw_groups2_f32",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 1, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/batch2_groups2_f32_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.5, 0.5, 2.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.25, -0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 1, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_rank3_equivalent_instance_norm_groups_equal_channels",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.Equivalent_InstanceNorm_G_C"
+      },
+      "attrs": { "num_groups": 3, "epsilon": 0.3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_rank3_equivalent_instance_norm_groups_equal_channels_input_x" }
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "bias": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "tolerance": 0.0001,
+          "data": {
+            "kind": "values",
+            "values": [-0.56495477, 1.48930046, -1.13334329, 0.20899761, 1.46688162, -0.98600774, -0.79911913, 0.31824524, 0.57370438, 0.42193634, 0.6525492, -1.64818992, -0.92380346, -0.60808484, 0.04711878, 1.48476953, -0.14644464, -0.82262872, -0.66852817, 1.63760153, -1.65898662, 0.27618144, 0.64840618, 0.734399]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_rank3_equivalent_instance_norm_groups_equal_channels_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.Equivalent_InstanceNorm_G_C",
+        "notes": "MLFloat16 typed variant of the ORT GroupNormalization instance-normalization case."
+      },
+      "attrs": { "num_groups": 3, "epsilon": 0.3 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_rank3_equivalent_instance_norm_groups_equal_channels_input_x" }
+          }
+        },
+        "scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "bias": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 4], "tolerance": 0.003 } }
+    },
+    {
+      "name": "ort_rank3_equivalent_layer_norm_one_group",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.Equivalent_LayerNorm_G_1"
+      },
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 1.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [0.5361, 1.1216, 1.7072, 1.2928, 1.8783, 2.4638] }
+        }
+      }
+    },
+    {
+      "name": "ort_rank3_equivalent_layer_norm_one_group_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.Equivalent_LayerNorm_G_1",
+        "notes": "MLFloat16 typed variant of ORT's GroupNormalization-as-LayerNorm case."
+      },
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [2.0, 1.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 2, 3],
+          "tolerance": 0.001,
+          "data": { "kind": "values", "values": [0.5361, 1.1216, 1.7072, 1.2928, 1.8783, 2.4638] }
+        }
+      }
+    },
+    {
+      "name": "ort_rank3_group_size_n",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.GroupSize_N"
+      },
+      "attrs": { "num_groups": 2, "epsilon": 0.3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_rank3_group_size_n_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [6], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float32", "shape": [6], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_rank3_group_size_n_output_y" } }
+        }
+      }
+    },
+    {
+      "name": "ort_rank3_group_size_n_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/group_norm_op_test.cc",
+        "test": "GroupNormalizationOpTest.GroupSize_N",
+        "notes": "MLFloat16 typed variant of ORT's six-channel, two-group normalization case; ORT uses 0.005 output tolerance for the fp16 typed test."
+      },
+      "attrs": { "num_groups": 2, "epsilon": 0.3 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_rank3_group_size_n_input_x" } }
+        },
+        "scale": { "dtype": "float16", "shape": [6], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float16", "shape": [6], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [2, 6, 4],
+          "tolerance": 0.005,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_rank3_group_size_n_output_y" } }
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_group_normalization_example",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_group_normalization_example" },
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_group_normalization_example_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 2, 2], "tolerance": 0.00005 } }
+    },
+    {
+      "name": "onnx_backend_group_normalization_epsilon",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_group_normalization_epsilon" },
+      "attrs": { "num_groups": 2, "epsilon": 0.009999999776482582 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_group_normalization_example_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 2, 2], "tolerance": 0.00005 } }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 8, 16, 16], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [8], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float32", "shape": [8], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 8, 16, 16], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_dim_f16",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [0, 8, 16, 16], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float16", "shape": [8], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float16", "shape": [8], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [0, 8, 16, 16], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_channels",
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 0, 4, 4], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "bias": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 0, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_instance_norm_4d_vec4_1x16x8x8_g16",
+      "attrs": { "num_groups": 16, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 16, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.19, "scale": 1.5 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.27, "cosStep": 0.13, "scale": 0.5, "offset": 1.0 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 16, 8, 8], "tolerance": 0.02 } }
+    },
+    {
+      "name": "f16_groupnorm_4d_vec4_cpg4_1x32x14x14_g8",
+      "attrs": { "num_groups": 8, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 32, 14, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.23, "scale": 2.0 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "scale": 0.4, "offset": 1.0 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.29, "scale": 0.3 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 32, 14, 14], "tolerance": 0.03 } }
+    },
+    {
+      "name": "rank6_vec4_g2_distinct_affine",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21, "scale": 1.5 }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 3.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, -1.0, 2.0, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 2, 2, 2], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "vec4_layernorm_equiv_g1_distinct_per_channel_affine",
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 2.0 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 3.0, -1.5, 2.5, -0.5, 4.0] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 1.0, -0.5, 2.0, -1.0, 0.5, 3.0, -2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 4, 4], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "scalar_layernorm_equiv_g1_distinct_affine_odd_spatial",
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 5, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.29, "scale": 1.75 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 3.0, -1.5, 2.5, -0.5, 4.0] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 1.0, -0.5, 2.0, -1.0, 0.5, 3.0, -2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 5, 5], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "vec4_instance_equiv_gEqC_distinct_scale_per_channel",
+      "attrs": { "num_groups": 6, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 6, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.07, "scale": 1.25 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 3.0, -1.5, 2.5] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, -0.5, 2.0, -1.0, 0.5] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 6, 4, 4], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "group_splitk_layernorm_equiv_65536",
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        },
+        "scale": { "dtype": "float32", "shape": [16], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float32", "shape": [16], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 16, 64, 64], "tolerance": 0.0002 } }
+    },
+    {
+      "name": "rank7_spatial",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 2, 1, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 2.0 }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.5, 1.5, 2.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.25, -0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 1, 2, 1, 2, 2], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "onnx21_f16_stash_float16_rounding",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc GroupNormalization-21 function body",
+        "notes": "TensorProto FLOAT16 stash (10). Sequential f16 stage arithmetic produces outputs separated by multiple f16 ULPs from the default float32-stash route."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "num_groups": 1, "epsilon": 0.00001, "stash_type": 10 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 1, 2],
+          "data": { "kind": "values", "values": [-3.650390625, -2.51171875, -3.216796875, -2.568359375] }
+        },
+        "scale": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [1.6484375, 3.365234375] } },
+        "bias": {
+          "dtype": "float16",
+          "shape": [2],
+          "data": { "kind": "values", "values": [-0.26123046875, 1.74609375] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 2, 1, 2],
+          "tolerance": 0.002,
+          "relTolerance": 0,
+          "data": { "kind": "values", "values": [-2.59375, 1.41796875, 0.1025390625, 4.765625] }
+        }
+      }
+    },
+    {
+      "name": "onnx21_f32_stash_float16_rounding",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc GroupNormalization-21 function body",
+        "notes": "TensorProto FLOAT16 stash (10) with float32 T. Non-f16 input and affine values expose both the cast into float16 stage one and the required cast back to float32 before scale and bias."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "num_groups": 1, "epsilon": 0.00001, "stash_type": 10 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2],
+          "data": { "kind": "values", "values": [-3.6501, -2.5113, -3.2172, -2.5687] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.6483, 3.3657] } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-0.2612, 1.7463] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2],
+          "tolerance": 0.000001,
+          "relTolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [-2.593608856201172, 1.417683720588684, 0.10207010805606842, 4.7668843269348145]
+          }
+        }
+      }
+    },
+    {
+      "name": "rank8_spatial",
+      "attrs": { "num_groups": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 2, 1, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.27 }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.5, -1.0, 2.0] } },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.25, -0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 1, 2, 1, 2, 2, 2], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "group_splitk_layernorm_equiv_65536_f16",
+      "provenance": {
+        "notes": "float16 twin of the split-K group reduction. Only f32 cases reached the split-K partial and apply shaders, so neither had ever emitted its f16 declaration."
+      },
+      "attrs": { "num_groups": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 16, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        },
+        "scale": { "dtype": "float16", "shape": [16], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float16", "shape": [16], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 16, 64, 64], "tolerance": 0.02 } }
+    }
+  ]
+}