sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +74 -0
build/webgpu/bench.json +200 -0
build/webgpu/layer-normalization.wgsl.jinja +192 -0
build/webgpu/manifest.json +975 -0
build/webgpu/metadata.json +19 -0
build/webgpu/norm-row-stats.wgsl.jinja +155 -0
build/webgpu/test.json +1786 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.LayerNormalization
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 17
+## Description
+Normalizes a tensor along a suffix of axes starting at `axis` by subtracting the mean and dividing by the square root of the variance plus `epsilon`, then scales and optionally shifts the result with learnable `Scale` and `B` tensors. The output `Y` has the same shape as `X`; optional outputs `Mean` and `InvStdDev` expose the per-normalization-group statistics computed during normalization.
+See the [ONNX `LayerNormalization` spec](https://onnx.ai/onnx/operators/onnx__LayerNormalization.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Tensor to be normalized. | required |
+| `Scale` | `scale` | `T` | — | — | Scale tensor applied after normalization. | required |
+| `B` | `b` | `T` | — | — | Optional bias tensor added after scaling. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | same as `X` | Normalized and scaled output tensor; same shape as X. | required |
+| `Mean` | `mean` | `float32` | same as `X` | — | Per-normalization-group mean in the ONNX broadcastable keepdims shape: dimensions before `axis` are preserved and dimensions from `axis` onward are 1. | optional |
+| `InvStdDev` | `invStdDev` | `float32` | same as `X` | — | Per-normalization-group reciprocal standard deviation `1 / sqrt(variance + epsilon)`, returned in the same ONNX broadcastable keepdims shape as `Mean`. | optional |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `-1` | The first axis of the normalization range; all axes from `axis` to the last are normalized together. Negative values count from the end; the default `-1` normalizes only the last dimension. |
+| `epsilon` | `0.00001` | Small constant added to the variance before taking the square root to avoid division by zero. |
+| `stash_type` | `1` | TensorProto element type used for the normalization stage and optional statistics; the implemented ONNX route supports the standard float32 value (`1`). |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`layer-normalization.wgsl.jinja`](build/webgpu/layer-normalization.wgsl.jinja)
+- [`norm-row-stats.wgsl.jinja`](build/webgpu/norm-row-stats.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.LayerNormalization", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 4] },
+  scale: { data: scaleData, shape: [4] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,200 @@

+{
+  "op": "ai.onnx.LayerNormalization",
+  "cases": [
+    {
+      "name": "layernorm-f32-256x1024",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 256, "dim": 1024 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [256, 1024], "dtype": "float32", "dist": "normal", "seed": 720, "scale": 0.5 },
+        "scale": { "shape": [1024], "dtype": "float32", "dist": "uniform", "seed": 721, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [1024], "dtype": "float32", "dist": "normal", "seed": 722, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [256, 1024], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-4096x4096",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 4096, "dim": 4096 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 730, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float32", "dist": "uniform", "seed": 731, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 732, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [4096, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f16-4096x4096",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "rows": 4096, "dim": 4096 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 4096], "dtype": "float16", "dist": "normal", "seed": 733, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float16", "dist": "uniform", "seed": 734, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [4096], "dtype": "float16", "dist": "normal", "seed": 735, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [4096, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-4096x2050-unaligned-cliff",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 4096, "dim": 2050 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 2050], "dtype": "float32", "dist": "normal", "seed": 740, "scale": 0.5 },
+        "scale": { "shape": [2050], "dtype": "float32", "dist": "uniform", "seed": 741, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [2050], "dtype": "float32", "dist": "normal", "seed": 742, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [4096, 2050], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-4096x1022-subgroup-nonvec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 4096, "dim": 1022 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 1022], "dtype": "float32", "dist": "normal", "seed": 743, "scale": 0.5 },
+        "scale": { "shape": [1022], "dtype": "float32", "dist": "uniform", "seed": 744, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [1022], "dtype": "float32", "dist": "normal", "seed": 745, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [4096, 1022], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f16-1x4096-rows1-decode",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "rows": 1, "dim": 4096 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [1, 4096], "dtype": "float16", "dist": "normal", "seed": 746, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float16", "dist": "uniform", "seed": 747, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [4096], "dtype": "float16", "dist": "normal", "seed": 748, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-512x4096-suffix-axis1-rank3",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "outer": 512, "hidden": 4096 },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": { "shape": [512, 256, 16], "dtype": "float32", "dist": "normal", "seed": 749, "scale": 0.5 },
+        "scale": { "shape": [256, 16], "dtype": "float32", "dist": "uniform", "seed": 750, "scale": 0.25, "offset": 1 },
+        "b": { "shape": [256, 16], "dtype": "float32", "dist": "normal", "seed": 751, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [512, 256, 16], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.outer * args.hidden * 2 + args.hidden * 2) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-8192x4096-broadcast-scale-generic-fallback",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "rows": 8192, "dim": 4096 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [8192, 4096], "dtype": "float32", "dist": "normal", "seed": 760, "scale": 0.5 },
+        "scale": {
+          "shape": [8192, 4096],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 761,
+          "scale": 0.25,
+          "offset": 1
+        }
+      },
+      "outputs": { "y": { "shape": [8192, 4096], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.rows * args.dim) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-65536x64-broadcast-scale-narrow-hidden-launchbound",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "rows": 65536, "dim": 64 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [65536, 64], "dtype": "float32", "dist": "normal", "seed": 762, "scale": 0.5 },
+        "scale": {
+          "shape": [65536, 64],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 763,
+          "scale": 0.25,
+          "offset": 1
+        }
+      },
+      "outputs": { "y": { "shape": [65536, 64], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.rows * args.dim) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "layernorm-f32-4096x4096-broadcast-bias-generic-fallback",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "rows": 4096, "dim": 4096 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 764, "scale": 0.5 },
+        "scale": {
+          "shape": [4096, 4096],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 765,
+          "scale": 0.25,
+          "offset": 1
+        },
+        "b": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 766, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [4096, 4096], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "(args.rows * args.dim * 2 + args.rows * args.dim * 2) * dtypeBytes(args.dtype)"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/layer-normalization.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,192 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
+fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
+{% if out_numel == 0 %}
+  return 0u;
+{% elif op_numel == 1 %}
+  return 0u;
+{% elif op_same %}
+  return out_index;
+{% else %}
+  var offset = 0u;
+{% for axis in range(outRank) %}
+{% set op_axis = axis - (outRank - opRank) %}
+{% if op_axis >= 0 and opShape[op_axis] != 1 %}
+{% set c_stride = namespace(value=1) %}
+{% for j in range(axis + 1, outRank) %}
+{% set c_stride.value = c_stride.value * outShape[j] %}
+{% endfor %}
+{% set op_stride = namespace(value=1) %}
+{% for j in range(op_axis + 1, opRank) %}
+{% set op_stride.value = op_stride.value * opShape[j] %}
+{% endfor %}
+{% if c_stride.value == 1 %}
+  let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
+{% else %}
+  let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
+{% endif %}
+{% if op_stride.value == 1 %}
+  offset = offset + coord{{ axis }};
+{% else %}
+  offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
+{% endif %}
+{% endif %}
+{% endfor %}
+  return offset;
+{% endif %}
+}
+{%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
+{% set op_numel = namespace(value=1) %}
+{% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
+{% set out_numel = namespace(value=1) %}
+{% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
+{{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
+{%- endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const EPSILON: f32 = {{ epsilon }};
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> partial: array<f32, WG>;
+var<workgroup> row_mean: f32;
+var<workgroup> row_inv: f32;
+{% set xNumel = namespace(value=1) %}
+{% for dim in source.xShape %}
+{% set xNumel.value = xNumel.value * dim %}
+{% endfor %}
+{% set scaleNumel = namespace(value=1) %}
+{% for dim in source.scaleShape %}
+{% set scaleNumel.value = scaleNumel.value * dim %}
+{% endfor %}
+{% if scaleNumel.value != 1 %}
+{{ offset_fn("scale_offset", source.scaleShape, source.scaleShape | length, source.scaleShape == source.xShape, scaleNumel.value, source.xShape, source.xShape | length, xNumel.value) }}
+{% endif %}
+{% if hasBias %}
+{% set biasNumel = namespace(value=1) %}
+{% for dim in source.biasShape %}
+{% set biasNumel.value = biasNumel.value * dim %}
+{% endfor %}
+{% if biasNumel.value != 1 %}
+{{ offset_fn("bias_offset", source.biasShape, source.biasShape | length, source.biasShape == source.xShape, biasNumel.value, source.xShape, source.xShape | length, xNumel.value) }}
+{% endif %}
+{% endif %}
+{% 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 %}
+// Reusing partial after this reduction requires a barrier between the read of
+// partial[0] and the next write, or the next round can race the prior readers.
+{% set trailingBarrier = trailingBarrier is defined and trailingBarrier %}
+fn reduce_sum(value: f32, tid: u32) -> f32 {
+  partial[tid] = value;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["partial"], idx="tid", wg="WG", form="head") }}
+{% if trailingBarrier %}
+  let total = partial[0];
+  workgroupBarrier();
+  return total;
+{% else %}
+  return partial[0];
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row = wg.x + wg.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let tid = lid.x;
+  let base = row * HIDDEN;
+  var local_sum = 0.0;
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let value = f32(x[base + d]);
+    local_sum = local_sum + value;
+  }
+  let sum = reduce_sum(local_sum, tid);
+  if (tid == 0u) {
+    row_mean = sum / f32(HIDDEN);
+  }
+  workgroupBarrier();
+  var local_var_sum = 0.0;
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let diff = f32(x[base + d]) - row_mean;
+    local_var_sum = local_var_sum + diff * diff;
+  }
+  let var_sum = reduce_sum(local_var_sum, tid);
+  if (tid == 0u) {
+    let variance = var_sum / f32(HIDDEN);
+    row_inv = inverseSqrt(variance + EPSILON);
+{% if writeMean %}
+    mean_out[row] = row_mean;
+{% endif %}
+{% if writeInvStdDev %}
+    inv_std_out[row] = row_inv;
+{% endif %}
+  }
+  workgroupBarrier();
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let index = base + d;
+    let normalized = (f32(x[index]) - row_mean) * row_inv;
+    var value = normalized * f32(scale[{% if scaleNumel.value == 1 %}0u{% else %}{{ broadcast_offset_call("scale_offset", source.scaleShape, source.xShape, "index") }}{% endif %}]);
+{% if hasBias %}
+    value = value + f32(bias[{% if biasNumel.value == 1 %}0u{% else %}{{ broadcast_offset_call("bias_offset", source.biasShape, source.xShape, "index") }}{% endif %}]);
+{% endif %}
+    y[index] = {{ scalar }}(value);
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,975 @@

+{
+  "domain": "ai.onnx",
+  "name": "LayerNormalization",
+  "sinceVersion": 17,
+  "description": "Normalizes a tensor along a suffix of axes starting at `axis` by subtracting the mean and dividing by the square root of the variance plus `epsilon`, then scales and optionally shifts the result with learnable `Scale` and `B` tensors. The output `Y` has the same shape as `X`; optional outputs `Mean` and `InvStdDev` expose the per-normalization-group statistics computed during normalization.",
+  "inputs": [
+    { "role": "X", "dtype": "T", "description": "Tensor to be normalized." },
+    { "role": "Scale", "dtype": "T", "description": "Scale tensor applied after normalization." },
+    { "role": "B", "dtype": "T", "optional": true, "description": "Optional bias tensor added after scaling." }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.X",
+      "description": "Normalized and scaled output tensor; same shape as X.",
+      "shape": "shapes.X"
+    },
+    {
+      "role": "Mean",
+      "dtype": "float32",
+      "optional": true,
+      "description": "Per-normalization-group mean in the ONNX broadcastable keepdims shape: dimensions before `axis` are preserved and dimensions from `axis` onward are 1.",
+      "rank": "ranks.X"
+    },
+    {
+      "role": "InvStdDev",
+      "dtype": "float32",
+      "optional": true,
+      "description": "Per-normalization-group reciprocal standard deviation `1 / sqrt(variance + epsilon)`, returned in the same ONNX broadcastable keepdims shape as `Mean`.",
+      "rank": "ranks.X"
+    }
+  ],
+  "attributes": { "axis": -1, "epsilon": 0.00001, "stash_type": 1 },
+  "attributeDescriptions": {
+    "axis": "The first axis of the normalization range; all axes from `axis` to the last are normalized together. Negative values count from the end; the default `-1` normalizes only the last dimension.",
+    "epsilon": "Small constant added to the variance before taking the square root to avoid division by zero.",
+    "stash_type": "TensorProto element type used for the normalization stage and optional statistics; the implemented ONNX route supports the standard float32 value (`1`)."
+  },
+  "attributeConstraints": { "stash_type": { "values": [1] } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "scale": { "kind": "tensor", "semantic": "Scale", "role": "input" },
+    "b": { "kind": "tensor", "semantic": "B", "role": "input", "required": false },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" },
+    "mean": { "kind": "tensor", "semantic": "Mean", "role": "output", "required": false },
+    "invStdDev": { "kind": "tensor", "semantic": "InvStdDev", "role": "output", "required": false }
+  },
+  "tunables": { "MAX_WORKGROUP_SIZE": 256, "SCALAR_FAST_MAX_HIDDEN": 1024 },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "normWorkgroupCap": "min(tunables.MAX_WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "hasSubgroupId": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "lastAxisWg": "min(normWorkgroupCap, pow2ceil(dim(shapes.X, -1)))",
+    "lastAxisWgVec4": "min(normWorkgroupCap, pow2ceil(dim(shapes.X, -1) / 4))",
+    "lastAxisContractOk": "ranks.X >= 1 and ranks.Y == ranks.X and numel(shapes.X) == numel(shapes.Y) and (attrs.axis == -1 or attrs.axis == ranks.X - 1)",
+    "suffixAxisContractOk": "ranks.X >= 2 and ranks.Y == ranks.X and numel(shapes.X) == numel(shapes.Y) and attrs.axis + ranks.X >= 0 and attrs.axis < ranks.X and not (attrs.axis == -1 or attrs.axis == ranks.X - 1)",
+    "axisNorm": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.X",
+    "normRows": "numel(shapes.X) / max(1, dim(shapes.X, -1)) if lastAxisContractOk else outer(shapes.X, axisNorm)",
+    "normRowStride": "max(1, min(normRows, device.limits.maxComputeWorkgroupsPerDimension))",
+    "suffixAxisSize": "numel(shapes.X) / max(1, outer(shapes.X, axisNorm))",
+    "suffixAxisWg": "min(normWorkgroupCap, pow2ceil(suffixAxisSize))",
+    "suffixAxisWgVec4": "min(normWorkgroupCap, pow2ceil(suffixAxisSize / 4))",
+    "genericHiddenSize": "dim(shapes.X, -1) if lastAxisContractOk else suffixAxisSize",
+    "genericWorkgroupSize": "lastAxisWg if lastAxisContractOk else suffixAxisWg",
+    "scaleExactOk": "ranks.X >= 1 and ranks.Scale >= 1 and numel(shapes.Scale) == dim(shapes.X, -1) and dim(shapes.Scale, -1) == dim(shapes.X, -1)",
+    "scaleBroadcastOk": "ranks.Scale >= 0 and ranks.Scale <= ranks.X and broadcastable(shapes.Scale, shapes.X)",
+    "biasExactOk": "present.b and ranks.X >= 1 and ranks.B >= 1 and numel(shapes.B) == dim(shapes.X, -1) and dim(shapes.B, -1) == dim(shapes.X, -1)",
+    "biasBroadcastOk": "present.b and ranks.B >= 0 and ranks.B <= ranks.X and broadcastable(shapes.B, shapes.X)",
+    "suffixScaleExactOk": "suffixAxisContractOk and scaleBroadcastOk and numel(shapes.Scale) == suffixAxisSize",
+    "suffixBiasExactOk": "present.b and suffixAxisContractOk and biasBroadcastOk and numel(shapes.B) == suffixAxisSize",
+    "lastAxisExactScaleOk": "lastAxisContractOk and scaleExactOk",
+    "lastAxisBroadcastScaleOk": "lastAxisContractOk and scaleBroadcastOk",
+    "suffixAxisBroadcastScaleOk": "suffixAxisContractOk and scaleBroadcastOk",
+    "suffixAxisExactAffineOk": "suffixAxisContractOk and suffixScaleExactOk and suffixBiasExactOk",
+    "lastAxisScalarFastOk": "dtypes.T == \"f16\" or dim(shapes.X, -1) <= tunables.SCALAR_FAST_MAX_HIDDEN",
+    "noStatsOutputs": "not present.mean and not present.invStdDev",
+    "meanOnlyOutputs": "present.mean and not present.invStdDev",
+    "invStdOnlyOutputs": "not present.mean and present.invStdDev",
+    "fullStatsOutputs": "present.mean and present.invStdDev",
+    "statsRowsOk": "fullStatsOutputs and ranks.X >= 1 and numel(shapes.Mean) == normRows and numel(shapes.InvStdDev) == normRows",
+    "meanRowsOk": "present.mean and ranks.X >= 1 and numel(shapes.Mean) == normRows",
+    "invStdRowsOk": "present.invStdDev and ranks.X >= 1 and numel(shapes.InvStdDev) == normRows",
+    "statsOuterOk": "fullStatsOutputs and ranks.X >= 2 and numel(shapes.Mean) == normRows and numel(shapes.InvStdDev) == normRows"
+  },
+  "bindingSets": {
+    "vec4Affine": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "vec4AffineBias": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "b",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "vec4AffineStats": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      { "name": "mean_out", "arg": "mean", "semantic": "Mean", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "inv_std_out",
+        "arg": "invStdDev",
+        "semantic": "InvStdDev",
+        "buffer": { "type": "storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "vec4AffineBiasStats": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "b",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
+      { "name": "mean_out", "arg": "mean", "semantic": "Mean", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "inv_std_out",
+        "arg": "invStdDev",
+        "semantic": "InvStdDev",
+        "buffer": { "type": "storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "scalarAffineMean": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      { "name": "mean_out", "arg": "mean", "semantic": "Mean", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "scalarAffineInvStd": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "inv_std_out",
+        "arg": "invStdDev",
+        "semantic": "InvStdDev",
+        "buffer": { "type": "storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "scalarAffineBiasMean": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "bias",
+        "arg": "b",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      { "name": "mean_out", "arg": "mean", "semantic": "Mean", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ],
+    "scalarAffineBiasInvStd": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "Scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "bias",
+        "arg": "b",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "inv_std_out",
+        "arg": "invStdDev",
+        "semantic": "InvStdDev",
+        "buffer": { "type": "storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normRows" },
+            { "name": "rowStride", "type": "u32", "value": "normRowStride" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "last_axis_row_vec4",
+      "priority": 110,
+      "when": ["f16Ok(dtypes.T)", "not present.b and noStatsOutputs", "lastAxisExactScaleOk", "dim(shapes.X, -1) % 4 == 0"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRowVec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": true,
+              "hasBias": false,
+              "writeStats": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWgVec4",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": "dim(shapes.X, -1) / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4Affine",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_row",
+      "priority": 100,
+      "when": ["not present.b and noStatsOutputs", "lastAxisExactScaleOk", "f16Ok(dtypes.T)"],
+      "demoteWhen": ["not lastAxisScalarFastOk"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRow",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": false,
+              "hasBias": false,
+              "writeStats": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWg",
+              "epsilon": "attrs.epsilon",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4Affine",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_bias_row_vec4",
+      "priority": 111,
+      "when": ["f16Ok(dtypes.T)", "present.b and noStatsOutputs and biasExactOk", "lastAxisExactScaleOk", "dim(shapes.X, -1) % 4 == 0"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRowVec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": true,
+              "hasBias": true,
+              "writeStats": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWgVec4",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": "dim(shapes.X, -1) / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineBias",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_bias_row",
+      "priority": 101,
+      "when": ["present.b and noStatsOutputs and biasExactOk", "lastAxisExactScaleOk", "f16Ok(dtypes.T)"],
+      "demoteWhen": ["not lastAxisScalarFastOk"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRow",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": false,
+              "hasBias": true,
+              "writeStats": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWg",
+              "epsilon": "attrs.epsilon",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineBias",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_stats_row_vec4",
+      "priority": 112,
+      "when": ["f16Ok(dtypes.T)", "not present.b and fullStatsOutputs and statsRowsOk", "lastAxisExactScaleOk", "dim(shapes.X, -1) % 4 == 0"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRowVec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": true,
+              "hasBias": false,
+              "writeStats": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWgVec4",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": "dim(shapes.X, -1) / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_stats_row",
+      "priority": 102,
+      "when": ["not present.b and fullStatsOutputs and statsRowsOk", "lastAxisExactScaleOk", "f16Ok(dtypes.T)"],
+      "demoteWhen": ["not lastAxisScalarFastOk"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRow",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": false,
+              "hasBias": false,
+              "writeStats": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWg",
+              "epsilon": "attrs.epsilon",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_bias_stats_row_vec4",
+      "priority": 113,
+      "when": ["f16Ok(dtypes.T)", "present.b and fullStatsOutputs and biasExactOk and statsRowsOk", "lastAxisExactScaleOk", "dim(shapes.X, -1) % 4 == 0"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRowVec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": true,
+              "hasBias": true,
+              "writeStats": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWgVec4",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": "dim(shapes.X, -1) / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineBiasStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_bias_stats_row",
+      "priority": 103,
+      "when": ["present.b and fullStatsOutputs and biasExactOk and statsRowsOk", "lastAxisExactScaleOk", "f16Ok(dtypes.T)"],
+      "demoteWhen": ["not lastAxisScalarFastOk"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.LastAxisRow",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": false,
+              "hasBias": true,
+              "writeStats": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "lastAxisWg",
+              "epsilon": "attrs.epsilon",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineBiasStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis_bias_exact_row_vec4",
+      "priority": 121,
+      "when": ["f16Ok(dtypes.T)", "present.b", "noStatsOutputs", "suffixAxisExactAffineOk", "suffixAxisSize % 4 == 0"],
+      "constants": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.SuffixAxisRowVec4",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"layer\"",
+              "vec4": true,
+              "hasBias": true,
+              "writeStats": false,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "suffixAxisSize",
+              "wg": "suffixAxisWgVec4",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": "suffixAxisSize / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "vec4AffineBias",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis",
+      "priority": 0,
+      "when": ["not present.b", "noStatsOutputs", "lastAxisBroadcastScaleOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": false,
+        "writeMean": false,
+        "writeInvStdDev": false,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "dim(shapes.X, -1)",
+        "workgroupSize": "lastAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale" }
+          },
+          "bindings": "vec4Affine",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_bias",
+      "priority": 10,
+      "when": ["present.b", "noStatsOutputs", "lastAxisBroadcastScaleOk", "biasBroadcastOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": true,
+        "writeMean": false,
+        "writeInvStdDev": false,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "dim(shapes.X, -1)",
+        "workgroupSize": "lastAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale", "biasShape": "shapes.B" }
+          },
+          "bindings": "vec4AffineBias",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_stats",
+      "priority": 20,
+      "when": ["not present.b", "fullStatsOutputs", "lastAxisBroadcastScaleOk", "statsRowsOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": false,
+        "writeMean": true,
+        "writeInvStdDev": true,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "dim(shapes.X, -1)",
+        "workgroupSize": "lastAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale" }
+          },
+          "bindings": "vec4AffineStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_bias_stats",
+      "priority": 30,
+      "when": ["present.b", "fullStatsOutputs", "lastAxisBroadcastScaleOk", "biasBroadcastOk", "statsRowsOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": true,
+        "writeMean": true,
+        "writeInvStdDev": true,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "dim(shapes.X, -1)",
+        "workgroupSize": "lastAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale", "biasShape": "shapes.B" }
+          },
+          "bindings": "vec4AffineBiasStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis",
+      "priority": 40,
+      "when": ["not present.b", "noStatsOutputs", "suffixAxisBroadcastScaleOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": false,
+        "writeMean": false,
+        "writeInvStdDev": false,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "suffixAxisSize",
+        "workgroupSize": "suffixAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.SuffixAxis",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale" }
+          },
+          "bindings": "vec4Affine",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis_bias",
+      "priority": 50,
+      "when": ["present.b", "noStatsOutputs", "suffixAxisBroadcastScaleOk", "biasBroadcastOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": true,
+        "writeMean": false,
+        "writeInvStdDev": false,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "suffixAxisSize",
+        "workgroupSize": "suffixAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.SuffixAxisBias",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale", "biasShape": "shapes.B" }
+          },
+          "bindings": "vec4AffineBias",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis_stats",
+      "priority": 45,
+      "when": ["not present.b", "fullStatsOutputs", "suffixAxisBroadcastScaleOk", "statsOuterOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": false,
+        "writeMean": true,
+        "writeInvStdDev": true,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "suffixAxisSize",
+        "workgroupSize": "suffixAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.SuffixAxisStats",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale" }
+          },
+          "bindings": "vec4AffineStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis_bias_stats",
+      "priority": 55,
+      "when": ["present.b", "fullStatsOutputs", "suffixAxisBroadcastScaleOk", "biasBroadcastOk", "statsOuterOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": true,
+        "writeMean": true,
+        "writeInvStdDev": true,
+        "scalar": "dtypes.T",
+        "vectorScalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "suffixAxisSize",
+        "workgroupSize": "suffixAxisWg",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.SuffixAxisBiasStats",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale", "biasShape": "shapes.B" }
+          },
+          "bindings": "vec4AffineBiasStats",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "mean_only",
+      "priority": 31,
+      "when": ["not present.b and meanOnlyOutputs and meanRowsOk", "lastAxisBroadcastScaleOk or suffixAxisBroadcastScaleOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": false,
+        "writeMean": true,
+        "writeInvStdDev": false,
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "genericHiddenSize",
+        "workgroupSize": "genericWorkgroupSize",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.MeanOnly",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale" }
+          },
+          "bindings": "scalarAffineMean",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "bias_mean_only",
+      "priority": 32,
+      "when": ["present.b and meanOnlyOutputs and biasBroadcastOk and meanRowsOk", "lastAxisBroadcastScaleOk or suffixAxisBroadcastScaleOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": true,
+        "writeMean": true,
+        "writeInvStdDev": false,
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "genericHiddenSize",
+        "workgroupSize": "genericWorkgroupSize",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.BiasMeanOnly",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale", "biasShape": "shapes.B" }
+          },
+          "bindings": "scalarAffineBiasMean",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "inv_std_dev_only",
+      "priority": 33,
+      "when": ["not present.b and invStdOnlyOutputs and invStdRowsOk", "lastAxisBroadcastScaleOk or suffixAxisBroadcastScaleOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": false,
+        "writeMean": false,
+        "writeInvStdDev": true,
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "genericHiddenSize",
+        "workgroupSize": "genericWorkgroupSize",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.InvStdDevOnly",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale" }
+          },
+          "bindings": "scalarAffineInvStd",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    },
+    {
+      "id": "bias_inv_std_dev_only",
+      "priority": 34,
+      "when": ["present.b and invStdOnlyOutputs and biasBroadcastOk and invStdRowsOk", "lastAxisBroadcastScaleOk or suffixAxisBroadcastScaleOk", "f16Ok(dtypes.T)"],
+      "constants": {
+        "hasBias": true,
+        "writeMean": false,
+        "writeInvStdDev": true,
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "hiddenSize": "genericHiddenSize",
+        "workgroupSize": "genericWorkgroupSize",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "LayerNormalization.BiasInvStdDevOnly",
+          "source": {
+            "shader": "layer-normalization.wgsl.jinja",
+            "inputs": { "xShape": "shapes.X", "scaleShape": "shapes.Scale", "biasShape": "shapes.B" }
+          },
+          "bindings": "scalarAffineBiasInvStd",
+          "dispatch": { "workgroups": "normRows" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "name": "ai.onnx.LayerNormalization",
+  "id": "_ai_onnx_layernormalization_webgpu_af1721d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "fLkWXUeB9lfy8s03WLvcd2LRTBWce01Yzzv5lBHMvig=",
+      "layer-normalization.wgsl.jinja": "3BXN4VIAbNR6se39KXKK5jV514N8v756GeH37Tsuzi0=",
+      "manifest.json": "j8ogYGGt7G1Pr0St9Gyhpff035GfQKHmuz6IBktIrV8=",
+      "norm-row-stats.wgsl.jinja": "82e5r5vFGd0ylf/r3n+fucFbRPPMH3II1VBaTTKSkp4=",
+      "test.json": "hjj+HEnuQ0NKto1gCH5+81D7bWaH+hgQWy5OlySa7C0="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.LayerNormalization" }
+}

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

	@@ -0,0 +1,155 @@

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

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1786 @@

+{
+  "op": "ai.onnx.LayerNormalization",
+  "fixtureArrays": {
+    "onnx_backend_layer_normalization_3d_input_x": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358],
+    "onnx_backend_layer_normalization_input_x": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902, -0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834, -0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235, 1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314, -0.8707971572875977, -0.5788496732711792, -0.3115525245666504, 0.056165341287851334, -1.1651498079299927, 0.9008265137672424, 0.4656624495983124, -1.5362436771392822, 1.4882521629333496, 1.895889163017273, 1.1787796020507812, -0.1799248307943344, -1.0707526206970215, 1.0544517040252686, -0.4031769335269928, 1.222445011138916, 0.2082749754190445, 0.9766390323638916, 0.3563663959503174, 0.7065731883049011, 0.01050002034753561, 1.7858705520629883, 0.12691208720207214, 0.4019893705844879, 1.8831506967544556, -1.3477590084075928, -1.2704850435256958, 0.969396710395813, -1.1731233596801758, 1.9436211585998535, -0.4136189818382263, -0.747454822063446, 1.922942042350769, 1.4805147647857666, 1.8675589561462402, 0.9060446619987488, -0.8612256646156311, 1.910064935684204, -0.26800337433815, 0.8024563789367676, 0.9472519755363464, -0.15501008927822113, 0.6140793561935425, 0.922206699848175]
+  },
+  "cases": [
+    {
+      "name": "subgroup_vec4_stats_no_bias_2x512",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.11, "scale": 0.5 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 512], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "subgroup_vec4_bias_2x512",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.19, "scale": 0.25 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 512], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "subgroup_vec4_f16_4x32",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.27 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4, 32], "tolerance": 0.005 } }
+    },
+    {
+      "name": "subgroup_vec4_f16_stats_4x32",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.19 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.17, "scale": 0.45 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [4, 32], "tolerance": 0.006 },
+        "mean": { "dtype": "float32", "shape": [4, 1], "tolerance": 0.002 },
+        "invStdDev": { "dtype": "float32", "shape": [4, 1], "tolerance": 0.02 }
+      }
+    },
+    {
+      "name": "subgroup_vec4_f16_bias_4x32",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.25 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.27, "cosStep": 0.09, "scale": 0.45 }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.21, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4, 32], "tolerance": 0.006 } }
+    },
+    {
+      "name": "subgroup_vec4_f16_bias_stats_4x32",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.12, "cosStep": 0.29 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.25, "cosStep": 0.11, "scale": 0.4 }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [4, 32], "tolerance": 0.006 },
+        "mean": { "dtype": "float32", "shape": [4, 1], "tolerance": 0.002 },
+        "invStdDev": { "dtype": "float32", "shape": [4, 1], "tolerance": 0.02 }
+      }
+    },
+    {
+      "name": "f32_subnormal_scale_preserves_tiny_outputs_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal per-element scale collapses so tiny normalized outputs cannot be preserved."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormalization",
+        "notes": "A valid subnormal scale should produce finite subnormal normalized outputs rather than being flushed to zero."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0, 2.0] } },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, 2e-40, -3e-40, 4e-40] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "f32_subnormal_scale_preserves_tiny_outputs_odd_hidden_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal per-element scale collapses so tiny normalized outputs cannot be preserved (odd hidden size)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormalization",
+        "notes": "Odd hidden-size companion for subnormal scale values; this exercises the non-vec4 last-axis path."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [-1.0, 0.0, 2.0] } },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -2e-40, 3e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3], "tolerance": 1e-44 } }
+    },
+    {
+      "name": "last_axis_y_only",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "scale": 0.4 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "last_axis_bias_and_stats",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.35 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.17, "scale": 0.2 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 7], "tolerance": 0.00002 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.001 }
+      }
+    },
+    {
+      "name": "f32_tiny_variance_epsilon_zero_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal variance collapses to zero so inverseSqrt produces Infinity instead of a finite large InvStdDev."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm17_opset",
+        "notes": "Valid epsilon=0 edge: normal inputs produce subnormal variance but finite order-one normalized outputs and finite large InvStdDev."
+      },
+      "attrs": { "epsilon": 0, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [1e-20, -1e-20, 2e-20, -2e-20] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.00001 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1], "tolerance": 1000000000000000, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "last_axis_mean_only_no_bias",
+      "provenance": {
+        "source": "ONNX LayerNormalization-17 optional-output contract",
+        "notes": "Requests Mean while independently omitting both B and InvStdDev."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.5, -2.0, 0.25, 4.0, 10.0, 10.5, 9.75, 10.25] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, -0.5, 2.0, 0.25] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "suffix_axis_inv_std_dev_only_no_bias",
+      "provenance": {
+        "source": "ONNX LayerNormalization-17 optional-output contract",
+        "notes": "Requests InvStdDev while independently omitting the earlier optional Mean output and B."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [-3.0, -1.0, 2.0, 4.0, 5.0, 8.0, 10.0, 8.0, 7.0, 3.0, 1.0, -2.0] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.5, -1.0, 2.0, -0.25, 0.75] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "suffix_axis_bias_inv_std_dev_only",
+      "provenance": {
+        "source": "ONNX LayerNormalization-17 optional-output contract",
+        "notes": "Requests B and InvStdDev while independently omitting the earlier optional Mean output."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [-2.0, 0.0, 1.0, 3.0, 7.0, 9.0, 12.0, 9.0, 6.0, 4.0, 2.0, -1.0] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.5, -1.0, 1.5, 0.25, 2.0, -0.75] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.1, -0.2, 0.3, -0.4, 0.5, -0.6] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "last_axis_mean_only_output",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_default_axis",
+        "notes": "ONNX optional outputs may be requested as a prefix; this asks for Y and Mean without InvStdDev."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.5, -2.0, 0.25, 4.0, 10.0, 10.5, 9.75, 10.25] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, -0.5, 2.0, 0.25] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.1, -0.2, 0.5] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_values_centered_variance_regression",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_LargeValues_NoNaN"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [40000.0, 40001.0, 40002.0, 40003.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 0.00001,
+          "data": {
+            "kind": "values",
+            "values": [-1.3416354656219482, -0.4472118020057678, 0.4472118020057678, 1.3416354656219482]
+          }
+        },
+        "mean": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [40001.5] }
+        },
+        "invStdDev": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "tolerance": 0.00001,
+          "data": { "kind": "values", "values": [0.8944236040115356] }
+        }
+      }
+    },
+    {
+      "name": "ort_very_large_values_no_nan",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_VeryLargeValues_NoNaN"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [80000.0, 80001.0, 80002.0, 80003.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 0.00001,
+          "data": {
+            "kind": "values",
+            "values": [-1.3416354656219482, -0.4472118020057678, 0.4472118020057678, 1.3416354656219482]
+          }
+        },
+        "mean": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [80001.5] }
+        },
+        "invStdDev": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "tolerance": 0.00001,
+          "data": { "kind": "values", "values": [0.8944236040115356] }
+        }
+      }
+    },
+    {
+      "name": "zero_variance_bias",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 4], "data": { "kind": "constant", "value": 5.0 } },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -0.25, 1.0, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_weights_large_input",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_ConstantWeights_LargeInput"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [2396.814, 2396.814, 2396.814, 2396.814] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.1, 0.1, 0.1, 0.1] }
+        }
+      }
+    },
+    {
+      "name": "ort_norm_size1_no_bias_stats",
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_NormSize1_NoBias"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 3, 1],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 0.0, 0.0] }
+        },
+        "mean": {
+          "dtype": "float32",
+          "shape": [2, 3, 1],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "invStdDev": {
+          "dtype": "float32",
+          "shape": [2, 3, 1],
+          "tolerance": 0.001,
+          "data": {
+            "kind": "values",
+            "values": [316.2277526855469, 316.2277526855469, 316.2277526855469, 316.2277526855469, 316.2277526855469, 316.2277526855469]
+          }
+        }
+      }
+    },
+    {
+      "name": "zero_rows_noop",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 4], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "axis1_rank3_scale_matrix",
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.35 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "axis1_rank3_scale_matrix_stats",
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.25 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.09, "scale": 0.3 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 },
+        "mean": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "axis1_rank3_scale_bias_matrix",
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.15, "scale": 0.25 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.11, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_opset17_float16_last_axis_stats",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm17_opset"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [1, 2, 3], "tolerance": 0.002 },
+        "mean": { "dtype": "float32", "shape": [1, 2, 1], "tolerance": 0.002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 2, 1], "tolerance": 0.01 }
+      }
+    },
+    {
+      "name": "ort_negative_suffix_axis_rank3",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_ValidScaleBias_Broadcast"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "data": { "kind": "values", "values": [1.2416, 0.946123, 13.1685, 0.36423, 21.145, 0.03941] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [-0.6953, 5.1824, -0.25, 0.75, 1.5, -2.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [0.6435, -0.3964, 0.25, -0.75, 1.0, 2.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "axis0_scalar_stats_whole_tensor",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [10.0, -10.0, 5.0, -5.0, 2.0, -2.0] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.5, -1.0, 2.0, -0.25, 1.5] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, -1.0, 0.5, -0.5, 2.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "ort_axis2_scalar_scale_no_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Scalar_NoBias_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scalar_scale_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Bias_Scalar_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.5] } },
+        "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_broadcast_dim0_scale_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Bias_Broadcast_Dim0"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [-1.0, 2.0, -10.264, 8.6453, 3.0, -4.0, 43.1561, -0.641239, -5.0, 6.0, -8.2164, 0.11412, 7.0, 8.0, 41.3156, 3.0458]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [-0.1, 1.7, -0.6953, 5.1824] }
+        },
+        "b": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [-2.0, 0.3, 0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_broadcast_dim1_scale_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Bias_Broadcast_Dim1"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 2],
+          "data": {
+            "kind": "values",
+            "values": [-1.0, 2.0, 3.0, -4.0, -5.0, 6.0, 7.0, 8.0, -10.264, 8.6453, 43.1561, -0.641239, -8.2164, 0.11412, 41.3156, 3.0458]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "values", "values": [-0.1, 1.7, -0.6953, 5.1824] }
+        },
+        "b": { "dtype": "float32", "shape": [2, 1, 2], "data": { "kind": "values", "values": [-2.0, 0.3, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4, 2],
+          "tolerance": 0.0001,
+          "data": {
+            "kind": "values",
+            "values": [-1.9, 2.0, -2.1, -1.4, -1.9, 2.0, -1.9, 2.0, 0.6953, 5.1824, -0.6953, -5.1824, 0.6953, 5.1824, -0.6953, -5.1824]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_axis1_inner_broadcast_scale",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Broadcast_Inner_Mixed"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [1.0, 0.5, 1.0, 0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis1_scale_bias_inner_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_ValidScaleBias_Broadcast"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "data": { "kind": "values", "values": [1.2416, 0.946123, 13.1685, 0.36423, 21.145, 0.03941] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-0.6953, 5.1824] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.6435, -0.3964] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [1.063606, -3.716114, 0.042961, -4.087264, -0.639629, -4.294445] }
+        }
+      }
+    },
+    {
+      "name": "ort_axis3_scale_bias_outer_inner_broadcast_4d",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Bias_4D_OuterInnerBroadcast_Axis3"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 1.1, 1.2, 1.3] }
+        },
+        "b": { "dtype": "float32", "shape": [1, 2, 1, 2], "data": { "kind": "values", "values": [0.0, 0.5, 1.0, 1.5] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [-1.0, 1.6, -1.0, 1.6, -0.2, 2.8, -0.2, 2.8] }
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_2d_axis1",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_2d_axis1" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.001 }
+      },
+      "attrs": { "axis": 1 }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_2d_axis_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_2d_axis_negative_1"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.001 }
+      },
+      "attrs": { "axis": -1 }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_3d_axis2_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_3d_axis2_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [0.9101788997650146, 0.31721821427345276, 0.7863279581069946, -0.4664191007614136, -0.9444462656974792]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.410049706697464, -0.017020413652062416, 0.37915173172950745, 2.2593090534210205, -0.0422571524977684]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.001 }
+      },
+      "attrs": { "axis": 2, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_3d_axis_negative_1_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_3d_axis_negative_1_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.9559450149536133, -0.34598177671432495, -0.463595986366272, 0.4814814627170563, -1.5407969951629639]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [0.06326199322938919, 0.15650653839111328, 0.23218104243278503, -0.5973160862922668, -0.23792172968387604]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 1], "tolerance": 0.001 }
+      },
+      "attrs": { "axis": -1, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis3",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis3" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-1.2986867427825928, 1.2760753631591797, 1.3250139951705933, 0.20533256232738495, 0.04513401538133621]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [2.3396248817443848, -0.2764328420162201, -0.2595769762992859, 0.36448124051094055, 1.471321940422058]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 4, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 4, 1], "tolerance": 0.001 }
+      },
+      "attrs": { "axis": 3 }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis_negative_1"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [1.5927706956863403, -0.25857263803482056, 0.3083312511444092, -1.3780834674835205, -0.3119761049747467]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.840290367603302, -1.0068317651748657, 1.6815767288208008, -0.7922866344451904, -0.5316058993339539]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 4, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 4, 1], "tolerance": 0.001 }
+      },
+      "attrs": { "axis": -1 }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_default_axis",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_default_axis"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [0.37642553448677063, -1.0994007587432861, 0.29823818802833557, 1.3263858556747437, -0.694567859172821]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.14963454008102417, -0.4351535439491272, 1.8492637872695923, 0.6722947359085083, 0.40746182203292847]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 4, 1], "tolerance": 0.00001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 4, 1], "tolerance": 0.001 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_2d_axis0",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_2d_axis0" },
+      "attrs": { "axis": 0, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_2d_axis_negative_2",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_2d_axis_negative_2"
+      },
+      "attrs": { "axis": -2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [-1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_3d_axis0_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_3d_axis0_epsilon"
+      },
+      "attrs": { "axis": 0, "epsilon": 0.10000000149011612 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834, -0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235, 1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314, -0.8707971572875977, -0.5788496732711792, -0.3115525245666504, 0.056165341287851334, -1.1651498079299927, 0.9008265137672424, 0.4656624495983124, -1.5362436771392822, 1.4882521629333496, 1.895889163017273, 1.1787796020507812, -0.1799248307943344, -1.0707526206970215, 1.0544517040252686]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_3d_axis1_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_3d_axis1_epsilon"
+      },
+      "attrs": { "axis": 1, "epsilon": 0.10000000149011612 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.06824160367250443, 1.7133426666259766, -0.7447548508644104, -0.8264385461807251, -0.09845252335071564, -0.6634783148765564, 1.1266359090805054, -1.0799314975738525, -1.1474686861038208, -0.43782004714012146, -0.49803245067596436, 1.9295320510864258, 0.9494208097457886, 0.08755124360322952, -1.225435495376587]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.8443629741668701, -1.0002152919769287, -1.5447710752487183, 1.1880297660827637, 0.31694260239601135, 0.9208588004112244, 0.31872764229774475, 0.8568305969238281, -0.6510255932807922, -1.034242868423462, 0.6815944910049438, -0.8034096360206604, -0.6895498037338257, -0.4555324912071228, 0.01747915893793106]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_3d_axis_negative_2_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_3d_axis_negative_2_epsilon"
+      },
+      "attrs": { "axis": -2, "epsilon": 0.10000000149011612 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.3539939224720001, -1.3749512434005737, -0.6436184048652649, -2.223403215408325, 0.6252314448356628, -1.602057695388794, -1.1043833494186401, 0.05216507986187935, -0.73956298828125, 1.543014645576477, -1.2928569316864014, 0.2670508623123169, -0.039282817393541336, -1.1680934429168701, 0.523276686668396]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.1715463250875473, 0.7717905640602112, 0.8235041499137878, 2.163235902786255, 1.336527943611145, -0.3691818416118622, -0.2393791824579239, 1.0996595621109009, 0.6552637219429016, 0.6401315331459045, -1.6169559955596924, -0.024326125159859657, -0.7380309104919434, 0.279924601316452, -0.09815038740634918]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_3d_axis_negative_3_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_3d_axis_negative_3_epsilon"
+      },
+      "attrs": { "axis": -3, "epsilon": 0.10000000149011612 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.4031769335269928, 1.222445011138916, 0.2082749754190445, 0.9766390323638916, 0.3563663959503174, 0.7065731883049011, 0.01050002034753561, 1.7858705520629883, 0.12691208720207214, 0.4019893705844879, 1.8831506967544556, -1.3477590084075928, -1.2704850435256958, 0.969396710395813, -1.1731233596801758, 1.9436211585998535, -0.4136189818382263, -0.747454822063446, 1.922942042350769, 1.4805147647857666, 1.8675589561462402, 0.9060446619987488, -0.8612256646156311, 1.910064935684204, -0.26800337433815, 0.8024563789367676, 0.9472519755363464, -0.15501008927822113, 0.6140793561935425, 0.922206699848175]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.37642553448677063, -1.0994007587432861, 0.29823818802833557, 1.3263858556747437, -0.694567859172821, -0.14963454008102417, -0.4351535439491272, 1.8492637872695923, 0.6722947359085083, 0.40746182203292847, -0.7699160575866699, 0.5392491817474365, -0.6743326783180237, 0.0318305566906929, -0.6358460783958435, 0.676433265209198, 0.5765908360481262, -0.20829875767230988, 0.39600670337677, -1.0930615663528442, -1.4912575483322144, 0.43939170241355896, 0.16667349636554718, 0.6350314617156982, 2.3831448554992676, 0.9444794654846191, -0.9128222465515137, 1.117016315460205, -1.31590735912323, -0.46158459782600403]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis0",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis0" },
+      "attrs": { "axis": 0, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.37642553448677063, -1.0994007587432861, 0.29823818802833557, 1.3263858556747437, -0.694567859172821, -0.14963454008102417, -0.4351535439491272, 1.8492637872695923, 0.6722947359085083, 0.40746182203292847, -0.7699160575866699, 0.5392491817474365, -0.6743326783180237, 0.0318305566906929, -0.6358460783958435, 0.676433265209198, 0.5765908360481262, -0.20829875767230988, 0.39600670337677, -1.0930615663528442, -1.4912575483322144, 0.43939170241355896, 0.16667349636554718, 0.6350314617156982, 2.3831448554992676, 0.9444794654846191, -0.9128222465515137, 1.117016315460205, -1.31590735912323, -0.46158459782600403, -0.06824160367250443, 1.7133426666259766, -0.7447548508644104, -0.8264385461807251, -0.09845252335071564, -0.6634783148765564, 1.1266359090805054, -1.0799314975738525, -1.1474686861038208, -0.43782004714012146, -0.49803245067596436, 1.9295320510864258, 0.9494208097457886, 0.08755124360322952, -1.225435495376587, 0.8443629741668701, -1.0002152919769287, -1.5447710752487183, 1.1880297660827637, 0.31694260239601135, 0.9208588004112244, 0.31872764229774475, 0.8568305969238281, -0.6510255932807922, -1.034242868423462, 0.6815944910049438, -0.8034096360206604, -0.6895498037338257, -0.4555324912071228, 0.01747915893793106, -0.3539939224720001, -1.3749512434005737, -0.6436184048652649, -2.223403215408325, 0.6252314448356628, -1.602057695388794, -1.1043833494186401, 0.05216507986187935, -0.73956298828125, 1.543014645576477, -1.2928569316864014, 0.2670508623123169, -0.039282817393541336, -1.1680934429168701, 0.523276686668396, -0.1715463250875473, 0.7717905640602112, 0.8235041499137878, 2.163235902786255, 1.336527943611145, -0.3691818416118622, -0.2393791824579239, 1.0996595621109009, 0.6552637219429016, 0.6401315331459045, -1.6169559955596924, -0.024326125159859657, -0.7380309104919434, 0.279924601316452, -0.09815038740634918, 0.9101788997650146, 0.31721821427345276, 0.7863279581069946, -0.4664191007614136, -0.9444462656974792, -0.410049706697464, -0.017020413652062416, 0.37915173172950745, 2.2593090534210205, -0.0422571524977684, -0.9559450149536133, -0.34598177671432495, -0.463595986366272, 0.4814814627170563, -1.5407969951629639, 0.06326199322938919, 0.15650653839111328, 0.23218104243278503, -0.5973160862922668, -0.23792172968387604, -1.4240609407424927, -0.49331986904144287, -0.5428614616394043, 0.4160500466823578, -1.1561824083328247, 0.7811980843544006, 1.494484543800354, -2.0699849128723145, 0.42625874280929565, 0.676908016204834]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.6374370455741882, -0.3972718119621277, -0.1328805834054947, -0.29779088497161865, -0.3090129792690277, -1.6760038137435913, 1.1523315906524658, 1.0796185731887817, -0.8133642673492432, -1.4664243459701538, 0.5210648775100708, -0.5757879614830017, 0.14195317029953003, -0.3193284273147583, 0.6915387511253357, 0.694749116897583, -0.7255973815917969, -1.383363962173462, -1.5829384326934814, 0.6103793978691101, -1.188859224319458, -0.5068163275718689, -0.596314013004303, -0.05256729573011398, -1.9362797737121582, 0.1887785941362381, 0.523891031742096, 0.08842208981513977, -0.3108861744403839, 0.09740016609430313, 0.3990463316440582, -2.772592782974243, 1.9559123516082764, 0.3900933265686035, -0.6524085998535156, -0.3909533619880676, 0.4937417805194855, -0.11610393971204758, -2.030684471130371, 2.06449294090271, -0.11054065823554993, 1.0201727151870728, -0.6920498609542847, 1.5363770723342896, 0.28634369373321533, 0.6088438630104065, -1.0452533960342407, 1.211145281791687, 0.6898181438446045, 1.3018462657928467, -0.62808758020401, -0.48102712631225586, 2.3039166927337646, -1.0600157976150513, -0.13594970107078552, 1.1368913650512695, 0.09772496670484543, 0.582953691482544, -0.39944902062416077, 0.3700558841228485, -1.3065268993377686, 1.6581306457519531, -0.1181640475988388, -0.6801782250404358, 0.6663830876350403, -0.4607197940349579, -1.3342584371566772, -1.3467174768447876, 0.6937731504440308, -0.15957343578338623, -0.13370156288146973, 1.0777437686920166, -1.1268258094787598, -0.7306777238845825, -0.38487979769706726, 0.09435158967971802, -0.042171452194452286, -0.28688719868659973, -0.06162640079855919, -0.10730527341365814, -0.7196043729782104, -0.8129929900169373, 0.27451634407043457, -0.8909150958061218, -1.1573553085327148, -0.3122922480106354, -0.15766701102256775, 2.256723403930664, -0.7047002911567688, 0.9432607293128967, 0.7471883296966553, -1.188944935798645, 0.7732529640197754, -1.1838806867599487, -2.659172296524048, 0.6063195466995239, -1.7558906078338623, 0.4509344696998596, -0.684010922908783, 1.6595507860183716, 1.0685093402862549, -0.4533858001232147, -0.6878376007080078, -1.214077353477478, -0.4409226179122925, -0.2803554832935333, -0.3646935522556305, 0.1567038595676422, 0.5785214900970459, 0.3496544659137726, -0.7641439437866211, -1.4377914667129517, 1.3645318746566772, -0.6894491910934448, -0.6522936224937439, -0.5211893320083618, -1.8430695533752441, -0.477973997592926, -0.47965580224990845, 0.6203582882881165]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [1, 1, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis1",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis1" },
+      "attrs": { "axis": 1, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-1.550429344177246, 0.41731882095336914, -0.9443684816360474, 0.23810315132141113, -1.4059629440307617, -0.5900576710700989, -0.11048940569162369, -1.6606998443603516, 0.11514787375926971, -0.37914755940437317, -1.7423561811447144, -1.303242802619934, 0.605120062828064, 0.8955559730529785, -0.13190864026546478, 0.40476182103157043, 0.2238435596227646, 0.3296229839324951, 1.2859840393066406, -1.5069984197616577, 0.6764607429504395, -0.38200896978378296, -0.2242589294910431, -0.3022497296333313, -0.37514710426330566, -1.2261961698532104, 0.1833391934633255, 1.670943021774292, -0.05613302066922188, -0.0013850427931174636, -0.6872990131378174, -0.11747454851865768, 0.4661664366722107, -0.3702424466609955, -0.45380404591560364, 0.4032645523548126, -0.9180047512054443, 0.2524966299533844, 0.820321798324585, 1.3599485158920288, -0.09038200974464417, 1.367597222328186, 1.0344098806381226, -0.9962126612663269, -1.2179385423660278, -0.3049636483192444, 1.0289355516433716, -0.07228700816631317, -0.60065758228302, 1.5522432327270508, 0.2869044840335846, -2.320594310760498, 0.31716063618659973, 0.5200406312942505, 0.22560866177082062, 0.4497120976448059, -0.06727560609579086, -1.3183958530426025, -0.3707039952278137, -0.9456157684326172]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.9327409267425537, -1.2630683183670044, 0.45248907804489136, 0.09789614379405975, -0.4481653571128845, -0.6493379473686218, -0.023423105478286743, 1.0791947841644287, -2.004215717315674, 0.3768765330314636, -0.5457119941711426, -1.8845858573913574, -1.9457030296325684, -0.9127835035324097, 0.219509556889534, 0.39306291937828064, -0.9389815926551819, 1.0170209407806396, 1.4229835271835327, 0.39608657360076904, -0.5914026498794556, 1.1244192123413086, 0.7553957104682922, 0.8674074411392212, -0.6564636826515198, -2.834554433822632, 2.116791009902954, -1.610878348350525, -0.03576807305216789, 2.3807454109191895, 0.33057674765586853, 0.9492464661598206, -1.502396583557129, -1.7776669263839722, -0.5327028036117554, 1.090749740600586, -0.34624946117401123, -0.7946363091468811, 0.1979672908782959, 1.081935167312622, -1.444940209388733, -1.2105430364608765, -0.788669228553772, 1.0946383476257324, 0.23482152819633484, 2.1321535110473633, 0.9364457130432129, -0.035095177590847015, 1.265077829360962, 0.2114970088005066, -0.7049213647842407, 0.6799748539924622, -0.6963266730308533, -0.29039710760116577, 1.3277827501296997, -0.10128148645162582, -0.8031414151191711, -0.46433767676353455, 1.021790623664856, -0.55254065990448]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [2, 1, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis2",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis2" },
+      "attrs": { "axis": 2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.23958276212215424, -0.3698011636734009, 0.9725357890129089, 2.1338682174682617, 0.40641549229621887, -0.19317670166492462, 0.7557402849197388, -0.5391326546669006, -0.7496903538703918, 0.03280874714255333, -2.582796573638916, -1.1539503335952759, -0.3479618430137634, -1.3533889055252075, -1.0326430797576904, -0.43674832582473755, -1.642965316772461, -0.4060717821121216, -0.5352701544761658, 0.02540520764887333]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [4, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.1541839838027954, 0.17250441014766693, 0.02106202207505703, 0.09945445507764816, 0.22739277780056, -1.0167386531829834, -0.11477532237768173, 0.30875125527381897, -1.3707599639892578, 0.8656529188156128, 1.081376075744629, -0.6313759684562683, -0.24133779108524323, -0.8781903386116028, 0.6993804574012756, -1.0612223148345947, -0.2224770039319992, -0.8589199185371399, 0.05095427855849266, -1.79422926902771]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis_negative_2",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis_negative_2"
+      },
+      "attrs": { "axis": -2, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.326461672782898, -0.9646064043045044, 0.059894684702157974, -0.21252304315567017, -0.7621145248413086, -0.8877801299095154, 0.9363985657691956, -0.525640606880188, 0.2711701989173889, -0.8014968633651733, -0.6471814513206482, 0.4722471535205841, 0.9304084777832031, -0.17531640827655792, -1.421919822692871, 1.9979560375213623, -0.8565493226051331, -1.5415873527526855, 2.5944244861602783, -0.40403228998184204]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-1.4617327451705933, -0.6834397912025452, 0.36754488945007324, 0.19031155109405518, -0.8517292141914368, 1.822723627090454, -0.5215796828269958, -1.184686541557312, 0.9606934189796448, 1.329062819480896, -0.8174930810928345, -1.401347279548645, 1.0304383039474487, -2.047323703765869, -1.2266216278076172, 0.9674461483955383, -0.05535254627466202, -0.2639373540878296, 0.3528166115283966, -0.15277442336082458]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [2, 3, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 3, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis_negative_3",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis_negative_3"
+      },
+      "attrs": { "axis": -3, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [-0.3868708610534668, -0.5102927684783936, 0.18392549455165863, -0.3854897618293762, -1.601836085319519, -0.8871809244155884, -0.9327890276908875, 1.2433193922042847, 0.8126740455627441, 0.5872593522071838, -0.5053583383560181, -0.8157915472984314, -0.5075175762176514, -1.051880121231079, 2.4972004890441895, -2.245321750640869, 0.5640085339546204, -1.2845523357391357, -0.1043434888124466, -0.9880019426345825, -1.177628993988037, -1.1401963233947754, 1.7549861669540405, -0.13298842310905457, -0.7657021880149841, 0.5557869672775269, 0.010349314659833908, 0.7200337648391724, -1.8242566585540771, 0.3036039173603058, 0.7726948261260986, -1.6615983247756958, 0.44819527864456177, 1.6961815357208252, -0.014857703819870949, 0.8214059472084045, 0.670570433139801, -0.7075057029724121, 0.03976673632860184, -1.5669946670532227, -0.45130303502082825, 0.2656879723072052, 0.723100483417511, 0.024612125009298325, 0.7199837565422058, -1.1029062271118164, -0.10169727355241776, 0.01927938498556614, 1.8495912551879883, -0.21416665613651276, -0.4990166425704956, 0.021351223811507225, -0.9191134572029114, 0.1927538514137268, -0.3650552034378052, -1.7913275957107544, -0.05858655273914337, -0.317543089389801, -1.6324232816696167, -0.06713415682315826]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.4893559217453003, 0.5213037729263306, 0.6119272112846375, -1.3414967060089111, 0.4768983721733093, 0.1484495848417282, 0.5290452241897583, 0.42262861132621765, -1.3597806692123413, -0.041400812566280365, -0.7578708529472351, -0.05008409544825554, -0.8974009156227112, 1.3124703168869019, -0.8589723706245422, -0.8989421725273132, 0.0745864063501358, -1.077099084854126, -0.4246633052825928, -0.8299645781517029, 1.4111720323562622, 0.7858038544654846, -0.05746951699256897, -0.39121705293655396, 0.9409176111221313, 0.40520408749580383, 0.4980524182319641, -0.026192236691713333, -1.688230037689209, -0.1124659851193428, -0.5324898958206177, 0.6450552940368652, 1.0118424892425537, -0.6579510569572449, 0.4683852195739746, 1.7358789443969727, -0.6677127480506897, 1.6819217205047607, -0.8525858521461487, 0.0229597557336092, -0.011145612224936485, 0.011498900130391121, -0.8376780152320862, -0.5911831259727478, -0.6677202582359314, 0.32696259021759033, 0.33003512024879456, 2.2259442806243896, 1.3709889650344849, -0.5098432302474976, 0.3248696029186249, 0.9971179962158203, 0.030601823702454567, -0.06964157521724701, 0.0515749417245388, 0.8672766089439392, -0.8483205437660217, -0.32566946744918823, 0.47043314576148987, 0.3114470839500427]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [2, 1, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "onnx_backend_layer_normalization_4d_axis_negative_4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_layer_normalization_4d_axis_negative_4"
+      },
+      "attrs": { "axis": -4, "epsilon": 0.00001 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_layer_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.6984571218490601, 0.003770889015868306, 0.9318483471870422, 0.33996498584747314, -0.01568211242556572, 0.1609281748533249, -0.1906534880399704, -0.3948495090007782, -0.26773354411125183, -1.1280113458633423, 0.2804417014122009, -0.9931235909461975, 0.841631293296814, -0.24945858120918274, 0.04949498176574707, 0.49383679032325745, 0.6433144807815552, -1.5706233978271484, -0.20690368115901947, 0.8801789283752441, -1.698105812072754, 0.3872804641723633, -2.2555642127990723, -1.022506833076477, 0.03863055258989334, -1.6567151546478271, -0.9855107665061951, -1.4718350172042847, 1.6481349468231201, 0.16422775387763977, 0.5672903060913086, -0.22267509996891022, -0.353431761264801, -1.6164741516113281, -0.2918373644351959, -0.7614921927452087, 0.8579239249229431, 1.1411018371582031, 1.466578722000122, 0.8525519371032715, -0.5986539125442505, -1.1158969402313232, 0.7666631937026978, 0.35629281401634216, -1.768538475036621, 0.3554818034172058, 0.8145198225975037, 0.05892558768391609, -0.18505367636680603, -0.8076484799385071, -1.4465347528457642, 0.8002979755401611, -0.3091144561767578, -0.23346665501594543, 1.732721209526062, 0.6845011115074158, 0.37082499265670776, 0.14206179976463318, 1.519994854927063, 1.719589352607727, 0.9295051097869873, 0.5822246074676514, -2.0946030616760254, 0.12372191250324249, -0.13010695576667786, 0.09395322948694229, 0.9430460929870605, -2.7396771907806396, -0.5693120360374451, 0.26990434527397156, -0.4668455421924591, -1.4169061183929443, 0.8689634799957275, 0.276871919631958, -0.9711045622825623, 0.3148171901702881, 0.8215857148170471, 0.005292646121233702, 0.8005648255348206, 0.07826017588376999, -0.3952289819717407, -1.1594204902648926, -0.08593076467514038, 0.19429293274879456, 0.875832736492157, -0.1151074692606926, 0.4574156105518341, -0.9646120071411133, -0.782629132270813, -0.11038929969072342, -1.0546284914016724, 0.8202478289604187, 0.46313032507896423, 0.27909576892852783, 0.33890411257743835, 2.021043539047241, -0.46886420249938965, -2.2014412879943848, 0.19930019974708557, -0.05060354247689247, -0.5175190567970276, -0.9788298606872559, -0.43918952345848083, 0.18133842945098877, -0.5028166770935059, 2.4124536514282227, -0.9605043530464172, -0.793117344379425, -2.2886199951171875, 0.2514844238758087, -2.016406536102295, -0.5394546389579773, -0.27567052841186523, -0.709727942943573, 1.7388726472854614, 0.9943943619728088, 1.3191368579864502, -0.8824188113212585, 1.1285940408706665, 0.4960009455680847]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.7714059352874756, 1.029438853263855, -0.9087632298469543, -0.4243176281452179, 0.8625960350036621, -2.6556191444396973, 1.5133280754089355, 0.5531320571899414, -0.045703962445259094, 0.2205076515674591, -1.0299352407455444, -0.3499433696269989, 1.1002843379974365, 1.2980220317840576, 2.6962239742279053, -0.07392466813325882, -0.6585529446601868, -0.5142339468002319, -1.0180418491363525, -0.07785475254058838, 0.38273242115974426, -0.03424227982759476, 1.0963468551635742, -0.23421579599380493, -0.3474506437778473, -0.5812684893608093, -1.6326345205307007, -1.5677677392959595, -1.1791579723358154, 1.3014280796051025, 0.8952602744102478, 1.3749641180038452, -1.3322116136550903, -1.9686247110366821, -0.6600562930107117, 0.17581894993782043, 0.49869027733802795, 1.0479722023010254, 0.2842796742916107, 1.7426687479019165, -0.22260567545890808, -0.9130792021751404, -1.6812182664871216, -0.8889713287353516, 0.2421179562807083, -0.8887202739715576, 0.9367424845695496, 1.4123276472091675, -2.369586944580078, 0.8640522956848145, -2.2396039962768555, 0.40149906277656555, 1.2248705625534058, 0.06485610455274582, -1.2796891927719116, -0.5854312181472778, -0.26164543628692627, -0.18224477767944336, -0.2028968334197998, -0.10988277941942215, 0.21348005533218384, -1.2085736989974976, -0.24201983213424683, 1.518261194229126, -0.3846454322338104, -0.4438360929489136, 1.0781973600387573, -2.5591845512390137, 1.1813786029815674, -0.6319037675857544, 0.1639285683631897, 0.09632135927677155, 0.9424681067466736, -0.26759475469589233, -0.6780257821083069, 1.2978458404541016, -2.3641738891601562, 0.020334182307124138, -1.3479254245758057, -0.7615733742713928, 2.011256694793701, -0.04459542781114578, 0.19506970047950745, -1.7815628051757812, -0.7290446758270264, 0.1965574026107788, 0.354757696390152, 0.616886556148529, 0.00862789899110794, 0.5270041823387146, 0.4537819027900696, -1.8297404050827026, 0.037005722522735596, 0.7679024338722229, 0.5898798108100891, -0.36385881900787354, -0.8056265115737915, -1.118311882019043, -0.13105401396751404, 1.1330798864364624, -1.9518041610717773, -0.6598917245864868, -1.1398024559020996, 0.7849575281143188, -0.554309606552124, -0.47063764929771423, -0.21694956719875336, 0.44539326429367065, -0.3923889994621277, -3.046143054962158, 0.5433118939399719, 0.43904295563697815, -0.2195410281419754, -1.0840365886688232, 0.35178011655807495, 0.37923553586006165, -0.47003287076950073, -0.2167314738035202, -0.9301565289497375, -0.1785890907049179]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.00005 },
+        "mean": { "dtype": "float32", "shape": [1, 1, 1, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [1, 1, 1, 1], "tolerance": 0.0005 }
+      }
+    },
+    {
+      "name": "ort_last_axis_scale_only_small",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [-10.264, 8.6453, 43.1561, -0.641239, -8.2164, 0.11412, 41.3156, 3.0458]
+          }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-0.6953, 5.1824] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [0.6953, 5.1824, -0.6953, -5.1824, 0.6953, 5.1824, -0.6953, -5.1824] }
+        }
+      }
+    },
+    {
+      "name": "ort_last_axis_scale_bias_small",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Bias"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "data": { "kind": "values", "values": [1.2416, 0.946123, 13.1685, 0.36423, 21.145, 0.03941] }
+        },
+        "scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-0.6953, 5.1824] } },
+        "b": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.6435, -0.3964] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [-0.0516, -5.5776, -0.0518, -5.5788, -0.0518, -5.5788] }
+        }
+      }
+    },
+    {
+      "name": "ort_last_axis_full_shape_scale_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/layer_norm_op_test.cc",
+        "test": "LayerNormTest.LayerNorm_Scale_Bias_NoBroadcast"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 2.0, 3.0, -4.0, -10.264, 8.6453, 43.1561, -0.641239] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [-0.1, 1.7, -0.6953, 5.1824, -0.1, 1.7, -0.6953, 5.1824] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [-2.0, 0.3, 0.0, 0.0, -2.0, 0.3, 0.0, 0.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [-1.9, 2.0, -0.6953, -5.1824, -1.9, 2.0, -0.6953, -5.1824] }
+        }
+      }
+    },
+    {
+      "name": "f16_realistic_hidden2048_bias_stats",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.021, "cosStep": 0.011, "scale": 0.25, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [8, 2048], "tolerance": 0.005 },
+        "mean": { "dtype": "float32", "shape": [8, 1], "tolerance": 0.005 },
+        "invStdDev": { "dtype": "float32", "shape": [8, 1], "tolerance": 0.02 }
+      }
+    },
+    {
+      "name": "f16_broadcast_scale_bias_generic",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.5, -0.5, 1.0, -1.0, 0.25, 0.75, -0.25, 0.6, 0.9, 0.1, -0.4, 0.4, 0.2, -0.8, 0.7, -0.3]
+          }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [1.2, 0.8, 1.0, 1.1] }
+        },
+        "b": { "dtype": "float16", "shape": [1, 2, 2], "data": { "kind": "values", "values": [0.1, -0.1, 0.05, 0.2] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4, 2, 2], "tolerance": 0.005 } }
+    },
+    {
+      "name": "f16_suffix_axis_bias_stats",
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 4, 8],
+          "data": {
+            "kind": "values",
+            "values": [0.4, 0.4413, 0.4613, 0.4606, 0.4411, 0.4055, 0.3572, 0.3001, 0.2386, 0.1768, 0.1189, 0.0685, 0.0287, 0.0017, -0.0111, -0.0095, 0.0058, 0.0334, 0.0707, 0.1146, 0.1613, 0.207, 0.2474, 0.2788, 0.2977, 0.3012, 0.2873, 0.2549, 0.2039, 0.1353, 0.0512, -0.0455, -0.1509, -0.2605, -0.3694, -0.4726, -0.5651, -0.6421, -0.6995, -0.734, -0.7433, -0.7261, -0.6823, -0.6131, -0.5207, -0.4086, -0.2809, -0.1427, 0.0005, 0.1431, 0.2794, 0.404, 0.5123, 0.6001, 0.6645, 0.7035, 0.7164, 0.7036, 0.6666, 0.6081, 0.5317, 0.4416, 0.3425, 0.2394]
+          }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [4, 8],
+          "data": {
+            "kind": "values",
+            "values": [1.15, 1.1804, 1.2075, 1.2303, 1.2474, 1.258, 1.2613, 1.2569, 1.2447, 1.2247, 1.1975, 1.1638, 1.1246, 1.0811, 1.0347, 0.9869, 0.9393, 0.8935, 0.8509, 0.813, 0.781, 0.7558, 0.738, 0.7282, 0.7265, 0.7325, 0.7459, 0.7659, 0.7914, 0.8214, 0.8544, 0.8891]
+          }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [4, 8],
+          "data": {
+            "kind": "values",
+            "values": [0.05, 0.0564, 0.0585, 0.0567, 0.0514, 0.0436, 0.0342, 0.0243, 0.0148, 0.0069, 0.001, -0.0022, -0.0026, -0.0003, 0.0041, 0.0102, 0.0168, 0.0232, 0.0284, 0.0314, 0.0315, 0.0283, 0.0216, 0.0116, -0.0014, -0.0164, -0.0324, -0.0484, -0.0631, -0.0752, -0.0837, -0.0878]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [2, 4, 8], "tolerance": 0.005 },
+        "mean": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.005 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1, 1], "tolerance": 0.02 }
+      }
+    },
+    {
+      "name": "per_row_scale_bias_no_outer_broadcast_generic_offset",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.07, "scale": 0.4, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.17, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.000002 } }
+    },
+    {
+      "name": "vec4_stats_subgroup_wide_hidden2048_f32_multi_subgroup_combine",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 1.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.003, "cosStep": 0.005, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 2048], "tolerance": 0.00001 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.00002 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "f32_last_axis_unaligned_hidden2050_bias",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 1.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.003, "cosStep": 0.005, "scale": 0.25, "offset": 1.0 }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.007, "cosStep": 0.011, "scale": 0.1 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2050], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "zero_variance_constant_row_vec4_width_exact_zeros",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 8], "data": { "kind": "constant", "value": 5.0 } },
+        "scale": { "dtype": "float32", "shape": [8], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 8],
+          "tolerance": 1e-7,
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "empty_last_axis_zero_hidden_noop",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [3, 0], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 0], "data": { "kind": "values", "values": [] } } }
+    },
+    {
+      "name": "last_axis_broadcast_scale_stats_generic_path",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.21 }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "last_axis_broadcast_scale_bias_stats_generic_path",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.07 }
+        },
+        "scale": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.5 } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.25 } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000002 },
+        "mean": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 },
+        "invStdDev": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.00001 }
+      }
+    }
+  ]
+}