sync 6fdf6301e2bb

Browse files

Files changed (8) hide show

README.md +99 -0
build/webgpu/bench.json +504 -0
build/webgpu/group-norm-combine.wgsl.jinja +50 -0
build/webgpu/group-norm-fused.wgsl.jinja +190 -0
build/webgpu/group-norm-slab.wgsl.jinja +277 -0
build/webgpu/manifest.json +633 -0
build/webgpu/metadata.json +32 -0
build/webgpu/test.json +1664 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.SkipGroupNorm
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Adds a residual and an optional bias to the input, then group-normalizes the sum with an optional fused SiLU: `S = X + skip + bias`, `Y = gamma * (S - mean) / sqrt(variance + epsilon) + beta`. Statistics are taken per batch item over each group of `C / groups` channels and every spatial position, under either layout. `skip` is either exactly `X`'s shape or a per-channel `(N, C)` / `(N, 1, 1, C)` residual. The sum is rounded to the tensor type before the statistics see it, so `Y` does not depend on whether `S` was requested. `X` and `gamma`/`beta` carry independent element types.
+See the [ONNX Runtime `SkipGroupNorm` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.SkipGroupNorm) for the reference semantics.
+## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `X` | `T` | `4` | — | Input image tensor, `(N, H, W, C)` when `channels_last` is 1 and `(N, C, H, W)` otherwise. | required |
+| `gamma` | `M` | `1` | — | Per-channel affine scale of shape `(C)`. | required |
+| `beta` | `M` | `1` | — | Per-channel affine offset of shape `(C)`. | required |
+| `skip` | `T` | — | — | Residual added to `X`. Either exactly `X`'s shape, or `(N, C)` / `(N, 1, 1, C)` to broadcast one value per batch item and channel. A `(N, 1, 1, C)` tensor at `H = W = 1` is an exact match and takes the non-broadcast route, which is what upstream does. | required |
+| `bias` | `T` | `1` | — | Optional per-channel bias of shape `(C)` added to `X + skip` before normalization. | optional |
+## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `Y` | `T` | `4` | same as `X` | Normalized tensor with the same shape and element type as `X`. | required |
+| `S` | `T` | `4` | same as `X` | Optional second output holding `X + skip + bias` at the tensor's element type, the values the statistics are taken over. | optional |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `activation` | — | Activation applied after the affine transform: `0` for none, `1` for SiLU. The request must supply it; the operator declares no default. |
+| `channels_last` | `1` | `1` when `X` and `Y` are `(N, H, W, C)`, `0` when they are `(N, C, H, W)`. Defaults to `1`. Any non-zero value selects the channels-last layout. |
+| `epsilon` | `0.00001` | Value added to the group variance before the reciprocal square root. Defaults to `1e-5`. |
+| `groups` | — | Number of channel groups the statistics are taken over. It must divide `C`. The request must supply it; the operator declares no default. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+| `M` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `split_bias_sum` — Three dispatches: partial group statistics, combine mean and reciprocal standard deviation, then normalize and apply affine weights. Channels-last statistics use four, two, or one aligned values per lane; group blocks and spatial rows follow device limits. Fuses skip, optional bias, and optional residual output.
+- `fused_bias_sum` — One workgroup per batch item and group for a request with a bias and the residual sum output, walking the group twice inside a single dispatch: once to form the residual sum and its moments, once to write the affine result.
+- `split_bias_plain` — Three dispatches: partial group statistics, combine mean and reciprocal standard deviation, then normalize and apply affine weights. Channels-last statistics use four, two, or one aligned values per lane; group blocks and spatial rows follow device limits. Fuses skip, optional bias, and optional residual output.
+- `fused_bias_plain` — One workgroup per batch item and group for a request with a bias and no residual sum output, walking the group twice inside a single dispatch: once to form the residual sum and its moments, once to write the affine result.
+- `split_sum` — Three dispatches: partial group statistics, combine mean and reciprocal standard deviation, then normalize and apply affine weights. Channels-last statistics use four, two, or one aligned values per lane; group blocks and spatial rows follow device limits. Fuses skip, optional bias, and optional residual output.
+- `fused_sum` — One workgroup per batch item and group for a request with no bias and the residual sum output, walking the group twice inside a single dispatch: once to form the residual sum and its moments, once to write the affine result.
+- `split_plain` — Three dispatches: partial group statistics, combine mean and reciprocal standard deviation, then normalize and apply affine weights. Channels-last statistics use four, two, or one aligned values per lane; group blocks and spatial rows follow device limits. Fuses skip, optional bias, and optional residual output.
+- `fused_plain` — One workgroup per batch item and group for a request with neither a bias nor the residual sum output, walking the group twice inside a single dispatch: once to form the residual sum and its moments, once to write the affine result.
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
+- [`group-norm-combine.wgsl.jinja`](build/webgpu/group-norm-combine.wgsl.jinja)
+- [`group-norm-fused.wgsl.jinja`](build/webgpu/group-norm-fused.wgsl.jinja)
+- [`group-norm-slab.wgsl.jinja`](build/webgpu/group-norm-slab.wgsl.jinja)
+## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.SkipGroupNorm", { version: 1 });
+const { Y } = await kernel({
+  X: { data: XData, shape: [2, 3, 2, 8] },
+  gamma: { data: gammaData, shape: [8] },
+  beta: { data: betaData, shape: [8] },
+  skip: { data: skipData, shape: [2, 8] },
+}, {
+  attrs: { groups: 2, activation: 0 },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,504 @@

+{
+  "tunableSpace": { "WORKGROUP_SIZE": [128, 256] },
+  "cases": [
+    {
+      "name": "flux-vae-mid-1x64x64x512-g32-f16-sum",
+      "preset": "smoke",
+      "provenance": { "note": "Flux.2 VAE decoder resnet: the residual feeds the next block, so S is requested" },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 2097152,
+        "channels": 512,
+        "streams": 3,
+        "skipElements": 2097152
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1, "epsilon": 0.000001 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "normal", "seed": 400, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 401, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 402, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "normal", "seed": 403, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [512], "dist": "normal", "seed": 404, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ],
+        "primary": true
+      }
+    },
+    {
+      "name": "flux-vae-mid-1x64x64x512-g32-f16-nosum",
+      "preset": "smoke",
+      "provenance": {
+        "note": "The same node with S not requested: the sum is round-tripped through Y instead of a second output"
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 2097152,
+        "channels": 512,
+        "streams": 2,
+        "skipElements": 2097152
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1, "epsilon": 0.000001 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "normal", "seed": 410, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 411, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 412, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "normal", "seed": 413, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [512], "dist": "normal", "seed": 414, "scale": 0.1 }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 64, 64, 512], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "sd15-unet-2x64x64x320-g32-f16-sum",
+      "preset": "smoke",
+      "provenance": {
+        "note": "Stable-Diffusion-1.5 UNet resnet at the classifier-free-guidance batch; 10 channels per group, no bias"
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 2621440,
+        "channels": 320,
+        "streams": 3,
+        "skipElements": 2621440
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 64, 64, 320], "dist": "normal", "seed": 420, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [320], "dist": "uniform", "seed": 421, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [320], "dist": "normal", "seed": 422, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 64, 64, 320], "dist": "normal", "seed": 423, "scale": 0.5 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 64, 64, 320], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [2, 64, 64, 320], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "sd15-unet-2x32x32x640-g32-f16-nosum",
+      "preset": "smoke",
+      "provenance": {
+        "note": "Stable-Diffusion-1.5 UNet mid-resolution resnet without the activation or the second output"
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 1310720,
+        "channels": 640,
+        "streams": 2,
+        "skipElements": 1310720
+      },
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 32, 32, 640], "dist": "normal", "seed": 430, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [640], "dist": "uniform", "seed": 431, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [640], "dist": "normal", "seed": 432, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 32, 32, 640], "dist": "normal", "seed": 433, "scale": 0.5 }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [2, 32, 32, 640], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "sdxl-1x144x112x640-g32-f16-sum",
+      "preset": "smoke",
+      "provenance": { "note": "An SDXL non-square aspect grid, where a spatial-divisor block split degenerates" },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 10321920,
+        "channels": 640,
+        "streams": 3,
+        "skipElements": 10321920
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 144, 112, 640], "dist": "normal", "seed": 440, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [640], "dist": "uniform", "seed": 441, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [640], "dist": "normal", "seed": 442, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 144, 112, 640], "dist": "normal", "seed": 443, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [640], "dist": "normal", "seed": 444, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 144, 112, 640], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [1, 144, 112, 640], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "broadcast-skip-rank2-1x32x32x320-g32-f16",
+      "preset": "smoke",
+      "provenance": { "note": "Per-channel residual spelled (N, C)" },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 327680,
+        "channels": 320,
+        "streams": 3,
+        "skipElements": 320
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 32, 32, 320], "dist": "normal", "seed": 450, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [320], "dist": "uniform", "seed": 451, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [320], "dist": "normal", "seed": 452, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 320], "dist": "normal", "seed": 453, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [320], "dist": "normal", "seed": 454, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 32, 32, 320], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [1, 32, 32, 320], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "broadcast-skip-rank4-1x32x32x320-g32-f16",
+      "preset": "smoke",
+      "provenance": {
+        "note": "The same per-channel residual and the same data spelled (N, 1, 1, C), to measure the two spellings against each other"
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 327680,
+        "channels": 320,
+        "streams": 3,
+        "skipElements": 320
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 32, 32, 320], "dist": "normal", "seed": 450, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [320], "dist": "uniform", "seed": 451, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [320], "dist": "normal", "seed": 452, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 1, 1, 320], "dist": "normal", "seed": 453, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [320], "dist": "normal", "seed": 454, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 32, 32, 320], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [1, 32, 32, 320], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nchw-2x320x64x64-g32-f16-sum",
+      "preset": "smoke",
+      "provenance": { "note": "Channels-first Stable-Diffusion resnet width with the fused SiLU" },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 2621440,
+        "channels": 320,
+        "streams": 3,
+        "skipElements": 2621440
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 0 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 320, 64, 64], "dist": "normal", "seed": 460, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [320], "dist": "uniform", "seed": 461, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [320], "dist": "normal", "seed": 462, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 320, 64, 64], "dist": "normal", "seed": 463, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [320], "dist": "normal", "seed": 464, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 320, 64, 64], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [2, 320, 64, 64], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "micro-2x2x2x16-g4-f16-sum",
+      "preset": "smoke",
+      "provenance": {
+        "note": "ORT: onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc SkipGroupNormTest.SkipGroupNorm_with_bias geometry"
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 128,
+        "channels": 16,
+        "streams": 3,
+        "skipElements": 128
+      },
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 2, 2, 16], "dist": "normal", "seed": 470, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [16], "dist": "uniform", "seed": 471, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [16], "dist": "normal", "seed": 472, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 2, 2, 16], "dist": "normal", "seed": 473, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [16], "dist": "normal", "seed": 474, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 2, 2, 16], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [2, 2, 2, 16], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "flux-vae-up-1x256x256x256-g32-f16-sum",
+      "preset": "scale",
+      "provenance": { "note": "Flux.2 VAE decoder after the second 2x upsample" },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 16777216,
+        "channels": 256,
+        "streams": 3,
+        "skipElements": 16777216
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1, "epsilon": 0.000001 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 256, 256, 256], "dist": "normal", "seed": 480, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [256], "dist": "uniform", "seed": 481, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [256], "dist": "normal", "seed": 482, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 256, 256, 256], "dist": "normal", "seed": 483, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [256], "dist": "normal", "seed": 484, "scale": 0.1 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 256, 256, 256], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [1, 256, 256, 256], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "flux-vae-convout-1x512x512x128-g32-f16-nosum",
+      "preset": "model",
+      "provenance": { "note": "Flux.2 VAE conv_norm_out geometry: the largest single dispatch in the decoder" },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 33554432,
+        "channels": 128,
+        "streams": 2,
+        "skipElements": 33554432
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1, "epsilon": 0.000001 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [1, 512, 512, 128], "dist": "normal", "seed": 490, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [128], "dist": "uniform", "seed": 491, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [128], "dist": "normal", "seed": 492, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [1, 512, 512, 128], "dist": "normal", "seed": 493, "scale": 0.5 },
+        "bias": { "dtype": "float16", "shape": [128], "dist": "normal", "seed": 494, "scale": 0.1 }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 512, 512, 128], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nhwc-even-groups-2x64x64x320-g32-f32",
+      "preset": "smoke",
+      "provenance": {
+        "note": "Even channel-group widths spanning one, multiple, and capacity-limited channel rows per workgroup, including float32 activations. Same distributions as the authored large NHWC case."
+      },
+      "vars": {
+        "dtype": "float32",
+        "weightDtype": "float32",
+        "count": 2621440,
+        "channels": 320,
+        "streams": 3,
+        "skipElements": 2621440
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [2, 64, 64, 320], "dist": "normal", "seed": 420, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [320], "dist": "uniform", "seed": 421, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [320], "dist": "normal", "seed": 422, "scale": 0.1 },
+        "skip": { "dtype": "float32", "shape": [2, 64, 64, 320], "dist": "normal", "seed": 423, "scale": 0.5 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 64, 64, 320], "dist": "empty" },
+        "S": { "dtype": "float32", "shape": [2, 64, 64, 320], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nhwc-even-groups-2x64x64x64-g32-f16",
+      "preset": "smoke",
+      "provenance": {
+        "note": "Even channel-group widths spanning one, multiple, and capacity-limited channel rows per workgroup, including float32 activations. Same distributions as the authored large NHWC case."
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 524288,
+        "channels": 64,
+        "streams": 3,
+        "skipElements": 524288
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 64, 64, 64], "dist": "normal", "seed": 420, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [64], "dist": "uniform", "seed": 421, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 422, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 64, 64, 64], "dist": "normal", "seed": 423, "scale": 0.5 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 64, 64, 64], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [2, 64, 64, 64], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nhwc-even-groups-2x64x64x192-g32-f16",
+      "preset": "smoke",
+      "provenance": {
+        "note": "Even channel-group widths spanning one, multiple, and capacity-limited channel rows per workgroup, including float32 activations. Same distributions as the authored large NHWC case."
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 1572864,
+        "channels": 192,
+        "streams": 3,
+        "skipElements": 1572864
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 64, 64, 192], "dist": "normal", "seed": 420, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [192], "dist": "uniform", "seed": 421, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [192], "dist": "normal", "seed": 422, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 64, 64, 192], "dist": "normal", "seed": 423, "scale": 0.5 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 64, 64, 192], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [2, 64, 64, 192], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nhwc-even-groups-2x64x64x448-g32-f16",
+      "preset": "smoke",
+      "provenance": {
+        "note": "Even channel-group widths spanning one, multiple, and capacity-limited channel rows per workgroup, including float32 activations. Same distributions as the authored large NHWC case."
+      },
+      "vars": {
+        "dtype": "float16",
+        "weightDtype": "float32",
+        "count": 3670016,
+        "channels": 448,
+        "streams": 3,
+        "skipElements": 3670016
+      },
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": { "dtype": "float16", "shape": [2, 64, 64, 448], "dist": "normal", "seed": 420, "scale": 0.5 },
+        "gamma": { "dtype": "float32", "shape": [448], "dist": "uniform", "seed": 421, "scale": 0.25, "offset": 1 },
+        "beta": { "dtype": "float32", "shape": [448], "dist": "normal", "seed": 422, "scale": 0.1 },
+        "skip": { "dtype": "float16", "shape": [2, 64, 64, 448], "dist": "normal", "seed": 423, "scale": 0.5 }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 64, 64, 448], "dist": "empty" },
+        "S": { "dtype": "float16", "shape": [2, 64, 64, 448], "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.count * args.streams + args.skipElements) * dtypeBytes(args.dtype) + args.channels * 2 * dtypeBytes(args.weightDtype)"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/group-norm-combine.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,50 @@

+{% set comp0 = ".x" if vecSpec > 1 else "" %}
+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// The partials hold sum(value - shift) and sum((value - shift)^2) over disjoint blocks
+// of the group, all taken against the same `shift` — the group's first element — so
+// they add directly. Recovering the mean adds the shift back; the variance is formed
+// from the shifted second moment and clamped at zero, which is where a group whose
+// values are all equal would otherwise land a few ulp below it.
+const GROUPS: u32 = {{ groupsSpec }}u;
+const CPG: u32 = {{ cpgSpec }}u;
+const ROWS: u32 = {{ normRowsSpec }}u;
+const PARTIALS_PER_ROW: u32 = {{ partialsPerRowSpec }}u;
+const GROUP_SIZE: f32 = {{ groupSizeSpec }}.0;
+const EPSILON: f32 = {{ epsilonSpec }};
+const WG: u32 = {{ combineWgSpec }}u;
+{% if phaseLayout %}
+const HW: u32 = {{ dimHWSpec }}u;
+const CHUNKS: u32 = {{ chunksSpec }}u;
+const VEC: u32 = {{ vecSpec }}u;
+{% else %}
+const CHANNELS: u32 = {{ dimCSpec }}u;
+const PLANE_V: u32 = {{ planeVecsSpec }}u;
+{% endif %}
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("WG", "row", "ROWS") }}
+  let n = row / GROUPS;
+  let g = row % GROUPS;
+{% if phaseLayout %}
+  let shift = f32(src[n * HW * CHUNKS + g * (CPG / VEC)]{{ comp0 }});
+{% else %}
+  let shift = f32(src[(n * CHANNELS + g * CPG) * PLANE_V]{{ comp0 }});
+{% endif %}
+  var total = vec2<f32>(0.0);
+  for (var k = 0u; k < PARTIALS_PER_ROW; k = k + 1u) {
+    total = total + partials[row * PARTIALS_PER_ROW + k];
+  }
+  let mean_d = total.x / GROUP_SIZE;
+  let variance = max(total.y / GROUP_SIZE - mean_d * mean_d, 0.0);
+  stats[row] = vec2<f32>(shift + mean_d, inverseSqrt(variance + EPSILON));
+}

build/webgpu/group-norm-fused.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,190 @@

+{% set comp0 = ".x" if useVec4 else "" %}
+{% macro pre_norm(idx, sidx) %}
+{% set x = vecF32 ~ "(X[" ~ idx ~ "])" %}
+{% set skipTerm = "skip_v" if broadcastSkip else vecF32 ~ "(skip[" ~ sidx ~ "])" %}
+{{ vecF32 ~ "(" ~ storeElement ~ "(" ~ x ~ " + " ~ skipTerm ~ " + bias_v))" if hasSkip else x }}{% endmacro %}
+{% macro shift_at(idx, sidx, ch) %}
+{% set x = "f32(X[" ~ idx ~ "]" ~ comp0 ~ ")" %}
+{% set skipTerm = "f32(skip[" ~ sidx ~ "]" ~ ("" if broadcastSkip else comp0) ~ ")" %}
+{% set biasTerm = " + f32(bias[" ~ ch ~ "])" if hasBias else "" %}
+{{ "f32(" ~ tScalar ~ "(" ~ x ~ " + " ~ skipTerm ~ biasTerm ~ "))" if hasSkip else x }}{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// Group normalization in one dispatch. The workgroup accumulates the group's shifted
+// moments, folds them through workgroup memory, then walks the same elements again to
+// write `gamma * (value - mean) * inv_std + beta`. Every intermediate is f32 whatever
+// the tensor type; only the store narrows. The reduction is a plain halving tree over
+// a power-of-two-padded scratch array, so it is deterministic and needs no subgroup
+// support.
+const CPG: u32 = {{ cpgSpec }}u;
+const GROUPS: u32 = {{ groupsSpec }}u;
+const ROWS: u32 = {{ normRowsSpec }}u;
+const GROUP_SIZE: f32 = {{ groupSizeSpec }}.0;
+const EPSILON: f32 = {{ epsilonSpec }};
+const WG: u32 = {{ wgSpec }}u;
+const LANES: u32 = {{ ldsSpec }}u;
+{% if phaseLayout %}
+const HW: u32 = {{ dimHWSpec }}u;
+const VEC: u32 = {{ vecSpec }}u;
+const CHUNKS: u32 = {{ chunksSpec }}u;
+const SEGMENT: u32 = {{ fusedChunksSpec }}u;
+const SUB_ROWS: u32 = {{ fusedRowsSpec }}u;
+{% if hasSkip and broadcastSkip %}
+const CHANNELS: u32 = {{ dimCSpec }}u;
+{% endif %}
+{% else %}
+const CHANNELS: u32 = {{ dimCSpec }}u;
+const PLANE_V: u32 = {{ planeVecsSpec }}u;
+const SPAN_V: u32 = {{ spanVecsSpec }}u;
+{% endif %}
+// LANES is WG rounded up to a power of two; the lanes past WG are zeroed so the
+// halving tree below stays uniform.
+var<workgroup> lane_sums: array<vec2<f32>, LANES>;
+{% if activationSpec == 1 %}
+fn silu_e(v: {{ vecF32 }}) -> {{ vecF32 }} {
+  return v / ({{ vecF32 }}(1.0) + exp(-v));
+}
+{% endif %}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  // 2D-folded row index: wg.y carries the high bits past the dispatch's per-axis
+  // workgroup fold width.
+  let row = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
+  if (row >= ROWS) {
+    return;
+  }
+  let n = row / GROUPS;
+  let g = row - n * GROUPS;
+{% if phaseLayout %}
+  let chunk = lid.x % SEGMENT;
+  let sub_row = lid.x / SEGMENT;
+  let c0 = g * CPG + chunk * VEC;
+  let base_v = n * HW * CHUNKS + g * SEGMENT;
+  let shift_index = base_v;
+{% if broadcastSkip %}
+  let shift_skip = n * CHANNELS + g * CPG;
+{% else %}
+  let skip_base = base_v;
+  let shift_skip = base_v;
+{% endif %}
+{% else %}
+  let base_v = (n * CHANNELS + g * CPG) * PLANE_V;
+  let shift_index = base_v;
+{% if broadcastSkip %}
+  let shift_skip = n * CHANNELS + g * CPG;
+{% else %}
+  let skip_base = base_v;
+  let shift_skip = base_v;
+{% endif %}
+{% endif %}
+{% if hasBias %}
+{% if phaseLayout %}
+{% if useVec4 %}
+  let bias_v = vec4<f32>(f32(bias[c0]), f32(bias[c0 + 1u]), f32(bias[c0 + 2u]), f32(bias[c0 + 3u]));
+{% else %}
+  let bias_v = f32(bias[c0]);
+{% endif %}
+{% else %}
+  // Channels first: the channel changes as this workgroup's walk crosses a plane, so
+  // the per-channel bias is refreshed inside the loop rather than hoisted.
+  var bias_v = {{ vecF32 }}(0.0);
+{% endif %}
+{% else %}
+  let bias_v = {{ vecF32 }}(0.0);
+{% endif %}
+{% if broadcastSkip %}
+{% if phaseLayout %}
+{% set sb = "n * CHANNELS + c0" %}
+{% if useVec4 %}
+  let skip_v = vec4<f32>(f32(skip[{{ sb }}]), f32(skip[{{ sb }} + 1u]), f32(skip[{{ sb }} + 2u]), f32(skip[{{ sb }} + 3u]));
+{% else %}
+  let skip_v = f32(skip[{{ sb }}]);
+{% endif %}
+{% else %}
+  // Channels first: this workgroup owns a whole group, so the channel — and with it
+  // the broadcast skip value — changes as the walk crosses a channel plane.
+  var skip_v = {{ vecF32 }}(0.0);
+{% endif %}
+{% endif %}
+  let shift = {{ shift_at("shift_index", "shift_skip", "g * CPG") }};
+  var acc = vec2<f32>(0.0);
+{% if phaseLayout %}
+  for (var hw = sub_row; hw < HW; hw = hw + SUB_ROWS) {
+    let idx = base_v + hw * CHUNKS + chunk;
+    let v = {{ pre_norm("idx", "skip_base + hw * CHUNKS + chunk") }};
+{% else %}
+  for (var i = lid.x; i < SPAN_V; i = i + WG) {
+    let idx = base_v + i;
+{% if hasSkip and (hasBias or broadcastSkip) %}
+    let channel = g * CPG + i / PLANE_V;
+{% if hasBias %}
+    bias_v = {{ vecF32 }}(f32(bias[channel]));
+{% endif %}
+{% if broadcastSkip %}
+    skip_v = {{ vecF32 }}(f32(skip[n * CHANNELS + channel]));
+{% endif %}
+{% endif %}
+    let v = {{ pre_norm("idx", "skip_base + i") }};
+{% endif %}
+    {{ fusedSrc }}[idx] = {{ storeElement }}(v);
+    let d = v - {{ vecF32 }}(shift);
+{% if useVec4 %}
+    acc = acc + vec2<f32>(d.x + d.y + d.z + d.w, dot(d, d));
+{% else %}
+    acc = acc + vec2<f32>(d, d * d);
+{% endif %}
+  }
+  lane_sums[lid.x] = acc;
+  if (lid.x + WG < LANES) {
+    lane_sums[lid.x + WG] = vec2<f32>(0.0);
+  }
+  workgroupBarrier();
+  var stride = LANES / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (lid.x < stride) {
+      lane_sums[lid.x] = lane_sums[lid.x] + lane_sums[lid.x + stride];
+    }
+    workgroupBarrier();
+    stride = stride / 2u;
+  }
+  let total = lane_sums[0];
+  let mean_d = total.x / GROUP_SIZE;
+  let variance = max(total.y / GROUP_SIZE - mean_d * mean_d, 0.0);
+  let mean = {{ vecF32 }}(shift + mean_d);
+  let inv_std = {{ vecF32 }}(inverseSqrt(variance + EPSILON));
+{% if phaseLayout %}
+{% if useVec4 %}
+  let gamma_v = vec4<f32>(f32(gamma[c0]), f32(gamma[c0 + 1u]), f32(gamma[c0 + 2u]), f32(gamma[c0 + 3u]));
+  let beta_v = vec4<f32>(f32(beta[c0]), f32(beta[c0 + 1u]), f32(beta[c0 + 2u]), f32(beta[c0 + 3u]));
+{% else %}
+  let gamma_v = f32(gamma[c0]);
+  let beta_v = f32(beta[c0]);
+{% endif %}
+  for (var hw = sub_row; hw < HW; hw = hw + SUB_ROWS) {
+    let idx = base_v + hw * CHUNKS + chunk;
+{% else %}
+  for (var i = lid.x; i < SPAN_V; i = i + WG) {
+    let idx = base_v + i;
+    let channel = g * CPG + i / PLANE_V;
+{% if useVec4 %}
+    let gamma_v = vec4<f32>(f32(gamma[channel]));
+    let beta_v = vec4<f32>(f32(beta[channel]));
+{% else %}
+    let gamma_v = f32(gamma[channel]);
+    let beta_v = f32(beta[channel]);
+{% endif %}
+{% endif %}
+    let normalized = ({{ vecF32 }}({{ fusedSrc }}[idx]) - mean) * inv_std;
+{% if activationSpec == 1 %}
+    Y[idx] = {{ storeElement }}(silu_e(gamma_v * normalized + beta_v));
+{% else %}
+    Y[idx] = {{ storeElement }}(gamma_v * normalized + beta_v);
+{% endif %}
+  }
+}

build/webgpu/group-norm-slab.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,277 @@

+{% set blockedStats = phase == "stats" and phaseLayout and (statsBlockedSpec | default(false)) %}
+{% set useVector = vecSpec > 1 %}
+{% set comp0 = ".x" if useVector else "" %}
+{% macro channel_values(buffer, channel) %}
+{{ vecF32 }}({% for i in range(vecSpec) %}f32({{ buffer }}[{{ channel }}{% if i > 0 %} + {{ i }}u{% endif %}]){% if not loop.last %}, {% endif %}{% endfor %}){% endmacro %}
+{% macro component_sum(value) %}
+{% for i in range(vecSpec) %}{{ value }}.{{ "xyzw"[i] }}{% if not loop.last %} + {% endif %}{% endfor %}{% endmacro %}
+{% macro pre_norm(idx, sidx) %}
+{% set x = vecF32 ~ "(X[" ~ idx ~ "])" %}
+{% set skipTerm = "skip_v" if broadcastSkip else vecF32 ~ "(skip[" ~ sidx ~ "])" %}
+{{ vecF32 ~ "(" ~ storeElement ~ "(" ~ x ~ " + " ~ skipTerm ~ " + bias_v))" if hasSkip else x }}{% endmacro %}
+{% macro shift_at(idx, sidx, ch) %}
+{% set x = "f32(X[" ~ idx ~ "]" ~ comp0 ~ ")" %}
+{% set skipTerm = "f32(skip[" ~ sidx ~ "]" ~ ("" if broadcastSkip else comp0) ~ ")" %}
+{% set biasTerm = " + f32(bias[" ~ ch ~ "])" if hasBias else "" %}
+{{ "f32(" ~ tScalar ~ "(" ~ x ~ " + " ~ skipTerm ~ biasTerm ~ "))" if hasSkip else x }}{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// One phase of group normalization over a 4-D image tensor.
+//
+// A group is every spatial position of `CPG` adjacent channels of one batch item. The
+// statistics phase accumulates the sum and the sum of squares of `value - shift`,
+// where `shift` is the group's first element: subtracting one member of the group
+// removes the common offset before the second moment is formed, so the variance the
+// combine pass derives does not cancel against a large mean. The apply phase reads the
+// combined mean and reciprocal standard deviation back and writes
+// `gamma * (value - mean) * inv_std + beta`, optionally through SiLU.
+//
+// Every accumulation and the whole affine are f32 whatever the tensor type; only the
+// store narrows.
+const CPG: u32 = {{ cpgSpec }}u;
+const GROUPS: u32 = {{ groupsSpec }}u;
+const WG: u32 = {{ wgSpec }}u;
+const UNITS: u32 = {{ unitsSpec }}u;
+{% if phaseLayout %}
+const HW: u32 = {{ dimHWSpec }}u;
+const VEC: u32 = {{ vecSpec }}u;
+const CHUNKS: u32 = {{ chunksSpec }}u;
+const ROWS_PER_STEP: u32 = {{ rowsPerStepSpec }}u;
+const ROWS_PER_SLAB: u32 = {{ rowsPerSlabSpec }}u;
+const SLABS: u32 = {{ splitSpec }}u;
+{% if blockedStats %}
+const GROUP_BLOCKS: u32 = {{ statsGroupBlocksSpec }}u;
+const GROUPS_PER_BLOCK: u32 = {{ statsGroupsPerBlockSpec }}u;
+const TILE_CHUNKS: u32 = {{ statsTileChunksSpec }}u;
+{% endif %}
+{% if hasSkip and broadcastSkip and phase == "stats" %}
+const CHANNELS: u32 = {{ dimCSpec }}u;
+{% endif %}
+{% else %}
+const CHANNELS: u32 = {{ dimCSpec }}u;
+const PLANE_V: u32 = {{ planeVecsSpec }}u;
+const CHUNK_V: u32 = {{ chunkVecsSpec }}u;
+const SPLITS: u32 = {{ splitSpec }}u;
+{% endif %}
+{% if phase == "stats" %}
+{% if not phaseLayout %}
+const PARTIALS_PER_ROW: u32 = {{ partialsPerRowSpec }}u;
+{% endif %}
+// Per-lane moments, combined below: the channels-last arm folds one contiguous
+// segment per group out of this array, the channels-first arm folds all of it.
+var<workgroup> lane_sums: array<vec2<f32>, WG>;
+{% endif %}
+{% if phase == "apply" and activationSpec == 1 %}
+fn silu_e(v: {{ vecF32 }}) -> {{ vecF32 }} {
+  return v / ({{ vecF32 }}(1.0) + exp(-v));
+}
+{% endif %}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  // 2D-folded work-unit index: wg.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+  let unit = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
+  if (unit >= UNITS) {
+    return;
+  }
+{% if phaseLayout %}
+  // Channels last. The workgroup owns one batch item and one slab of spatial rows and
+  // covers every channel of those rows: lane `lid.x` keeps channel chunk
+  // `lid.x % CHUNKS` and steps through the slab `ROWS_PER_STEP` rows at a time, so the
+  // workgroup reads whole contiguous channel rows and each lane's group is fixed.
+{% if blockedStats %}
+  let spatial_unit = unit / GROUP_BLOCKS;
+  let group_begin = (unit % GROUP_BLOCKS) * GROUPS_PER_BLOCK;
+  let n = spatial_unit / SLABS;
+  let slab = spatial_unit - n * SLABS;
+  let chunk = group_begin * (CPG / VEC) + lid.x % TILE_CHUNKS;
+  let sub_row = lid.x / TILE_CHUNKS;
+{% else %}
+  let n = unit / SLABS;
+  let slab = unit - n * SLABS;
+  let chunk = lid.x % CHUNKS;
+  let sub_row = lid.x / CHUNKS;
+{% endif %}
+  let c0 = chunk * VEC;
+{% if phase == "stats" or groupUniform %}
+  let g = c0 / CPG;
+{% endif %}
+  let image_base = n * HW * CHUNKS;
+  let hw_begin = slab * ROWS_PER_SLAB;
+  let hw_end = min(hw_begin + ROWS_PER_SLAB, HW);
+{% else %}
+  // Channels first. The workgroup owns one contiguous chunk of one (batch item,
+  // channel) plane, so every lane shares the channel and the group and the loads are
+  // contiguous across the workgroup. Adjacent units are adjacent chunks of one plane.
+  let plane = unit / SPLITS;
+  let part = unit - plane * SPLITS;
+  let n = plane / CHANNELS;
+  let c = plane % CHANNELS;
+  let g = c / CPG;
+  let base_v = plane * PLANE_V;
+  let chunk_begin = part * CHUNK_V;
+  let chunk_end = min(chunk_begin + CHUNK_V, PLANE_V);
+{% endif %}
+{% if phase == "stats" %}
+{% if blockedStats %}
+  var acc = vec2<f32>(0.0);
+  // Keep tail lanes participating in the barrier without reading outside the group.
+  if (g < GROUPS) {
+{% endif %}
+{% if phaseLayout %}
+  let shift_index = image_base + g * (CPG / VEC);
+{% else %}
+  let shift_index = (n * CHANNELS + g * CPG) * PLANE_V;
+{% endif %}
+{% if broadcastSkip %}
+  let shift_skip = n * CHANNELS + g * CPG;
+{% elif phaseLayout %}
+  let skip_base = image_base;
+  let shift_skip = shift_index;
+{% else %}
+  let skip_base = base_v;
+  let shift_skip = shift_index;
+{% endif %}
+{% if hasBias %}
+{% if useVector %}
+{% if phaseLayout %}
+  let bias_v = {{ channel_values("bias", "c0") }};
+{% else %}
+  let bias_v = {{ vecF32 }}(f32(bias[c]));
+{% endif %}
+{% else %}
+{% if phaseLayout %}
+  let bias_v = f32(bias[c0]);
+{% else %}
+  let bias_v = f32(bias[c]);
+{% endif %}
+{% endif %}
+{% else %}
+  let bias_v = {{ vecF32 }}(0.0);
+{% endif %}
+{% if broadcastSkip %}
+{% if phaseLayout %}
+{% set sb = "n * CHANNELS + c0" %}
+{% if useVector %}
+  let skip_v = {{ channel_values("skip", sb) }};
+{% else %}
+  let skip_v = f32(skip[{{ sb }}]);
+{% endif %}
+{% else %}
+  // A broadcast skip is one value per (batch item, channel); this whole workgroup
+  // shares a channel, so it is read once and splatted.
+  let skip_v = {{ vecF32 }}(f32(skip[n * CHANNELS + c]));
+{% endif %}
+{% endif %}
+  let shift = {{ shift_at("shift_index", "shift_skip", "g * CPG") }};
+{% if not blockedStats %}
+  var acc = vec2<f32>(0.0);
+{% endif %}
+{% if phaseLayout %}
+  for (var hw = hw_begin + sub_row; hw < hw_end; hw = hw + ROWS_PER_STEP) {
+    let idx = image_base + hw * CHUNKS + chunk;
+    let v = {{ pre_norm("idx", "skip_base + hw * CHUNKS + chunk") }};
+{% else %}
+  for (var i = chunk_begin + lid.x; i < chunk_end; i = i + WG) {
+    let idx = base_v + i;
+    let v = {{ pre_norm("idx", "skip_base + i") }};
+{% endif %}
+    sum_dst[idx] = {{ storeElement }}(v);
+    let d = v - {{ vecF32 }}(shift);
+{% if useVector %}
+    acc = acc + vec2<f32>({{ component_sum("d") }}, dot(d, d));
+{% else %}
+    acc = acc + vec2<f32>(d, d * d);
+{% endif %}
+  }
+{% if blockedStats %}
+  }
+{% endif %}
+  lane_sums[lid.x] = acc;
+  workgroupBarrier();
+{% if phaseLayout %}
+  // One aligned segment of `CPG / VEC` lanes per group per sub-row. Each group's
+  // partial is summed in a fixed order by a single lane, so the value does not depend
+  // on how the workgroup was scheduled.
+  let segment = CPG / VEC;
+  for (var sel = lid.x; sel < {% if blockedStats %}min(GROUPS_PER_BLOCK, GROUPS - group_begin){% else %}GROUPS{% endif %}; sel = sel + WG) {
+    var total = vec2<f32>(0.0);
+    for (var r = 0u; r < ROWS_PER_STEP; r = r + 1u) {
+      let seg_base = r * {% if blockedStats %}TILE_CHUNKS{% else %}CHUNKS{% endif %} + sel * segment;
+      for (var t = 0u; t < segment; t = t + 1u) {
+        total = total + lane_sums[seg_base + t];
+      }
+    }
+    partials[(n * GROUPS + {% if blockedStats %}group_begin + {% endif %}sel) * SLABS + slab] = total;
+  }
+{% else %}
+  // WG is a power of two on this arm, so the halving stride is uniform across the
+  // workgroup and every lane reaches every barrier.
+  var stride = WG / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (lid.x < stride) {
+      lane_sums[lid.x] = lane_sums[lid.x] + lane_sums[lid.x + stride];
+    }
+    workgroupBarrier();
+    stride = stride / 2u;
+  }
+  if (lid.x == 0u) {
+    partials[(n * GROUPS + g) * PARTIALS_PER_ROW + (c - g * CPG) * SPLITS + part] = lane_sums[0];
+  }
+{% endif %}
+{% else %}
+{% if groupUniform %}
+  let st = stats[n * GROUPS + g];
+  let mean = {{ vecF32 }}(st.x);
+  let inv_std = {{ vecF32 }}(st.y);
+{% else %}
+  // The four channels this lane holds do not all belong to one group, which is what
+  // `channels_per_group % 4 != 0` means. The statistics are still read once per lane,
+  // as a component-wise gather over at most two adjacent groups.
+  let row_base = n * GROUPS;
+  let s0 = stats[row_base + c0 / CPG];
+  let s1 = stats[row_base + (c0 + 1u) / CPG];
+  let s2 = stats[row_base + (c0 + 2u) / CPG];
+  let s3 = stats[row_base + (c0 + 3u) / CPG];
+  let mean = vec4<f32>(s0.x, s1.x, s2.x, s3.x);
+  let inv_std = vec4<f32>(s0.y, s1.y, s2.y, s3.y);
+{% endif %}
+{% if useVector %}
+{% if phaseLayout %}
+  let gamma_v = {{ channel_values("gamma", "c0") }};
+  let beta_v = {{ channel_values("beta", "c0") }};
+{% else %}
+  let gamma_v = {{ vecF32 }}(f32(gamma[c]));
+  let beta_v = {{ vecF32 }}(f32(beta[c]));
+{% endif %}
+{% else %}
+{% if phaseLayout %}
+  let gamma_v = f32(gamma[c0]);
+  let beta_v = f32(beta[c0]);
+{% else %}
+  let gamma_v = f32(gamma[c]);
+  let beta_v = f32(beta[c]);
+{% endif %}
+{% endif %}
+{% if phaseLayout %}
+  for (var hw = hw_begin + sub_row; hw < hw_end; hw = hw + ROWS_PER_STEP) {
+    let idx = image_base + hw * CHUNKS + chunk;
+{% else %}
+  for (var i = chunk_begin + lid.x; i < chunk_end; i = i + WG) {
+    let idx = base_v + i;
+{% endif %}
+    let normalized = ({{ vecF32 }}({{ applySrc }}[idx]) - mean) * inv_std;
+{% if activationSpec == 1 %}
+    Y[idx] = {{ storeElement }}(silu_e(gamma_v * normalized + beta_v));
+{% else %}
+    Y[idx] = {{ storeElement }}(gamma_v * normalized + beta_v);
+{% endif %}
+  }
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,633 @@

+{
+  "domain": "com.microsoft",
+  "name": "SkipGroupNorm",
+  "sinceVersion": 1,
+  "inputs": {
+    "X": { "dtype": "T", "rank": 4 },
+    "gamma": { "dtype": "M", "rank": 1 },
+    "beta": { "dtype": "M", "rank": 1 },
+    "skip": { "dtype": "T" },
+    "bias": { "dtype": "T", "rank": 1, "optional": true }
+  },
+  "outputs": {
+    "Y": { "dtype": "T", "rank": 4, "shape": "shapes.X" },
+    "S": { "dtype": "T", "rank": 4, "optional": true, "shape": "shapes.X" }
+  },
+  "attributes": { "activation": {}, "groups": {}, "channels_last": { "default": 1 }, "epsilon": { "default": 0.00001 } },
+  "attributeConstraints": { "activation": { "required": true, "values": [0, 1] }, "groups": { "required": true } },
+  "typeConstraints": { "T": ["float32", "float16"], "M": ["float32", "float16"] },
+  "tunables": {
+    "WORKGROUP_SIZE": { "default": 256 },
+    "MAX_SPLITS": { "default": 256 },
+    "VALUES_PER_SPLIT": { "default": 4096 },
+    "SPLIT_MIN_ELEMENTS": { "default": 262144 },
+    "APPLY_SLAB_ELEMENTS": { "default": 8192 },
+    "FUSED_MIN_ROWS": { "default": 64 },
+    "FUSED_MAX_GROUP_BYTES": { "default": 131072 }
+  },
+  "derive": {
+    "attrsPresent": "has(attrs, \"groups\") and has(attrs, \"activation\")",
+    "numGroups": "attrs.groups if attrsPresent else 0",
+    "activationKind": "attrs.activation if attrsPresent else 0",
+    "channelsLast": "attrs.channels_last != 0",
+    "rankOk": "ranks.X == 4 and ranks.Y == 4 and ranks.gamma == 1 and ranks.beta == 1",
+    "dimN": "dim(shapes.X, 0) if rankOk else 0",
+    "dimC": "(dim(shapes.X, 3) if channelsLast else dim(shapes.X, 1)) if rankOk else 0",
+    "dimHW": "(dim(shapes.X, 1) * dim(shapes.X, 2) if channelsLast else dim(shapes.X, 2) * dim(shapes.X, 3)) if rankOk else 0",
+    "groupsOk": "numGroups >= 1 and dimC >= 1 and dimC % numGroups == 0",
+    "channelsPerGroup": "dimC / numGroups if groupsOk else 0",
+    "planeCount": "dimN * dimC",
+    "groupSize": "channelsPerGroup * dimHW",
+    "normRows": "dimN * numGroups",
+    "elementCount": "numel(shapes.X) if rankOk else 0",
+    "shapeOk": "rankOk and groupsOk and dimN > 0 and dimHW > 0 and groupSize > 0 and sameShape(shapes.Y, shapes.X) and dim(shapes.gamma, 0) == dimC and dim(shapes.beta, 0) == dimC",
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+    "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "statsVec4Ok": "(channelsPerGroup % 4 == 0 and dimC % 4 == 0) if channelsLast else (dimHW % 4 == 0)",
+    "applyVec4Ok": "(dimC % 4 == 0) if channelsLast else (dimHW % 4 == 0)",
+    "fusedVecWidth": "4 if statsVec4Ok else 1",
+    "statsVecWidth": "4 if statsVec4Ok else (2 if channelsLast and channelsPerGroup % 2 == 0 else 1)",
+    "applyVecWidth": "4 if applyVec4Ok else 1",
+    "workgroupCap": "max(1, floor(pow2ceil(min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap) + 1) / 2))",
+    "statsChunks": "max(1, floor(dimC / statsVecWidth))",
+    "fusedRowChunks": "max(1, floor(dimC / fusedVecWidth))",
+    "statsSegment": "max(1, floor(channelsPerGroup / statsVecWidth))",
+    "statsGroupsPerBlock": "max(1, min(numGroups, floor(workgroupCap / statsSegment))) if channelsLast and statsChunks > deviceWorkgroupCap and statsSegment <= workgroupCap else max(1, numGroups)",
+    "statsGroupBlocks": "ceilDiv(max(1, numGroups), statsGroupsPerBlock)",
+    "statsTileChunks": "statsGroupsPerBlock * statsSegment if statsGroupBlocks > 1 else statsChunks",
+    "applyChunks": "max(1, floor(dimC / applyVecWidth))",
+    "statsRowsPerStep": "max(1, floor(workgroupCap / statsTileChunks))",
+    "applyRowsPerStep": "max(1, floor(workgroupCap / applyChunks))",
+    "planeVecs": "max(1, floor(dimHW / statsVecWidth))",
+    "slabCount": "ceilDiv(dimHW, ceilDiv(dimHW, max(1, min(tunables.MAX_SPLITS, dimHW, ceilDiv(dimHW * dimC, tunables.VALUES_PER_SPLIT)))))",
+    "planeSplits": "max(1, min(tunables.MAX_SPLITS, planeVecs, ceilDiv(dimHW, tunables.VALUES_PER_SPLIT)))",
+    "splitCount": "slabCount if channelsLast else planeSplits",
+    "partialsPerRow": "slabCount if channelsLast else max(1, channelsPerGroup) * planeSplits",
+    "partialElements": "max(1, normRows * partialsPerRow)",
+    "statsElements": "max(1, normRows)",
+    "statsWorkgroup": "(statsTileChunks * statsRowsPerStep) if channelsLast else min(workgroupCap, max(32, pow2ceil(ceilDiv(planeVecs, planeSplits))))",
+    "applyWorkgroup": "(applyChunks * applyRowsPerStep) if channelsLast else statsWorkgroup",
+    "applySlabCount": "ceilDiv(dimHW, ceilDiv(dimHW, max(1, min(dimHW, min(device.limits.maxComputeWorkgroupsPerDimension, 65535), ceilDiv(elementCount, tunables.APPLY_SLAB_ELEMENTS)))))",
+    "applySplit": "applySlabCount if channelsLast else planeSplits",
+    "applyUnits": "(dimN * applySlabCount) if channelsLast else (planeCount * planeSplits)",
+    "workUnits": "(dimN * slabCount * statsGroupBlocks) if channelsLast else (planeCount * planeSplits)",
+    "splitFits": "statsWorkgroup <= deviceWorkgroupCap and applyWorkgroup <= deviceWorkgroupCap and statsWorkgroup * 8 <= device.limits.maxComputeWorkgroupStorageSize and partialElements * 8 <= storageBufferLimit and workUnits > 0 and workUnits <= foldedDispatchCapacity and applyUnits <= foldedDispatchCapacity",
+    "fusedChunks": "max(1, floor(channelsPerGroup / fusedVecWidth)) if channelsLast else 1",
+    "fusedWorkgroup": "(fusedChunks * max(1, floor(workgroupCap / fusedChunks))) if channelsLast else min(workgroupCap, max(32, pow2ceil(max(1, floor(groupSize / fusedVecWidth)))))",
+    "fusedLanes": "pow2ceil(fusedWorkgroup)",
+    "fusedFits": "fusedWorkgroup <= deviceWorkgroupCap and fusedLanes * 8 <= device.limits.maxComputeWorkgroupStorageSize and normRows <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "combineWorkgroup": "min(workgroupCap, 64)",
+    "statsElement": "(\"vec\" ~ statsVecWidth ~ \"<\" ~ dtypes.T ~ \">\") if statsVecWidth > 1 else dtypes.T",
+    "fusedElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if statsVec4Ok else dtypes.T",
+    "applyElement": "(\"vec4<\" ~ dtypes.T ~ \">\") if applyVec4Ok else dtypes.T",
+    "mScalar": "dtypes.M",
+    "tScalar": "dtypes.T",
+    "skipExact": "rankOk and ranks.skip == 4 and sameShape(shapes.skip, shapes.X)",
+    "skipBroadcast2": "rankOk and ranks.skip == 2 and dim(shapes.skip, 0) == dimN and dim(shapes.skip, 1) == dimC",
+    "skipBroadcast4": "rankOk and ranks.skip == 4 and dim(shapes.skip, 0) == dimN and dim(shapes.skip, 1) == 1 and dim(shapes.skip, 2) == 1 and dim(shapes.skip, 3) == dimC",
+    "broadcastSkipSpec": "(not skipExact) and (skipBroadcast2 or skipBroadcast4)",
+    "skipOk": "skipExact or broadcastSkipSpec",
+    "biasOk": "(not present.bias) or (ranks.bias == 1 and dim(shapes.bias, 0) == dimC)",
+    "sumOk": "(not present.S) or (ranks.S == 4 and sameShape(shapes.S, shapes.X))",
+    "contractOk": "shapeOk and skipOk and biasOk and sumOk and f16Ok(dtypes.T) and f16Ok(dtypes.M)",
+    "skipElement": "dtypes.T if broadcastSkipSpec else statsElement"
+  },
+  "when": ["contractOk"],
+  "bindings": {
+    "xIn": { "arg": "X", "name": "X", "elementType": "$statsElement" },
+    "gamma": { "elementType": "$mScalar" },
+    "beta": { "elementType": "$mScalar" },
+    "yOut": { "arg": "Y", "name": "Y", "elementType": "$applyElement" },
+    "partialsW": { "name": "partials", "buffer": "storage", "elementType": "vec2<f32>" },
+    "partialsR": { "name": "partials", "buffer": "read-only-storage", "elementType": "vec2<f32>" },
+    "statsW": { "name": "stats", "buffer": "storage", "elementType": "vec2<f32>" },
+    "statsR": { "name": "stats", "buffer": "read-only-storage", "elementType": "vec2<f32>" },
+    "yFused": { "arg": "Y", "name": "Y", "elementType": "$statsElement" },
+    "skipIn": { "arg": "skip", "name": "skip", "elementType": "$skipElement" },
+    "biasIn": { "arg": "bias", "name": "bias", "elementType": "$tScalar" },
+    "sumDstS": { "arg": "S", "name": "sum_dst", "elementType": "$statsElement" },
+    "sumDstY": { "arg": "Y", "name": "sum_dst", "elementType": "$statsElement" },
+    "srcCombineS": { "arg": "S", "name": "src", "buffer": "read-only-storage", "elementType": "$statsElement" },
+    "srcCombineY": { "arg": "Y", "name": "src", "buffer": "read-only-storage", "elementType": "$statsElement" },
+    "srcApplyS": { "arg": "S", "name": "src", "buffer": "read-only-storage", "elementType": "$applyElement" }
+  },
+  "variants": [
+    {
+      "id": "split_bias_sum",
+      "priority": 20,
+      "when": ["splitFits", "present.bias", "present.S"],
+      "demoteWhen": ["elementCount < tunables.SPLIT_MIN_ELEMENTS * statsGroupBlocks", "(not channelsLast) and normRows >= tunables.FUSED_MIN_ROWS and groupSize * dtypeBytes(tensorDtypes.X) <= tunables.FUSED_MAX_GROUP_BYTES"],
+      "derive": {
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "true",
+        "broadcastSkip": "broadcastSkipSpec",
+        "vecSpec": "statsVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "statsChunks",
+        "rowsPerStepSpec": "statsRowsPerStep",
+        "splitSpec": "splitCount",
+        "rowsPerSlabSpec": "ceilDiv(dimHW, slabCount)",
+        "planeVecsSpec": "planeVecs",
+        "chunkVecsSpec": "ceilDiv(planeVecs, planeSplits)",
+        "partialsPerRowSpec": "partialsPerRow",
+        "wgSpec": "statsWorkgroup",
+        "combineWgSpec": "combineWorkgroup",
+        "vecF32": "(\"vec\" ~ statsVecWidth ~ \"<f32>\") if statsVecWidth > 1 else \"f32\"",
+        "storeElement": "statsElement",
+        "applySrc": "\"src\"",
+        "groupUniform": "true"
+      },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[partialElements, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[statsElements, 2]" }
+      ],
+      "passes": [
+        {
+          "id": "stats",
+          "name": "GroupNorm.Partials",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"stats\"",
+            "unitsSpec": "workUnits",
+            "statsBlockedSpec": "channelsLast and statsGroupBlocks > 1",
+            "statsGroupsPerBlockSpec": "statsGroupsPerBlock",
+            "statsGroupBlocksSpec": "statsGroupBlocks",
+            "statsTileChunksSpec": "statsTileChunks"
+          },
+          "bindings": ["xIn", "skipIn", "biasIn", "sumDstS", "partialsW"],
+          "dispatch": { "x": "min(workUnits, 65535)", "y": "ceilDiv(workUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "combine",
+          "name": "GroupNorm.Combine",
+          "shader": "group-norm-combine.wgsl.jinja",
+          "bindings": ["srcCombineS", "partialsR", "statsW"],
+          "dispatch": {
+            "x": "min(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "y": "ceilDiv(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "apply",
+          "name": "GroupNorm.Apply",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"apply\"",
+            "vecSpec": "applyVecWidth",
+            "chunksSpec": "applyChunks",
+            "rowsPerStepSpec": "applyRowsPerStep",
+            "wgSpec": "applyWorkgroup",
+            "splitSpec": "applySplit",
+            "rowsPerSlabSpec": "ceilDiv(dimHW, applySlabCount)",
+            "groupUniform": "(channelsPerGroup % applyVecWidth == 0) if channelsLast else true",
+            "storeElement": "applyElement",
+            "vecF32": "\"vec4<f32>\" if applyVec4Ok else \"f32\"",
+            "unitsSpec": "applyUnits",
+            "applySrc": "\"src\""
+          },
+          "bindings": ["srcApplyS", "gamma", "beta", "statsR", "yOut"],
+          "dispatch": { "x": "min(applyUnits, 65535)", "y": "ceilDiv(applyUnits, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "fused_bias_sum",
+      "priority": 10,
+      "when": ["fusedFits", "present.bias", "present.S"],
+      "derive": {
+        "statsElement": "fusedElement",
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "true",
+        "broadcastSkip": "broadcastSkipSpec",
+        "useVec4": "statsVec4Ok",
+        "vecSpec": "fusedVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "fusedRowChunks",
+        "fusedChunksSpec": "fusedChunks",
+        "fusedRowsSpec": "max(1, floor(fusedWorkgroup / fusedChunks))",
+        "spanVecsSpec": "max(1, floor(groupSize / fusedVecWidth))",
+        "wgSpec": "fusedWorkgroup",
+        "ldsSpec": "fusedLanes",
+        "vecF32": "\"vec4<f32>\" if statsVec4Ok else \"f32\"",
+        "storeElement": "fusedElement",
+        "fusedSrc": "\"sum_dst\"",
+        "planeVecsSpec": "planeVecs",
+        "skipElement": "tScalar if broadcastSkipSpec else fusedElement"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNorm.Fused",
+          "shader": "group-norm-fused.wgsl.jinja",
+          "bindings": ["xIn", "skipIn", "biasIn", "gamma", "beta", "sumDstS", "yFused"],
+          "dispatch": { "x": "min(normRows, 65535)", "y": "ceilDiv(normRows, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "split_bias_plain",
+      "priority": 20,
+      "when": ["splitFits", "present.bias", "not present.S"],
+      "demoteWhen": ["elementCount < tunables.SPLIT_MIN_ELEMENTS * statsGroupBlocks", "(not channelsLast) and normRows >= tunables.FUSED_MIN_ROWS and groupSize * dtypeBytes(tensorDtypes.X) <= tunables.FUSED_MAX_GROUP_BYTES"],
+      "derive": {
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "true",
+        "broadcastSkip": "broadcastSkipSpec",
+        "vecSpec": "statsVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "statsChunks",
+        "rowsPerStepSpec": "statsRowsPerStep",
+        "splitSpec": "splitCount",
+        "rowsPerSlabSpec": "ceilDiv(dimHW, slabCount)",
+        "planeVecsSpec": "planeVecs",
+        "chunkVecsSpec": "ceilDiv(planeVecs, planeSplits)",
+        "partialsPerRowSpec": "partialsPerRow",
+        "wgSpec": "statsWorkgroup",
+        "combineWgSpec": "combineWorkgroup",
+        "vecF32": "(\"vec\" ~ statsVecWidth ~ \"<f32>\") if statsVecWidth > 1 else \"f32\"",
+        "storeElement": "statsElement",
+        "applySrc": "\"Y\"",
+        "groupUniform": "true"
+      },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[partialElements, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[statsElements, 2]" }
+      ],
+      "passes": [
+        {
+          "id": "stats",
+          "name": "GroupNorm.Partials",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"stats\"",
+            "unitsSpec": "workUnits",
+            "statsBlockedSpec": "channelsLast and statsGroupBlocks > 1",
+            "statsGroupsPerBlockSpec": "statsGroupsPerBlock",
+            "statsGroupBlocksSpec": "statsGroupBlocks",
+            "statsTileChunksSpec": "statsTileChunks"
+          },
+          "bindings": ["xIn", "skipIn", "biasIn", "sumDstY", "partialsW"],
+          "dispatch": { "x": "min(workUnits, 65535)", "y": "ceilDiv(workUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "combine",
+          "name": "GroupNorm.Combine",
+          "shader": "group-norm-combine.wgsl.jinja",
+          "bindings": ["srcCombineY", "partialsR", "statsW"],
+          "dispatch": {
+            "x": "min(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "y": "ceilDiv(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "apply",
+          "name": "GroupNorm.Apply",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"apply\"",
+            "vecSpec": "applyVecWidth",
+            "chunksSpec": "applyChunks",
+            "rowsPerStepSpec": "applyRowsPerStep",
+            "wgSpec": "applyWorkgroup",
+            "splitSpec": "applySplit",
+            "rowsPerSlabSpec": "ceilDiv(dimHW, applySlabCount)",
+            "groupUniform": "(channelsPerGroup % applyVecWidth == 0) if channelsLast else true",
+            "storeElement": "applyElement",
+            "vecF32": "\"vec4<f32>\" if applyVec4Ok else \"f32\"",
+            "unitsSpec": "applyUnits",
+            "applySrc": "\"Y\""
+          },
+          "bindings": ["gamma", "beta", "statsR", "yOut"],
+          "dispatch": { "x": "min(applyUnits, 65535)", "y": "ceilDiv(applyUnits, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "fused_bias_plain",
+      "priority": 10,
+      "when": ["fusedFits", "present.bias", "not present.S"],
+      "derive": {
+        "statsElement": "fusedElement",
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "true",
+        "broadcastSkip": "broadcastSkipSpec",
+        "useVec4": "statsVec4Ok",
+        "vecSpec": "fusedVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "fusedRowChunks",
+        "fusedChunksSpec": "fusedChunks",
+        "fusedRowsSpec": "max(1, floor(fusedWorkgroup / fusedChunks))",
+        "spanVecsSpec": "max(1, floor(groupSize / fusedVecWidth))",
+        "wgSpec": "fusedWorkgroup",
+        "ldsSpec": "fusedLanes",
+        "vecF32": "\"vec4<f32>\" if statsVec4Ok else \"f32\"",
+        "storeElement": "fusedElement",
+        "fusedSrc": "\"Y\"",
+        "planeVecsSpec": "planeVecs",
+        "skipElement": "tScalar if broadcastSkipSpec else fusedElement"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNorm.Fused",
+          "shader": "group-norm-fused.wgsl.jinja",
+          "bindings": ["xIn", "skipIn", "biasIn", "gamma", "beta", "yFused"],
+          "dispatch": { "x": "min(normRows, 65535)", "y": "ceilDiv(normRows, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "split_sum",
+      "priority": 20,
+      "when": ["splitFits", "not present.bias", "present.S"],
+      "demoteWhen": ["elementCount < tunables.SPLIT_MIN_ELEMENTS * statsGroupBlocks", "(not channelsLast) and normRows >= tunables.FUSED_MIN_ROWS and groupSize * dtypeBytes(tensorDtypes.X) <= tunables.FUSED_MAX_GROUP_BYTES"],
+      "derive": {
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "false",
+        "broadcastSkip": "broadcastSkipSpec",
+        "vecSpec": "statsVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "statsChunks",
+        "rowsPerStepSpec": "statsRowsPerStep",
+        "splitSpec": "splitCount",
+        "rowsPerSlabSpec": "ceilDiv(dimHW, slabCount)",
+        "planeVecsSpec": "planeVecs",
+        "chunkVecsSpec": "ceilDiv(planeVecs, planeSplits)",
+        "partialsPerRowSpec": "partialsPerRow",
+        "wgSpec": "statsWorkgroup",
+        "combineWgSpec": "combineWorkgroup",
+        "vecF32": "(\"vec\" ~ statsVecWidth ~ \"<f32>\") if statsVecWidth > 1 else \"f32\"",
+        "storeElement": "statsElement",
+        "applySrc": "\"src\"",
+        "groupUniform": "true"
+      },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[partialElements, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[statsElements, 2]" }
+      ],
+      "passes": [
+        {
+          "id": "stats",
+          "name": "GroupNorm.Partials",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"stats\"",
+            "unitsSpec": "workUnits",
+            "statsBlockedSpec": "channelsLast and statsGroupBlocks > 1",
+            "statsGroupsPerBlockSpec": "statsGroupsPerBlock",
+            "statsGroupBlocksSpec": "statsGroupBlocks",
+            "statsTileChunksSpec": "statsTileChunks"
+          },
+          "bindings": ["xIn", "skipIn", "sumDstS", "partialsW"],
+          "dispatch": { "x": "min(workUnits, 65535)", "y": "ceilDiv(workUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "combine",
+          "name": "GroupNorm.Combine",
+          "shader": "group-norm-combine.wgsl.jinja",
+          "bindings": ["srcCombineS", "partialsR", "statsW"],
+          "dispatch": {
+            "x": "min(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "y": "ceilDiv(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "apply",
+          "name": "GroupNorm.Apply",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"apply\"",
+            "vecSpec": "applyVecWidth",
+            "chunksSpec": "applyChunks",
+            "rowsPerStepSpec": "applyRowsPerStep",
+            "wgSpec": "applyWorkgroup",
+            "splitSpec": "applySplit",
+            "rowsPerSlabSpec": "ceilDiv(dimHW, applySlabCount)",
+            "groupUniform": "(channelsPerGroup % applyVecWidth == 0) if channelsLast else true",
+            "storeElement": "applyElement",
+            "vecF32": "\"vec4<f32>\" if applyVec4Ok else \"f32\"",
+            "unitsSpec": "applyUnits",
+            "applySrc": "\"src\""
+          },
+          "bindings": ["srcApplyS", "gamma", "beta", "statsR", "yOut"],
+          "dispatch": { "x": "min(applyUnits, 65535)", "y": "ceilDiv(applyUnits, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "fused_sum",
+      "priority": 10,
+      "when": ["fusedFits", "not present.bias", "present.S"],
+      "derive": {
+        "statsElement": "fusedElement",
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "false",
+        "broadcastSkip": "broadcastSkipSpec",
+        "useVec4": "statsVec4Ok",
+        "vecSpec": "fusedVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "fusedRowChunks",
+        "fusedChunksSpec": "fusedChunks",
+        "fusedRowsSpec": "max(1, floor(fusedWorkgroup / fusedChunks))",
+        "spanVecsSpec": "max(1, floor(groupSize / fusedVecWidth))",
+        "wgSpec": "fusedWorkgroup",
+        "ldsSpec": "fusedLanes",
+        "vecF32": "\"vec4<f32>\" if statsVec4Ok else \"f32\"",
+        "storeElement": "fusedElement",
+        "fusedSrc": "\"sum_dst\"",
+        "planeVecsSpec": "planeVecs",
+        "skipElement": "tScalar if broadcastSkipSpec else fusedElement"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNorm.Fused",
+          "shader": "group-norm-fused.wgsl.jinja",
+          "bindings": ["xIn", "skipIn", "gamma", "beta", "sumDstS", "yFused"],
+          "dispatch": { "x": "min(normRows, 65535)", "y": "ceilDiv(normRows, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "split_plain",
+      "priority": 20,
+      "when": ["splitFits", "not present.bias", "not present.S"],
+      "demoteWhen": ["elementCount < tunables.SPLIT_MIN_ELEMENTS * statsGroupBlocks", "(not channelsLast) and normRows >= tunables.FUSED_MIN_ROWS and groupSize * dtypeBytes(tensorDtypes.X) <= tunables.FUSED_MAX_GROUP_BYTES"],
+      "derive": {
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "false",
+        "broadcastSkip": "broadcastSkipSpec",
+        "vecSpec": "statsVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "statsChunks",
+        "rowsPerStepSpec": "statsRowsPerStep",
+        "splitSpec": "splitCount",
+        "rowsPerSlabSpec": "ceilDiv(dimHW, slabCount)",
+        "planeVecsSpec": "planeVecs",
+        "chunkVecsSpec": "ceilDiv(planeVecs, planeSplits)",
+        "partialsPerRowSpec": "partialsPerRow",
+        "wgSpec": "statsWorkgroup",
+        "combineWgSpec": "combineWorkgroup",
+        "vecF32": "(\"vec\" ~ statsVecWidth ~ \"<f32>\") if statsVecWidth > 1 else \"f32\"",
+        "storeElement": "statsElement",
+        "applySrc": "\"Y\"",
+        "groupUniform": "true"
+      },
+      "intermediates": [
+        { "id": "partials", "dtype": "float32", "shape": "[partialElements, 2]" },
+        { "id": "stats", "dtype": "float32", "shape": "[statsElements, 2]" }
+      ],
+      "passes": [
+        {
+          "id": "stats",
+          "name": "GroupNorm.Partials",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"stats\"",
+            "unitsSpec": "workUnits",
+            "statsBlockedSpec": "channelsLast and statsGroupBlocks > 1",
+            "statsGroupsPerBlockSpec": "statsGroupsPerBlock",
+            "statsGroupBlocksSpec": "statsGroupBlocks",
+            "statsTileChunksSpec": "statsTileChunks"
+          },
+          "bindings": ["xIn", "skipIn", "sumDstY", "partialsW"],
+          "dispatch": { "x": "min(workUnits, 65535)", "y": "ceilDiv(workUnits, 65535)", "z": 1 }
+        },
+        {
+          "id": "combine",
+          "name": "GroupNorm.Combine",
+          "shader": "group-norm-combine.wgsl.jinja",
+          "bindings": ["srcCombineY", "partialsR", "statsW"],
+          "dispatch": {
+            "x": "min(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "y": "ceilDiv(ceilDiv((normRows), (combineWorkgroup)), 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "apply",
+          "name": "GroupNorm.Apply",
+          "shader": "group-norm-slab.wgsl.jinja",
+          "derive": {
+            "phase": "\"apply\"",
+            "vecSpec": "applyVecWidth",
+            "chunksSpec": "applyChunks",
+            "rowsPerStepSpec": "applyRowsPerStep",
+            "wgSpec": "applyWorkgroup",
+            "splitSpec": "applySplit",
+            "rowsPerSlabSpec": "ceilDiv(dimHW, applySlabCount)",
+            "groupUniform": "(channelsPerGroup % applyVecWidth == 0) if channelsLast else true",
+            "storeElement": "applyElement",
+            "vecF32": "\"vec4<f32>\" if applyVec4Ok else \"f32\"",
+            "unitsSpec": "applyUnits",
+            "applySrc": "\"Y\""
+          },
+          "bindings": ["gamma", "beta", "statsR", "yOut"],
+          "dispatch": { "x": "min(applyUnits, 65535)", "y": "ceilDiv(applyUnits, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "fused_plain",
+      "priority": 10,
+      "when": ["fusedFits", "not present.bias", "not present.S"],
+      "derive": {
+        "statsElement": "fusedElement",
+        "phaseLayout": "channelsLast",
+        "hasSkip": "true",
+        "hasBias": "false",
+        "broadcastSkip": "broadcastSkipSpec",
+        "useVec4": "statsVec4Ok",
+        "vecSpec": "fusedVecWidth",
+        "dimCSpec": "dimC",
+        "dimHWSpec": "dimHW",
+        "cpgSpec": "channelsPerGroup",
+        "groupsSpec": "numGroups",
+        "groupSizeSpec": "groupSize",
+        "normRowsSpec": "normRows",
+        "epsilonSpec": "attrs.epsilon",
+        "activationSpec": "activationKind",
+        "chunksSpec": "fusedRowChunks",
+        "fusedChunksSpec": "fusedChunks",
+        "fusedRowsSpec": "max(1, floor(fusedWorkgroup / fusedChunks))",
+        "spanVecsSpec": "max(1, floor(groupSize / fusedVecWidth))",
+        "wgSpec": "fusedWorkgroup",
+        "ldsSpec": "fusedLanes",
+        "vecF32": "\"vec4<f32>\" if statsVec4Ok else \"f32\"",
+        "storeElement": "fusedElement",
+        "fusedSrc": "\"Y\"",
+        "planeVecsSpec": "planeVecs",
+        "skipElement": "tScalar if broadcastSkipSpec else fusedElement"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "GroupNorm.Fused",
+          "shader": "group-norm-fused.wgsl.jinja",
+          "bindings": ["xIn", "skipIn", "gamma", "beta", "yFused"],
+          "dispatch": { "x": "min(normRows, 65535)", "y": "ceilDiv(normRows, 65535)", "z": 1 }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,32 @@

+{
+  "name": "com.microsoft.SkipGroupNorm",
+  "id": "_com_microsoft_skipgroupnorm_webgpu_92ea3fc",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "VviePAedFI1tXjO4aRxUsDVvypknA6GNCHfB/ySq07c=",
+      "group-norm-combine.wgsl.jinja": "weqcoPzdY5lQ8K/lZVEF03/sdbu1c/2jwdrW/KT9xr4=",
+      "group-norm-fused.wgsl.jinja": "523MK9Q6GT26VjwUdiq/KsvpQdE6BH38iuS9ftCxDeY=",
+      "group-norm-slab.wgsl.jinja": "gsFbtBGWwxpdm/76plQtm7psUEN6HoFrM05U0y+LwBM=",
+      "manifest.json": "KRDZNaGuWK334a4ewL0oA0L3/sAdyt9pwqbJb8Ms4O8=",
+      "test.json": "w/W+fF5md2sOc+Q6Db/HPoF8z1qw48sJpvma6ph27Bs="
+    }
+  },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.1",
+    "variants": {
+      "split_bias_sum": ["group-norm-combine.wgsl.jinja", "group-norm-slab.wgsl.jinja"],
+      "fused_bias_sum": ["group-norm-fused.wgsl.jinja"],
+      "split_bias_plain": ["group-norm-combine.wgsl.jinja", "group-norm-slab.wgsl.jinja"],
+      "fused_bias_plain": ["group-norm-fused.wgsl.jinja"],
+      "split_sum": ["group-norm-combine.wgsl.jinja", "group-norm-slab.wgsl.jinja"],
+      "fused_sum": ["group-norm-fused.wgsl.jinja"],
+      "split_plain": ["group-norm-combine.wgsl.jinja", "group-norm-slab.wgsl.jinja"],
+      "fused_plain": ["group-norm-fused.wgsl.jinja"]
+    }
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1664 @@

+{
+  "fixtureArrays": {
+    "wb_input": [-0.768555, 1.575195, -0.698242, 1.587891, 0.371826, -0.280029, -1.328125, 0.127197, -0.197144, 0.982422, -0.671387, -1.925781, 1.800781, -0.020218, -0.782227, 1.291992, -0.935059, 1.782227, -0.674316, -1.943359, -0.218994, 0.054138, -1.539062, -0.546387, -2.160156, 1.195312, 1.65332, -0.674316, 0.224731, -0.093262, 1.160156, -0.389404, 1.748047, 0.766113, 0.234375, 0.011177, -0.055847, -0.930664, -0.490234, -0.655762, -0.382568, -0.554688, 0.910645, -0.227295, 1.6875, 0.028397, -0.241699, -0.480957, -0.355713, -2.095703, -0.443359, -0.126221, -0.815918, 0.792969, -0.450439, -0.952148, -1.174805, 0.242798, 0.13855, -0.237061, -0.994141, 0.346436, 0.147705, 0.125854, -0.51709, 0.253906, 0.400146, -0.540039, -0.788574, 0.146606, -0.409668, 0.281982, 1.444336, 0.044434, -0.366699, 2.25, -0.453613, -0.652344, 1.828125, -0.244751, 0.307129, -0.051361, 0.106384, 0.844727, 1.648438, -0.904785, -0.35376, 0.510742, 0.074829, -0.311279, 0.274902, 1.594727, 1.367188, 0.098755, 0.043304, -0.207397, 0.068298, -0.601074, 0.083008, 0.264893, -0.65918, -0.216797, -0.086548, -0.683594, -0.964844, -2.591797, -0.817383, -0.461914, -1.84082, -0.712402, -0.052094, -0.583008, 1.114258, 0.190308, 1.087891, 0.005146, 1.041992, 1.363281, -0.273682, -0.465576, -0.027618, 1.345703, 0.789551, -0.015991, 0.401611, 0.726562, 0.598633, 0.133667],
+    "wb_gamma": [0.241255, 0.55666, -0.835532, 0.564596, -1.338308, -0.278924, 0.357326, -1.745484, 0.277184, 0.101415, -0.018637, -0.526188, -0.011698, -2.349411, 0.206578, 0.357679],
+    "wb_beta": [-1.194839, 0.209146, -0.677225, -0.547338, 1.275685, -1.099577, 0.470916, 0.293907, -1.094209, 2.350204, -1.633769, 0.248753, -0.180166, 0.365134, -0.555731, 1.843083],
+    "wb_skip": [0.892578, -0.471924, -0.423096, 1.277344, 0.25708, -1.366211, 1.552734, 0.441406, -0.033142, -0.059418, 1.536133, -0.225464, 1.472656, 0.591309, -0.38623, -2.197266, 0.0896, -0.256592, -1.873047, 0.916992, 0.39209, 0.015526, -0.949219, 0.566895, -0.220459, 1.262695, -0.437744, -2.283203, -0.264893, -0.660156, 2.353516, 1.992188, 0.865723, -0.854004, -1.014648, 0.899414, -1.041016, 1.378906, -0.075073, -2.541016, -0.883789, -0.428711, 0.981934, -0.072754, 2.214844, 0.658203, 0.170166, -1.727539, -0.672363, -1.373047, 0.318115, 0.422363, 0.260742, -0.547852, 0.545898, -0.155762, 0.679688, 2.861328, -0.300781, -0.504883, 1.548828, 0.35376, -0.387695, -1.595703, -0.170166, -0.002897, 0.273193, -0.383545, -1.082031, -0.894043, -1.048828, -0.044708, 0.049286, 0.220215, 0.272705, -0.853027, -0.489258, 0.513672, 0.977051, 0.310547, -0.577148, -0.479004, 0.838867, 0.872559, -0.510254, 0.101807, -0.299805, -1.179688, -1.555664, 0.668457, 0.939453, 0.118103, -0.376709, 0.735352, -0.214233, -1.987305, -0.931152, 1.268555, 1.427734, -0.757812, -1.324219, 0.375488, 1.364258, -1.708008, 0.976562, -0.037659, -1.779297, -0.196655, 1.636719, 0.69043, 0.941895, -1.882812, 0.431641, 0.203857, 1.306641, -0.126343, 1.408203, 1.188477, 0.432861, -2.296875, -0.475342, 1.517578, -0.824219, 1.288086, -0.028244, 1.918945, 0.352295, 0.693359],
+    "wb_bias": [-0.537598, 0.500488, -0.252441, -0.460693, -1.640625, -1.298828, 0.331787, -1.588867, 1.000977, 1.458984, 0.702637, 0.147827, 1.143555, 0.533691, -0.07251, 0.51123],
+    "wb_norm": [-1.213867, 0.856445, -0.119141, 0.386475, 0.714355, -0.804688, 1.048828, -0.42627, -1.091797, 2.435547, -1.641602, 0.989746, -0.200928, 0.267334, -0.800781, 1.577148, -1.357422, 1.000977, 0.613281, -0.963867, 1.179688, -1.169922, 0.30835, 0.304199, -1.396484, 2.513672, -1.644531, 1.206055, -0.180664, 1.896484, -0.294678, 2.046875, -0.844238, 0.448486, -0.294189, -0.291504, 2.480469, -1.250977, 0.833008, 4.59375, -1.238281, 2.335938, -1.651367, 0.491943, -0.204834, 0.12561, -0.682129, 1.333984, -1.384766, -0.708008, -0.630859, -0.504883, 1.924805, -1.208008, 1.013672, 1.80957, -1.128906, 2.546875, -1.631836, 0.61084, -0.184326, 0.110046, -0.700195, 1.47168, -1.511719, 0.492188, -0.847168, -1.373047, 2.837891, -0.998047, 0.521484, 0.262207, -0.810547, 2.400391, -1.628906, 0.049896, -0.174927, 1.076172, -0.252197, 1.78418, -1.418945, 0.09082, -1.056641, 0.002945, 0.627441, -0.989746, 0.679199, 1.130859, -1.371094, 2.408203, -1.645508, -0.062988, -0.192017, -0.655762, -0.718262, 1.170898, -1.550781, 0.706055, -1.492188, -1.148438, 2.921875, -1.136719, 1.058594, 2.78125, -1.089844, 2.201172, -1.597656, 0.785645, -0.181396, 0.868164, -0.552246, 1.097656, -1.015625, 0.56543, -2.173828, -0.955078, -0.336426, -1.503906, 0.838867, 3.136719, -1.186523, 2.580078, -1.629883, 0.094604, -0.186523, -3.884766, -0.54248, 1.990234],
+    "wb_add_out": [-0.414062, 1.604492, -1.374023, 2.404297, -1.011719, -2.945312, 0.556641, -1.020508, 0.770508, 2.382812, 1.567383, -2.003906, 4.417969, 1.105469, -1.240234, -0.394531, -1.382812, 2.027344, -2.800781, -1.487305, -1.466797, -1.229492, -2.15625, -1.568359, -1.379883, 3.917969, 1.917969, -2.808594, 1.103516, -0.219727, 3.441406, 2.113281, 2.076172, 0.412598, -1.033203, 0.449951, -2.738281, -0.851562, -0.233521, -4.785156, -0.265625, 0.475586, 2.595703, -0.152222, 5.046875, 1.220703, -0.144043, -1.697266, -1.566406, -2.96875, -0.377686, -0.164551, -2.195312, -1.053711, 0.427246, -2.697266, 0.505859, 4.5625, 0.540527, -0.594238, 1.698242, 1.233398, -0.3125, -0.958496, -1.224609, 0.751465, 0.420898, -1.384766, -3.511719, -2.046875, -1.126953, -1.351562, 2.494141, 1.724609, 0.608398, 1.544922, 0.200684, 0.39502, 2.732422, 0.577148, -0.807617, -0.029785, 0.692871, 1.256836, -0.502441, -2.101562, -0.321777, -2.257812, -0.479492, 1.816406, 1.916992, 1.860352, 2.134766, 1.367188, -0.243408, -1.683594, -1.400391, 1.167969, 1.257812, -0.953613, -3.625, -1.140625, 1.609375, -3.980469, 1.012695, -1.170898, -1.894531, -0.510742, 0.939453, 0.511719, 0.817383, -1.955078, 1.007812, 0.894531, 2.142578, -0.582031, 0.80957, 1.25293, 0.490967, -4.351562, 0.497803, 4.320312, 0.667969, 1.419922, 1.516602, 3.179688, 0.878906, 1.337891],
+    "bc_input": [0.588867, 0.896484, -0.213623, 0.803223, 0.65918, -0.216187, 1.197266, -0.486084, -0.71875, 0.332031, -0.364746, -0.831543, -0.031219, -1.05957, 0.161621, 1.519531, 0.169312, 1.048828, 1.330078, 0.450195, -2.867188, -1.456055, 0.708496, -1.120117, -1.208984, -1.199219, -1.505859, -0.549316, 0.505371, 0.723145, -0.359131, -0.250977, -0.879883, -0.305664, 0.709473, 0.81543, 0.617676, -0.638672, 0.066772, -2.330078, -1.316406, 1.744141, 1.12207, -0.633789, -1.802734, -0.825684, 0.622559, -0.481689, -1.364258, -0.536621, -0.464111, 0.247437, -0.213989, 0.384521, 0.556641, -0.303711, -0.160034, 0.882324, -0.212036, -0.796387, 0.153076, -1.311523, 2.212891, 0.685059],
+    "bc_gamma": [0.789682, 0.869051, -0.010169, -0.021685, 0.506611, 1.267444, -0.312695, 0.877844, 0.598637, 0.598314, -1.721544, -0.593328, 0.986705, -0.419391, -0.852584, -0.572351, 0.912797, -0.586863, 0.477761, -0.484418, -0.193835, 0.347757, 0.327637, -1.100304, 1.233108, -0.272569, -0.688656, 0.687245, 0.398386, 0.888089, -0.792587, -0.769029, -0.427778, 0.100768, -2.18706, 1.279301, 1.109054, 0.375992, 1.514775, 1.271436, 0.822896, -0.47675, 0.475507, -1.011297, 1.177197, 1.58654, -1.059944, -0.145351, 0.841555, -2.014113, -0.230498, 0.302128, -0.180508, 0.980534, -0.126871, 0.203151, -0.754841, 0.42057, -1.085798, 1.335042, -0.67493, 2.453507, 2.139259, 1.087436],
+    "bc_beta": [-0.064518, -0.262683, 0.827528, -0.960938, 1.062519, 2.417941, 0.212789, -1.63843, 1.875453, -0.883058, -0.006704, 0.424894, -0.869972, 0.727008, 0.879303, -3.024141, -2.610873, 1.269641, 0.883006, 0.804167, -1.510324, 2.258091, -0.00675, -1.553668, -1.659453, 0.579603, 0.652358, 0.007077, 0.09918, 0.418658, -0.273778, -1.036199, -1.128691, -0.296022, -0.224056, 1.476306, 0.577624, -0.372049, -0.581659, -1.841807, -0.361721, 0.05116, -0.749332, -2.634807, 0.562719, -0.738667, 0.024864, -1.135937, -1.368144, -1.458886, -0.946683, 1.953936, -1.198661, 0.166648, 0.447206, -0.45814, -0.553395, 0.1129, 0.255989, -0.184551, 1.254163, -0.260479, -1.232429, 1.902575],
+    "bc_skip": [0.952148, 1.342773, -0.172974, -0.395264, 1.119141, 0.330566, 0.281494, 0.4729, -0.692871, -0.634766, 0.013504, -1.866211, -0.428223, 0.669922, -0.323486, 0.713867, -0.350586, 0.65918, -0.288574, 0.324219, -0.300781, -0.789551, -0.216431, -0.221436, -0.08667, 0.366211, -0.643555, -0.977051, 0.001021, 0.415527, -0.271729, 0.836426, 0.03537, -0.806152, 0.936035, -0.021332, -1.095703, 0.97168, 1.648438, 0.84082, 0.837402, 0.60791, -1.894531, 0.666016, -0.171143, 1.625977, -0.620117, -0.039581, 1.702148, -2.410156, 1.56543, -0.756348, 1.446289, 0.583496, -0.497559, -0.271729, -0.956055, -1.642578, 0.833496, -1.136719, 1.248047, -2.515625, 0.080383, 0.376221],
+    "bc_norm": [0.494873, 1.017578, 0.841797, -0.949219, 1.552734, 1.333984, 0.012703, -2.511719, 1.424805, -0.818359, -0.128418, 1.462891, -0.882812, 0.709961, 0.693848, -4.210938, -2.505859, 0.513184, 1.300781, 0.460938, -1.172852, 1.851562, 0.167969, -0.885254, -2.535156, 0.656738, 1.683594, -0.627441, 0.478271, 1.782227, -0.196777, -1.824219, -0.791016, -0.398682, -3.197266, 2.275391, 0.052704, -0.286865, 1.567383, -3.552734, -0.646973, -0.927734, -1.032227, -2.722656, -1.337891, 0.432129, -0.040253, -1.080078, -1.118164, 3.123047, -1.15332, 1.84375, -1.378906, 0.941406, 0.437256, -0.542969, -0.218872, 0.006115, -0.265869, -1.356445, 0.649902, -4.882812, 1.696289, 2.679688],
+    "bc_add_out": [1.541016, 2.238281, -0.386719, 0.407959, 1.77832, 0.11438, 1.478516, -0.013184, -1.412109, -0.302734, -0.351318, -2.697266, -0.459473, -0.389648, -0.161865, 2.234375, -0.181274, 1.708008, 1.041016, 0.774414, -3.167969, -2.246094, 0.492188, -1.341797, -1.295898, -0.833008, -2.148438, -1.526367, 0.506348, 1.138672, -0.630859, 0.585449, -0.844727, -1.111328, 1.645508, 0.793945, -0.478027, 0.333008, 1.714844, -1.489258, -0.479004, 2.351562, -0.772461, 0.032227, -1.973633, 0.800293, 0.002441, -0.521484, 0.337891, -2.947266, 1.101562, -0.508789, 1.232422, 0.967773, 0.059082, -0.575195, -1.116211, -0.760254, 0.621582, -1.933594, 1.401367, -3.828125, 2.292969, 1.061523],
+    "nhwc_group_block_tail_c320_float32_input_X": [-0.5, -0.4375, -0.375, -0.3125, -0.25, -0.1875, -0.125, -0.0625, 0, 0.0625, 0.125, 0.1875, 0.25, 0.3125, 0.375, 0.4375, 0.5]
+  },
+  "cases": [
+    {
+      "name": "ort_skip_group_norm_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_with_bias",
+        "note": "explicit channels_last = 1"
+      },
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_skip" } }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_bias" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        },
+        "S": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_add_out" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_with_bias_default_layout",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_with_bias",
+        "note": "channels_last left at its default, the gtest's channels_last = -1 instantiation"
+      },
+      "attrs": { "groups": 4, "activation": 0, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_skip" } }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_bias" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        },
+        "S": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_add_out" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_with_bias_no_sum_output",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_with_bias",
+        "note": "the same table with the optional S output not requested; Y is unchanged because the sum is round-tripped at the tensor type either way"
+      },
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_skip" } }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_bias" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 16],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/wb_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_broadcast_skip_rank2_with_sum",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_no_bias_broadcast_skip",
+        "note": "skip spelled (N, C), which broadcasts; S requested"
+      },
+      "attrs": { "groups": 8, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_skip" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        },
+        "S": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_add_out" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_broadcast_skip_rank2_no_sum",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_no_bias_broadcast_skip",
+        "note": "skip spelled (N, C), which broadcasts; S not requested"
+      },
+      "attrs": { "groups": 8, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_skip" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_broadcast_skip_rank4_with_sum",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_no_bias_broadcast_skip",
+        "note": "skip spelled (N, 1, 1, C), which at H = W = 1 is an exact shape match and takes the non-broadcast route; S requested"
+      },
+      "attrs": { "groups": 8, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_skip" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        },
+        "S": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_add_out" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_broadcast_skip_rank4_no_sum",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skip_group_norm_op_test.cc",
+        "test": "SkipGroupNormTest.SkipGroupNorm_no_bias_broadcast_skip",
+        "note": "skip spelled (N, 1, 1, C), which at H = W = 1 is an exact shape match and takes the non-broadcast route; S not requested"
+      },
+      "attrs": { "groups": 8, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_input" } }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_gamma" } }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_beta" } }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_skip" } }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float16",
+          "shape": [1, 1, 1, 64],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/bc_norm" } },
+          "tolerance": 0.02,
+          "relTolerance": 0.004
+        }
+      }
+    },
+    {
+      "name": "nhwc_vec4_bias_sum_silu_f16",
+      "attrs": { "groups": 8, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 32, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.28 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 32, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.42, "cosStep": 0.16 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 32, 32, 64], "tolerance": 0.01, "relTolerance": 0.004 },
+        "S": { "dtype": "float16", "shape": [1, 32, 32, 64], "tolerance": 0.01, "relTolerance": 0.004 }
+      }
+    },
+    {
+      "name": "nhwc_vec4_bias_nosum_f32",
+      "attrs": { "groups": 8, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 32, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.30000000000000004 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 32, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.44, "cosStep": 0.18 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 32, 32, 64], "tolerance": 0.0005, "relTolerance": 0.0005 } }
+    },
+    {
+      "name": "nhwc_vec4_nobias_sum_f32",
+      "attrs": { "groups": 4, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 8, 8, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.34 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 8, 8, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.22000000000000003 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 8, 8, 32], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 8, 8, 32], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_vec4_nobias_nosum_f16",
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 8, 8, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.34 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 8, 8, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.48, "cosStep": 0.22000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [2, 8, 8, 32], "tolerance": 0.01, "relTolerance": 0.004 } }
+    },
+    {
+      "name": "nhwc_scalar_odd_cpg_bias_sum",
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 3, 3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 3, 3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 3, 3, 96], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 3, 3, 96], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_scalar_odd_cpg_nobias_nosum_f16",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [2, 3, 3, 96], "tolerance": 0.01, "relTolerance": 0.004 } }
+    },
+    {
+      "name": "nhwc_broadcast_skip_rank2_bias_sum",
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 4, 4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.4 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.54, "cosStep": 0.28 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 4, 4, 32], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 4, 4, 32], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_broadcast_skip_rank4_nobias_sum_f16",
+      "attrs": { "groups": 4, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 4, 4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.4 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 1, 1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.54, "cosStep": 0.28 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 4, 4, 32], "tolerance": 0.01, "relTolerance": 0.004 },
+        "S": { "dtype": "float16", "shape": [2, 4, 4, 32], "tolerance": 0.01, "relTolerance": 0.004 }
+      }
+    },
+    {
+      "name": "nhwc_broadcast_skip_scalar_cpg_nosum",
+      "attrs": { "groups": 4, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 5, 5, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.45999999999999996 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6, "cosStep": 0.33999999999999997 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 5, 5, 12], "tolerance": 0.0005, "relTolerance": 0.0005 } }
+    },
+    {
+      "name": "nchw_vec4_bias_sum_f32",
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 16, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.48 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 16, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.62, "cosStep": 0.36 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 16, 32, 32], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 16, 32, 32], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nchw_vec4_nobias_sum_silu_f16",
+      "attrs": { "groups": 4, "activation": 1, "channels_last": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 16, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.48 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 16, 32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.62, "cosStep": 0.36 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 16, 32, 32], "tolerance": 0.01, "relTolerance": 0.004 },
+        "S": { "dtype": "float16", "shape": [1, 16, 32, 32], "tolerance": 0.01, "relTolerance": 0.004 }
+      }
+    },
+    {
+      "name": "nchw_scalar_hw9_bias_nosum",
+      "attrs": { "groups": 3, "activation": 1, "channels_last": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 6, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.54 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 6, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.6799999999999999, "cosStep": 0.42 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [2, 6, 3, 3], "tolerance": 0.0005, "relTolerance": 0.0005 } }
+    },
+    {
+      "name": "nchw_broadcast_skip_rank2_bias_sum",
+      "attrs": { "groups": 2, "activation": 0, "channels_last": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 8, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.58 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.72, "cosStep": 0.45999999999999996 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 8, 4, 4], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 8, 4, 4], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nchw_broadcast_skip_rank4_nobias_nosum_f16",
+      "attrs": { "groups": 2, "activation": 1, "channels_last": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 8, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.58 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 1, 1, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.72, "cosStep": 0.45999999999999996 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [2, 8, 4, 4], "tolerance": 0.01, "relTolerance": 0.004 } }
+    },
+    {
+      "name": "nhwc_f16_weights_bias_sum",
+      "attrs": { "groups": 2, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 8, 8, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.43, "cosStep": 0.6 }
+        },
+        "gamma": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 8, 8, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.74, "cosStep": 0.48 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 8, 8, 16], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 8, 8, 16], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_all_f16_nobias_sum",
+      "attrs": { "groups": 2, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 8, 8, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.47, "cosStep": 0.64 }
+        },
+        "gamma": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 8, 8, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.78, "cosStep": 0.52 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 8, 8, 16], "tolerance": 0.01, "relTolerance": 0.004 },
+        "S": { "dtype": "float16", "shape": [1, 8, 8, 16], "tolerance": 0.01, "relTolerance": 0.004 }
+      }
+    },
+    {
+      "name": "nhwc_vec4_large_epsilon_bias_nosum",
+      "attrs": { "groups": 2, "activation": 0, "channels_last": 1, "epsilon": 0.5 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 4, 4, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.7000000000000001 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 4, 4, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8400000000000001, "cosStep": 0.5800000000000001 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 4, 4, 8], "tolerance": 0.0005, "relTolerance": 0.0005 } }
+    },
+    {
+      "name": "nchw_multi_split_bias_sum_f32",
+      "attrs": { "groups": 4, "activation": 0, "channels_last": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 16, 96, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.59, "cosStep": 0.76 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 16, 96, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.8999999999999999, "cosStep": 0.64 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 16, 96, 96], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 16, 96, 96], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_multi_slab_nobias_sum_f16_silu",
+      "attrs": { "groups": 8, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 32, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.78 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 32, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.9199999999999999, "cosStep": 0.66 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 32, 32, 64], "tolerance": 0.01, "relTolerance": 0.004 },
+        "S": { "dtype": "float16", "shape": [1, 32, 32, 64], "tolerance": 0.01, "relTolerance": 0.004 }
+      }
+    },
+    {
+      "name": "nhwc_scalar_channels_not_multiple_of_four_bias_sum",
+      "attrs": { "groups": 3, "activation": 1, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 3, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.73, "cosStep": 0.9 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 3, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.17 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 3, 3, 6], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 3, 3, 6], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_scalar_channels_not_multiple_of_four_broadcast_nosum",
+      "attrs": { "groups": 5, "activation": 0, "channels_last": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 4, 4, 10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.79, "cosStep": 0.31 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.43, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [2, 4, 4, 10], "tolerance": 0.0005, "relTolerance": 0.0005 } }
+    },
+    {
+      "name": "nhwc_group_block_tail_c320_float32",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 320],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/nhwc_group_block_tail_c320_float32_input_X" }
+          }
+        },
+        "gamma": {
+          "dtype": "float16",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float16",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 320],
+          "data": {
+            "kind": "cycle",
+            "values": [-0.625, -0.5, -0.375, -0.25, -0.125, 0.0, 0.125, 0.25, 0.375, 0.5, 0.625]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 17, 19, 320], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 17, 19, 320], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_group_block_tail_c320_float16",
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 17, 19, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 17, 19, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 17, 19, 320], "tolerance": 0.01, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "nhwc_group_block_tail_c96_float32",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 96],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/nhwc_group_block_tail_c320_float32_input_X" }
+          }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 96],
+          "data": {
+            "kind": "cycle",
+            "values": [-0.625, -0.5, -0.375, -0.25, -0.125, 0.0, 0.125, 0.25, 0.375, 0.5, 0.625]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 17, 19, 96], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 17, 19, 96], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_group_block_tail_c384_float16",
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 17, 19, 384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [384],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 17, 19, 384], "tolerance": 0.01, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "nhwc_group_block_tail_c320_float32_nonexact_inputs",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float16",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float16",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [320],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 17, 19, 320], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 17, 19, 320], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_group_block_tail_c96_float32_nonexact_inputs",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 17, 19, 96], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 17, 19, 96], "tolerance": 0.0005, "relTolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "nhwc_pair_stats_c64_float32",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 17, 19, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 17, 19, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 17, 19, 64], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 17, 19, 64], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_pair_stats_c64_float16",
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "tunables": { "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 17, 19, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 17, 19, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 17, 19, 64], "tolerance": 0.01, "relTolerance": 0.002 },
+        "S": { "dtype": "float16", "shape": [2, 17, 19, 64], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_pair_stats_c192_float32",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 17, 19, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 17, 19, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 17, 19, 192], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 17, 19, 192], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_pair_stats_c192_float16",
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "tunables": { "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 17, 19, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.36 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 17, 19, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.5, "cosStep": 0.24 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 17, 19, 192], "tolerance": 0.01, "relTolerance": 0.002 },
+        "S": { "dtype": "float16", "shape": [2, 17, 19, 192], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_pair_stats_c448_float32",
+      "attrs": { "groups": 32, "activation": 0, "channels_last": 1 },
+      "tunables": { "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 17, 19, 448],
+          "data": { "kind": "cycle", "values": [-0.75, -0.5, -0.25, 0.0, 0.25, 0.5, 0.75] }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [448],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [448],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 448],
+          "data": { "kind": "cycle", "values": [-0.625, -0.375, -0.125, 0.125, 0.375, 0.625] }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [448],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [2, 17, 19, 448], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [2, 17, 19, 448], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_pair_stats_c448_float16",
+      "attrs": { "groups": 32, "activation": 1, "channels_last": 1 },
+      "tunables": { "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 17, 19, 448],
+          "data": { "kind": "cycle", "values": [-0.75, -0.5, -0.25, 0.0, 0.25, 0.5, 0.75] }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [448],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [448],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 1, 1, 448],
+          "data": { "kind": "cycle", "values": [-0.625, -0.375, -0.125, 0.125, 0.375, 0.625] }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [448],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 17, 19, 448], "tolerance": 0.01, "relTolerance": 0.002 },
+        "S": { "dtype": "float16", "shape": [2, 17, 19, 448], "tolerance": 0, "relTolerance": 0 }
+      }
+    },
+    {
+      "name": "nhwc_min_limit_group_tail_c288_g96_float32",
+      "attrs": { "groups": 96, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 288],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/nhwc_group_block_tail_c320_float32_input_X" }
+          }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [288],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [288],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 288],
+          "data": {
+            "kind": "cycle",
+            "values": [-0.625, -0.5, -0.375, -0.25, -0.125, 0.0, 0.125, 0.25, 0.375, 0.5, 0.625]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [288],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 17, 19, 288], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 17, 19, 288], "tolerance": 0, "relTolerance": 0 }
+      },
+      "provenance": {
+        "notes": "Complete-group statistics blocks with a partial final block under the standard 256-invocation WebGPU minimum tier. Odd groups retain scalar loads; even groups exercise vec2 blocks."
+      }
+    },
+    {
+      "name": "nhwc_min_limit_group_tail_c288_g96_float16",
+      "attrs": { "groups": 96, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 17, 19, 288],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/nhwc_group_block_tail_c320_float32_input_X" }
+          }
+        },
+        "gamma": {
+          "dtype": "float16",
+          "shape": [288],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float16",
+          "shape": [288],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 288],
+          "data": {
+            "kind": "cycle",
+            "values": [-0.625, -0.5, -0.375, -0.25, -0.125, 0.0, 0.125, 0.25, 0.375, 0.5, 0.625]
+          }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [288],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 17, 19, 288], "tolerance": 0.01, "relTolerance": 0.002 },
+        "S": { "dtype": "float16", "shape": [1, 17, 19, 288], "tolerance": 0, "relTolerance": 0 }
+      },
+      "provenance": {
+        "notes": "Complete-group statistics blocks with a partial final block under the standard 256-invocation WebGPU minimum tier. Odd groups retain scalar loads; even groups exercise vec2 blocks."
+      }
+    },
+    {
+      "name": "nhwc_min_limit_group_tail_c640_g64_float32",
+      "attrs": { "groups": 64, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 17, 19, 640],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/nhwc_group_block_tail_c320_float32_input_X" }
+          }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [640],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [640],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [1, 640],
+          "data": {
+            "kind": "cycle",
+            "values": [-0.625, -0.5, -0.375, -0.25, -0.125, 0.0, 0.125, 0.25, 0.375, 0.5, 0.625]
+          }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [640],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float32", "shape": [1, 17, 19, 640], "tolerance": 0.0005, "relTolerance": 0.0005 },
+        "S": { "dtype": "float32", "shape": [1, 17, 19, 640], "tolerance": 0, "relTolerance": 0 }
+      },
+      "provenance": {
+        "notes": "Complete-group statistics blocks with a partial final block under the standard 256-invocation WebGPU minimum tier. Odd groups retain scalar loads; even groups exercise vec2 blocks."
+      }
+    },
+    {
+      "name": "nhwc_min_limit_group_tail_c640_g64_float16",
+      "attrs": { "groups": 64, "activation": 0, "channels_last": 1 },
+      "tunables": { "WORKGROUP_SIZE": 64, "SPLIT_MIN_ELEMENTS": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [1, 17, 19, 640],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/nhwc_group_block_tail_c320_float32_input_X" }
+          }
+        },
+        "gamma": {
+          "dtype": "float16",
+          "shape": [640],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "beta": {
+          "dtype": "float16",
+          "shape": [640],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [1, 640],
+          "data": {
+            "kind": "cycle",
+            "values": [-0.625, -0.5, -0.375, -0.25, -0.125, 0.0, 0.125, 0.25, 0.375, 0.5, 0.625]
+          }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [640],
+          "data": { "kind": "cycle", "values": [-0.375, 0.125, 0.25, -0.5, 0.0625, 0.1875, 0.3125] }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [1, 17, 19, 640], "tolerance": 0.01, "relTolerance": 0.002 },
+        "S": { "dtype": "float16", "shape": [1, 17, 19, 640], "tolerance": 0, "relTolerance": 0 }
+      },
+      "provenance": {
+        "notes": "Complete-group statistics blocks with a partial final block under the standard 256-invocation WebGPU minimum tier. Odd groups retain scalar loads; even groups exercise vec2 blocks."
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_webgpu_cpg2_c6_g3_broadcast_n11c_bias_sum",
+      "attrs": { "groups": 3, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [2, 3, 2, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13, "offset": 1.0, "scale": 0.5 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23, "scale": 0.5 }
+        },
+        "skip": {
+          "dtype": "float16",
+          "shape": [2, 1, 1, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.53, "scale": 0.5 }
+        }
+      },
+      "outputs": {
+        "Y": { "dtype": "float16", "shape": [2, 3, 2, 6], "tolerance": 0.01, "relTolerance": 0.004 },
+        "S": { "dtype": "float16", "shape": [2, 3, 2, 6], "tolerance": 0.01, "relTolerance": 0.004 }
+      }
+    },
+    {
+      "name": "ort_skip_group_norm_webgpu_cpg4_c8_g2_broadcast_nc_nobias_nosum",
+      "attrs": { "groups": 2, "activation": 0, "channels_last": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41 }
+        },
+        "gamma": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13, "offset": 1.0, "scale": 0.5 }
+        },
+        "beta": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23, "scale": 0.5 }
+        },
+        "skip": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [2, 3, 2, 8], "tolerance": 0.0005, "relTolerance": 0.0005 } }
+    }
+  ]
+}