sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +84 -0
build/webgpu/bench.json +108 -0
build/webgpu/embed-mask-index.wgsl.jinja +28 -0
build/webgpu/embed-normalize.wgsl.jinja +114 -0
build/webgpu/embed-sum.wgsl.jinja +81 -0
build/webgpu/manifest.json +1708 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +1325 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.EmbedLayerNormalization
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+BERT embedding fusion: looks up word and position tables, optionally adds a segment table, then applies layer normalization. A segment table without IDs uses row 0. `embedding_sum` is the pre-normalization sum. `mask_index` is the first zero or the sequence length; without `mask`, it is zero. Batch and sequence dimensions must be non-empty.
+See the [ONNX Runtime `EmbedLayerNormalization` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.EmbedLayerNormalization) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `input_ids` | `inputIdsT` | `T1` | `int32` | `2` | — | Word ids of shape `(batch_size, sequence_length)`. | required |
+| `segment_ids` | `segmentIdsT` | `T1` | `int32` | `2` | — | Segment ids `(batch_size, sequence_length)`. Requires `segment_embedding`; when omitted with that table present, every token uses row 0. Values must be valid non-negative table-row indices. | optional |
+| `word_embedding` | `wordEmbeddingT` | `T` | same as logical dtype | `2` | — | Non-empty word embedding table `(vocab, hidden_size)`. Every `input_ids` value must be a valid non-negative row index. | required |
+| `position_embedding` | `positionEmbeddingT` | `T` | same as logical dtype | `2` | — | Non-empty position embedding table `(max_positions, hidden_size)`. Without `position_ids`, it must contain at least `sequence_length` rows. | required |
+| `segment_embedding` | `segmentEmbeddingT` | `T` | same as logical dtype | `2` | — | Non-empty segment embedding table `(segments, hidden_size)`. If `segment_ids` is absent, row 0 is used for every token. | optional |
+| `gamma` | `gammaT` | `T` | same as logical dtype | `1` | — | Layer-normalization scale of shape `(hidden_size)`. | required |
+| `beta` | `betaT` | `T` | same as logical dtype | `1` | — | Layer-normalization bias of shape `(hidden_size)`. | required |
+| `mask` | `maskT` | `T1` | `int32` | `2` | — | Attention mask of shape `(batch_size, sequence_length)`. Only used to produce `mask_index`. | optional |
+| `position_ids` | `positionIdsT` | `T1` | `int32` | `2` | — | Position ids `(batch_size, sequence_length)`, or `(1, sequence_length)` to share one row across the batch. Values must be valid non-negative table-row indices; absent uses the position within the sequence. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `output` | `outputT` | `T` | same as logical dtype | `3` | derived; see description | Normalized embeddings of shape `(batch_size, sequence_length, hidden_size)`. | required |
+| `mask_index` | `maskIndexT` | `T1` | `int32` | `1` | `[input_ids[0]]` | Position of the first zero in each mask row, or `sequence_length` when no zero exists; shape `(batch_size)`. It is zero when the optional mask input is absent. | optional |
+| `embedding_sum` | `embeddingSumT` | `T` | same as logical dtype | `3` | derived; see description | The summed embeddings before normalization, including the segment term when present. Float16 uses staged `(word + segment) + position`; float32 uses `(word + position) + segment`. | optional |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `epsilon` | `9.999999960041972e-13` | Non-negative epsilon added to the layer-normalization variance before taking the square root. |
+| `mask_index_type` | — | Optional shape-inference hint for the `mask_index` output type. The schema's `T1` constraint fixes the runtime tensor type to int32. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+| `T1` | `int32` |
+## 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
+- [`embed-mask-index.wgsl.jinja`](build/webgpu/embed-mask-index.wgsl.jinja)
+- [`embed-normalize.wgsl.jinja`](build/webgpu/embed-normalize.wgsl.jinja)
+- [`embed-sum.wgsl.jinja`](build/webgpu/embed-sum.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/com.microsoft.EmbedLayerNormalization", { version: 1 });
+const { outputT } = await kernel({
+  inputIdsT: { data: inputIdsTData, shape: [1, 2] },
+  wordEmbeddingT: { data: wordEmbeddingTData, shape: [2, 2] },
+  positionEmbeddingT: { data: positionEmbeddingTData, shape: [2, 2] },
+  gammaT: { data: gammaTData, shape: [2] },
+  betaT: { data: betaTData, shape: [2] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,108 @@

+{
+  "op": "com.microsoft.EmbedLayerNormalization",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "eln-bert-base-b8-s384-h768",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "inputs": {
+        "inputIdsT": { "shape": [8, 384], "dtype": "int32", "dist": "randint", "seed": 9301, "min": 0, "max": 30521 },
+        "wordEmbeddingT": { "shape": [30522, 768], "dtype": "float32", "dist": "normal", "seed": 9302, "scale": 0.02 },
+        "positionEmbeddingT": { "shape": [512, 768], "dtype": "float32", "dist": "normal", "seed": 9303, "scale": 0.02 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 9304, "scale": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 9305, "scale": 0.1 }
+      },
+      "outputs": { "outputT": { "shape": [8, 384, 768], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "8 * 384 * 768 * 3 * 4" }] }
+    },
+    {
+      "name": "eln-bert-base-b1-s512-h768",
+      "preset": "model",
+      "provenance": {
+        "notes": "BERT class defaults (vocab_size 30522, hidden_size 768, max_position_embeddings 512, type_vocab_size 2) at a single full-length sequence."
+      },
+      "vars": { "batch": 1, "seq": 512, "hidden": 768, "vocab": 30522 },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputIdsT": { "shape": [1, 512], "dtype": "int32", "dist": "randint", "seed": 5100, "min": 0, "max": 30521 },
+        "segmentIdsT": { "shape": [1, 512], "dtype": "int32", "dist": "randint", "seed": 5101, "min": 0, "max": 1 },
+        "wordEmbeddingT": { "shape": [30522, 768], "dtype": "float32", "dist": "normal", "seed": 5102, "scale": 0.1 },
+        "positionEmbeddingT": { "shape": [512, 768], "dtype": "float32", "dist": "normal", "seed": 5103, "scale": 0.1 },
+        "segmentEmbeddingT": { "shape": [2, 768], "dtype": "float32", "dist": "normal", "seed": 5104, "scale": 0.1 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 5105, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 5106, "scale": 0.05 },
+        "maskT": { "shape": [1, 512], "dtype": "int32", "dist": "randint", "seed": 5107, "min": 0, "max": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 512, 768], "dtype": "float32" },
+        "maskIndexT": { "shape": [1], "dtype": "int32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.seq * args.hidden * 3 * 4" }] }
+    },
+    {
+      "name": "eln-bert-base-b32-s128-h768",
+      "preset": "model",
+      "provenance": {
+        "notes": "BERT class defaults at a 32-sequence serving batch of 128 tokens, the shape short-text classification actually runs."
+      },
+      "vars": { "batch": 32, "seq": 128, "hidden": 768, "vocab": 30522 },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputIdsT": { "shape": [32, 128], "dtype": "int32", "dist": "randint", "seed": 5200, "min": 0, "max": 30521 },
+        "segmentIdsT": { "shape": [32, 128], "dtype": "int32", "dist": "randint", "seed": 5201, "min": 0, "max": 1 },
+        "wordEmbeddingT": { "shape": [30522, 768], "dtype": "float32", "dist": "normal", "seed": 5202, "scale": 0.1 },
+        "positionEmbeddingT": { "shape": [512, 768], "dtype": "float32", "dist": "normal", "seed": 5203, "scale": 0.1 },
+        "segmentEmbeddingT": { "shape": [2, 768], "dtype": "float32", "dist": "normal", "seed": 5204, "scale": 0.1 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 5205, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 5206, "scale": 0.05 },
+        "maskT": { "shape": [32, 128], "dtype": "int32", "dist": "randint", "seed": 5207, "min": 0, "max": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [32, 128, 768], "dtype": "float32" },
+        "maskIndexT": { "shape": [32], "dtype": "int32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.seq * args.hidden * 3 * 4" }] }
+    },
+    {
+      "name": "eln-distilbert-b8-s512-h768",
+      "preset": "model",
+      "provenance": {
+        "notes": "DistilBERT class defaults (vocab_size 30522, dim 768, max_position_embeddings 512); DistilBERT has no token-type embedding, so the segment inputs are absent."
+      },
+      "vars": { "batch": 8, "seq": 512, "hidden": 768, "vocab": 30522 },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputIdsT": { "shape": [8, 512], "dtype": "int32", "dist": "randint", "seed": 5300, "min": 0, "max": 30521 },
+        "wordEmbeddingT": { "shape": [30522, 768], "dtype": "float32", "dist": "normal", "seed": 5302, "scale": 0.1 },
+        "positionEmbeddingT": { "shape": [512, 768], "dtype": "float32", "dist": "normal", "seed": 5303, "scale": 0.1 },
+        "gammaT": { "shape": [768], "dtype": "float32", "dist": "uniform", "seed": 5305, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [768], "dtype": "float32", "dist": "normal", "seed": 5306, "scale": 0.05 },
+        "maskT": { "shape": [8, 512], "dtype": "int32", "dist": "randint", "seed": 5307, "min": 0, "max": 1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [8, 512, 768], "dtype": "float32" },
+        "maskIndexT": { "shape": [8], "dtype": "int32" }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.seq * args.hidden * 3 * 4" }] }
+    },
+    {
+      "name": "eln-deberta-v2-b8-s512-h1536",
+      "preset": "model",
+      "provenance": {
+        "notes": "DeBERTa-v2 class defaults (vocab_size 128100, hidden_size 1536, max_position_embeddings 512, type_vocab_size 0) -- the widest embedding table in this set."
+      },
+      "vars": { "batch": 8, "seq": 512, "hidden": 1536, "vocab": 128100 },
+      "attrs": { "epsilon": 1e-12 },
+      "inputs": {
+        "inputIdsT": { "shape": [8, 512], "dtype": "int32", "dist": "randint", "seed": 5400, "min": 0, "max": 128099 },
+        "wordEmbeddingT": { "shape": [128100, 1536], "dtype": "float32", "dist": "normal", "seed": 5402, "scale": 0.1 },
+        "positionEmbeddingT": { "shape": [512, 1536], "dtype": "float32", "dist": "normal", "seed": 5403, "scale": 0.1 },
+        "gammaT": { "shape": [1536], "dtype": "float32", "dist": "uniform", "seed": 5405, "scale": 0.1, "offset": 1 },
+        "betaT": { "shape": [1536], "dtype": "float32", "dist": "normal", "seed": 5406, "scale": 0.05 }
+      },
+      "outputs": { "outputT": { "shape": [8, 512, 1536], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.seq * args.hidden * 3 * 4" }] }
+    }
+  ]
+}

build/webgpu/embed-mask-index.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,28 @@

+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.EmbedLayerNormalization, mask_index pass. With a mask, return
+// the first zero position or the sequence length when every position is set.
+// Without a mask, initialize the optional output to zero.
+{% if hasMask %}
+const SEQUENCE: u32 = {{ sequenceLength }}u;
+{% endif %}
+@compute @workgroup_size({{ maskWorkgroupSize }}, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let batch = gid.x;
+  if (batch >= params.batch) {
+    return;
+  }
+  var first_zero: i32 = 0;
+{% if hasMask %}
+  first_zero = i32(SEQUENCE);
+  let base = batch * SEQUENCE;
+  for (var s: u32 = 0u; s < SEQUENCE; s = s + 1u) {
+    if (mask[base + s] == 0) {
+      first_zero = i32(s);
+      break;
+    }
+  }
+{% endif %}
+  mask_index[batch] = first_zero;
+}

build/webgpu/embed-normalize.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.EmbedLayerNormalization, normalization pass.
+//   output = (sum - mean) / sqrt(variance + epsilon) * gamma + beta
+// One workgroup owns one (batch, position) row of the summed embedding the
+// previous pass left in `output`, and normalizes it in place. The statistics
+// accumulate in f32 over the stored tensor-type values.
+const HIDDEN: u32 = {{ hidden }}u;
+const EPSILON: f32 = {{ epsilon }};
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> partial: array<f32, WG>;
+{% 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(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  // 2D-folded row index: wg.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+  let token = wg.x + wg.y * nwg.x;
+  if (token >= params.tokens) {
+    return;
+  }
+  let tid = lid.x;
+  let base = token * HIDDEN;
+  var local_sum = 0.0;
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    local_sum = local_sum + f32(output[base + i]);
+  }
+  let mean = reduce_sum(local_sum, tid) / f32(HIDDEN);
+  // Separates the mean reduction's read of partial[0] from the variance
+  // reduction's writes to the same workgroup array.
+  workgroupBarrier();
+  var local_sq = 0.0;
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let centred = f32(output[base + i]) - mean;
+    local_sq = local_sq + centred * centred;
+  }
+  let deviation = sqrt(reduce_sum(local_sq, tid) / f32(HIDDEN) + EPSILON);
+  workgroupBarrier();
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    let centred = f32(output[base + i]) - mean;
+    output[base + i] = {{ scalar }}(centred / deviation * f32(gamma[i]) + f32(beta[i]));
+  }
+}

build/webgpu/embed-sum.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,81 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.EmbedLayerNormalization, embedding-sum pass.
+//   sum = word_embedding[input_ids[t]] + position_embedding[p(t)]
+//         (+ segment_embedding[segment_ids[t] or 0])
+// One workgroup writes one token's row of the summed embedding into `output`,
+// where the normalization pass reads it back. The row bases -- which word,
+// position and segment rows this token gathers -- are the same for every element
+// of the row, so they are resolved once per workgroup.
+// The sum is materialized in the tensor type -- it is also the optional
+// `embedding_sum` output -- so the statistics the next pass computes describe
+// the values this op emits rather than an unobservable wider intermediate.
+const HIDDEN: u32 = {{ hidden }}u;
+{% if not hasPositionIds or broadcastPositionIds %}
+// Read only where the position row comes from the position within the sequence:
+// a per-token position_ids indexes directly instead.
+const SEQUENCE: u32 = {{ sequenceLength }}u;
+{% endif %}
+const WG: u32 = {{ workgroupSize }}u;
+// Defensive row clamps keep an invalid id from reading outside an embedding
+// table. Request validation rejects such ids before a valid dispatch.
+const WORD_ROWS: u32 = {{ wordRows }}u;
+const POSITION_ROWS: u32 = {{ positionRows }}u;
+{% if hasSegment and hasSegmentIds %}
+const SEGMENT_ROWS: u32 = {{ segmentRows }}u;
+{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  // 2D-folded row index: wg.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+  let token = wg.x + wg.y * nwg.x;
+  if (token >= params.tokens) {
+    return;
+  }
+  let tid = lid.x;
+  let base = token * HIDDEN;
+  let word_row = min(u32(max(input_ids[token], 0)), WORD_ROWS - 1u) * HIDDEN;
+{% if hasPositionIds %}
+{% if broadcastPositionIds %}
+  // A position_ids of shape (1, sequence_length) is shared by every batch, so
+  // the row index drops back to the position within the sequence.
+  let position_id = position_ids[token % SEQUENCE];
+{% else %}
+  let position_id = position_ids[token];
+{% endif %}
+  let position_row = min(u32(max(position_id, 0)), POSITION_ROWS - 1u) * HIDDEN;
+{% else %}
+  let position_row = min(token % SEQUENCE, POSITION_ROWS - 1u) * HIDDEN;
+{% endif %}
+{% if hasSegment %}
+{% if hasSegmentIds %}
+  let segment_row = min(u32(max(segment_ids[token], 0)), SEGMENT_ROWS - 1u) * HIDDEN;
+{% else %}
+  let segment_row = 0u;
+{% endif %}
+{% endif %}
+  for (var i = tid; i < HIDDEN; i = i + WG) {
+    // Materialize each addition in T; association is observable for f16.
+{% if hasSegment and scalar == "f16" %}
+    var stored = {{ scalar }}(f32(word_embedding[word_row + i]) + f32(segment_embedding[segment_row + i]));
+    stored = {{ scalar }}(f32(stored) + f32(position_embedding[position_row + i]));
+{% else %}
+    var stored = {{ scalar }}(f32(word_embedding[word_row + i]) + f32(position_embedding[position_row + i]));
+{% if hasSegment %}
+    stored = {{ scalar }}(f32(stored) + f32(segment_embedding[segment_row + i]));
+{% endif %}
+{% endif %}
+{% if writeEmbeddingSum %}
+    embedding_sum[base + i] = stored;
+{% endif %}
+    output[base + i] = stored;
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1708 @@

+{
+  "domain": "com.microsoft",
+  "name": "EmbedLayerNormalization",
+  "sinceVersion": 1,
+  "description": "BERT embedding fusion: looks up word and position tables, optionally adds a segment table, then applies layer normalization. A segment table without IDs uses row 0. `embedding_sum` is the pre-normalization sum. `mask_index` is the first zero or the sequence length; without `mask`, it is zero. Batch and sequence dimensions must be non-empty.",
+  "inputs": [
+    {
+      "role": "input_ids",
+      "dtype": "T1",
+      "rank": 2,
+      "description": "Word ids of shape `(batch_size, sequence_length)`."
+    },
+    {
+      "role": "segment_ids",
+      "dtype": "T1",
+      "rank": 2,
+      "optional": true,
+      "description": "Segment ids `(batch_size, sequence_length)`. Requires `segment_embedding`; when omitted with that table present, every token uses row 0. Values must be valid non-negative table-row indices."
+    },
+    {
+      "role": "word_embedding",
+      "dtype": "T",
+      "rank": 2,
+      "description": "Non-empty word embedding table `(vocab, hidden_size)`. Every `input_ids` value must be a valid non-negative row index."
+    },
+    {
+      "role": "position_embedding",
+      "dtype": "T",
+      "rank": 2,
+      "description": "Non-empty position embedding table `(max_positions, hidden_size)`. Without `position_ids`, it must contain at least `sequence_length` rows."
+    },
+    {
+      "role": "segment_embedding",
+      "dtype": "T",
+      "rank": 2,
+      "optional": true,
+      "description": "Non-empty segment embedding table `(segments, hidden_size)`. If `segment_ids` is absent, row 0 is used for every token."
+    },
+    { "role": "gamma", "dtype": "T", "rank": 1, "description": "Layer-normalization scale of shape `(hidden_size)`." },
+    { "role": "beta", "dtype": "T", "rank": 1, "description": "Layer-normalization bias of shape `(hidden_size)`." },
+    {
+      "role": "mask",
+      "dtype": "T1",
+      "rank": 2,
+      "optional": true,
+      "description": "Attention mask of shape `(batch_size, sequence_length)`. Only used to produce `mask_index`."
+    },
+    {
+      "role": "position_ids",
+      "dtype": "T1",
+      "rank": 2,
+      "optional": true,
+      "description": "Position ids `(batch_size, sequence_length)`, or `(1, sequence_length)` to share one row across the batch. Values must be valid non-negative table-row indices; absent uses the position within the sequence."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": 3,
+      "shape": "[dim(shapes.inputIdsT, 0), dim(shapes.inputIdsT, 1), hidden]",
+      "description": "Normalized embeddings of shape `(batch_size, sequence_length, hidden_size)`."
+    },
+    {
+      "role": "mask_index",
+      "dtype": "T1",
+      "rank": 1,
+      "optional": true,
+      "shape": "[dim(shapes.inputIdsT, 0)]",
+      "description": "Position of the first zero in each mask row, or `sequence_length` when no zero exists; shape `(batch_size)`. It is zero when the optional mask input is absent."
+    },
+    {
+      "role": "embedding_sum",
+      "dtype": "T",
+      "rank": 3,
+      "optional": true,
+      "shape": "[dim(shapes.inputIdsT, 0), dim(shapes.inputIdsT, 1), hidden]",
+      "description": "The summed embeddings before normalization, including the segment term when present. Float16 uses staged `(word + segment) + position`; float32 uses `(word + position) + segment`."
+    }
+  ],
+  "attributes": { "epsilon": 9.999999960041972e-13 },
+  "attributeDescriptions": {
+    "epsilon": "Non-negative epsilon added to the layer-normalization variance before taking the square root.",
+    "mask_index_type": "Optional shape-inference hint for the `mask_index` output type. The schema's `T1` constraint fixes the runtime tensor type to int32."
+  },
+  "attributeConstraints": { "mask_index_type": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32", "float16"], "T1": ["int32"] },
+  "args": {
+    "inputIdsT": { "kind": "tensor", "semantic": "input_ids", "role": "input", "dtype": "int32" },
+    "segmentIdsT": { "kind": "tensor", "semantic": "segment_ids", "role": "input", "dtype": "int32", "required": false },
+    "wordEmbeddingT": { "kind": "tensor", "semantic": "word_embedding", "role": "weights" },
+    "positionEmbeddingT": { "kind": "tensor", "semantic": "position_embedding", "role": "weights" },
+    "segmentEmbeddingT": { "kind": "tensor", "semantic": "segment_embedding", "role": "weights", "required": false },
+    "gammaT": { "kind": "tensor", "semantic": "gamma", "role": "weights" },
+    "betaT": { "kind": "tensor", "semantic": "beta", "role": "weights" },
+    "maskT": { "kind": "tensor", "semantic": "mask", "role": "input", "dtype": "int32", "required": false },
+    "positionIdsT": {
+      "kind": "tensor",
+      "semantic": "position_ids",
+      "role": "input",
+      "dtype": "int32",
+      "required": false
+    },
+    "outputT": { "kind": "tensor", "semantic": "output", "role": "output" },
+    "maskIndexT": { "kind": "tensor", "semantic": "mask_index", "role": "output", "dtype": "int32", "required": false },
+    "embeddingSumT": { "kind": "tensor", "semantic": "embedding_sum", "role": "output", "required": false }
+  },
+  "tunables": { "WORKGROUP_SIZE": 128, "MASK_WORKGROUP_SIZE": 64 },
+  "derive": {
+    "batchSize": "dim(shapes.inputIdsT, 0)",
+    "sequenceLength": "dim(shapes.inputIdsT, 1)",
+    "tokens": "batchSize * sequenceLength",
+    "hidden": "dim(shapes.wordEmbeddingT, 1)",
+    "epsilonValue": "attrs.epsilon",
+    "epsilonOk": "epsilonValue >= 0",
+    "tableShapeOk": "ranks.wordEmbeddingT == 2 and dim(shapes.wordEmbeddingT, 0) > 0 and ranks.positionEmbeddingT == 2 and dim(shapes.positionEmbeddingT, 0) > 0 and dim(shapes.positionEmbeddingT, 1) == hidden and ranks.gammaT == 1 and ranks.betaT == 1 and dim(shapes.gammaT, 0) == hidden and dim(shapes.betaT, 0) == hidden and hidden > 0",
+    "segmentContract": "(not present.segmentIdsT or present.segmentEmbeddingT) and (ranks.segmentIdsT == 2 and sameShape(shapes.segmentIdsT, shapes.inputIdsT) if present.segmentIdsT else true) and (ranks.segmentEmbeddingT == 2 and dim(shapes.segmentEmbeddingT, 0) > 0 and dim(shapes.segmentEmbeddingT, 1) == hidden if present.segmentEmbeddingT else true)",
+    "positionIdsContract": "(ranks.positionIdsT == 2 and dim(shapes.positionIdsT, 1) == sequenceLength and (dim(shapes.positionIdsT, 0) == batchSize or dim(shapes.positionIdsT, 0) == 1) if present.positionIdsT else dim(shapes.positionEmbeddingT, 0) >= sequenceLength)",
+    "broadcastPositionIds": "dim(shapes.positionIdsT, 0) == 1 if present.positionIdsT else false",
+    "maskContract": "ranks.maskT == 2 and sameShape(shapes.maskT, shapes.inputIdsT) if present.maskT else true",
+    "maskIndexTypeOk": "not has(attrs, \"mask_index_type\") or attrs.mask_index_type == 0 or attrs.mask_index_type == 1",
+    "ioShapeOk": "ranks.inputIdsT == 2 and ranks.outputT == 3 and dim(shapes.outputT, 0) == batchSize and dim(shapes.outputT, 1) == sequenceLength and dim(shapes.outputT, 2) == hidden and tensorDtypes.outputT == tensorDtypes.wordEmbeddingT and tensorDtypes.positionEmbeddingT == tensorDtypes.wordEmbeddingT and tensorDtypes.gammaT == tensorDtypes.wordEmbeddingT and tensorDtypes.betaT == tensorDtypes.wordEmbeddingT and f16Ok(tensorDtypes.wordEmbeddingT)",
+    "embeddingSumContract": "ranks.embeddingSumT == 3 and sameShape(shapes.embeddingSumT, shapes.outputT) and tensorDtypes.embeddingSumT == tensorDtypes.wordEmbeddingT if present.embeddingSumT else true",
+    "maskIndexShapeOk": "ranks.maskIndexT == 1 and dim(shapes.maskIndexT, 0) == batchSize if present.maskIndexT else true",
+    "embedContractOk": "epsilonOk and tableShapeOk and segmentContract and positionIdsContract and maskContract and maskIndexTypeOk and ioShapeOk and embeddingSumContract and maskIndexShapeOk and batchSize > 0 and sequenceLength > 0",
+    "dispatchFits": "tunables.WORKGROUP_SIZE <= device.limits.maxComputeInvocationsPerWorkgroup and tunables.MASK_WORKGROUP_SIZE <= device.limits.maxComputeInvocationsPerWorkgroup"
+  },
+  "constants": {
+    "aScalar": "dtypes.T",
+    "scalar": "dtypes.T",
+    "usesF16": "dtypes.T == \"f16\"",
+    "hidden": "hidden",
+    "sequenceLength": "sequenceLength",
+    "epsilon": "epsilonValue",
+    "workgroupSize": "tunables.WORKGROUP_SIZE",
+    "maskWorkgroupSize": "tunables.MASK_WORKGROUP_SIZE",
+    "wordRows": "dim(shapes.wordEmbeddingT, 0)",
+    "positionRows": "dim(shapes.positionEmbeddingT, 0)",
+    "segmentRows": "dim(shapes.segmentEmbeddingT, 0) if present.segmentEmbeddingT else 1",
+    "hasSegment": "present.segmentEmbeddingT",
+    "hasSegmentIds": "present.segmentIdsT",
+    "hasPositionIds": "present.positionIdsT",
+    "broadcastPositionIds": "broadcastPositionIds",
+    "writeEmbeddingSum": "present.embeddingSumT",
+    "hasMask": "present.maskT",
+    "HIDDEN_LEN": "hidden"
+  },
+  "bindingSets": {
+    "embed_noseg_nopos_nosum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_noseg_nopos_sum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "embedding_sum",
+        "arg": "embeddingSumT",
+        "semantic": "embedding_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_noseg_posids_nosum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "position_ids",
+        "arg": "positionIdsT",
+        "semantic": "position_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_noseg_posids_sum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "position_ids",
+        "arg": "positionIdsT",
+        "semantic": "position_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "embedding_sum",
+        "arg": "embeddingSumT",
+        "semantic": "embedding_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_segdefault_nopos_nosum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_segdefault_nopos_sum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "embedding_sum",
+        "arg": "embeddingSumT",
+        "semantic": "embedding_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_segdefault_posids_nosum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "position_ids",
+        "arg": "positionIdsT",
+        "semantic": "position_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_segdefault_posids_sum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "position_ids",
+        "arg": "positionIdsT",
+        "semantic": "position_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "embedding_sum",
+        "arg": "embeddingSumT",
+        "semantic": "embedding_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_seg_nopos_nosum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "segment_ids",
+        "arg": "segmentIdsT",
+        "semantic": "segment_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_seg_nopos_sum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "segment_ids",
+        "arg": "segmentIdsT",
+        "semantic": "segment_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "embedding_sum",
+        "arg": "embeddingSumT",
+        "semantic": "embedding_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_seg_posids_nosum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "segment_ids",
+        "arg": "segmentIdsT",
+        "semantic": "segment_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "position_ids",
+        "arg": "positionIdsT",
+        "semantic": "position_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "embed_seg_posids_sum": [
+      {
+        "name": "input_ids",
+        "arg": "inputIdsT",
+        "semantic": "input_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "segment_ids",
+        "arg": "segmentIdsT",
+        "semantic": "segment_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "position_ids",
+        "arg": "positionIdsT",
+        "semantic": "position_ids",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "word_embedding",
+        "arg": "wordEmbeddingT",
+        "semantic": "word_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "position_embedding",
+        "arg": "positionEmbeddingT",
+        "semantic": "position_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "segment_embedding",
+        "arg": "segmentEmbeddingT",
+        "semantic": "segment_embedding",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "embedding_sum",
+        "arg": "embeddingSumT",
+        "semantic": "embedding_sum",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "normalizeSet": [
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "gamma",
+        "arg": "gammaT",
+        "semantic": "gamma",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "beta",
+        "arg": "betaT",
+        "semantic": "beta",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$HIDDEN_LEN"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "tokens", "type": "u32", "value": "tokens" }] }
+      }
+    ],
+    "maskIndexSet": [
+      {
+        "name": "mask",
+        "arg": "maskT",
+        "semantic": "mask",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "mask_index",
+        "arg": "maskIndexT",
+        "semantic": "mask_index",
+        "buffer": { "type": "storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "batch", "type": "u32", "value": "batchSize" }] }
+      }
+    ],
+    "zeroMaskIndexSet": [
+      {
+        "name": "mask_index",
+        "arg": "maskIndexT",
+        "semantic": "mask_index",
+        "buffer": { "type": "storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "batch", "type": "u32", "value": "batchSize" }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "noseg_nopos_nosum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "not present.positionIdsT", "not present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_nopos_nosum_mask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "not present.positionIdsT", "not present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_nopos_sum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "not present.positionIdsT", "present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_nopos_sum_mask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "not present.positionIdsT", "present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_posids_nosum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "present.positionIdsT", "not present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_posids_nosum_mask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "present.positionIdsT", "not present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_posids_sum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "present.positionIdsT", "present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_posids_sum_mask",
+      "when": ["embedContractOk", "dispatchFits", "not present.segmentEmbeddingT", "present.positionIdsT", "present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_nopos_nosum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "not present.positionIdsT", "not present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "seg_nopos_nosum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "not present.positionIdsT", "not present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_nopos_sum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "not present.positionIdsT", "present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "seg_nopos_sum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "not present.positionIdsT", "present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_posids_nosum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "present.positionIdsT", "not present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "seg_posids_nosum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "present.positionIdsT", "not present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_posids_sum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "present.positionIdsT", "present.embeddingSumT", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "seg_posids_sum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "present.segmentIdsT", "present.positionIdsT", "present.embeddingSumT", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_nopos_nosum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_nopos_nosum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_nopos_sum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_nopos_sum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_posids_nosum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_posids_nosum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_posids_sum_nomask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "not present.maskIndexT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_posids_sum_mask",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT", "not present.segmentIdsT", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.MaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "maskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_nopos_nosum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"noseg\" != \"noseg\")", "present.segmentIdsT == (\"noseg\" == \"seg\")", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_nopos_sum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"noseg\" != \"noseg\")", "present.segmentIdsT == (\"noseg\" == \"seg\")", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_posids_nosum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"noseg\" != \"noseg\")", "present.segmentIdsT == (\"noseg\" == \"seg\")", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "noseg_posids_sum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"noseg\" != \"noseg\")", "present.segmentIdsT == (\"noseg\" == \"seg\")", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_noseg_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_nopos_nosum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"segdefault\" != \"noseg\")", "present.segmentIdsT == (\"segdefault\" == \"seg\")", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_nopos_sum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"segdefault\" != \"noseg\")", "present.segmentIdsT == (\"segdefault\" == \"seg\")", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_posids_nosum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"segdefault\" != \"noseg\")", "present.segmentIdsT == (\"segdefault\" == \"seg\")", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "segdefault_posids_sum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"segdefault\" != \"noseg\")", "present.segmentIdsT == (\"segdefault\" == \"seg\")", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_segdefault_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_nopos_nosum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"seg\" != \"noseg\")", "present.segmentIdsT == (\"seg\" == \"seg\")", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_nopos_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_nopos_sum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"seg\" != \"noseg\")", "present.segmentIdsT == (\"seg\" == \"seg\")", "present.positionIdsT == (\"nopos\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_nopos_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_posids_nosum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"seg\" != \"noseg\")", "present.segmentIdsT == (\"seg\" == \"seg\")", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"nosum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_posids_nosum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "seg_posids_sum_mask_without_input",
+      "when": ["embedContractOk", "dispatchFits", "present.segmentEmbeddingT == (\"seg\" != \"noseg\")", "present.segmentIdsT == (\"seg\" == \"seg\")", "present.positionIdsT == (\"posids\" == \"posids\")", "present.embeddingSumT == (\"sum\" == \"sum\")", "present.maskIndexT", "not present.maskT"],
+      "passes": [
+        {
+          "id": "sum",
+          "name": "EmbedLayerNormalization.EmbeddingSum",
+          "shader": "embed-sum.wgsl.jinja",
+          "bindings": "embed_seg_posids_sum",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "normalize",
+          "name": "EmbedLayerNormalization.Normalize",
+          "shader": "embed-normalize.wgsl.jinja",
+          "bindings": "normalizeSet",
+          "dispatch": { "workgroups": "tokens" }
+        },
+        {
+          "id": "maskIndex",
+          "name": "EmbedLayerNormalization.ZeroMaskIndex",
+          "shader": "embed-mask-index.wgsl.jinja",
+          "bindings": "zeroMaskIndexSet",
+          "dispatch": { "threads": "batchSize", "workgroupSize": "tunables.MASK_WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,20 @@

+{
+  "name": "com.microsoft.EmbedLayerNormalization",
+  "id": "_com_microsoft_embedlayernormalization_webgpu_b297617",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "LyhXY2RrY86niUBQGOouS5wCSVNuVu0RylHfXPXwZRc=",
+      "embed-mask-index.wgsl.jinja": "7wT7/LzdrOnanlTU5kmneEYpGb6D8v0H9/9mra8eF5Q=",
+      "embed-normalize.wgsl.jinja": "YObvmdAvDuyx8f4pwTmVxphDSivMh4Y2oovmT7eGm1s=",
+      "embed-sum.wgsl.jinja": "cCsjv9CmNWgQLtMOAvEQOveXEO6jpeX7NMD4PiaaIW0=",
+      "manifest.json": "C/rakYKSOpdBjI3xKiqWRE18Dg002pKoEVcA/yRtprM=",
+      "test.json": "UE2iD8tAA2HxgRcTihjDe6LYfamqEWUopaGGER7vNQc="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.EmbedLayerNormalization" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1325 @@

+{
+  "op": "com.microsoft.EmbedLayerNormalization",
+  "cases": [
+    {
+      "name": "segment_embedding_without_segment_ids_uses_row_zero",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/cpu/bert/embed_layer_norm.cc",
+        "notes": "The public schema makes both segment inputs independently optional. The pinned CPU and CUDA providers use segment row 0 when the table is present without segment_ids."
+      },
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, -1.0, 0.0, 1.0, 2.0] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8] }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [0.25, -0.5, 0.75, -1.0] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } },
+        "betaT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00001 },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "segment_default_nopos_nosum_nomask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1, 4] } }
+    },
+    {
+      "name": "segment_default_nopos_nosum_mask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] },
+        "maskT": { "dtype": "int32", "shape": [1, 1] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] }
+      }
+    },
+    {
+      "name": "segment_default_nopos_sum_mask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] },
+        "maskT": { "dtype": "int32", "shape": [1, 1] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 1, 4] }
+      }
+    },
+    {
+      "name": "segment_default_posids_nosum_nomask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1, 4] } }
+    },
+    {
+      "name": "segment_default_posids_nosum_mask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] },
+        "maskT": { "dtype": "int32", "shape": [1, 1] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] }
+      }
+    },
+    {
+      "name": "segment_default_posids_sum_nomask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 1, 4] }
+      }
+    },
+    {
+      "name": "segment_default_posids_sum_mask",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] },
+        "maskT": { "dtype": "int32", "shape": [1, 1] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 1, 4] }
+      }
+    },
+    {
+      "name": "segment_default_nopos_nosum_mask_without_input",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] }
+      }
+    },
+    {
+      "name": "segment_default_nopos_sum_mask_without_input",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 1, 4] }
+      }
+    },
+    {
+      "name": "segment_default_posids_nosum_mask_without_input",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] }
+      }
+    },
+    {
+      "name": "segment_default_posids_sum_mask_without_input",
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 1] },
+        "wordEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "positionEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "segmentEmbeddingT": { "dtype": "float32", "shape": [1, 4] },
+        "gammaT": { "dtype": "float32", "shape": [4] },
+        "betaT": { "dtype": "float32", "shape": [4] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 4] },
+        "maskIndexT": { "dtype": "int32", "shape": [1] },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 1, 4] }
+      }
+    },
+    {
+      "name": "noseg_nopos_nosum_mask_index_without_mask",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/cpu/bert/embed_layer_norm.cc",
+        "notes": "Pins ORT's independently optional mask_index output: when mask is absent, the output is initialized to zeros."
+      },
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } }
+      }
+    },
+    {
+      "name": "noseg_nopos_nosum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] }
+      }
+    },
+    {
+      "name": "noseg_nopos_sum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "noseg_nopos_sum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "noseg_posids_nosum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } }
+      }
+    },
+    {
+      "name": "noseg_posids_nosum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] }
+      }
+    },
+    {
+      "name": "noseg_posids_sum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "noseg_posids_sum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "seg_nopos_nosum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } }
+      }
+    },
+    {
+      "name": "seg_nopos_nosum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] }
+      }
+    },
+    {
+      "name": "seg_nopos_sum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "seg_nopos_sum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "seg_posids_nosum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } }
+      }
+    },
+    {
+      "name": "seg_posids_nosum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] }
+      }
+    },
+    {
+      "name": "seg_posids_sum_mask_index_without_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "seg_posids_sum_mask",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "positionIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 2, 4, 0, 2, 4, 0, 2] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] },
+        "embeddingSumT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "noseg_nopos_sum_mask_index_omitted",
+      "provenance": { "notes": "Compact optional-output regression: embedding_sum does not imply mask_index." },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.2, -0.4, 0.7, 0.1] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.3, 0.6, -0.2, 0.5] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.1, 0.9] } },
+        "betaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.05, -0.1] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "noseg_posids_sum_mask_index_omitted",
+      "provenance": {
+        "notes": "Compact optional-output regression: requesting embedding_sum does not require mask_index."
+      },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [1, 0] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.2, -0.4, 0.7, 0.1] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.3, 0.6, -0.2, 0.5] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.1, 0.9] } },
+        "betaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.05, -0.1] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "seg_nopos_nosum_mask_index_omitted",
+      "provenance": { "notes": "Compact optional-output regression: segment embeddings do not require mask_index." },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "segmentIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.2, -0.4, 0.7, 0.1] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.3, 0.6, -0.2, 0.5] }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.1, 0.2, -0.3, 0.4] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.1, 0.9] } },
+        "betaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.05, -0.1] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "seg_posids_nosum_mask_index_omitted",
+      "provenance": {
+        "notes": "Compact optional-output regression: segment and position ids do not require mask_index."
+      },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "segmentIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [1, 0] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.2, -0.4, 0.7, 0.1] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.3, 0.6, -0.2, 0.5] }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.1, 0.2, -0.3, 0.4] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.1, 0.9] } },
+        "betaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.05, -0.1] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "seg_posids_sum_mask_index_omitted",
+      "provenance": {
+        "notes": "Compact optional-output regression: embedding_sum remains independently requestable with the full embedding input set."
+      },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "segmentIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [1, 0] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.2, -0.4, 0.7, 0.1] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.3, 0.6, -0.2, 0.5] }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.1, 0.2, -0.3, 0.4] }
+        },
+        "gammaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.1, 0.9] } },
+        "betaT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.05, -0.1] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "position_ids_broadcast_row",
+      "provenance": {
+        "notes": "position_ids of shape (1, sequence_length) is shared by every batch, so the row index drops back to the position within the sequence."
+      },
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "positionIdsT": { "dtype": "int32", "shape": [1, 4], "data": { "kind": "values", "values": [0, 2, 4, 0] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "mask_full_row_and_partial_row",
+      "provenance": {
+        "notes": "Pins the public first-zero contract with one full row and one non-left-aligned row. The latter distinguishes first-zero (1) from the CPU provider's divergent count-ones result (3)."
+      },
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 0, 1, 1] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "maskIndexT": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [4, 1] } }
+      }
+    },
+    {
+      "name": "mask_input_without_mask_index",
+      "provenance": {
+        "notes": "Pins the independent optionality in the opposite direction: a mask input is valid when mask_index is not requested."
+      },
+      "attrs": { "mask_index_type": 0 },
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 0, 1, 1, 0, 1, 0, 1] } }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "explicit_epsilon",
+      "attrs": { "epsilon": 0.00001 },
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 4, 8], "tolerance": 0.00001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "hidden_wider_than_workgroup_stride",
+      "provenance": {
+        "notes": "A hidden size that is not a multiple of the workgroup, so the strided loops cover a partial final step."
+      },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [3, 10] } },
+        "segmentIdsT": { "dtype": "int32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "wordEmbeddingT": {
+          "dtype": "float32",
+          "shape": [12, 40],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float32",
+          "shape": [4, 40],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float32",
+          "shape": [2, 40],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float32",
+          "shape": [40],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float32",
+          "shape": [40],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 40], "tolerance": 0.00001, "relTolerance": 0.0001 },
+        "embeddingSumT": { "dtype": "float32", "shape": [1, 2, 40], "tolerance": 0.00001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "f16_full",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "segmentIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0, 0, 0] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float16",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float16",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float16",
+          "shape": [3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.53, "cosStep": 0.19, "scale": 0.4 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "maskT": { "dtype": "int32", "shape": [2, 4], "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 0, 0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 4, 8], "tolerance": 0.004, "relTolerance": 0.01 },
+        "maskIndexT": { "dtype": "int32", "shape": [2] },
+        "embeddingSumT": { "dtype": "float16", "shape": [2, 4, 8], "tolerance": 0.004, "relTolerance": 0.01 }
+      }
+    },
+    {
+      "name": "f16_cuda_association_embedding_sum_exact",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/cuda/bert/embed_layer_norm_impl.cu",
+        "notes": "The first lane distinguishes staged (word + segment) + position from (word + position) + segment by one float16 ULP."
+      },
+      "inputs": {
+        "inputIdsT": { "dtype": "int32", "shape": [1, 1], "data": { "kind": "values", "values": [0] } },
+        "segmentIdsT": { "dtype": "int32", "shape": [1, 1], "data": { "kind": "values", "values": [0] } },
+        "wordEmbeddingT": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [-0.9599609375, 1.0, -1.0, 0.0] }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [3.134765625, 2.0, 1.0, 0.5] }
+        },
+        "segmentEmbeddingT": {
+          "dtype": "float16",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [0.2060546875, 3.0, 2.0, -0.25] }
+        },
+        "gammaT": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } },
+        "betaT": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 4], "tolerance": 0.004, "relTolerance": 0.01 },
+        "embeddingSumT": {
+          "dtype": "float16",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [2.380859375, 6.0, 2.0, 0.25] },
+          "tolerance": 0,
+          "relTolerance": 0
+        }
+      }
+    },
+    {
+      "name": "f16_minimal",
+      "inputs": {
+        "inputIdsT": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [3, 10, 17, 4, 11, 18, 5, 12] }
+        },
+        "wordEmbeddingT": {
+          "dtype": "float16",
+          "shape": [20, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.9 }
+        },
+        "positionEmbeddingT": {
+          "dtype": "float16",
+          "shape": [6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.11, "scale": 0.6 }
+        },
+        "gammaT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.3, "offset": 1.0 }
+        },
+        "betaT": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.61, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float16", "shape": [2, 4, 8], "tolerance": 0.004, "relTolerance": 0.01 } }
+    }
+  ]
+}