sync 2e7068faf55e

Browse files

Files changed (10) hide show

README.md +75 -0
build/webgpu/bench.json +173 -0
build/webgpu/manifest.json +369 -0
build/webgpu/metadata.json +22 -0
build/webgpu/norm-row-stats.wgsl.jinja +124 -0
build/webgpu/rms-normalization-splitk-normalize.wgsl.jinja +93 -0
build/webgpu/rms-normalization-splitk-partials.wgsl.jinja +93 -0
build/webgpu/rms-normalization-stash-f16-serial.wgsl.jinja +140 -0
build/webgpu/rms-normalization.wgsl.jinja +144 -0
build/webgpu/test.json +1590 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.RMSNormalization
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 23
+## Description
+Computes RMS normalization over the suffix dimensions of `X` starting at `axis`: `Y = X / sqrt(mean(X^2) + epsilon) * scale`. The normalization stage supports TensorProto `stash_type` values `1` (float32) and `10` (float16), and is cast back to the dtype of `X` before `scale` is applied. The input type `T` and scale/output type `V` may independently be float16 or float32; ONNX's bfloat16 and double cases are not yet implemented.
+See the [ONNX `RMSNormalization` spec](https://onnx.ai/onnx/operators/onnx__RMSNormalization.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Input tensor to be normalized; the RMS is taken over the last dimensions starting at `axis`. | required |
+| `scale` | `scale` | `V` | — | — | Scale tensor, unidirectionally broadcastable to `X`; its dtype `V` may differ from the input dtype `T`. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `V` | same as `X` | same as `X` | Normalized and scaled output tensor; same shape as `X` and same dtype `V` as `scale`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `-1` | The first dimension of the normalization suffix; negative values count from the end, so the default `-1` normalizes over only the last dimension. |
+| `epsilon` | `0.00001` | Small constant added to the mean square before taking the square root to avoid division by zero. |
+| `stash_type` | `1` | TensorProto element type used for normalization: `1` computes in float32, while `10` computes in float16. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+| `V` | `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
+- [`norm-row-stats.wgsl.jinja`](build/webgpu/norm-row-stats.wgsl.jinja)
+- [`rms-normalization-splitk-normalize.wgsl.jinja`](build/webgpu/rms-normalization-splitk-normalize.wgsl.jinja)
+- [`rms-normalization-splitk-partials.wgsl.jinja`](build/webgpu/rms-normalization-splitk-partials.wgsl.jinja)
+- [`rms-normalization-stash-f16-serial.wgsl.jinja`](build/webgpu/rms-normalization-stash-f16-serial.wgsl.jinja)
+- [`rms-normalization.wgsl.jinja`](build/webgpu/rms-normalization.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.RMSNormalization", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 2, 3] },
+  scale: { data: scaleData, shape: [3] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,173 @@

+{
+  "op": "ai.onnx.RMSNormalization",
+  "cases": [
+    {
+      "name": "rmsnorm-f32-256x1024",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 256, "dim": 1024 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [256, 1024], "dtype": "float32", "dist": "normal", "seed": 710, "scale": 0.5 },
+        "scale": { "shape": [1024], "dtype": "float32", "dist": "uniform", "seed": 711, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [256, 1024], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f32-4096x4096",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "rows": 4096, "dim": 4096 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 712, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float32", "dist": "uniform", "seed": 713, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [4096, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f16-4096x4096",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "rows": 4096, "dim": 4096 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4096, 4096], "dtype": "float16", "dist": "normal", "seed": 714, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float16", "dist": "uniform", "seed": 715, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [4096, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f16-1x4096-rows1-decode",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "rows": 1, "dim": 4096 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [1, 4096], "dtype": "float16", "dist": "normal", "seed": 7200, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float16", "dist": "uniform", "seed": 7201, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [1, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f16-256x4096-aligned-vec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "rows": 256, "dim": 4096 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [256, 4096], "dtype": "float16", "dist": "normal", "seed": 7210, "scale": 0.5 },
+        "scale": { "shape": [4096], "dtype": "float16", "dist": "uniform", "seed": 7211, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [256, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f16-256x4094-unaligned-alignmentcliff",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "rows": 256, "dim": 4094 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [256, 4094], "dtype": "float16", "dist": "normal", "seed": 7220, "scale": 0.5 },
+        "scale": { "shape": [4094], "dtype": "float16", "dist": "uniform", "seed": 7221, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [256, 4094], "dtype": "float16" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f32-8x256x1024-suffixaxis1-generic",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "outer": 8, "normSize": 262144 },
+      "attrs": { "epsilon": 0.000001, "axis": 1 },
+      "inputs": {
+        "x": { "shape": [8, 256, 1024], "dtype": "float32", "dist": "normal", "seed": 7230, "scale": 0.5 },
+        "scale": {
+          "shape": [256, 1024],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 7231,
+          "scale": 0.25,
+          "offset": 1
+        }
+      },
+      "outputs": { "y": { "shape": [8, 256, 1024], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.outer * args.normSize * 2 + args.normSize) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "rmsnorm-f32-65535x512-undercliff",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "rows": 65535, "dim": 512 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [65535, 512], "dtype": "float32", "dist": "normal", "seed": 7240, "scale": 0.5 },
+        "scale": { "shape": [512], "dtype": "float32", "dist": "uniform", "seed": 7241, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [65535, 512], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f32-70000x512-rows-dispatchcliff",
+      "preset": "stress",
+      "provenance": { "notes": "Stress-only dispatch/capacity case: declared tensors occupy 273 MiB of GPU storage." },
+      "vars": { "dtype": "float32", "rows": 70000, "dim": 512 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [70000, 512], "dtype": "float32", "dist": "normal", "seed": 7250, "scale": 0.5 },
+        "scale": { "shape": [512], "dtype": "float32", "dist": "uniform", "seed": 7251, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [70000, 512], "dtype": "float32" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "rmsnorm-f16-8x262144-suffixaxis0-nosplitk-occupancy-starve",
+      "preset": "stress",
+      "vars": { "dtype": "float16", "outer": 8, "normSize": 262144 },
+      "attrs": { "epsilon": 0.000001, "axis": 0 },
+      "inputs": {
+        "x": { "shape": [8, 262144], "dtype": "float16", "dist": "normal", "seed": 7260, "scale": 0.5 },
+        "scale": { "shape": [262144], "dtype": "float16", "dist": "uniform", "seed": 7261, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [8, 262144], "dtype": "float16", "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "(args.outer * args.normSize * 2 + args.normSize) * dtypeBytes(args.dtype)" }
+        ]
+      }
+    },
+    {
+      "name": "rmsnorm-f32-4x262144-lastaxis-scalarscale-fallback-starve",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "rows": 4, "dim": 262144 },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "shape": [4, 262144], "dtype": "float32", "dist": "normal", "seed": 7270, "scale": 0.5 },
+        "scale": { "shape": [1], "dtype": "float32", "dist": "uniform", "seed": 7271, "scale": 0.25, "offset": 1 }
+      },
+      "outputs": { "y": { "shape": [4, 262144], "dtype": "float32", "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "(args.rows * args.dim * 2 + args.dim) * dtypeBytes(args.dtype)" }]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,369 @@

+{
+  "domain": "ai.onnx",
+  "name": "RMSNormalization",
+  "sinceVersion": 23,
+  "description": "Computes RMS normalization over the suffix dimensions of `X` starting at `axis`: `Y = X / sqrt(mean(X^2) + epsilon) * scale`. The normalization stage supports TensorProto `stash_type` values `1` (float32) and `10` (float16), and is cast back to the dtype of `X` before `scale` is applied. The input type `T` and scale/output type `V` may independently be float16 or float32; ONNX's bfloat16 and double cases are not yet implemented.",
+  "inputs": [
+    {
+      "role": "X",
+      "dtype": "T",
+      "description": "Input tensor to be normalized; the RMS is taken over the last dimensions starting at `axis`."
+    },
+    {
+      "role": "scale",
+      "dtype": "V",
+      "description": "Scale tensor, unidirectionally broadcastable to `X`; its dtype `V` may differ from the input dtype `T`."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "V",
+      "rank": "ranks.X",
+      "shape": "shapes.X",
+      "description": "Normalized and scaled output tensor; same shape as `X` and same dtype `V` as `scale`."
+    }
+  ],
+  "attributes": { "axis": -1, "epsilon": 0.00001, "stash_type": 1 },
+  "attributeDescriptions": {
+    "axis": "The first dimension of the normalization suffix; negative values count from the end, so the default `-1` normalizes over only the last dimension.",
+    "epsilon": "Small constant added to the mean square before taking the square root to avoid division by zero.",
+    "stash_type": "TensorProto element type used for normalization: `1` computes in float32, while `10` computes in float16."
+  },
+  "attributeConstraints": { "stash_type": { "values": [1, 10] } },
+  "typeConstraints": { "T": ["float32", "float16"], "V": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "scale": { "kind": "tensor", "semantic": "scale", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "SPLIT_MAX_ROWS": 256,
+    "SPLIT_MIN_HIDDEN": 16384,
+    "SPLIT_TARGET_ELEMENTS": 4096,
+    "MAX_SPLITS": 64
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
+    "reportedNonWave32Adapter": "not wave32Adapter and (has(device.adapterInfo, \"subgroupMinSize\") or has(device.adapterInfo, \"subgroupMaxSize\"))",
+    "normMaxWorkgroup": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "hasSubgroupId": "device.features.has(\"subgroups\") and device.wgslLanguageFeatures.has(\"subgroup_id\")",
+    "axisNorm": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.X",
+    "normalizedRows": "outer(shapes.X, axisNorm)",
+    "normalizedHidden": "dim(shapes.X, axisNorm) * inner(shapes.X, axisNorm)",
+    "normalizedDispatchRows": "0 if normalizedHidden == 0 else normalizedRows",
+    "normalizedWorkgroupHidden": "max(1, normalizedHidden)",
+    "normalizationShapeOk": "ranks.X >= 1 and ranks.scale >= 0 and ranks.scale <= ranks.X and sameShape(shapes.Y, shapes.X) and attrs.axis + ranks.X >= 0 and attrs.axis < ranks.X and broadcastable(shapes.scale, shapes.X) and f16Ok(dtypes.T) and f16Ok(dtypes.V)",
+    "baseOk": "normalizationShapeOk and attrs.stash_type == onnxDtypeCode(\"float32\")",
+    "stashF16Ok": "normalizationShapeOk and attrs.stash_type == onnxDtypeCode(\"float16\")",
+    "lastAxisOk": "baseOk and (attrs.axis == -1 or attrs.axis == ranks.X - 1)",
+    "suffixAxisOk": "baseOk and ranks.X >= 2 and not (attrs.axis == -1 or attrs.axis == ranks.X - 1)"
+  },
+  "bindingSets": {
+    "rows": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xElement"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$ioElement" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "normalizedRows" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(normalizedRows, device.limits.maxComputeWorkgroupsPerDimension))"
+            }
+          ]
+        }
+      }
+    ],
+    "splitPartials": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xElement"
+      },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "splitRows" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(splitRows, device.limits.maxComputeWorkgroupsPerDimension))"
+            }
+          ]
+        }
+      }
+    ],
+    "splitNormalize": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "X",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xElement"
+      },
+      {
+        "name": "scale",
+        "arg": "scale",
+        "semantic": "scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$ioElement"
+      },
+      { "name": "partials", "semantic": "partials", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$ioElement" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "splitRows" },
+            {
+              "name": "rowStride",
+              "type": "u32",
+              "value": "max(1, min(splitRows, device.limits.maxComputeWorkgroupsPerDimension))"
+            }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "stash_f16_serial",
+      "priority": 1000,
+      "when": "stashF16Ok",
+      "constants": {
+        "scalar": "dtypes.V",
+        "xElement": "dtypes.T",
+        "ioElement": "dtypes.V",
+        "usesF16": "dtypes.T == \"f16\" or dtypes.V == \"f16\"",
+        "hiddenSize": "normalizedHidden",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RMSNormalization.StashF16Serial",
+          "source": {
+            "shader": "rms-normalization-stash-f16-serial.wgsl.jinja",
+            "inputs": {
+              "xShape": "shapes.X",
+              "scaleShape": "shapes.scale",
+              "xRank": "ranks.X",
+              "scaleRank": "ranks.scale"
+            }
+          },
+          "bindings": "rows",
+          "dispatch": { "workgroups": "normalizedDispatchRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis_splitk",
+      "priority": 15,
+      "derive": {
+        "splitRows": "normalizedRows",
+        "splitHidden": "normalizedHidden",
+        "split": "min(tunables.MAX_SPLITS, pow2ceil(ceilDiv(splitHidden, tunables.SPLIT_TARGET_ELEMENTS)))"
+      },
+      "when": ["baseOk", "ranks.X >= 2", "normalizedRows <= tunables.SPLIT_MAX_ROWS", "normalizedHidden >= tunables.SPLIT_MIN_HIDDEN", "min(tunables.MAX_SPLITS, pow2ceil(ceilDiv(normalizedHidden, tunables.SPLIT_TARGET_ELEMENTS))) <= device.limits.maxComputeWorkgroupsPerDimension", "normalizedRows * min(tunables.MAX_SPLITS, pow2ceil(ceilDiv(normalizedHidden, tunables.SPLIT_TARGET_ELEMENTS))) * 4 <= device.limits.maxStorageBufferBindingSize", "normalizedRows * min(tunables.MAX_SPLITS, pow2ceil(ceilDiv(normalizedHidden, tunables.SPLIT_TARGET_ELEMENTS))) * 4 <= device.limits.maxBufferSize"],
+      "demoteWhen": ["reportedNonWave32Adapter"],
+      "constants": {
+        "scalar": "dtypes.V",
+        "xElement": "dtypes.T",
+        "ioElement": "dtypes.V",
+        "usesF16": "dtypes.T == \"f16\" or dtypes.V == \"f16\"",
+        "hiddenSize": "splitHidden",
+        "workgroupSize": "normMaxWorkgroup",
+        "split": "split",
+        "epsilon": "attrs.epsilon"
+      },
+      "intermediates": [{ "id": "partials", "dtype": "float32", "shape": "[splitRows * split]" }],
+      "passes": [
+        {
+          "id": "partials",
+          "name": "RMSNormalization.SplitKPartials",
+          "source": { "shader": "rms-normalization-splitk-partials.wgsl.jinja", "inputs": {} },
+          "bindings": "splitPartials",
+          "dispatch": { "workgroups": "splitRows", "z": "split" }
+        },
+        {
+          "id": "normalize",
+          "name": "RMSNormalization.SplitKNormalize",
+          "source": {
+            "shader": "rms-normalization-splitk-normalize.wgsl.jinja",
+            "inputs": {
+              "xShape": "shapes.X",
+              "scaleShape": "shapes.scale",
+              "xRank": "ranks.X",
+              "scaleRank": "ranks.scale",
+              "writeStats": false,
+              "rmsScaleAfterCast": true
+            }
+          },
+          "bindings": "splitNormalize",
+          "dispatch": { "workgroups": "splitRows", "z": "split" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis",
+      "priority": 0,
+      "when": "lastAxisOk",
+      "constants": {
+        "scalar": "dtypes.V",
+        "xElement": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\" or dtypes.V == \"f16\"",
+        "ioElement": "dtypes.V",
+        "hiddenSize": "normalizedHidden",
+        "workgroupSize": "min(normMaxWorkgroup, pow2ceil(normalizedWorkgroupHidden))",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RMSNormalization",
+          "source": {
+            "shader": "rms-normalization.wgsl.jinja",
+            "inputs": {
+              "xShape": "shapes.X",
+              "scaleShape": "shapes.scale",
+              "xRank": "ranks.X",
+              "scaleRank": "ranks.scale",
+              "writeStats": false,
+              "rmsScaleAfterCast": true
+            }
+          },
+          "bindings": "rows",
+          "dispatch": { "workgroups": "normalizedDispatchRows" }
+        }
+      ]
+    },
+    {
+      "id": "suffix_axis",
+      "priority": 10,
+      "when": "suffixAxisOk",
+      "constants": {
+        "scalar": "dtypes.V",
+        "xElement": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\" or dtypes.V == \"f16\"",
+        "ioElement": "dtypes.V",
+        "hiddenSize": "normalizedHidden",
+        "workgroupSize": "min(normMaxWorkgroup, pow2ceil(normalizedWorkgroupHidden))",
+        "epsilon": "attrs.epsilon"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RMSNormalization.SuffixAxis",
+          "source": {
+            "shader": "rms-normalization.wgsl.jinja",
+            "inputs": {
+              "xShape": "shapes.X",
+              "scaleShape": "shapes.scale",
+              "xRank": "ranks.X",
+              "scaleRank": "ranks.scale",
+              "writeStats": false,
+              "rmsScaleAfterCast": true
+            }
+          },
+          "bindings": "rows",
+          "dispatch": { "workgroups": "normalizedDispatchRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_row_vec4",
+      "priority": 110,
+      "when": ["lastAxisOk", "dtypes.T == dtypes.V", "ranks.scale >= 1", "numel(shapes.scale) == dim(shapes.X, -1)", "dim(shapes.scale, -1) == dim(shapes.X, -1)", "dim(shapes.X, -1) % 4 == 0"],
+      "constants": { "xElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "ioElement": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RMSNormalization.LastAxisRow",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"rms\"",
+              "vec4": true,
+              "writeStats": false,
+              "rmsScaleAfterCast": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "min(normMaxWorkgroup, pow2ceil(max(1, dim(shapes.X, -1) / 4)))",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": "dim(shapes.X, -1) / 4",
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "rows",
+          "dispatch": { "workgroups": "normalizedDispatchRows" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_row",
+      "priority": 100,
+      "when": ["lastAxisOk", "dtypes.T == dtypes.V", "ranks.scale >= 1", "numel(shapes.scale) == dim(shapes.X, -1)", "dim(shapes.scale, -1) == dim(shapes.X, -1)", "true"],
+      "constants": { "xElement": "dtypes.T", "ioElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RMSNormalization.LastAxisRow",
+          "source": {
+            "shader": "norm-row-stats.wgsl.jinja",
+            "inputs": {
+              "mode": "\"rms\"",
+              "vec4": false,
+              "writeStats": false,
+              "rmsScaleAfterCast": true,
+              "scalar": "dtypes.T",
+              "usesF16": "dtypes.T == \"f16\"",
+              "hidden": "dim(shapes.X, -1)",
+              "wg": "min(normMaxWorkgroup, pow2ceil(max(1, dim(shapes.X, -1))))",
+              "epsilon": "attrs.epsilon",
+              "hiddenVec": 1,
+              "vecType": "\"vec4<\" ~ dtypes.T ~ \">\"",
+              "combineSubgroups": "hasSubgroupId"
+            }
+          },
+          "subgroupCollectivesWidth": "portable",
+          "bindings": "rows",
+          "dispatch": { "workgroups": "normalizedDispatchRows" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.RMSNormalization",
+  "id": "_ai_onnx_rmsnormalization_webgpu_78878c2",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "aQ8RZqdCxZw7VI+7n2XQneyE8kM1Qzire3++ZtvURxg=",
+      "manifest.json": "Nl952QCnY+Pw6VgYi/lKfYZEh3PpBL2LijyCQqtMTTg=",
+      "norm-row-stats.wgsl.jinja": "uxWNeifDOH5P8jfrDfE95ow3QGhewBtScBqM76ps2YI=",
+      "rms-normalization-splitk-normalize.wgsl.jinja": "TXhEDbPnIGP5Ba5SckUHyuuAdCWCip2BH5sJLZ43DdE=",
+      "rms-normalization-splitk-partials.wgsl.jinja": "GbYu7Cj/LLO9y2Rvhw30RJAUAduI0xfN5YS3H/4Z1k8=",
+      "rms-normalization-stash-f16-serial.wgsl.jinja": "tzxVAVThZ5QRnu0dtHoYwqYErZwth2dViidXHaqsVGo=",
+      "rms-normalization.wgsl.jinja": "2YtdapcEAWQUL/wotiMPs8i6wfRV3VG+y+CPhdu47nM=",
+      "test.json": "09ASLvCprHAheaBj+hX9jc5xDu/05mUj7mOoFwJWWyk="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.RMSNormalization" }
+}

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

	@@ -0,0 +1,124 @@

+{% if source.usesF16 %}
+enable f16;
+{% endif %}
+{% set combineSubgroups = source.combineSubgroups %}
+{% set scalarIo = source.scalarIo if source.scalarIo is defined else false %}
+{% set rmsScaleVec = "vec4<f32>(scale[i])" %}
+{% set rmsScaleScalar = "f32(scale[i])" %}
+{% 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.
+//
+// RMS mode uses sum_sq / HIDDEN without computing or subtracting a mean.
+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<f32, WG>;
+fn reduce_scalar(value: f32{{ reduceThreadParameters }}) -> f32 {
+  let s = subgroupAdd(value);
+  if (num_sg == 1u) {
+    return s;
+  }
+  if (sg_lane == 0u) {
+    sg_partials[sg_id] = s;
+  }
+  workgroupBarrier();
+  var total = 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>;
+fn reduce_scalar(value: f32, tid: u32) -> f32 {
+  tr0[tid] = value;
+  workgroupBarrier();
+  var stride: u32 = WG / 2u;
+  loop {
+    if (stride == 0u) { break; }
+    if (tid < stride) { tr0[tid] = tr0[tid] + tr0[tid + stride]; }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  let reduced = tr0[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 %}
+  var acc = 0.0;
+{% if source.vec4 %}
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+    let v = vec4<f32>(x[base + i]);
+    acc = acc + dot(v, v);
+  }
+{% else %}
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let v = f32(x[base + i]);
+    acc = acc + v * v;
+  }
+{% endif %}
+  let total = reduce_scalar(acc{{ reduceThreadArguments }});
+  let inv = inverseSqrt(total / f32(HIDDEN) + EPSILON);
+{% if source.vec4 %}
+  for (var i = tid; i < HIDDEN_V; i = i + WG) {
+    let idx = base + i;
+    let v = vec4<f32>(x[idx]);
+    y[idx] = {{ source.vecType }}(v * inv) * {{ source.vecType }}({{ rmsScaleVec }});
+  }
+{% else %}
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let idx = base + i;
+    let v = f32(x[idx]);
+    y[idx] = {{ source.scalar }}(v * inv) * {{ source.scalar }}({{ rmsScaleScalar }});
+  }
+{% endif %}
+}

build/webgpu/rms-normalization-splitk-normalize.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,93 @@

+// Split-K normalize pass. Each workgroup (row = wg.x, split index = wg.z)
+// folds the SPLIT per-row partial sums of squares into the RMS scale, then
+// normalizes its HIDDEN/SPLIT slice. SPLIT is small (<=64), so the serial fold
+// avoids a third combine pass. Scale offsets follow the suffix-axis broadcast
+// contract.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const EPSILON: f32 = {{ epsilon }};
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+{% if source.scaleRank > 0 %}
+const X_RANK: u32 = {{ source.xRank }}u;
+const SCALE_RANK: u32 = {{ source.scaleRank }}u;
+const X_SHAPE: array<u32, {{ source.xRank }}> = array<u32, {{ source.xRank }}>({% for d in source.xShape %}{{ d }}u{% if not loop.last %}, {% endif %}{% endfor %});
+const SCALE_SHAPE: array<u32, {{ source.scaleRank }}> = array<u32, {{ source.scaleRank }}>({% for d in source.scaleShape %}{{ d }}u{% if not loop.last %}, {% endif %}{% endfor %});
+fn x_stride(axis: u32) -> u32 {
+  var stride = 1u;
+  for (var i = axis + 1u; i < X_RANK; i += 1u) {
+    stride *= X_SHAPE[i];
+  }
+  return stride;
+}
+fn scale_stride(axis: u32) -> u32 {
+  var stride = 1u;
+  for (var i = axis + 1u; i < SCALE_RANK; i += 1u) {
+    stride *= SCALE_SHAPE[i];
+  }
+  return stride;
+}
+{% endif %}
+fn scale_offset({% if source.scaleRank > 0 %}out_index: u32{% endif %}) -> u32 {
+{% if source.scaleRank == 0 %}
+  return 0u;
+{% else %}
+  var rem = out_index;
+  var offset = 0u;
+  for (var axis = 0u; axis < X_RANK; axis += 1u) {
+    let stride = x_stride(axis);
+    let coord = rem / stride;
+    rem %= stride;
+    let scale_axis = i32(axis) - i32(X_RANK - SCALE_RANK);
+    if (scale_axis >= 0) {
+      let s_axis = u32(scale_axis);
+      if (SCALE_SHAPE[s_axis] != 1u) {
+        offset += coord * scale_stride(s_axis);
+      }
+    }
+  }
+  return offset;
+{% 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 k = wg.z;
+  let tid = lid.x;
+  var total = 0.0;
+  for (var i = 0u; i < SPLIT; i = i + 1u) {
+    total = total + partials[row * SPLIT + i];
+  }
+  let inv = inverseSqrt(total / f32(HIDDEN) + EPSILON);
+  let chunk = (HIDDEN + SPLIT - 1u) / SPLIT;
+  let start = k * chunk;
+  var end = start + chunk;
+  if (end > HIDDEN) { end = HIDDEN; }
+  let base = row * HIDDEN;
+  var d = start + tid;
+  loop {
+    if (d >= end) { break; }
+    let index = base + d;
+    // Preserve the ONNX stage boundary: round Normalized to X's dtype before
+    // the affine scale is applied.
+    let normalized = {{ xElement }}(f32(x[index]) * inv);
+    let value = f32(normalized) * f32(scale[scale_offset({% if source.scaleRank > 0 %}index{% endif %})]);
+    y[index] = {{ scalar }}(value);
+    d = d + WG;
+  }
+}

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

	@@ -0,0 +1,93 @@

+{% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" %}
+{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
+{%- else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
+{%- endif %}
+{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
+  var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
+  loop {
+{% if form == "head" %}
+{% if breakInline %}
+    if ({{ svar }} == 0u) { break; }
+{% else %}
+    if ({{ svar }} == 0u) {
+      break;
+    }
+{% endif %}
+{% endif %}
+{% if bodyInline %}
+    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
+{% else %}
+    if ({{ idx }} < {{ svar }}) {
+{% for a in arrays %}
+      {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
+{% endfor %}
+    }
+{% endif %}
+{% if form == "head" %}
+{% if barrierFirst %}
+    workgroupBarrier();
+    {{ svar }} = {{ svar }} / 2u;
+{% else %}
+    {{ svar }} = {{ svar }} / 2u;
+    workgroupBarrier();
+{% endif %}
+{% else %}
+    workgroupBarrier();
+    if ({{ svar }} == 1u) {
+      break;
+    }
+    {{ svar }} = {{ svar }} / 2u;
+{% endif %}
+  }
+{%- endmacro %}
+/* Split-K partial sum-of-squares for tensors with few rows and a large hidden
+   dimension. A workgroup-per-row kernel exposes too little parallelism in this
+   regime, so this pass splits each row across SPLIT workgroups
+   (row = wg.x, split index = wg.z). Each workgroup accumulates a partial
+   sum-of-squares over its HIDDEN/SPLIT slice and writes one partial to scratch.
+   The normalize pass folds the SPLIT partials per row. Split-K reassociates the
+   f32 sum, so this route is not bit-identical to the unsplit reduction. */
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const WG: u32 = {{ workgroupSize }}u;
+const SPLIT: u32 = {{ split }}u;
+var<workgroup> red: array<f32, WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let row = wg.x + wg.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let k = wg.z;
+  let tid = lid.x;
+  let chunk = (HIDDEN + SPLIT - 1u) / SPLIT;
+  let start = k * chunk;
+  var end = start + chunk;
+  if (end > HIDDEN) { end = HIDDEN; }
+  let base = row * HIDDEN;
+  var acc = 0.0;
+  var d = start + tid;
+  loop {
+    if (d >= end) { break; }
+    let v = f32(x[base + d]);
+    acc = acc + v * v;
+    d = d + WG;
+  }
+  red[tid] = acc;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["red"], idx="tid", wg="WG", typed=true, form="head", breakInline=true, bodyInline=true) }}
+  if (tid == 0u) {
+    partials[row * SPLIT + k] = red[0];
+  }
+}

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

	@@ -0,0 +1,140 @@

+{% if usesF16 %}enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const EPSILON: f32 = {{ epsilon }};
+fn round_f16_bits_rte(value: f32) -> u32 {
+  let bits = bitcast<u32>(value);
+  let sign = (bits >> 16u) & 0x8000u;
+  let exponent_f32 = (bits >> 23u) & 0xffu;
+  let mantissa_f32 = bits & 0x7fffffu;
+  if (exponent_f32 == 0xffu) {
+    if (mantissa_f32 != 0u) {
+      return 0x7e00u;
+    }
+    return sign | 0x7c00u;
+  }
+  var exponent_f16 = i32(exponent_f32) - 127 + 15;
+  if (exponent_f16 >= 0x1f) {
+    return sign | 0x7c00u;
+  }
+  if (exponent_f16 <= 0) {
+    if (exponent_f16 < -10) {
+      return sign;
+    }
+    let significand = mantissa_f32 | 0x800000u;
+    let shift = u32(14 - exponent_f16);
+    let halfway = 1u << (shift - 1u);
+    let discarded = significand & ((1u << shift) - 1u);
+    var fraction = significand >> shift;
+    if (discarded > halfway || (discarded == halfway && (fraction & 1u) == 1u)) {
+      fraction = fraction + 1u;
+    }
+    return sign | fraction;
+  }
+  let halfway = 1u << 12u;
+  let discarded = mantissa_f32 & 0x1fffu;
+  var mantissa_f16 = mantissa_f32 >> 13u;
+  if (discarded > halfway || (discarded == halfway && (mantissa_f16 & 1u) == 1u)) {
+    mantissa_f16 = mantissa_f16 + 1u;
+    if (mantissa_f16 == 0x400u) {
+      mantissa_f16 = 0u;
+      exponent_f16 = exponent_f16 + 1;
+    }
+  }
+  if (exponent_f16 >= 0x1f) {
+    return sign | 0x7c00u;
+  }
+  return sign | (u32(exponent_f16) << 10u) | mantissa_f16;
+}
+fn widen_f16_bits(value: u32) -> f32 {
+  return unpack2x16float(value & 0xffffu).x;
+}
+{% if source.scaleRank > 0 %}
+const X_RANK: u32 = {{ source.xRank }}u;
+const SCALE_RANK: u32 = {{ source.scaleRank }}u;
+const X_SHAPE: array<u32, {{ source.xRank }}> = array<u32, {{ source.xRank }}>({% for d in source.xShape %}{{ d }}u{% if not loop.last %}, {% endif %}{% endfor %});
+const SCALE_SHAPE: array<u32, {{ source.scaleRank }}> = array<u32, {{ source.scaleRank }}>({% for d in source.scaleShape %}{{ d }}u{% if not loop.last %}, {% endif %}{% endfor %});
+fn x_stride(axis: u32) -> u32 {
+  var stride = 1u;
+  for (var i = axis + 1u; i < X_RANK; i += 1u) {
+    stride *= X_SHAPE[i];
+  }
+  return stride;
+}
+fn scale_stride(axis: u32) -> u32 {
+  var stride = 1u;
+  for (var i = axis + 1u; i < SCALE_RANK; i += 1u) {
+    stride *= SCALE_SHAPE[i];
+  }
+  return stride;
+}
+{% endif %}
+fn scale_offset({% if source.scaleRank > 0 %}out_index: u32{% endif %}) -> u32 {
+{% if source.scaleRank == 0 %}
+  return 0u;
+{% else %}
+  var rem = out_index;
+  var offset = 0u;
+  for (var axis = 0u; axis < X_RANK; axis += 1u) {
+    let stride = x_stride(axis);
+    let coord = rem / stride;
+    rem %= stride;
+    let scale_axis = i32(axis) - i32(X_RANK - SCALE_RANK);
+    if (scale_axis >= 0) {
+      let s_axis = u32(scale_axis);
+      if (SCALE_SHAPE[s_axis] != 1u) {
+        offset += coord * scale_stride(s_axis);
+      }
+    }
+  }
+  return offset;
+{% endif %}
+}
+// ONNX RMSNormalization-23 expresses the stash_type=FLOAT16 stage as a graph
+// of f16 tensor operators. Keeping one normalization row on one invocation
+// preserves f16 accumulation and every stage boundary.
+@compute @workgroup_size(1, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>) {
+  let row = wg.x + wg.y * params.rowStride;
+  if (row >= params.rows) {
+    return;
+  }
+  let base = row * HIDDEN;
+  // Preserve the reduction state as bits so native compilers cannot reassociate
+  // the loop into a wider-precision sum across binary16 tensor boundaries.
+  var square_sum = round_f16_bits_rte(0.0);
+  for (var d = 0u; d < HIDDEN; d = d + 1u) {
+    let value = round_f16_bits_rte(f32(x[base + d]));
+    let squared = round_f16_bits_rte(widen_f16_bits(value) * widen_f16_bits(value));
+    square_sum = round_f16_bits_rte(widen_f16_bits(square_sum) + widen_f16_bits(squared));
+  }
+  let mean_square = round_f16_bits_rte(widen_f16_bits(square_sum) / f32(HIDDEN));
+  let epsilon_f16 = round_f16_bits_rte(EPSILON);
+  let mean_square_epsilon = round_f16_bits_rte(widen_f16_bits(mean_square) + widen_f16_bits(epsilon_f16));
+  let rms = round_f16_bits_rte(sqrt(widen_f16_bits(mean_square_epsilon)));
+  for (var d = 0u; d < HIDDEN; d = d + 1u) {
+    let index = base + d;
+    let value_f16 = round_f16_bits_rte(f32(x[index]));
+    let normalized = round_f16_bits_rte(widen_f16_bits(value_f16) / widen_f16_bits(rms));
+    let value = widen_f16_bits(normalized)
+      * f32(scale[scale_offset({% if source.scaleRank > 0 %}index{% endif %})]);
+    y[index] = {{ scalar }}(value);
+  }
+}

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

	@@ -0,0 +1,144 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const HIDDEN: u32 = {{ hiddenSize }}u;
+const EPSILON: f32 = {{ epsilon }};
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> partial: array<f32, WG>;
+{% if source.scaleRank > 0 %}
+const X_RANK: u32 = {{ source.xRank }}u;
+const SCALE_RANK: u32 = {{ source.scaleRank }}u;
+const X_SHAPE: array<u32, {{ source.xRank }}> = array<u32, {{ source.xRank }}>({% for d in source.xShape %}{{ d }}u{% if not loop.last %}, {% endif %}{% endfor %});
+const SCALE_SHAPE: array<u32, {{ source.scaleRank }}> = array<u32, {{ source.scaleRank }}>({% for d in source.scaleShape %}{{ d }}u{% if not loop.last %}, {% endif %}{% endfor %});
+fn x_stride(axis: u32) -> u32 {
+  var stride = 1u;
+  for (var i = axis + 1u; i < X_RANK; i += 1u) {
+    stride *= X_SHAPE[i];
+  }
+  return stride;
+}
+fn scale_stride(axis: u32) -> u32 {
+  var stride = 1u;
+  for (var i = axis + 1u; i < SCALE_RANK; i += 1u) {
+    stride *= SCALE_SHAPE[i];
+  }
+  return stride;
+}
+{% endif %}
+fn scale_offset({% if source.scaleRank > 0 %}out_index: u32{% endif %}) -> u32 {
+{% if source.scaleRank == 0 %}
+  return 0u;
+{% else %}
+  var rem = out_index;
+  var offset = 0u;
+  for (var axis = 0u; axis < X_RANK; axis += 1u) {
+    let stride = x_stride(axis);
+    let coord = rem / stride;
+    rem %= stride;
+    let scale_axis = i32(axis) - i32(X_RANK - SCALE_RANK);
+    if (scale_axis >= 0) {
+      let s_axis = u32(scale_axis);
+      if (SCALE_SHAPE[s_axis] != 1u) {
+        offset += coord * scale_stride(s_axis);
+      }
+    }
+  }
+  return offset;
+{% 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_sq = 0.0;
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let value = f32(x[base + d]);
+    local_sq = local_sq + value * value;
+  }
+  let inv = inverseSqrt(reduce_sum(local_sq, tid) / f32(HIDDEN) + EPSILON);
+  for (var d = tid; d < HIDDEN; d = d + WG) {
+    let index = base + d;
+    // ONNX stage one ends by casting Normalized back to X's dtype; Scale is
+    // applied only after that rounding point.
+    let normalized = {{ xElement }}(f32(x[index]) * inv);
+    let value = f32(normalized) * f32(scale[scale_offset({% if source.scaleRank > 0 %}index{% endif %})]);
+    y[base + d] = {{ scalar }}(value);
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1590 @@

+{
+  "op": "ai.onnx.RMSNormalization",
+  "fixtureArrays": {
+    "f16_scalar_cast_x": [-1.1103515625, 2.982421875, 1.248046875, -1.8544921875],
+    "f16_scalar_cast_scale": [2.015625],
+    "f16_scalar_cast_y": [-1.150390625, 3.091796875, 1.29296875, -1.921875],
+    "onnx_backend_rms_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_rms_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],
+    "ort_axis2_vector3_scale_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29],
+    "ort_f16_axis1_outer_inner_broadcast_scale_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23]
+  },
+  "cases": [
+    {
+      "name": "subgroup_vec4_2x512",
+      "attrs": { "epsilon": 0.000001, "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 } }
+    },
+    {
+      "name": "f32_default_epsilon_small_magnitude",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 0.001 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07, "scale": 0.4 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000001 } },
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc",
+        "test": "RMSNormalization-23 schema",
+        "notes": "Pins the default-epsilon path: no epsilon attribute is passed, so manifest and oracle defaults must both match the ONNX schema default 1e-5. Small-magnitude rows make epsilon dominate the mean-square, so a wrong default (e.g. 1e-6) diverges by >2x."
+      }
+    },
+    {
+      "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": "last_axis_3x8",
+      "attrs": { "epsilon": 0.000001, "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_rank3",
+      "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 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_tiny_rms_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 RMS denominator collapses to zero so normalization is non-finite."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm",
+        "notes": "Valid epsilon=0 edge: the RMS denominator is positive subnormal, so tiny normal inputs normalize to finite order-one values."
+      },
+      "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] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_subnormal_scale_last_axis_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero in f32; the subnormal affine scale collapses to zero, losing the tiny output (vec4 path)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale",
+        "notes": "Subnormal scale values are valid; with a nonzero RMS denominator they should produce subnormal outputs, not zeros."
+      },
+      "attrs": { "epsilon": 0, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 4], "data": { "kind": "values", "values": [-1.0, 1.0, 2.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_last_axis_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 affine scale collapses to zero (odd hidden size)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale",
+        "notes": "Odd hidden-size companion for valid subnormal scale outputs."
+      },
+      "attrs": { "epsilon": 0, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [-1.0, 1.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": "ort_basic_1x2x3",
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [0.4629, 0.9258, 1.3887, 0.7895, 0.9869, 1.1843] }
+        }
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm"
+      }
+    },
+    {
+      "name": "mixed_x_f32_scale_y_f16",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc",
+        "test": "RMSNormalization-23 schema",
+        "notes": "Exercises the standard's independent T and V variables with float32 X and float16 scale/output."
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.75, -1.25, 1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 3], "tolerance": 0.002 } }
+    },
+    {
+      "name": "mixed_x_f16_scale_y_f32",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc",
+        "test": "RMSNormalization-23 schema",
+        "notes": "Exercises the standard's independent T and V variables with float16 X and float32 scale/output, including the cast-to-T stage boundary before scaling."
+      },
+      "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": "float32", "shape": [3], "data": { "kind": "values", "values": [0.75, -1.25, 1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "zero_rows_noop",
+      "attrs": { "epsilon": 0.000001, "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.000001, "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": "ort_basic_1x2x3_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_float16"
+      },
+      "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,
+          "data": { "kind": "values", "values": [0.4629, 0.9258, 1.3887, 0.7895, 0.9869, 1.1843] }
+        }
+      }
+    },
+    {
+      "name": "ort_scale_2x2x2_f32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_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.7521, 4.7215, -0.9832, -0.1089, 0.9832, 0.1018, -0.9806, 0.5388] }
+        }
+      }
+    },
+    {
+      "name": "ort_scale_2x2x2_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Float16"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "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": "float16", "shape": [2], "data": { "kind": "values", "values": [-0.6953, 5.1824] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [2, 2, 2],
+          "tolerance": 0.01,
+          "data": { "kind": "values", "values": [0.7521, 4.7215, -0.9832, -0.1089, 0.9832, 0.1018, -0.9806, 0.5388] }
+        }
+      }
+    },
+    {
+      "name": "ort_axis2_vector3_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Vector3_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 5, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axis2_vector3_scale_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.5, 1.5, 1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scalar_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Scalar_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 5, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axis2_vector3_scale_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_batch_outer_broadcast_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Bx1x3_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 1, 3],
+          "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.2, 1.2, 1.2, 1.4, 1.4, 1.4] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_negative_axis_outer_inner_broadcast_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xSx1xW_AxisNeg2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": -2 },
+      "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.2, 1.4, 1.6] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_outer_inner_broadcast_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xSx1xW_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "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.2, 1.4, 1.6] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [0.0, 0.6414, 1.069, 1.9243, 0.9978, 1.4254, 1.4967, 1.9956] }
+        }
+      }
+    },
+    {
+      "name": "ort_f16_axis1_outer_inner_broadcast_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Float16_OuterInnerBroadcast_Axis1"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_f16_axis1_outer_inner_broadcast_scale_input_x" }
+          }
+        },
+        "scale": { "dtype": "float16", "shape": [1, 3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 4], "tolerance": 0.02 } }
+    },
+    {
+      "name": "ort_axis2_ranked_scalar_scale_1x1x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1x1x1_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": [1, 1, 1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_trailing_rank3_scale_1x1x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1x1x3_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 5, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axis2_vector3_scale_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "constant", "value": 1.5 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_outer_sequence_scale_1x4x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xSx3_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_f16_axis1_outer_inner_broadcast_scale_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 4, 3],
+          "data": { "kind": "values", "values": [1.1, 1.1, 1.1, 1.2, 1.2, 1.2, 1.3, 1.3, 1.3, 1.4, 1.4, 1.4] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_full_scale_no_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_NoBroadcast_BxSx3_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 5, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_axis2_vector3_scale_input_x" } }
+        },
+        "scale": { "dtype": "float32", "shape": [2, 5, 3], "data": { "kind": "constant", "value": 1.5 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis1_nchw_channel_scale_1x4x1x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xCx1x1_Axis1"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0]
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "values", "values": [1.1, 1.2, 1.3, 1.4] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis1_rank3_scale_1x3x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xCx1_Axis1"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [1, 3, 1], "data": { "kind": "values", "values": [1.0, 1.2, 1.4] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis1_rank4_scale_1x3x2x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1x3x2x1_Axis1"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 1.1, 1.2, 1.3, 1.4, 1.5] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scale_1x1xHx1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1x1xHx1_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "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, 1, 2, 1], "data": { "kind": "values", "values": [1.0, 1.3] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scale_1x1x1xW",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1x1x1xW_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [1, 1, 1, 3], "data": { "kind": "values", "values": [1.0, 1.2, 1.4] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scale_1xSx1x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xSx1x1_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [1, 3, 1, 1], "data": { "kind": "values", "values": [1.0, 1.2, 1.4] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scale_Bx1x1xW",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Bx1x1xW_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 1, 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": [2, 1, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 1.1, 1.3, 1.4] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis2_scale_1x1xHxW",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1x1xHxW_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 1.1, 1.2, 1.3, 1.4, 1.5] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis3_rank5_scale_1xSx1x1xC",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_1xSx1x1xC_Axis3"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_f16_axis1_outer_inner_broadcast_scale_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 1, 3],
+          "data": { "kind": "values", "values": [1.0, 1.1, 1.2, 1.3, 1.4, 1.5] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2, 2, 3], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_f16_axis2_outer_broadcast_scale",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_Scale_Float16_OuterBroadcast_BxSx1_Axis2"
+      },
+      "attrs": { "epsilon": 0.00001, "axis": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0] }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [2, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 3], "tolerance": 0.02 } }
+    },
+    {
+      "name": "ort_axis1_inner_mixed_scale_1x4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/rms_norm_op_test.cc",
+        "test": "RMSNormalizationOpTest.RMSNorm_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": "onnx_backend_rms_normalization_2d_axis0",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_2d_axis0" },
+      "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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.0001 } },
+      "attrs": { "axis": 0 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_2d_axis1",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_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": [1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.0001 } },
+      "attrs": { "axis": 1 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_2d_axis_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_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": [-1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.0001 } },
+      "attrs": { "axis": -1 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_2d_axis_negative_2",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_2d_axis_negative_2"
+      },
+      "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": [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.0001 } },
+      "attrs": { "axis": -2 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_3d_axis0_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_3d_axis0_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 0, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_3d_axis1_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_3d_axis1_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 1, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_3d_axis2_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_3d_axis2_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [0.37642553448677063, -1.0994007587432861, 0.29823818802833557, 1.3263858556747437, -0.694567859172821]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 2, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_3d_axis_negative_1_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_3d_axis_negative_1_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "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, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": -1, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_3d_axis_negative_2_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_3d_axis_negative_2_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": -2, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_3d_axis_negative_3_epsilon",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_3d_axis_negative_3_epsilon"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_3d_input_x" }
+          }
+        },
+        "scale": {
+          "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.0001 } },
+      "attrs": { "axis": -3, "epsilon": 0.10000000149011612 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis0",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis0" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 0 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis1",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis1" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 1 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis2",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis2" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 2 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis3",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis3" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.22260567545890808, -0.9130792021751404, -1.6812182664871216, -0.8889713287353516, 0.2421179562807083]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": 3 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis_negative_1"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [-0.8887202739715576, 0.9367424845695496, 1.4123276472091675, -2.369586944580078, 0.8640522956848145]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0002 } },
+      "attrs": { "axis": -1 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis_negative_2",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis_negative_2"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4, 5],
+          "data": {
+            "kind": "values",
+            "values": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": -2 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis_negative_3",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis_negative_3"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } },
+      "attrs": { "axis": -3 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_4d_axis_negative_4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_4d_axis_negative_4"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "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.0001 } },
+      "attrs": { "axis": -4 }
+    },
+    {
+      "name": "onnx_backend_rms_normalization_default_axis",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_rms_normalization_default_axis" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_rms_normalization_input_x" } }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": {
+            "kind": "values",
+            "values": [0.37642553448677063, -1.0994007587432861, 0.29823818802833557, 1.3263858556747437, -0.694567859172821]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 4, 5], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "f16_last_axis_vec4_hidden2048_multisubgroup",
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [64, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.021, "scale": 0.4 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.011, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [64, 2048], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_last_axis_unaligned_hidden2050_alignment_cliff",
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [32, 2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.017, "scale": 0.4 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [2050],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.013, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [32, 2050], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_suffix_axis1_nchw_channel_norm",
+      "attrs": { "epsilon": 0.000001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 64, 16, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.019, "scale": 0.4 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [1, 64, 1, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.027, "cosStep": 0.011, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 16, 16], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f32_last_axis_vec4_hidden1024_scale_offset_rank3",
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 8, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.021, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.009, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 8, 1024], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_suffix_axis_splitk_2x64x512",
+      "attrs": { "epsilon": 0.000001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 64, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.023, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [64, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.009, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 512], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "f32_suffix_axis_splitk_scalar_scale_1x32x512",
+      "attrs": { "epsilon": 0.000001, "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.013, "scale": 0.5 }
+        },
+        "scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.75] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 512], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "empty_outer_0x3x4_suffix_axis",
+      "attrs": { "epsilon": 0.000001, "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 3, 4], "data": { "kind": "values", "values": [] } },
+        "scale": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.009, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [0, 3, 4],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [] }
+        }
+      }
+    },
+    {
+      "name": "empty_hidden_2x0_lastaxis_array0_const",
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "scale": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 0], "tolerance": 0.000001, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "rows_fold_65537x8_lastaxis_vec4_dispatch2d",
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [65537, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.021, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.009, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [65537, 8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "lastaxis_broadcast_scale1_fallback_hidden8",
+      "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": [1], "data": { "kind": "values", "values": [0.75] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "all_zero_row_epsilon_floor_finite_not_nan",
+      "attrs": { "epsilon": 0.01, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 2.0, -2.0, 2.0, -2.0] }
+        },
+        "scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "f16_lastaxis_unaligned_hidden4094",
+      "provenance": {
+        "notes": "Compact correctness sibling for the f16 unaligned-hidden benchmark cliff; dim=4094 exercises the scalar tail path next to the vec4/subgroup path."
+      },
+      "attrs": { "epsilon": 0.000001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 4094],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.023, "scale": 0.5 }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [4094],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.009, "scale": 0.25, "offset": 1.0 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 4094], "tolerance": 0.02 } }
+    },
+    {
+      "name": "onnx23_f16_cast_before_scale_row_vec4_exact",
+      "provenance": {
+        "notes": "ONNX RMSNormalization-23 stage one casts Normalized back to X's dtype before Scale. This vector-scale case selects the optimized vec4 row kernel and differs by one f16 ULP if the cast is deferred until after Scale."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [-1.4765625, -1.861328125, 1.4697265625, 1.0703125] }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0.316162109375, -2.791015625, -0.0877685546875, 2.0859375] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "tolerance": 0,
+          "relTolerance": 0,
+          "data": { "kind": "values", "values": [-0.312255859375, 3.47265625, -0.08624267578125, 1.4921875] }
+        }
+      }
+    },
+    {
+      "name": "onnx23_f16_cast_before_scale_generic_scalar_exact",
+      "provenance": {
+        "notes": "ONNX RMSNormalization-23 stage-one cast boundary on the generic broadcast kernel. A scalar Scale prevents the specialized row variant; two outputs differ by one f16 ULP if normalization stays in float32 through the multiply."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_scalar_cast_x" } }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_scalar_cast_scale" } }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "tolerance": 0,
+          "relTolerance": 0,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_scalar_cast_y" } }
+        }
+      }
+    },
+    {
+      "name": "onnx23_f16_cast_before_scale_splitk_exact",
+      "provenance": {
+        "notes": "Forces the split-K shared kernel on the same exact f16 boundary case, pinning the standardized cast of Normalized back to X before Scale."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "tunables": { "SPLIT_MIN_HIDDEN": 1, "SPLIT_TARGET_ELEMENTS": 1 },
+      "attrs": { "epsilon": 0.00001, "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_scalar_cast_x" } }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_scalar_cast_scale" } }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "tolerance": 0,
+          "relTolerance": 0,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_scalar_cast_y" } }
+        }
+      }
+    },
+    {
+      "name": "onnx23_f16_stash_float16_rounding",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc RMSNormalization-23 function body",
+        "notes": "TensorProto FLOAT16 stash (10). Sequential f16 sum-of-squares and normalization produce outputs separated by multiple f16 ULPs from the default float32-stash route."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "epsilon": 0.00001, "axis": -1, "stash_type": 10 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [8.171875, 0.15771484375, -0.1475830078125, -2.125] }
+        },
+        "scale": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [3.640625, -3.275390625, -1.51953125, -0.85205078125] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "tolerance": 0.004,
+          "relTolerance": 0,
+          "data": { "kind": "values", "values": [7.0546875, -0.1224365234375, 0.053131103515625, 0.429443359375] }
+        }
+      }
+    },
+    {
+      "name": "onnx23_f32_stash_float16_rounding",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc RMSNormalization-23 function body",
+        "notes": "TensorProto FLOAT16 stash (10) with float32 X and V. Non-f16 input and scale values expose both the cast into float16 stage one and the cast of Normalized back to float32 before the affine multiply."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "epsilon": 0.00001, "axis": -1, "stash_type": 10 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [2.2851, -3.3284, 0.07473, 0.8293] }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [2.3753, 2.8352, 2.3861, 1.0394] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 0.000001,
+          "relTolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [2.635098457336426, -4.579512596130371, 0.08650777488946915, 0.41819608211517334]
+          }
+        }
+      }
+    },
+    {
+      "name": "onnx23_f32_stash_float16_scalar_scale",
+      "provenance": {
+        "source": "onnx/defs/nn/defs.cc RMSNormalization-23 function body",
+        "notes": "Float32 X with TensorProto FLOAT16 stash (10) and a scalar float32 Scale covers the scalar broadcast path while preserving the float16 normalization and float32 affine boundaries."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "epsilon": 0.00001, "axis": -1, "stash_type": 10 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [8.1716, 0.15772, -0.14759, -2.1252] }
+        },
+        "scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.234567] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "tolerance": 0.000001,
+          "relTolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [2.3919737339019775, 0.04615309461951256, -0.043176691979169846, -0.6221060752868652]
+          }
+        }
+      }
+    },
+    {
+      "name": "explicit_stash_type_float32",
+      "attrs": { "epsilon": 0.00001, "axis": -1, "stash_type": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.37, "scale": 0.5 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000002 } }
+    }
+  ]
+}