sync 6fdf6301e2bb

Browse files

Files changed (8) hide show

README.md +40 -3
build/webgpu/bench.json +121 -0
build/webgpu/manifest.json +1136 -248
build/webgpu/metadata.json +42 -13
build/webgpu/norm-skip-row-vec4.wgsl.jinja +6 -38
build/webgpu/norm-skip-row.wgsl.jinja +32 -44
build/webgpu/norm-stats-copy.wgsl.jinja +15 -0
build/webgpu/test.json +750 -2

README.md CHANGED Viewed

@@ -12,7 +12,7 @@ tags:
 ## Description
-Fuses skip addition with layer normalization for rank-2 or rank-3 input and a non-empty hidden axis. With exact-shape skip, float32 and float16 output-only paths support optional `beta`; adding `bias` requires `beta`. Returning the residual sum requires `beta`: float32 supports optional `bias` and arbitrary hidden sizes, while float16 requires `bias` and a hidden size divisible by four. Broadcast skip is supported for rank-3 float32 input, required `beta`, no `bias` or residual output, and a hidden size divisible by four. Bfloat16 and training statistics are not implemented.
 See the [ONNX Runtime `SkipLayerNormalization` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.SkipLayerNormalization) for the reference semantics.
@@ -31,6 +31,8 @@ See the [ONNX Runtime `SkipLayerNormalization` contrib-operator spec](https://gi
 | Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
 | `outputT` | `output` | `T` | same as `inputT` | same as `inputT` | Normalized output tensor with the same shape as `input`. | required |
 | `residualT` | `input_skip_bias_sum` | `T` | same as `inputT` | same as `inputT` | Sum of `input`, `skip`, and `bias` (when present) before normalization, with the same shape as `input`. | optional |
 ## Attributes
@@ -46,6 +48,40 @@ Default values (overridable per request):
 | Variable | Allowed dtypes |
 | --- | --- |
 | `T` | `float32`, `float16` |
 ## Device requirements
@@ -56,14 +92,15 @@ Some implementation variants require `shader-f16`. These are route-specific capa
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
-- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`norm-skip-row-vec4.wgsl.jinja`](build/webgpu/norm-skip-row-vec4.wgsl.jinja)
 - [`norm-skip-row.wgsl.jinja`](build/webgpu/norm-skip-row.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
-npm install --save-exact @huggingface/kernels@0.0.1-preview.2
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

 ## Description
+Fuses skip addition with layer normalization over a non-empty final hidden axis of rank-2 or rank-3 input. Exact-shape skip supports float32 and float16. Optional float32 mean and inverse-standard-deviation outputs expose row statistics. Broadcast skip supports rank-3 float32 input with beta, no bias or residual output, and a hidden size divisible by four. Output-only and residual-only paths have additional beta, bias and alignment requirements stated by their variants. Bfloat16 is not implemented.
 See the [ONNX Runtime `SkipLayerNormalization` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.SkipLayerNormalization) for the reference semantics.
 | Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
 | `outputT` | `output` | `T` | same as `inputT` | same as `inputT` | Normalized output tensor with the same shape as `input`. | required |
+| `meanT` | `mean` | `U` | same as `inputT` | derived | Per-row mean; zero for simplified RMS normalization. Shape matches the input with its final axis replaced by one. | optional |
+| `invStdT` | `inv_std_var` | `U` | same as `inputT` | derived | Per-row inverse standard deviation, or inverse RMS for simplified normalization. Shape matches the input with its final axis replaced by one. | optional |
 | `residualT` | `input_skip_bias_sum` | `T` | same as `inputT` | same as `inputT` | Sum of `input`, `skip`, and `bias` (when present) before normalization, with the same shape as `input`. | optional |
 ## Attributes
 | Variable | Allowed dtypes |
 | --- | --- |
 | `T` | `float32`, `float16` |
+| `U` | `float32` |
+## 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.
+- `hidden1_f32_no_beta` — Closed-form one-element rows bind only gamma, optional beta, the output, and row parameters. Each invocation writes one row without reading the unused residual inputs.
+- `hidden1_f32_beta` — Closed-form one-element rows bind only gamma, optional beta, the output, and row parameters. Each invocation writes one row without reading the unused residual inputs.
+- `hidden1_f16_no_beta` — Closed-form one-element rows bind only gamma, optional beta, the output, and row parameters. Each invocation writes one row without reading the unused residual inputs.
+- `hidden1_f16_beta` — Closed-form one-element rows bind only gamma, optional beta, the output, and row parameters. Each invocation writes one row without reading the unused residual inputs.
+- `stats_mean_plain` — Row normalization returning mean statistics with `plain` optional inputs and outputs.
+- `stats_mean_residual` — Row normalization returning mean statistics with `residual` optional inputs and outputs.
+- `stats_mean_beta` — Row normalization returning mean statistics with `beta` optional inputs and outputs.
+- `stats_mean_beta_residual` — Row normalization returning mean statistics with `beta_residual` optional inputs and outputs.
+- `stats_mean_bias` — Row normalization returning mean statistics with `bias` optional inputs and outputs.
+- `stats_mean_bias_residual` — Row normalization returning mean statistics with `bias_residual` optional inputs and outputs.
+- `stats_mean_bias_beta` — Row normalization returning mean statistics with `bias_beta` optional inputs and outputs.
+- `stats_mean_bias_beta_residual` — Row normalization returning mean statistics with `bias_beta_residual` optional inputs and outputs.
+- `stats_inv_plain` — Row normalization returning inv statistics with `plain` optional inputs and outputs.
+- `stats_inv_residual` — Row normalization returning inv statistics with `residual` optional inputs and outputs.
+- `stats_inv_beta` — Row normalization returning inv statistics with `beta` optional inputs and outputs.
+- `stats_inv_beta_residual` — Row normalization returning inv statistics with `beta_residual` optional inputs and outputs.
+- `stats_inv_bias` — Row normalization returning inv statistics with `bias` optional inputs and outputs.
+- `stats_inv_bias_residual` — Row normalization returning inv statistics with `bias_residual` optional inputs and outputs.
+- `stats_inv_bias_beta` — Row normalization returning inv statistics with `bias_beta` optional inputs and outputs.
+- `stats_inv_bias_beta_residual` — Row normalization returning inv statistics with `bias_beta_residual` optional inputs and outputs.
+- `stats_both_plain` — Row normalization returning both statistics with `plain` optional inputs and outputs.
+- `stats_both_residual` — Row normalization returning both statistics with `residual` optional inputs and outputs.
+- `stats_both_beta` — Row normalization returning both statistics with `beta` optional inputs and outputs.
+- `stats_both_beta_residual` — Row normalization returning both statistics with `beta_residual` optional inputs and outputs.
+- `stats_both_bias` — Row normalization returning both statistics with `bias` optional inputs and outputs.
+- `stats_both_bias_residual` — Row normalization returning both statistics with `bias_residual` optional inputs and outputs.
+- `stats_both_bias_beta` — Row normalization returning both statistics with `bias_beta` optional inputs and outputs.
+- `stats_both_bias_beta_residual` — Row normalization returning both statistics with `bias_beta_residual` optional inputs and outputs.
 ## Device requirements
 - [`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
 - [`norm-skip-row-vec4.wgsl.jinja`](build/webgpu/norm-skip-row-vec4.wgsl.jinja)
 - [`norm-skip-row.wgsl.jinja`](build/webgpu/norm-skip-row.wgsl.jinja)
+- [`norm-stats-copy.wgsl.jinja`](build/webgpu/norm-stats-copy.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.

build/webgpu/bench.json CHANGED Viewed

@@ -240,6 +240,127 @@
         "residualT": { "shape": [65535, 1], "dtype": "float32" }
       },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 3 * args.hidden)" }] }
     }
   ]
 }

         "residualT": { "shape": [65535, 1], "dtype": "float32" }
       },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "4 * (4 * args.rows * args.hidden + 3 * args.hidden)" }] }
+    },
+    {
+      "name": "independent_float32_rows257_option2",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [257, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float32", "shape": [257, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [257, 1] } }
+    },
+    {
+      "name": "independent_float32_rows65537_option0",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [65537, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float32", "shape": [65537, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 1.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [65537, 1] } }
+    },
+    {
+      "name": "independent_float32_rows65537_option2",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [65537, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float32", "shape": [65537, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [65537, 1] } }
+    },
+    {
+      "name": "independent_float32_rows65537_option3",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [65537, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float32", "shape": [65537, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [65537, 1] },
+        "residualT": { "dtype": "float32", "shape": [65537, 1] }
+      }
+    },
+    {
+      "name": "independent_float16_rows257_option2",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [257, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float16", "shape": [257, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [257, 1] } }
+    },
+    {
+      "name": "independent_float16_rows65537_option0",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [65537, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float16", "shape": [65537, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 1.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [65537, 1] } }
+    },
+    {
+      "name": "independent_float16_rows65537_option2",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [65537, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float16", "shape": [65537, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [65537, 1] } }
+    },
+    {
+      "name": "independent_float32_rows524289_wg8_fold",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [524289, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float32", "shape": [524289, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float32", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [524289, 1] },
+        "residualT": { "dtype": "float32", "shape": [524289, 1] }
+      },
+      "tunables": { "MAX_WORKGROUP_SIZE": 8 }
+    },
+    {
+      "name": "independent_float16_rows524289_wg8_fold",
+      "preset": "edge",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [524289, 1], "dist": "normal", "seed": 1301, "scale": 0.2 },
+        "skipT": { "dtype": "float16", "shape": [524289, 1], "dist": "normal", "seed": 1302, "scale": 0.2 },
+        "gammaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 1.125 },
+        "betaT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": -0.25 },
+        "biasT": { "dtype": "float16", "shape": [1], "dist": "constant", "value": 0.125 }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [524289, 1] } },
+      "tunables": { "MAX_WORKGROUP_SIZE": 8 }
     }
   ]
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -11,6 +11,20 @@
   },
   "outputs": {
     "outputT": { "onnx": "output", "dtype": "T", "rank": "ranks.inputT", "shape": "shapes.inputT" },
     "residualT": {
       "onnx": "input_skip_bias_sum",
       "dtype": "T",
@@ -20,36 +34,37 @@
     }
   },
   "attributes": { "epsilon": { "default": 9.999999960041972e-13 } },
-  "typeConstraints": { "T": ["float32", "float16"] },
   "tunables": { "MAX_WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "rowCount": "numel(shapes.inputT) / max(1, dim(shapes.inputT, -1))",
     "hiddenSize": "dim(shapes.inputT, -1)",
-    "skipWg": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(hiddenSize)))",
     "skipWgVec4": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(ceilDiv(hiddenSize, 4))))",
-    "portableWideExecution": "not has(device.adapterInfo, \"subgroupMinSize\") or device.adapterInfo.subgroupMinSize >= 32",
-    "broadcastRows": "dim(shapes.inputT, 0) * dim(shapes.inputT, 1)",
-    "broadcastHiddenSize": "dim(shapes.inputT, 2)",
-    "broadcastSkipWgVec4": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(ceilDiv(broadcastHiddenSize, 4))))",
     "rowDispatchFits": "rowCount <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
-    "broadcastDispatchFits": "broadcastRows <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "normResourcesFit": "skipWg * 8 <= device.limits.maxComputeWorkgroupStorageSize and skipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
-    "broadcastResourcesFit": "broadcastSkipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
     "epsilonOk": "attrs.epsilon >= 0",
     "coreContract": "epsilonOk and (ranks.inputT == 2 or ranks.inputT == 3) and ranks.skipT == ranks.inputT and ranks.gammaT == 1 and ranks.outputT == ranks.inputT and sameShape(shapes.inputT, shapes.skipT) and sameShape(shapes.outputT, shapes.inputT) and dim(shapes.inputT, -1) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, -1)",
     "residualOutputContract": "present.residualT and sameShape(shapes.residualT, shapes.inputT)",
     "outputOnlyContract": "not present.residualT",
-    "betaContract": "false if not present.betaT else (ranks.betaT == 1 and dim(shapes.betaT, 0) == dim(shapes.inputT, -1))",
-    "noBetaContract": "not present.betaT",
     "f32MainDtypes": "tensorDtypes.inputT == \"float32\" and tensorDtypes.skipT == \"float32\" and tensorDtypes.gammaT == \"float32\" and tensorDtypes.outputT == \"float32\"",
     "f16MainDtypes": "tensorDtypes.inputT == \"float16\" and tensorDtypes.skipT == \"float16\" and tensorDtypes.gammaT == \"float16\" and tensorDtypes.outputT == \"float16\"",
     "f32ResidualDtypes": "f32MainDtypes and tensorDtypes.residualT == \"float32\" if present.residualT else false",
     "f16ResidualDtypes": "f16MainDtypes and tensorDtypes.residualT == \"float16\" if present.residualT else false",
     "vec4Aligned": "dim(shapes.inputT, -1) % 4 == 0",
-    "broadcastSkipShapeOk": "(ranks.skipT == 2 and dim(shapes.skipT, 0) == dim(shapes.inputT, 1) and dim(shapes.skipT, 1) == dim(shapes.inputT, 2)) or (ranks.skipT == 3 and ((dim(shapes.skipT, 0) == 1 and dim(shapes.skipT, 1) == dim(shapes.inputT, 1) and dim(shapes.skipT, 2) == dim(shapes.inputT, 2)) or sameShape(shapes.skipT, shapes.inputT)))",
-    "broadcastOutputOnlyContract": "false if ranks.inputT != 3 or not present.betaT else (epsilonOk and not present.biasT and not present.residualT and dim(shapes.inputT, 2) % 4 == 0 and broadcastSkipShapeOk and ranks.gammaT == 1 and ranks.betaT == 1 and ranks.outputT == 3 and tensorDtypes.inputT == \"float32\" and tensorDtypes.skipT == \"float32\" and tensorDtypes.gammaT == \"float32\" and tensorDtypes.betaT == \"float32\" and tensorDtypes.outputT == \"float32\" and dim(shapes.inputT, 2) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, 2) and dim(shapes.betaT, 0) == dim(shapes.inputT, 2) and sameShape(shapes.outputT, shapes.inputT))",
     "hasSubgroups": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
     "hasF16": "device.features.has(\"shader-f16\")",
     "f32_beta_no_bias_residual_contract": "coreContract and residualOutputContract and betaContract and f32ResidualDtypes and not present.biasT and tensorDtypes.betaT == \"float32\"",
     "f32_beta_bias_residual_contract": "false if not present.biasT else (coreContract and residualOutputContract and betaContract and f32ResidualDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float32\" and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)",
     "f16_beta_bias_residual_contract": "false if not present.biasT else (hasF16 and coreContract and residualOutputContract and betaContract and f16ResidualDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float16\" and tensorDtypes.biasT == \"float16\" and dim(shapes.biasT, 0) == hiddenSize)",
@@ -58,19 +73,19 @@
     "f32_beta_bias_output_only_contract": "false if not present.biasT else (coreContract and outputOnlyContract and betaContract and f32MainDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float32\" and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)",
     "f16_no_beta_output_contract": "hasF16 and coreContract and outputOnlyContract and noBetaContract and f16MainDtypes and not present.biasT",
     "f16_beta_no_bias_output_only_contract": "hasF16 and coreContract and outputOnlyContract and betaContract and f16MainDtypes and not present.biasT and tensorDtypes.betaT == \"float16\"",
-    "f16_beta_bias_output_only_contract": "false if not present.biasT else (hasF16 and coreContract and outputOnlyContract and betaContract and f16MainDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float16\" and tensorDtypes.biasT == \"float16\" and dim(shapes.biasT, 0) == hiddenSize)"
   },
   "bindings": {
-    "input": { "arg": "inputT", "buffer": "read-only-storage", "elementType": "$vectorScalar" },
-    "skip": { "arg": "skipT", "buffer": "read-only-storage", "elementType": "$vectorScalar" },
-    "gamma": { "arg": "gammaT", "buffer": "read-only-storage", "elementType": "$vectorScalar", "length": "$HIDDEN_LEN" },
-    "beta": { "arg": "betaT", "buffer": "read-only-storage", "elementType": "$vectorScalar", "length": "$HIDDEN_LEN" },
-    "output": { "arg": "outputT", "buffer": "storage", "elementType": "$vectorScalar" },
-    "bias": { "arg": "biasT", "buffer": "read-only-storage", "elementType": "$vectorScalar", "length": "$HIDDEN_LEN" },
-    "input_skip_bias_sum": { "arg": "residualT", "buffer": "storage", "elementType": "$vectorScalar" },
-    "params_2": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "rows", "type": "u32", "value": "rowCount" },
         {
@@ -81,51 +96,149 @@
         { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
       ]
     },
-    "input_2": { "arg": "inputT", "name": "input", "buffer": "read-only-storage", "elementType": "$scalar" },
-    "skip_2": { "arg": "skipT", "name": "skip", "buffer": "read-only-storage", "elementType": "$scalar" },
-    "bias_2": {
-      "arg": "biasT",
-      "name": "bias",
-      "buffer": "read-only-storage",
-      "elementType": "$scalar",
-      "length": "$HIDDEN_LEN"
-    },
-    "gamma_2": {
-      "arg": "gammaT",
-      "name": "gamma",
-      "buffer": "read-only-storage",
-      "elementType": "$scalar",
-      "length": "$HIDDEN_LEN"
-    },
-    "beta_2": {
-      "arg": "betaT",
-      "name": "beta",
-      "buffer": "read-only-storage",
-      "elementType": "$scalar",
-      "length": "$HIDDEN_LEN"
-    },
-    "output_2": { "arg": "outputT", "name": "output", "buffer": "storage", "elementType": "$scalar" },
-    "input_skip_bias_sum_2": {
-      "arg": "residualT",
-      "name": "input_skip_bias_sum",
-      "buffer": "storage",
-      "elementType": "$scalar"
-    },
-    "params_3": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "rows", "type": "u32", "value": "rowCount" },
         { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
       ]
-    }
   },
   "variants": [
     {
       "id": "beta_output_only_vec4_broadcast",
       "priority": 19,
-      "when": ["broadcastOutputOnlyContract", "broadcastResourcesFit", "broadcastDispatchFits"],
-      "derive": { "scalar": "\"f32\"", "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "broadcastHiddenSize / 4" },
       "passes": [
         {
           "id": "main",
@@ -136,7 +249,6 @@
             "hasBias": false,
             "hasBeta": true,
             "writeResidualSum": false,
-            "usesF16Spec": false,
             "broadcastSkip": true,
             "hidden": "broadcastHiddenSize",
             "hiddenVec": "broadcastHiddenSize / 4",
@@ -164,22 +276,15 @@
               ]
             }
           ],
-          "dispatch": { "x": "min(broadcastRows, 65535)", "y": "ceilDiv(broadcastRows, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "beta_bias_vec4",
       "priority": 15,
-      "when": ["f32_beta_bias_residual_contract", "vec4Aligned", "hasSubgroups or \"bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
-      "derive": {
-        "scalar": "\"f32\"",
-        "vectorScalar": "\"vec4<f32>\"",
-        "hasBias": "\"bias\" == \"bias\"",
-        "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize / 4"
-      },
       "passes": [
         {
           "id": "normalize",
@@ -187,26 +292,23 @@
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": "\"bias\" == \"bias\"",
             "hasBeta": true,
             "writeResidualSum": true,
-            "usesF16Spec": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "bias", "gamma", "beta", "output", "input_skip_bias_sum", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "beta_bias_row",
       "priority": 5,
-      "when": ["f32_beta_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
         "useSubgroups": "hasSubgroups",
@@ -223,17 +325,77 @@
           "name": "SkipLayerNormalization.Row.Normalize",
           "shader": "norm-skip-row.wgsl.jinja",
           "derive": { "writeResidualSum": true },
-          "bindings": ["input_2", "skip_2", "bias_2", "gamma_2", "beta_2", "output_2", "input_skip_bias_sum_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "beta_bias_vec4_f16",
       "priority": 21,
-      "when": ["f16_beta_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
-      "derive": { "scalar": "\"f16\"", "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
@@ -244,24 +406,22 @@
             "hasBias": true,
             "hasBeta": true,
             "writeResidualSum": true,
-            "usesF16Spec": true,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "bias", "gamma", "beta", "output", "input_skip_bias_sum", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "no_beta_output_only_vec4",
       "priority": 20,
-      "when": ["f32_no_beta_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
-      "derive": { "scalar": "\"f32\"", "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
@@ -269,30 +429,28 @@
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": false,
-            "hasBeta": false,
             "writeResidualSum": false,
-            "usesF16Spec": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "gamma", "output", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "no_beta_output_only_row",
       "priority": 10,
-      "when": ["f32_no_beta_output_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
-        "hasBias": false,
-        "hasBeta": false,
         "writeResidualSum": false,
         "useSubgroups": "hasSubgroups",
         "scalar": "\"f32\"",
@@ -304,17 +462,20 @@
           "id": "main",
           "name": "SkipLayerNormalization.NoBetaOutputOnly.Row",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "gamma_2", "output_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "no_beta_output_only_vec4_f16",
       "priority": 20,
-      "when": ["f16_no_beta_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits"],
-      "derive": { "scalar": "\"f16\"", "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
@@ -322,34 +483,31 @@
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": false,
-            "hasBeta": false,
             "writeResidualSum": false,
-            "usesF16Spec": true,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "gamma", "output", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "no_beta_output_only_row_f16",
       "priority": 10,
-      "when": ["f16_no_beta_output_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
-        "hasBias": false,
-        "hasBeta": false,
         "writeResidualSum": false,
         "useSubgroups": "hasSubgroups",
         "scalar": "\"f16\"",
-        "usesF16": true,
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
       },
@@ -358,23 +516,20 @@
           "id": "main",
           "name": "SkipLayerNormalization.NoBetaOutputOnly.Row.F16",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "gamma_2", "output_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_no_bias_vec4",
       "priority": 20,
-      "when": ["f32_beta_no_bias_residual_contract", "vec4Aligned", "hasSubgroups or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
-      "derive": {
-        "scalar": "\"f32\"",
-        "vectorScalar": "\"vec4<f32>\"",
-        "hasBias": "\"no_bias\" == \"bias\"",
-        "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize / 4"
-      },
       "passes": [
         {
           "id": "main",
@@ -382,32 +537,30 @@
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": "\"no_bias\" == \"bias\"",
-            "hasBeta": true,
-            "writeResidualSum": true,
-            "usesF16Spec": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "gamma", "beta", "output", "input_skip_bias_sum", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_no_bias_row",
       "priority": 10,
-      "when": ["f32_beta_no_bias_residual_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
         "useSubgroups": "hasSubgroups",
-        "hasBeta": true,
-        "writeResidualSum": true,
-        "hasBias": "\"no_bias\" == \"bias\"",
         "scalar": "\"f32\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
@@ -417,82 +570,74 @@
           "id": "main",
           "name": "SkipLayerNormalization.Row",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "gamma_2", "beta_2", "output_2", "input_skip_bias_sum_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_no_bias_output_only_vec4",
       "priority": 20,
-      "when": ["f32_beta_no_bias_output_only_contract", "vec4Aligned", "hasSubgroups or \"no_bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
-      "derive": {
-        "scalar": "\"f32\"",
-        "vectorScalar": "\"vec4<f32>\"",
-        "hasBias": "\"no_bias\" == \"bias\"",
-        "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize / 4"
-      },
       "passes": [
         {
           "id": "main",
-          "name": "SkipLayerNormalization.Vec4",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": "\"no_bias\" == \"bias\"",
-            "hasBeta": true,
             "writeResidualSum": false,
-            "usesF16Spec": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
-            "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "gamma", "beta", "output", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_no_bias_output_only_row",
       "priority": 10,
-      "when": ["f32_beta_no_bias_output_only_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
-        "useSubgroups": "hasSubgroups",
-        "hasBeta": true,
         "writeResidualSum": false,
-        "hasBias": "\"no_bias\" == \"bias\"",
-        "scalar": "\"f32\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
       },
       "passes": [
         {
           "id": "main",
-          "name": "SkipLayerNormalization.Row",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "gamma_2", "beta_2", "output_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "beta_bias_output_only_vec4",
       "priority": 20,
-      "when": ["f32_beta_bias_output_only_contract", "vec4Aligned", "hasSubgroups or \"bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
-      "derive": {
-        "scalar": "\"f32\"",
-        "vectorScalar": "\"vec4<f32>\"",
-        "hasBias": "\"bias\" == \"bias\"",
-        "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize / 4"
-      },
       "passes": [
         {
           "id": "main",
@@ -500,32 +645,30 @@
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": "\"bias\" == \"bias\"",
-            "hasBeta": true,
             "writeResidualSum": false,
-            "usesF16Spec": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "gamma", "beta", "bias", "output", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
       "id": "beta_bias_output_only_row",
       "priority": 10,
-      "when": ["f32_beta_bias_output_only_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
-        "useSubgroups": "hasSubgroups",
-        "hasBeta": true,
         "writeResidualSum": false,
-        "hasBias": "\"bias\" == \"bias\"",
         "scalar": "\"f32\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
@@ -535,23 +678,20 @@
           "id": "main",
           "name": "SkipLayerNormalization.Row",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "bias_2", "gamma_2", "beta_2", "output_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_no_bias_output_only_vec4_f16",
       "priority": 20,
-      "when": ["f16_beta_no_bias_output_only_contract", "vec4Aligned", "hasSubgroups or \"no_bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
-      "derive": {
-        "scalar": "\"f16\"",
-        "vectorScalar": "\"vec4<f16>\"",
-        "hasBias": "\"no_bias\" == \"bias\"",
-        "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize / 4"
-      },
       "passes": [
         {
           "id": "main",
@@ -559,34 +699,31 @@
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
-            "hasBias": "\"no_bias\" == \"bias\"",
-            "hasBeta": true,
             "writeResidualSum": false,
-            "usesF16Spec": true,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
-          "bindings": ["input", "skip", "gamma", "beta", "output", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_no_bias_output_only_row_f16",
       "priority": 10,
-      "when": ["f16_beta_no_bias_output_only_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
-        "useSubgroups": "hasSubgroups",
-        "hasBeta": true,
         "writeResidualSum": false,
-        "hasBias": "\"no_bias\" == \"bias\"",
         "scalar": "\"f16\"",
-        "usesF16": true,
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
       },
@@ -595,71 +732,822 @@
           "id": "main",
           "name": "SkipLayerNormalization.Row.F16",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "gamma_2", "beta_2", "output_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
       ]
     },
     {
-      "id": "beta_bias_output_only_vec4_f16",
-      "priority": 20,
-      "when": ["f16_beta_bias_output_only_contract", "vec4Aligned", "hasSubgroups or \"bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits"],
       "derive": {
-        "scalar": "\"f16\"",
-        "vectorScalar": "\"vec4<f16>\"",
-        "hasBias": "\"bias\" == \"bias\"",
         "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize / 4"
       },
       "passes": [
         {
           "id": "main",
-          "name": "SkipLayerNormalization.Vec4.F16",
-          "shader": "norm-skip-row-vec4.wgsl.jinja",
-          "derive": {
-            "simplified": false,
-            "hasBias": "\"bias\" == \"bias\"",
-            "hasBeta": true,
-            "writeResidualSum": false,
-            "usesF16Spec": true,
-            "hidden": "hiddenSize",
-            "hiddenVec": "hiddenSize / 4",
-            "wg": "skipWgVec4",
-            "vecType": "\"vec4<f16>\"",
-            "useSubgroups": "hasSubgroups"
-          },
-          "bindings": ["input", "skip", "gamma", "beta", "bias", "output", "params_2"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
-      ]
     },
     {
-      "id": "beta_bias_output_only_row_f16",
-      "priority": 10,
-      "when": ["f16_beta_bias_output_only_contract", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
-        "useSubgroups": "hasSubgroups",
-        "hasBeta": true,
-        "writeResidualSum": false,
-        "hasBias": "\"bias\" == \"bias\"",
-        "scalar": "\"f16\"",
-        "usesF16": true,
         "workgroupSize": "skipWg",
-        "HIDDEN_LEN": "hiddenSize"
       },
       "passes": [
         {
           "id": "main",
-          "name": "SkipLayerNormalization.Row.F16",
           "shader": "norm-skip-row.wgsl.jinja",
-          "bindings": ["input_2", "skip_2", "bias_2", "gamma_2", "beta_2", "output_2", "params_3"],
-          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 },
-          "subgroupCollectivesWidth": "portable"
         }
-      ]
     }
   ]
 }

   },
   "outputs": {
     "outputT": { "onnx": "output", "dtype": "T", "rank": "ranks.inputT", "shape": "shapes.inputT" },
+    "meanT": {
+      "onnx": "mean",
+      "dtype": "U",
+      "optional": true,
+      "rank": "ranks.inputT",
+      "shape": "prefix(shapes.inputT, ranks.inputT - 1) + [1]"
+    },
+    "invStdT": {
+      "onnx": "inv_std_var",
+      "dtype": "U",
+      "optional": true,
+      "rank": "ranks.inputT",
+      "shape": "prefix(shapes.inputT, ranks.inputT - 1) + [1]"
+    },
     "residualT": {
       "onnx": "input_skip_bias_sum",
       "dtype": "T",
     }
   },
   "attributes": { "epsilon": { "default": 9.999999960041972e-13 } },
+  "typeConstraints": { "T": ["float32", "float16"], "U": ["float32"] },
   "tunables": { "MAX_WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "rowCount": "numel(shapes.inputT) / max(1, dim(shapes.inputT, -1))",
     "hiddenSize": "dim(shapes.inputT, -1)",
+    "skipWg": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)) if hiddenSize == 1 else (max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(hiddenSize))))",
     "skipWgVec4": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(ceilDiv(hiddenSize, 4))))",
     "rowDispatchFits": "rowCount <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "normResourcesFit": "skipWg * 8 <= device.limits.maxComputeWorkgroupStorageSize and skipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
     "epsilonOk": "attrs.epsilon >= 0",
     "coreContract": "epsilonOk and (ranks.inputT == 2 or ranks.inputT == 3) and ranks.skipT == ranks.inputT and ranks.gammaT == 1 and ranks.outputT == ranks.inputT and sameShape(shapes.inputT, shapes.skipT) and sameShape(shapes.outputT, shapes.inputT) and dim(shapes.inputT, -1) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, -1)",
     "residualOutputContract": "present.residualT and sameShape(shapes.residualT, shapes.inputT)",
     "outputOnlyContract": "not present.residualT",
     "f32MainDtypes": "tensorDtypes.inputT == \"float32\" and tensorDtypes.skipT == \"float32\" and tensorDtypes.gammaT == \"float32\" and tensorDtypes.outputT == \"float32\"",
     "f16MainDtypes": "tensorDtypes.inputT == \"float16\" and tensorDtypes.skipT == \"float16\" and tensorDtypes.gammaT == \"float16\" and tensorDtypes.outputT == \"float16\"",
     "f32ResidualDtypes": "f32MainDtypes and tensorDtypes.residualT == \"float32\" if present.residualT else false",
     "f16ResidualDtypes": "f16MainDtypes and tensorDtypes.residualT == \"float16\" if present.residualT else false",
     "vec4Aligned": "dim(shapes.inputT, -1) % 4 == 0",
     "hasSubgroups": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
     "hasF16": "device.features.has(\"shader-f16\")",
+    "statsRequested": "present.meanT or present.invStdT",
+    "portableWideExecution": "not has(device.adapterInfo, \"subgroupMinSize\") or device.adapterInfo.subgroupMinSize >= 32",
+    "broadcastRows": "dim(shapes.inputT, 0) * dim(shapes.inputT, 1)",
+    "broadcastHiddenSize": "dim(shapes.inputT, 2)",
+    "broadcastSkipWgVec4": "max(1, min(tunables.MAX_WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, pow2ceil(ceilDiv(broadcastHiddenSize, 4))))",
+    "broadcastDispatchFits": "broadcastRows <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "broadcastResourcesFit": "broadcastSkipWgVec4 * 8 <= device.limits.maxComputeWorkgroupStorageSize",
+    "betaContract": "false if not present.betaT else (ranks.betaT == 1 and dim(shapes.betaT, 0) == dim(shapes.inputT, -1))",
+    "noBetaContract": "not present.betaT",
+    "broadcastSkipShapeOk": "(ranks.skipT == 2 and dim(shapes.skipT, 0) == dim(shapes.inputT, 1) and dim(shapes.skipT, 1) == dim(shapes.inputT, 2)) or (ranks.skipT == 3 and ((dim(shapes.skipT, 0) == 1 and dim(shapes.skipT, 1) == dim(shapes.inputT, 1) and dim(shapes.skipT, 2) == dim(shapes.inputT, 2)) or sameShape(shapes.skipT, shapes.inputT)))",
+    "broadcastOutputOnlyContract": "false if ranks.inputT != 3 or not present.betaT else (epsilonOk and not present.biasT and not present.residualT and dim(shapes.inputT, 2) % 4 == 0 and broadcastSkipShapeOk and ranks.gammaT == 1 and ranks.betaT == 1 and ranks.outputT == 3 and tensorDtypes.inputT == \"float32\" and tensorDtypes.skipT == \"float32\" and tensorDtypes.gammaT == \"float32\" and tensorDtypes.betaT == \"float32\" and tensorDtypes.outputT == \"float32\" and dim(shapes.inputT, 2) > 0 and dim(shapes.gammaT, 0) == dim(shapes.inputT, 2) and dim(shapes.betaT, 0) == dim(shapes.inputT, 2) and sameShape(shapes.outputT, shapes.inputT))",
     "f32_beta_no_bias_residual_contract": "coreContract and residualOutputContract and betaContract and f32ResidualDtypes and not present.biasT and tensorDtypes.betaT == \"float32\"",
     "f32_beta_bias_residual_contract": "false if not present.biasT else (coreContract and residualOutputContract and betaContract and f32ResidualDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float32\" and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)",
     "f16_beta_bias_residual_contract": "false if not present.biasT else (hasF16 and coreContract and residualOutputContract and betaContract and f16ResidualDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float16\" and tensorDtypes.biasT == \"float16\" and dim(shapes.biasT, 0) == hiddenSize)",
     "f32_beta_bias_output_only_contract": "false if not present.biasT else (coreContract and outputOnlyContract and betaContract and f32MainDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float32\" and tensorDtypes.biasT == \"float32\" and dim(shapes.biasT, 0) == hiddenSize)",
     "f16_no_beta_output_contract": "hasF16 and coreContract and outputOnlyContract and noBetaContract and f16MainDtypes and not present.biasT",
     "f16_beta_no_bias_output_only_contract": "hasF16 and coreContract and outputOnlyContract and betaContract and f16MainDtypes and not present.biasT and tensorDtypes.betaT == \"float16\"",
+    "f16_beta_bias_output_only_contract": "false if not present.biasT else (hasF16 and coreContract and outputOnlyContract and betaContract and f16MainDtypes and ranks.biasT == 1 and tensorDtypes.betaT == \"float16\" and tensorDtypes.biasT == \"float16\" and dim(shapes.biasT, 0) == hiddenSize)",
+    "statsContract": "statsRequested and coreContract and f16Ok(dtypes.T) and (not present.biasT or (ranks.biasT == 1 and dim(shapes.biasT, 0) == hiddenSize)) and (not present.residualT or sameShape(shapes.residualT, shapes.inputT)) and (not present.betaT or (ranks.betaT == 1 and dim(shapes.betaT, 0) == hiddenSize))"
   },
   "bindings": {
+    "input": { "arg": "inputT", "elementType": "$vectorScalar" },
+    "skip": { "arg": "skipT", "elementType": "$vectorScalar" },
+    "gamma": { "arg": "gammaT", "elementType": "$vectorScalar", "length": "$HIDDEN_LEN" },
+    "beta": { "arg": "betaT", "elementType": "$vectorScalar", "length": "$HIDDEN_LEN" },
+    "output": { "arg": "outputT", "elementType": "$vectorScalar" },
+    "bias": { "arg": "biasT", "elementType": "$vectorScalar", "length": "$HIDDEN_LEN" },
+    "input_skip_bias_sum": { "arg": "residualT", "elementType": "$vectorScalar" },
+    "params_main": {
       "name": "params",
       "struct": [
         { "name": "rows", "type": "u32", "value": "rowCount" },
         {
         { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
       ]
     },
+    "input_main": { "arg": "inputT", "name": "input", "elementType": "$scalar" },
+    "skip_main": { "arg": "skipT", "name": "skip", "elementType": "$scalar" },
+    "bias_main": { "arg": "biasT", "name": "bias", "elementType": "$scalar", "length": "$HIDDEN_LEN" },
+    "gamma_main": { "arg": "gammaT", "name": "gamma", "elementType": "$scalar", "length": "$HIDDEN_LEN" },
+    "beta_main": { "arg": "betaT", "name": "beta", "elementType": "$scalar", "length": "$HIDDEN_LEN" },
+    "output_main": { "arg": "outputT", "name": "output", "elementType": "$scalar" },
+    "input_skip_bias_sum_main": { "arg": "residualT", "name": "input_skip_bias_sum", "elementType": "$scalar" },
+    "params__uniform": {
       "name": "params",
       "struct": [
         { "name": "rows", "type": "u32", "value": "rowCount" },
         { "name": "epsilon", "type": "f32", "value": "attrs.epsilon" }
       ]
+    },
+    "mean": { "arg": "meanT", "elementType": "f32" },
+    "inv_std_var": { "arg": "invStdT", "elementType": "f32" },
+    "row_stats": { "scratch": "rowStats", "elementType": "vec2<f32>" },
+    "row_stats_read": {
+      "scratch": "rowStats",
+      "name": "row_stats",
+      "buffer": "read-only-storage",
+      "elementType": "vec2<f32>"
+    },
+    "params_stats": { "name": "params", "struct": [{ "name": "rows", "type": "u32", "value": "rowCount" }] }
   },
   "variants": [
+    {
+      "id": "hidden1_f32_no_beta",
+      "priority": 20,
+      "when": ["f32_no_beta_output_contract", "hiddenSize == 1", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": {
+        "simplified": false,
+        "hasBias": false,
+        "hasBeta": "present.betaT",
+        "writeResidualSum": false,
+        "useSubgroups": false,
+        "scalar": "dtypes.T",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Hidden1",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["gamma_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (skipWg)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "hidden1_f32_beta",
+      "priority": 20,
+      "when": ["(f32_beta_no_bias_output_only_contract or f32_beta_bias_output_only_contract)", "hiddenSize == 1", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": {
+        "simplified": false,
+        "hasBias": false,
+        "hasBeta": "present.betaT",
+        "writeResidualSum": false,
+        "useSubgroups": false,
+        "scalar": "dtypes.T",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Hidden1",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["gamma_main", "beta_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (skipWg)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "hidden1_f16_no_beta",
+      "priority": 20,
+      "when": ["f16_no_beta_output_contract", "hiddenSize == 1", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": {
+        "simplified": false,
+        "hasBias": false,
+        "hasBeta": "present.betaT",
+        "writeResidualSum": false,
+        "useSubgroups": false,
+        "scalar": "dtypes.T",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Hidden1",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["gamma_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (skipWg)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "hidden1_f16_beta",
+      "priority": 20,
+      "when": ["(f16_beta_no_bias_output_only_contract or f16_beta_bias_output_only_contract)", "hiddenSize == 1", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": {
+        "simplified": false,
+        "hasBias": false,
+        "hasBeta": "present.betaT",
+        "writeResidualSum": false,
+        "useSubgroups": false,
+        "scalar": "dtypes.T",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Hidden1",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["gamma_main", "beta_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (skipWg)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
     {
       "id": "beta_output_only_vec4_broadcast",
       "priority": 19,
+      "when": ["broadcastOutputOnlyContract", "broadcastResourcesFit", "broadcastDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "broadcastHiddenSize / 4" },
       "passes": [
         {
           "id": "main",
             "hasBias": false,
             "hasBeta": true,
             "writeResidualSum": false,
             "broadcastSkip": true,
             "hidden": "broadcastHiddenSize",
             "hiddenVec": "broadcastHiddenSize / 4",
               ]
             }
           ],
+          "dispatch": { "x": "min(broadcastRows, 65535)", "y": "ceilDiv(broadcastRows, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "beta_bias_vec4",
       "priority": 15,
+      "when": ["f32_beta_bias_residual_contract", "vec4Aligned", "hasSubgroups or \"bias\" == \"bias\" or portableWideExecution", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f32>\"", "hasBias": "\"bias\" == \"bias\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "normalize",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
             "hasBeta": true,
             "writeResidualSum": true,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "bias", "gamma", "beta", "output", "input_skip_bias_sum", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "beta_bias_row",
       "priority": 5,
+      "when": ["f32_beta_bias_residual_contract", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
         "useSubgroups": "hasSubgroups",
           "name": "SkipLayerNormalization.Row.Normalize",
           "shader": "norm-skip-row.wgsl.jinja",
           "derive": { "writeResidualSum": true },
+          "bindings": ["input_main", "skip_main", "bias_main", "gamma_main", "beta_main", "output_main", "input_skip_bias_sum_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "beta_no_bias_vec4",
+      "priority": 20,
+      "when": ["f32_beta_no_bias_residual_contract", "vec4Aligned", "hasSubgroups or portableWideExecution", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": {
+        "vectorScalar": "\"vec4<f32>\"",
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "HIDDEN_LEN": "hiddenSize / 4"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Vec4",
+          "shader": "norm-skip-row-vec4.wgsl.jinja",
+          "derive": {
+            "simplified": false,
+            "hasBeta": true,
+            "writeResidualSum": true,
+            "hidden": "hiddenSize",
+            "hiddenVec": "hiddenSize / 4",
+            "wg": "skipWgVec4",
+            "vecType": "\"vec4<f32>\"",
+            "useSubgroups": "hasSubgroups"
+          },
+          "bindings": ["input", "skip", "gamma", "beta", "output", "input_skip_bias_sum", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
+        }
+      ]
+    },
+    {
+      "id": "beta_no_bias_row",
+      "priority": 10,
+      "when": ["f32_beta_no_bias_residual_contract", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": {
+        "simplified": false,
+        "useSubgroups": "hasSubgroups",
+        "hasBeta": true,
+        "writeResidualSum": true,
+        "hasBias": "\"no_bias\" == \"bias\"",
+        "scalar": "\"f32\"",
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SkipLayerNormalization.Row",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "output_main", "input_skip_bias_sum_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "beta_bias_vec4_f16",
       "priority": 21,
+      "when": ["f16_beta_bias_residual_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
             "hasBias": true,
             "hasBeta": true,
             "writeResidualSum": true,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "bias", "gamma", "beta", "output", "input_skip_bias_sum", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "no_beta_output_only_vec4",
       "priority": 20,
+      "when": ["f32_no_beta_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
+            "hasBias": "present.biasT",
+            "hasBeta": "present.betaT",
             "writeResidualSum": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "gamma", "output", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "no_beta_output_only_row",
       "priority": 10,
+      "when": ["f32_no_beta_output_contract", "normResourcesFit", "rowDispatchFits", "hiddenSize != 1", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
         "writeResidualSum": false,
         "useSubgroups": "hasSubgroups",
         "scalar": "\"f32\"",
           "id": "main",
           "name": "SkipLayerNormalization.NoBetaOutputOnly.Row",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "no_beta_output_only_vec4_f16",
       "priority": 20,
+      "when": ["f16_no_beta_output_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
+            "hasBias": "present.biasT",
+            "hasBeta": "present.betaT",
             "writeResidualSum": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "gamma", "output", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "no_beta_output_only_row_f16",
       "priority": 10,
+      "when": ["f16_no_beta_output_contract", "normResourcesFit", "rowDispatchFits", "hiddenSize != 1", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
         "writeResidualSum": false,
         "useSubgroups": "hasSubgroups",
         "scalar": "\"f16\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
       },
           "id": "main",
           "name": "SkipLayerNormalization.NoBetaOutputOnly.Row.F16",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
+      "id": "beta_no_bias_output_only_vec4",
       "priority": 20,
+      "when": ["f32_beta_no_bias_output_only_contract", "vec4Aligned", "hasSubgroups or portableWideExecution", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
+            "hasBias": "present.biasT",
+            "hasBeta": "present.betaT",
+            "writeResidualSum": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "gamma", "beta", "output", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
+      "id": "beta_no_bias_output_only_row",
       "priority": 10,
+      "when": ["f32_beta_no_bias_output_only_contract", "normResourcesFit", "rowDispatchFits", "hiddenSize != 1", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": false,
         "useSubgroups": "hasSubgroups",
         "scalar": "\"f32\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
           "id": "main",
           "name": "SkipLayerNormalization.Row",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
+      "id": "beta_no_bias_output_only_vec4_f16",
       "priority": 20,
+      "when": ["f16_beta_no_bias_output_only_contract", "vec4Aligned", "hasSubgroups or portableWideExecution", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
+          "name": "SkipLayerNormalization.Vec4.F16",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
+            "hasBias": "present.biasT",
+            "hasBeta": "present.betaT",
             "writeResidualSum": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
+            "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "gamma", "beta", "output", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
+      "id": "beta_no_bias_output_only_row_f16",
       "priority": 10,
+      "when": ["f16_beta_no_bias_output_only_contract", "normResourcesFit", "rowDispatchFits", "hiddenSize != 1", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
         "writeResidualSum": false,
+        "useSubgroups": "hasSubgroups",
+        "scalar": "\"f16\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
       },
       "passes": [
         {
           "id": "main",
+          "name": "SkipLayerNormalization.Row.F16",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "beta_bias_output_only_vec4",
       "priority": 20,
+      "when": ["f32_beta_bias_output_only_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f32>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
+            "hasBias": "present.biasT",
+            "hasBeta": "present.betaT",
             "writeResidualSum": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f32>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "gamma", "beta", "bias", "output", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "beta_bias_output_only_row",
       "priority": 10,
+      "when": ["f32_beta_bias_output_only_contract", "normResourcesFit", "rowDispatchFits", "hiddenSize != 1", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
         "writeResidualSum": false,
+        "useSubgroups": "hasSubgroups",
         "scalar": "\"f32\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
           "id": "main",
           "name": "SkipLayerNormalization.Row",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "bias_main", "gamma_main", "beta_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
+      "id": "beta_bias_output_only_vec4_f16",
       "priority": 20,
+      "when": ["f16_beta_bias_output_only_contract", "vec4Aligned", "normResourcesFit", "rowDispatchFits", "not present.meanT and not present.invStdT"],
+      "derive": { "vectorScalar": "\"vec4<f16>\"", "HIDDEN_LEN": "hiddenSize / 4" },
       "passes": [
         {
           "id": "main",
           "shader": "norm-skip-row-vec4.wgsl.jinja",
           "derive": {
             "simplified": false,
+            "hasBias": "present.biasT",
+            "hasBeta": "present.betaT",
             "writeResidualSum": false,
             "hidden": "hiddenSize",
             "hiddenVec": "hiddenSize / 4",
             "wg": "skipWgVec4",
             "vecType": "\"vec4<f16>\"",
             "useSubgroups": "hasSubgroups"
           },
+          "bindings": ["input", "skip", "gamma", "beta", "bias", "output", "params_main"],
+          "dispatch": { "x": "min(rowCount, 65535)", "y": "ceilDiv(rowCount, 65535)", "z": 1 }
         }
       ]
     },
     {
+      "id": "beta_bias_output_only_row_f16",
       "priority": 10,
+      "when": ["f16_beta_bias_output_only_contract", "normResourcesFit", "rowDispatchFits", "hiddenSize != 1", "not present.meanT and not present.invStdT"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
         "writeResidualSum": false,
+        "useSubgroups": "hasSubgroups",
         "scalar": "\"f16\"",
         "workgroupSize": "skipWg",
         "HIDDEN_LEN": "hiddenSize"
       },
           "id": "main",
           "name": "SkipLayerNormalization.Row.F16",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "bias_main", "gamma_main", "beta_main", "output_main", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount * (1 if hiddenSize == 1 else skipWg)), (skipWg)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
+      "id": "stats_mean_plain",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "not present.biasT", "not present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
       "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
         "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
       },
       "passes": [
         {
           "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
         }
+      ],
+      "intermediates": []
     },
     {
+      "id": "stats_mean_residual",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "not present.biasT", "not present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
       "derive": {
         "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
         "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
       },
       "passes": [
         {
           "id": "main",
           "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "input_skip_bias_sum_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
         }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_mean_beta",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "not present.biasT", "present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_mean_beta_residual",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "not present.biasT", "present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "input_skip_bias_sum_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_mean_bias",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "present.biasT", "not present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_mean_bias_residual",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "present.biasT", "not present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "input_skip_bias_sum_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_mean_bias_beta",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "present.biasT", "present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "beta_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_mean_bias_beta_residual",
+      "when": ["statsContract", "present.meanT", "not present.invStdT", "present.biasT", "present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "beta_main", "input_skip_bias_sum_main", "output_main", "mean", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_plain",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "not present.biasT", "not present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_residual",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "not present.biasT", "not present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "input_skip_bias_sum_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_beta",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "not present.biasT", "present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_beta_residual",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "not present.biasT", "present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "input_skip_bias_sum_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_bias",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "present.biasT", "not present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_bias_residual",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "present.biasT", "not present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "input_skip_bias_sum_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_bias_beta",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "present.biasT", "present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "beta_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_inv_bias_beta_residual",
+      "when": ["statsContract", "not present.meanT", "present.invStdT", "present.biasT", "present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "beta_main", "input_skip_bias_sum_main", "output_main", "inv_std_var", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": []
+    },
+    {
+      "id": "stats_both_plain",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "not present.biasT", "not present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_residual",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "not present.biasT", "not present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "input_skip_bias_sum_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_beta",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "not present.biasT", "present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_beta_residual",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "not present.biasT", "present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "beta_main", "input_skip_bias_sum_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_bias",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "present.biasT", "not present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_bias_residual",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "present.biasT", "not present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "input_skip_bias_sum_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_bias_beta",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "present.biasT", "present.betaT", "not present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "beta_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
+    },
+    {
+      "id": "stats_both_bias_beta_residual",
+      "when": ["statsContract", "present.meanT", "present.invStdT", "present.biasT", "present.betaT", "present.residualT", "normResourcesFit", "rowDispatchFits"],
+      "derive": {
+        "simplified": false,
+        "hasBias": "present.biasT",
+        "hasBeta": "present.betaT",
+        "writeResidualSum": "present.residualT",
+        "writeMean": "present.meanT",
+        "writeInvStd": "present.invStdT",
+        "scalar": "dtypes.T",
+        "useSubgroups": false,
+        "workgroupSize": "skipWg",
+        "HIDDEN_LEN": "hiddenSize",
+        "packedStatistics": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "norm-skip-row.wgsl.jinja",
+          "bindings": ["input_main", "skip_main", "gamma_main", "bias_main", "beta_main", "input_skip_bias_sum_main", "output_main", "row_stats", "params__uniform"],
+          "dispatch": {
+            "x": "min(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "y": "ceilDiv(ceilDiv(rowCount, skipWg) if hiddenSize == 1 else rowCount, 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "statistics",
+          "shader": "norm-stats-copy.wgsl.jinja",
+          "derive": { "workgroupSize": 64 },
+          "bindings": ["row_stats_read", "mean", "inv_std_var", "params_stats"],
+          "dispatch": {
+            "x": "min(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((rowCount), (workgroupSize)), 65535)",
+            "z": 1
+          }
+        }
+      ],
+      "intermediates": [{ "id": "rowStats", "dtype": "float32", "shape": "[rowCount, 2]" }]
     }
   ]
 }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,41 +1,70 @@
 {
   "name": "com.microsoft.SkipLayerNormalization",
-  "id": "_com_microsoft_skiplayernormalization_webgpu_28a934b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "dZADuGnHvG3kWGPS0BGgmikuwdi16tD/Sk4HbGC9UV0=",
-      "manifest.json": "G6ZofqPdSuhkAxEuuKIlW4fQ/fCsDwxEzafKwVqGEY8=",
-      "norm-skip-row-vec4.wgsl.jinja": "bH8L9BYJ/3XQXE2zP8kvZc4wlhAbcAolI4/p0xXskzM=",
-      "norm-skip-row.wgsl.jinja": "aERQfDDXx6lwKuyahwcemgFllRSq46egf2++7f6bNtA=",
-      "test.json": "K+KtXhWTSxbRlFwFfszdd1YKz9JnAOuDZ9ebLg7Y8w0="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "beta_output_only_vec4_broadcast": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_bias_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_bias_row": ["norm-skip-row.wgsl.jinja"],
       "beta_bias_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
       "no_beta_output_only_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
       "no_beta_output_only_row": ["norm-skip-row.wgsl.jinja"],
       "no_beta_output_only_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
       "no_beta_output_only_row_f16": ["norm-skip-row.wgsl.jinja"],
-      "beta_no_bias_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
-      "beta_no_bias_row": ["norm-skip-row.wgsl.jinja"],
       "beta_no_bias_output_only_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_no_bias_output_only_row": ["norm-skip-row.wgsl.jinja"],
-      "beta_bias_output_only_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
-      "beta_bias_output_only_row": ["norm-skip-row.wgsl.jinja"],
       "beta_no_bias_output_only_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_no_bias_output_only_row_f16": ["norm-skip-row.wgsl.jinja"],
       "beta_bias_output_only_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
-      "beta_bias_output_only_row_f16": ["norm-skip-row.wgsl.jinja"]
     }
   }
 }

 {
   "name": "com.microsoft.SkipLayerNormalization",
+  "id": "_com_microsoft_skiplayernormalization_webgpu_a30f1f2",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "cVSe9BmYyynkTmGaCfti3+3as2tY8s7SiHcDUGZXadQ=",
+      "manifest.json": "eBeukaRPaePto8cDuxf/pVi95UUFHU+jR4Tcc4/FFh4=",
+      "norm-skip-row-vec4.wgsl.jinja": "hpoa1+W36ADP2AlqvZAoWFHCZPTBrUIlgVIQXE98oTE=",
+      "norm-skip-row.wgsl.jinja": "LOw1CZFwolmaeaXpWSTi8ToiXzTfvJDoaapF2N6a0x4=",
+      "norm-stats-copy.wgsl.jinja": "4PNrRNFbMWP9csH7RrxSvJuAGSn5aDfZZPw/IZEVmjU=",
+      "test.json": "upshAVlCpYOJn3G1M4TGhznjHzx8ybo8uae8zkB932s="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
+      "hidden1_f32_no_beta": ["norm-skip-row.wgsl.jinja"],
+      "hidden1_f32_beta": ["norm-skip-row.wgsl.jinja"],
+      "hidden1_f16_no_beta": ["norm-skip-row.wgsl.jinja"],
+      "hidden1_f16_beta": ["norm-skip-row.wgsl.jinja"],
       "beta_output_only_vec4_broadcast": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_bias_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_bias_row": ["norm-skip-row.wgsl.jinja"],
+      "beta_no_bias_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
+      "beta_no_bias_row": ["norm-skip-row.wgsl.jinja"],
       "beta_bias_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
       "no_beta_output_only_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
       "no_beta_output_only_row": ["norm-skip-row.wgsl.jinja"],
       "no_beta_output_only_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
       "no_beta_output_only_row_f16": ["norm-skip-row.wgsl.jinja"],
       "beta_no_bias_output_only_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_no_bias_output_only_row": ["norm-skip-row.wgsl.jinja"],
       "beta_no_bias_output_only_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
       "beta_no_bias_output_only_row_f16": ["norm-skip-row.wgsl.jinja"],
+      "beta_bias_output_only_vec4": ["norm-skip-row-vec4.wgsl.jinja"],
+      "beta_bias_output_only_row": ["norm-skip-row.wgsl.jinja"],
       "beta_bias_output_only_vec4_f16": ["norm-skip-row-vec4.wgsl.jinja"],
+      "beta_bias_output_only_row_f16": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_plain": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_beta": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_beta_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_bias": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_bias_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_bias_beta": ["norm-skip-row.wgsl.jinja"],
+      "stats_mean_bias_beta_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_plain": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_beta": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_beta_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_bias": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_bias_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_bias_beta": ["norm-skip-row.wgsl.jinja"],
+      "stats_inv_bias_beta_residual": ["norm-skip-row.wgsl.jinja"],
+      "stats_both_plain": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_residual": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_beta": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_beta_residual": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_bias": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_bias_residual": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_bias_beta": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"],
+      "stats_both_bias_beta_residual": ["norm-skip-row.wgsl.jinja", "norm-stats-copy.wgsl.jinja"]
     }
   }
 }

build/webgpu/norm-skip-row-vec4.wgsl.jinja CHANGED Viewed

@@ -1,52 +1,20 @@
 {% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
-{% if op == "max" %}
-{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
-{%- else %}
-{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
-{%- endif %}
-{% endmacro %}
-{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
   var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
   loop {
-{% if form == "head" %}
-{% if breakInline %}
     if ({{ svar }} == 0u) { break; }
-{% else %}
-    if ({{ svar }} == 0u) {
-      break;
-    }
-{% endif %}
-{% endif %}
-{% if bodyInline %}
-    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
-{% else %}
     if ({{ idx }} < {{ svar }}) {
 {% for a in arrays %}
       {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
 {% endfor %}
     }
-{% endif %}
-{% if form == "head" %}
-{% if barrierFirst %}
-    workgroupBarrier();
-    {{ svar }} = {{ svar }} / 2u;
-{% else %}
     {{ svar }} = {{ svar }} / 2u;
     workgroupBarrier();
-{% endif %}
-{% else %}
-    workgroupBarrier();
-    if ({{ svar }} == 1u) {
-      break;
-    }
-    {{ svar }} = {{ svar }} / 2u;
-{% endif %}
-  }
-{%- endmacro %}{% set broadcastSkip = broadcastSkip is defined and broadcastSkip %}
-{% set useSubgroups = useSubgroups %}
-{% if usesF16Spec %}
-enable f16;
-{% endif %}
 {% if useSubgroups %}
 enable subgroups;
 {% endif %}

 {% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" or op == "min" %}
+{{ a }}[{{ idx }}] = {{ op }}({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);{% else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] {{ "*" if op == "prod" else "+" }} {{ a }}[{{ idx }} + {{ svar }}];{% endif %}{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false, reuse=false) %}
   var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
   loop {
     if ({{ svar }} == 0u) { break; }
     if ({{ idx }} < {{ svar }}) {
 {% for a in arrays %}
       {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
 {% endfor %}
     }
     {{ svar }} = {{ svar }} / 2u;
     workgroupBarrier();
+  }{% endmacro %}
+{% set broadcastSkip = broadcastSkip is defined and broadcastSkip %}
 {% if useSubgroups %}
 enable subgroups;
 {% endif %}

build/webgpu/norm-skip-row.wgsl.jinja CHANGED Viewed

@@ -1,61 +1,30 @@
 {% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
-{% if op == "max" %}
-{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
-{%- else %}
-{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
-{%- endif %}
-{% endmacro %}
-{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
   var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
   loop {
-{% if form == "head" %}
-{% if breakInline %}
-    if ({{ svar }} == 0u) { break; }
-{% else %}
     if ({{ svar }} == 0u) {
       break;
     }
-{% endif %}
-{% endif %}
-{% if bodyInline %}
-    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
-{% else %}
     if ({{ idx }} < {{ svar }}) {
 {% for a in arrays %}
       {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
 {% endfor %}
     }
-{% endif %}
-{% if form == "head" %}
-{% if barrierFirst %}
-    workgroupBarrier();
     {{ svar }} = {{ svar }} / 2u;
-{% else %}
-    {{ svar }} = {{ svar }} / 2u;
-    workgroupBarrier();
-{% endif %}
-{% else %}
     workgroupBarrier();
-    if ({{ svar }} == 1u) {
-      break;
-    }
-    {{ svar }} = {{ svar }} / 2u;
-{% endif %}
-  }
-{%- endmacro %}
-/* One workgroup normalizes each row of residual = input + skip, with an
- * optional bias. */
 {% set degenerateRow = (not simplified) and hiddenSize == 1 %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {% if useSubgroups and not degenerateRow %}
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% if not degenerateRow or writeResidualSum %}
 const HIDDEN: u32 = {{ hiddenSize }}u;
 {% endif %}
 const WG: u32 = {{ workgroupSize }}u;
@@ -88,8 +57,8 @@ fn reduce_pair(value: vec2<f32>, tid: u32) -> vec2<f32> {
 }
 {% endif %}
 {% endif %}
-{% if not degenerateRow or writeResidualSum %}
 fn residual_value(row: u32, d: u32) -> f32 {
   let index = row * HIDDEN + d;
   var value = f32(input[index]) + f32(skip[index]);
@@ -102,16 +71,19 @@ fn residual_value(row: u32, d: u32) -> f32 {
 @compute @workgroup_size(WG, 1, 1)
 fn main(
-  @builtin(workgroup_id) wg: vec3<u32>{% if not degenerateRow %},
   @builtin(local_invocation_id) lid: vec3<u32>{% endif %}{% if useSubgroups and not degenerateRow %},
   @builtin(subgroup_invocation_id) sg_lane: u32,
   @builtin(subgroup_id) sg_id: u32,
   @builtin(num_subgroups) num_sg: u32{% endif %}
 ) {
-  // 2D-folded row index: wg.y carries the high bits past the per-axis dispatch fold width.
-  // Reduces to wg.x when the dispatch does not fold;
-  // the row >= params.rows guard drops the over-dispatched tail.
   let row = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
   if (row >= params.rows) {
     return;
   }
@@ -124,6 +96,14 @@ fn main(
   // the variance are exactly zero and the output reduces to beta. The closed
   // form avoids computing that zero by subtracting two equal rounded values.
   let row_inv = inverseSqrt(params.epsilon);
 {% if writeResidualSum %}
   let residual = residual_value(row, 0u);
   input_skip_bias_sum[row] = {{ scalar }}(residual);
@@ -152,6 +132,14 @@ fn main(
   let row_mean = shift + mean_delta;
   let variance = max(totals.y / f32(HIDDEN) - mean_delta * mean_delta, 0.0);
   let row_inv = inverseSqrt(variance + params.epsilon);
   for (var d = tid; d < HIDDEN; d = d + WG) {
     let index = row * HIDDEN + d;
     let residual = residual_value(row, d);

 {% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" or op == "min" %}
+{{ a }}[{{ idx }}] = {{ op }}({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);{% else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] {{ "*" if op == "prod" else "+" }} {{ a }}[{{ idx }} + {{ svar }}];{% endif %}{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false, reuse=false) %}
   var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
   loop {
     if ({{ svar }} == 0u) {
       break;
     }
     if ({{ idx }} < {{ svar }}) {
 {% for a in arrays %}
       {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
 {% endfor %}
     }
     {{ svar }} = {{ svar }} / 2u;
     workgroupBarrier();
+  }{% endmacro %}
+/* Normalize residual = input + skip, with an optional bias. Reductions use
+ * one workgroup per row; closed-form one-element rows use one invocation. */
 {% set degenerateRow = (not simplified) and hiddenSize == 1 %}
 {% if useSubgroups and not degenerateRow %}
 enable subgroups;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
+{% if not degenerateRow or writeResidualSum or (writeMean is defined and writeMean) %}
 const HIDDEN: u32 = {{ hiddenSize }}u;
 {% endif %}
 const WG: u32 = {{ workgroupSize }}u;
 }
 {% endif %}
 {% endif %}
+{% if not degenerateRow or writeResidualSum or (writeMean is defined and writeMean) %}
 fn residual_value(row: u32, d: u32) -> f32 {
   let index = row * HIDDEN + d;
   var value = f32(input[index]) + f32(skip[index]);
 @compute @workgroup_size(WG, 1, 1)
 fn main(
+  @builtin({{ "global_invocation_id" if degenerateRow else "workgroup_id" }}) {{ "gid" if degenerateRow else "wg" }}: vec3<u32>{% if not degenerateRow %},
   @builtin(local_invocation_id) lid: vec3<u32>{% endif %}{% if useSubgroups and not degenerateRow %},
   @builtin(subgroup_invocation_id) sg_lane: u32,
   @builtin(subgroup_id) sg_id: u32,
   @builtin(num_subgroups) num_sg: u32{% endif %}
 ) {
+  // Fold the row grid across workgroups; independent rows also include the
+  // invocation offset. The bounds guard drops the final dispatch tail.
+{% if degenerateRow %}
+  let row = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
+{% else %}
   let row = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
+{% endif %}
   if (row >= params.rows) {
     return;
   }
   // the variance are exactly zero and the output reduces to beta. The closed
   // form avoids computing that zero by subtracting two equal rounded values.
   let row_inv = inverseSqrt(params.epsilon);
+{% if packedStatistics is defined and packedStatistics %}
+  row_stats[row] = vec2<f32>(residual_value(row, 0u), row_inv);
+{% elif writeMean is defined and writeMean %}
+  mean[row] = residual_value(row, 0u);
+{% endif %}
+{% if writeInvStd is defined and writeInvStd and not (packedStatistics is defined and packedStatistics) %}
+  inv_std_var[row] = row_inv;
+{% endif %}
 {% if writeResidualSum %}
   let residual = residual_value(row, 0u);
   input_skip_bias_sum[row] = {{ scalar }}(residual);
   let row_mean = shift + mean_delta;
   let variance = max(totals.y / f32(HIDDEN) - mean_delta * mean_delta, 0.0);
   let row_inv = inverseSqrt(variance + params.epsilon);
+{% if packedStatistics is defined and packedStatistics %}
+  if (tid == 0u) { row_stats[row] = vec2<f32>(row_mean, row_inv); }
+{% elif writeMean is defined and writeMean %}
+  if (tid == 0u) { mean[row] = row_mean; }
+{% endif %}
+{% if writeInvStd is defined and writeInvStd and not (packedStatistics is defined and packedStatistics) %}
+  if (tid == 0u) { inv_std_var[row] = row_inv; }
+{% endif %}
   for (var d = tid; d < HIDDEN; d = d + WG) {
     let index = row * HIDDEN + d;
     let residual = residual_value(row, d);

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

	@@ -0,0 +1,15 @@

+{% 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 }}
+// Separate optional outputs keep the reduction within eight storage bindings.
+@compute @workgroup_size({{ workgroupSize }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(workgroupSize, "row", "params.rows", guardInline=true) }}
+  let stats = row_stats[row];
+  mean[row] = stats.x;
+  inv_std_var[row] = stats.y;
+}

build/webgpu/test.json CHANGED Viewed

@@ -476,7 +476,7 @@
     {
       "name": "no_beta_output_only_hidden6_unaligned_row",
       "provenance": {
-        "notes": "Beta and optional outputs are omitted, and hidden size 6 is not divisible by four, selecting the scalar no-beta output-only row path on every tier."
       },
       "attrs": { "epsilon": 0.00001 },
       "inputs": {
@@ -531,7 +531,7 @@
     {
       "name": "beta_no_bias_output_only_hidden6_unaligned_row",
       "provenance": {
-        "notes": "Beta is present, bias and optional outputs are omitted, and hidden size 6 selects the scalar output-only row path on every tier."
       },
       "attrs": { "epsilon": 0.00001 },
       "inputs": {
@@ -1172,6 +1172,754 @@
         }
       },
       "outputs": { "outputT": { "dtype": "float16", "shape": [3, 8], "tolerance": 0.002 } }
     }
   ]
 }

     {
       "name": "no_beta_output_only_hidden6_unaligned_row",
       "provenance": {
+        "notes": "Beta and optional outputs are omitted; hidden size six exercises a non-four-aligned row."
       },
       "attrs": { "epsilon": 0.00001 },
       "inputs": {
     {
       "name": "beta_no_bias_output_only_hidden6_unaligned_row",
       "provenance": {
+        "notes": "Beta is present, bias and optional outputs are omitted, and hidden size 6 leaves a partial four-element storage group."
       },
       "attrs": { "epsilon": 0.00001 },
       "inputs": {
         }
       },
       "outputs": { "outputT": { "dtype": "float16", "shape": [3, 8], "tolerance": 0.002 } }
+    },
+    {
+      "name": "independent_rows_float32_257x1_option0_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_257x1_option1_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_257x1_option2_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_3x31x1_option2_wg1_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 1 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_3x31x1_option2_wg8_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 8 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_3x31x1_option2_wg64_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 64 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_3x31x1_option2_wg128_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 128 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float32_524289x1_option3_wg8_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 8 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [524289, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [524289, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [524289, 1], "tolerance": 0, "allowNaN": false },
+        "residualT": { "dtype": "float32", "shape": [524289, 1], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "independent_rows_float32_257x1_option2_wgdefault_epsilon0",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "independent_rows_float16_257x1_option0_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [257, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_257x1_option1_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [257, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_257x1_option2_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [257, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_3x31x1_option2_wg1_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 1 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_3x31x1_option2_wg8_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 8 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_3x31x1_option2_wg64_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 64 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_3x31x1_option2_wg128_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 128 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [3, 31, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [3, 31, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_524289x1_option2_wg8_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "tunables": { "MAX_WORKGROUP_SIZE": 8 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [524289, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [524289, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [524289, 1], "tolerance": 0, "allowNaN": false } }
+    },
+    {
+      "name": "independent_rows_float16_257x1_option2_wgdefault_epsilon0",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float16", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [257, 1], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "independent_rows_float32_257x1_option3_wgdefault_epsilon0.00001",
+      "provenance": {
+        "notes": "One-element normalization has an intrinsically uniform result. Uniform inputs pin its closed-form value, workgroup tails, optional outputs, and epsilon behavior; the GPU row-coverage test checks poisoned outputs and varying residuals separately."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "skipT": { "dtype": "float32", "shape": [257, 1], "data": { "kind": "constant", "value": 0.5 } },
+        "gammaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.125] } },
+        "betaT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.25] } },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0, "allowNaN": false },
+        "residualT": { "dtype": "float32", "shape": [257, 1], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "ort_skip_layer_norm_large_magnitude_row",
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [10000.0, 10001.0, 9999.0, 10000.0] }
+        },
+        "skipT": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 4], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "stats_mean_plain",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_mean_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_mean_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "betaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_mean_beta_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "betaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_mean_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_mean_bias_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_mean_bias_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_mean_bias_beta_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_inv_plain",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_inv_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_inv_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "betaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_inv_beta_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "betaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_inv_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_inv_bias_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_inv_bias_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_inv_bias_beta_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_both_plain",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_both_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_both_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "betaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_both_beta_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "betaT": { "dtype": "float32", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_both_bias",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_both_bias_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_both_bias_beta",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "stats_both_bias_beta_residual",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [5], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [2, 3, 5], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "ort_saved_statistics",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc",
+        "test": "SkipLayerNormStatistics"
+      },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [10000.0, 10001.0, 9999.0, 10000.0] }
+        },
+        "skipT": { "dtype": "float32", "shape": [1, 1, 4], "data": { "kind": "constant", "value": 0.0 } },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 1.0 } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.0, 1.4142135, -1.4142135, 0.0] }
+        },
+        "meanT": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "values", "values": [10000.0] } },
+        "invStdT": {
+          "dtype": "float32",
+          "shape": [1, 1, 1],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [1.4142135] }
+        }
+      }
+    },
+    {
+      "name": "stats_single_element_row",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [1, 1, 1], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [1, 1, 1], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "invStdT": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_single_element_mean",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [1, 1, 1], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [1, 1, 1], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.002, "relTolerance": 0.002 },
+        "meanT": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "stats_single_element_inverse",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputT": { "dtype": "float16", "shape": [1, 1, 1], "data": { "kind": "linspace", "start": -2.0, "end": 3.0 } },
+        "skipT": { "dtype": "float16", "shape": [1, 1, 1], "data": { "kind": "linspace", "start": 0.4, "end": -0.7 } },
+        "gammaT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": 0.5, "end": 1.5 } },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": -0.2, "end": 0.3 } },
+        "betaT": { "dtype": "float16", "shape": [1], "data": { "kind": "linspace", "start": 0.2, "end": -0.3 } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.002, "relTolerance": 0.002 },
+        "invStdT": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00001, "relTolerance": 0.00001 },
+        "residualT": { "dtype": "float16", "shape": [1, 1, 1], "tolerance": 0.002 }
+      }
     }
   ]
 }