sync 6fdf6301e2bb

Browse files

Files changed (6) hide show

README.md +81 -0
build/webgpu/bench.json +442 -0
build/webgpu/manifest.json +146 -0
build/webgpu/metadata.json +21 -0
build/webgpu/rotary-embedding-slices.wgsl.jinja +156 -0
build/webgpu/test.json +1181 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.RotaryEmbedding
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Rotary positional embedding (RoPE): each head's embedding vector is rotated with the `cos_cache` and `sin_cache` rows selected by `position_ids`, which is either a single base offset (token `s` reads row `position_ids[0] + s`) or a `(batch_size, sequence_length)` table. `input` is rank 3 `(batch_size, sequence_length, hidden_size)` or rank 4 `(batch_size, num_heads, sequence_length, head_size)`. `rotary_embedding_dim` rotates a prefix and copies the tail unchanged. Rotation arithmetic is float32 with one narrowing store. Bfloat16 and non-default `scale` are not implemented.
+See the [ONNX Runtime `RotaryEmbedding` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.RotaryEmbedding) for the reference semantics.
+## Inputs
+| Name | Upstream name | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `x` | `input` | `T` | same as logical dtype | — | — | Input token embeddings, shaped `(batch_size, sequence_length, hidden_size)` at rank 3 or `(batch_size, num_heads, sequence_length, head_size)` at rank 4. At rank 3 the head size comes from `num_heads` when that attribute is positive and from `2 * cos_cache.shape[1]` otherwise; at rank 4 both the head count and head size are read from the shape. | required |
+| `positionIds` | `position_ids` | `M` | `uint32` | — | — | Logical int64 cache-row selector in either upstream format: a scalar or one-element vector holding a base offset, so token `s` reads row `position_ids[0] + s`; or a `(batch_size, sequence_length)` table read per token. Valid positions are non-negative rows of the caches and use uint32 WebGPU storage, so a negative value is rejected at the host boundary. | required |
+| `cos` | `cos_cache` | `T` | same as logical dtype | `2` | — | Precomputed cosine values of shape `(max_sequence_length, rotary_dim / 2)`, where `rotary_dim` is `rotary_embedding_dim` when that is positive and the head size otherwise. | required |
+| `sin` | `sin_cache` | `T` | same as logical dtype | `2` | — | Precomputed sine values with the same shape and type as `cos_cache`. | required |
+## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `y` | `output` | `T` | same as `x` | same as `x` | Rotary-position-encoded tensor with the same shape and type as `input`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `interleaved` | `0` | Set to 1 to pair adjacent even/odd elements, or 0 to pair each element of the first half of the rotary window with the matching element of the second half. Default is 0. |
+| `is_packed_batching` | `0` | Ragged (packed) batch inputs. Its only upstream effect is to lift the `sequence_length <= max_sequence_length` bound, and this implementation never imposes that bound: every gathered row is required to index the caches whatever the sequence length. Default is 0. |
+| `num_heads` | `0` | Number of attention heads. Default is 0, which asks the rank-3 path to take the head size from the cache width instead; a positive value is required whenever `rotary_embedding_dim` is nonzero. At rank 4 the head count comes from the input shape and this attribute is not consulted. |
+| `rotary_embedding_dim` | `0` | Positive even count of leading head-dimension elements to rotate; the remaining tail is copied unchanged. Default is 0, meaning the whole head dimension, which must then be even. An odd value is rejected. |
+| `scale` | `1` | Declared scale for the gathered rotation. No ONNX Runtime provider applies it, so the default and only accepted value is 1.0 and a caller's other value is rejected rather than silently discarded. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+| `M` | `int64` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
+- [`rotary-embedding-slices.wgsl.jinja`](build/webgpu/rotary-embedding-slices.wgsl.jinja)
+## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.RotaryEmbedding", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 2, 18] },
+  positionIds: { data: positionIdsData, shape: [1, 2] },
+  cos: { data: cosData, shape: [4, 3] },
+  sin: { data: sinData, shape: [4, 3] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,442 @@

+{
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256, 512], "ITEMS_PER_LANE": [1, 2, 4] },
+  "cases": [
+    {
+      "name": "llama3-8b-prefill-s512-q-h32-d128-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "512-token prefill of the query projection: rank-3 packed heads with a rank-2 cache and a per-token position table. Read-once, write-once traffic.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float32", "elems": 2097152, "cacheElems": 32768 },
+      "inputs": {
+        "x": { "shape": [1, 512, 4096], "dtype": "float32", "dist": "normal", "seed": 4010, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4010,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4011, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4012, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 512, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-prefill-s512-q-h32-d128-f16",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "The float16 prefill of the same projection, the dtype shipped model graphs use.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float16", "elems": 2097152, "cacheElems": 32768 },
+      "inputs": {
+        "x": { "shape": [1, 512, 4096], "dtype": "float16", "dist": "normal", "seed": 4020, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4020,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 64], "dtype": "float16", "dist": "uniform", "seed": 4021, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float16", "dist": "uniform", "seed": 4022, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 512, 4096], "dtype": "float16" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-prefill-s512-q-h32-d128-interleaved-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 1 },
+      "provenance": {
+        "notes": "The interleaved pairing at the same shape. One of the two ONNX Runtime fusion paths emits interleaved=1, so both layouts occur in real exports.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float32", "elems": 2097152, "cacheElems": 32768 },
+      "inputs": {
+        "x": { "shape": [1, 512, 4096], "dtype": "float32", "dist": "normal", "seed": 4030, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4030,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4031, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4032, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 512, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-decode-q-h32-d128-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "One decoded token through a format-0 base offset, the shape a decode loop emits. 32 KB of traffic, so this measures launch latency.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float32", "elems": 4096, "cacheElems": 64 },
+      "inputs": {
+        "x": { "shape": [1, 1, 4096], "dtype": "float32", "dist": "normal", "seed": 4040, "scale": 1 },
+        "positionIds": { "shape": [1], "dtype": "uint32", "data": { "kind": "values", "values": [1024] } },
+        "cos": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4041, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4042, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 1, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-decode-k-h8-d128-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "The grouped-query key projection of the same decode step: eight heads, the smallest dispatch this operator sees per token.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float32", "elems": 1024, "cacheElems": 64 },
+      "inputs": {
+        "x": { "shape": [1, 1, 1024], "dtype": "float32", "dist": "normal", "seed": 4050, "scale": 1 },
+        "positionIds": { "shape": [1], "dtype": "uint32", "data": { "kind": "values", "values": [1024] } },
+        "cos": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4051, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4052, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 1, 1024], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "qwen3-4b-prefill-s512-rank4-h32-d128-f16",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "Rank-4 BNSH prefill, the transposed layout the ONNX Runtime attention fusion emits after a Transpose.",
+        "model": "Qwen/Qwen3-4B"
+      },
+      "vars": { "dtype": "float16", "elems": 2097152, "cacheElems": 32768 },
+      "inputs": {
+        "x": { "shape": [1, 32, 512, 128], "dtype": "float16", "dist": "normal", "seed": 4060, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4060,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 64], "dtype": "float16", "dist": "uniform", "seed": 4061, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float16", "dist": "uniform", "seed": 4062, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 32, 512, 128], "dtype": "float16" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "phi2-prefill-s512-rotary32-d80-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0, "rotary_embedding_dim": 32, "num_heads": 32 },
+      "provenance": {
+        "notes": "Partial rotary: 32 of 80 head elements rotate, so 60 percent of the traffic is an untouched tail the same invocations copy at full vector width.",
+        "model": "microsoft/phi-2"
+      },
+      "vars": { "dtype": "float32", "elems": 1310720, "cacheElems": 8192 },
+      "inputs": {
+        "x": { "shape": [1, 512, 2560], "dtype": "float32", "dist": "normal", "seed": 4070, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4070,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 16], "dtype": "float32", "dist": "uniform", "seed": 4071, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 16], "dtype": "float32", "dist": "uniform", "seed": 4072, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 512, 2560], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "phi3-mini-prefill-s512-d96-h32-f16",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "Head size 96 with a 48-wide cache: a non-power-of-two head that still clears the four-pair route's gates.",
+        "model": "microsoft/Phi-3-mini-4k-instruct"
+      },
+      "vars": { "dtype": "float16", "elems": 1572864, "cacheElems": 24576 },
+      "inputs": {
+        "x": { "shape": [1, 512, 3072], "dtype": "float16", "dist": "normal", "seed": 4080, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4080,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 48], "dtype": "float16", "dist": "uniform", "seed": 4081, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 48], "dtype": "float16", "dist": "uniform", "seed": 4082, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 512, 3072], "dtype": "float16" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "pair-route-prefill-s512-rotary126-d128-h32-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0, "rotary_embedding_dim": 126, "num_heads": 32 },
+      "provenance": {
+        "notes": "A rotary window of 126 gives 63 rotation pairs and leaves two unrotated elements in each 128-wide head, at the same geometry as the paired Llama prefill case."
+      },
+      "vars": { "dtype": "float32", "elems": 2097152, "cacheElems": 32256 },
+      "inputs": {
+        "x": { "shape": [1, 512, 4096], "dtype": "float32", "dist": "normal", "seed": 4090, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4090,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 63], "dtype": "float32", "dist": "uniform", "seed": 4091, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 63], "dtype": "float32", "dist": "uniform", "seed": 4092, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 512, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-prefill-s1536-q-h32-d128-f32",
+      "preset": "model",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "A 1536-token prefill, sized so its working set matches the elementwise control of the same working set.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float32", "elems": 6291456, "cacheElems": 98304 },
+      "inputs": {
+        "x": { "shape": [1, 1536, 4096], "dtype": "float32", "dist": "normal", "seed": 4100, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 1536],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4100,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4101, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4102, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 1536, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "pair-route-prefill-s2048-rotary126-d128-h32-f32",
+      "preset": "model",
+      "attrs": { "interleaved": 0, "rotary_embedding_dim": 126, "num_heads": 32 },
+      "provenance": {
+        "notes": "The single-pair route at the same 67 MB working set as the four-pair route's model case, so the two schedules are directly comparable."
+      },
+      "vars": { "dtype": "float32", "elems": 8388608, "cacheElems": 129024 },
+      "inputs": {
+        "x": { "shape": [1, 2048, 4096], "dtype": "float32", "dist": "normal", "seed": 4110, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 2048],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4110,
+          "step": 1,
+          "offset": 0,
+          "mod": 8192
+        },
+        "cos": { "shape": [8192, 63], "dtype": "float32", "dist": "uniform", "seed": 4111, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [8192, 63], "dtype": "float32", "dist": "uniform", "seed": 4112, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 2048, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "qwen3-4b-decode-rank4-h32-d128-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "One decoded token in the rank-4 layout through a format-0 base offset: 32 head blocks of one token each, the smallest rank-4 dispatch.",
+        "model": "Qwen/Qwen3-4B"
+      },
+      "vars": { "dtype": "float32", "elems": 4096, "cacheElems": 64 },
+      "inputs": {
+        "x": { "shape": [1, 32, 1, 128], "dtype": "float32", "dist": "normal", "seed": 4120, "scale": 1 },
+        "positionIds": { "shape": [1], "dtype": "uint32", "data": { "kind": "values", "values": [1024] } },
+        "cos": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4121, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 64], "dtype": "float32", "dist": "uniform", "seed": 4122, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 32, 1, 128], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "pair-route-rank4-prefill-s512-rotary126-f32",
+      "preset": "smoke",
+      "attrs": { "interleaved": 0, "rotary_embedding_dim": 126, "num_heads": 32 },
+      "provenance": {
+        "notes": "The single-pair route in the rank-4 layout, where the head block shares one gathered cache row across its heads exactly as the vector route does."
+      },
+      "vars": { "dtype": "float32", "elems": 2097152, "cacheElems": 32256 },
+      "inputs": {
+        "x": { "shape": [1, 32, 512, 128], "dtype": "float32", "dist": "normal", "seed": 4130, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 512],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4130,
+          "step": 1,
+          "offset": 0,
+          "mod": 4096
+        },
+        "cos": { "shape": [4096, 63], "dtype": "float32", "dist": "uniform", "seed": 4131, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [4096, 63], "dtype": "float32", "dist": "uniform", "seed": 4132, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 32, 512, 128], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-prefill-s4096-q-h32-d128-f16",
+      "preset": "model",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "A 4096-token prefill: a 67 MB working set, far past cache, so the bandwidth figure is a DRAM figure.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float16", "elems": 16777216, "cacheElems": 262144 },
+      "inputs": {
+        "x": { "shape": [1, 4096, 4096], "dtype": "float16", "dist": "normal", "seed": 4140, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 4096],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4140,
+          "step": 1,
+          "offset": 0,
+          "mod": 8192
+        },
+        "cos": { "shape": [8192, 64], "dtype": "float16", "dist": "uniform", "seed": 4141, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [8192, 64], "dtype": "float16", "dist": "uniform", "seed": 4142, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 4096, 4096], "dtype": "float16" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "llama3-8b-prefill-s2048-q-h32-d128-f32",
+      "preset": "model",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "The float32 counterpart at the same 67 MB working set.",
+        "model": "meta-llama/Meta-Llama-3-8B"
+      },
+      "vars": { "dtype": "float32", "elems": 8388608, "cacheElems": 131072 },
+      "inputs": {
+        "x": { "shape": [1, 2048, 4096], "dtype": "float32", "dist": "normal", "seed": 4150, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 2048],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4150,
+          "step": 1,
+          "offset": 0,
+          "mod": 8192
+        },
+        "cos": { "shape": [8192, 64], "dtype": "float32", "dist": "uniform", "seed": 4151, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [8192, 64], "dtype": "float32", "dist": "uniform", "seed": 4152, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 2048, 4096], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    },
+    {
+      "name": "qwen3-4b-prefill-s2048-rank4-h32-d128-f32",
+      "preset": "model",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "notes": "The rank-4 layout at the same 67 MB working set, where the pinned (batch, head) slice carries 2048 tokens.",
+        "model": "Qwen/Qwen3-4B"
+      },
+      "vars": { "dtype": "float32", "elems": 8388608, "cacheElems": 131072 },
+      "inputs": {
+        "x": { "shape": [1, 32, 2048, 128], "dtype": "float32", "dist": "normal", "seed": 4160, "scale": 1 },
+        "positionIds": {
+          "shape": [1, 2048],
+          "dtype": "uint32",
+          "dist": "linearMod",
+          "seed": 4160,
+          "step": 1,
+          "offset": 0,
+          "mod": 8192
+        },
+        "cos": { "shape": [8192, 64], "dtype": "float32", "dist": "uniform", "seed": 4161, "scale": 0.1, "offset": 0.9 },
+        "sin": { "shape": [8192, 64], "dtype": "float32", "dist": "uniform", "seed": 4162, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 32, 2048, 128], "dtype": "float32" } },
+      "bench": {
+        "primary": false,
+        "metrics": [{ "type": "bandwidth", "value": "(args.elems * 2 + args.cacheElems * 2) * dtypeBytes(args.dtype)" }]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,146 @@

+{
+  "domain": "com.microsoft",
+  "name": "RotaryEmbedding",
+  "sinceVersion": 1,
+  "inputs": {
+    "x": { "onnx": "input", "dtype": "T" },
+    "positionIds": { "onnx": "position_ids", "dtype": "M", "storage": "uint32", "narrowing": "checked" },
+    "cos": { "onnx": "cos_cache", "dtype": "T", "rank": 2 },
+    "sin": { "onnx": "sin_cache", "dtype": "T", "rank": 2 }
+  },
+  "outputs": { "y": { "onnx": "output", "dtype": "T", "rank": "ranks.x", "shape": "shapes.x" } },
+  "attributes": {
+    "interleaved": { "default": 0 },
+    "is_packed_batching": { "default": 0 },
+    "num_heads": { "default": 0 },
+    "rotary_embedding_dim": { "default": 0 },
+    "scale": { "default": 1 }
+  },
+  "attributeConstraints": {
+    "interleaved": { "values": [0, 1] },
+    "is_packed_batching": { "values": [0, 1] },
+    "scale": { "values": [1] }
+  },
+  "typeConstraints": { "T": ["float32", "float16"], "M": ["int64"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 }, "ITEMS_PER_LANE": { "default": 2 } },
+  "derive": {
+    "foldedDispatchCapacity": "min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "rank3": "ranks.x == 3",
+    "rank4": "ranks.x == 4",
+    "rankOk": "rank3 or rank4",
+    "cacheShapeOk": "ranks.cos == 2 and ranks.sin == 2 and sameShape(shapes.cos, shapes.sin)",
+    "cacheWidth": "dim(shapes.cos, 1) if cacheShapeOk else 0",
+    "batchSize": "dim(shapes.x, 0) if rankOk else 0",
+    "seqLength": "(dim(shapes.x, 1) if rank3 else dim(shapes.x, 2)) if rankOk else 0",
+    "hiddenSize": "dim(shapes.x, 2) if rank3 else 0",
+    "rank3HeadSize": "(hiddenSize / attrs.num_heads if attrs.num_heads > 0 and hiddenSize % max(1, attrs.num_heads) == 0 else 2 * cacheWidth) if rank3 else 0",
+    "headSize": "rank3HeadSize if rank3 else (dim(shapes.x, 3) if rank4 else 0)",
+    "numHeads": "(hiddenSize / headSize if headSize > 0 and hiddenSize % max(1, headSize) == 0 else 0) if rank3 else (dim(shapes.x, 1) if rank4 else 0)",
+    "rotaryDim": "attrs.rotary_embedding_dim if attrs.rotary_embedding_dim > 0 else headSize",
+    "halfRotaryDim": "rotaryDim / 2 if rotaryDim % 2 == 0 else 0",
+    "geometryOk": "rankOk and cacheShapeOk and headSize > 0 and numHeads > 0 and batchSize > 0 and seqLength > 0 and rotaryDim > 0 and rotaryDim % 2 == 0 and rotaryDim <= headSize and cacheWidth == halfRotaryDim",
+    "attrsOk": "attrs.num_heads >= 0 and attrs.rotary_embedding_dim >= 0 and (attrs.rotary_embedding_dim == 0 or attrs.num_heads > 0)",
+    "contractOk": "f16Ok(dtypes.T) and geometryOk and attrsOk and sameShape(shapes.x, shapes.y)",
+    "itemsPerLane": "tunables.ITEMS_PER_LANE if (rank3 or numHeads % max(1, tunables.ITEMS_PER_LANE) == 0) else 1",
+    "headBlocks": "ceilDiv(numHeads, max(1, itemsPerLane)) if rank4 else 1",
+    "sliceRows": "(seqLength * numHeads if rank3 else seqLength) if rankOk else 0",
+    "sliceCount": "(batchSize if rank3 else batchSize * headBlocks) if rankOk else 0",
+    "sliceDispatchOk": "sliceCount <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "workgroupSize": "max(1, min(tunables.WORKGROUP_SIZE, min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)))",
+    "laneUnits": "workgroupSize * itemsPerLane if rank3 else workgroupSize",
+    "quadHeadUnits": "headSize / 8 if headSize % 8 == 0 else 0",
+    "quadHeadStride": "headSize / 4 if headSize % 4 == 0 else 0",
+    "quadRotStart": "rotaryDim / 4 if rotaryDim % 4 == 0 else 0",
+    "quadRotUnits": "rotaryDim / 8 if rotaryDim % 8 == 0 else 0",
+    "quadHasTail": "quadHeadUnits > quadRotUnits",
+    "quadBlocks": "ceilDiv(sliceRows * quadHeadUnits, laneUnits)",
+    "quadOk": "headSize % 8 == 0 and rotaryDim % 8 == 0 and quadBlocks <= foldedDispatchCapacity and sliceDispatchOk",
+    "pairHeadUnits": "ceilDiv(headSize, 2)",
+    "pairHasTail": "pairHeadUnits > halfRotaryDim",
+    "pairTailGuard": "headSize % 2 == 1",
+    "pairBlocks": "ceilDiv(sliceRows * pairHeadUnits, laneUnits)",
+    "pairOk": "pairBlocks <= foldedDispatchCapacity and sliceDispatchOk",
+    "posOffsetOk": "ranks.positionIds <= 1 and numel(shapes.positionIds) == 1",
+    "posTableOk": "ranks.positionIds == 2 and dim(shapes.positionIds, 0) == batchSize and dim(shapes.positionIds, 1) == seqLength",
+    "interleaved": "attrs.interleaved != 0",
+    "scalarType": "dtypes.T",
+    "vectorType": "\"vec4<\" ~ dtypes.T ~ \">\""
+  },
+  "when": ["contractOk"],
+  "bindings": {
+    "x": { "elementType": "$vectorType" },
+    "position_ids": { "arg": "positionIds", "elementType": "$M" },
+    "cos_cache": { "arg": "cos", "elementType": "$vectorType" },
+    "sin_cache": { "arg": "sin", "elementType": "$vectorType" },
+    "y": { "elementType": "$vectorType" },
+    "x_scalar": { "arg": "x", "name": "x", "elementType": "$T" },
+    "cos_scalar": { "arg": "cos", "name": "cos_cache", "elementType": "$T" },
+    "sin_scalar": { "arg": "sin", "name": "sin_cache", "elementType": "$T" },
+    "y_scalar": { "arg": "y", "name": "y", "elementType": "$T" },
+    "params": { "struct": [{ "name": "sliceRows", "type": "u32", "value": "sliceRows" }] }
+  },
+  "variants": [
+    {
+      "id": "quad",
+      "priority": 10,
+      "when": ["posTableOk or posOffsetOk", "quadOk"],
+      "derive": {
+        "rank": "ranks.x",
+        "posOffset": "posOffsetOk",
+        "useVec4": "true",
+        "headUnits": "quadHeadUnits",
+        "headStride": "quadHeadStride",
+        "rotStart": "quadRotStart",
+        "rotUnits": "quadRotUnits",
+        "hasTail": "quadHasTail",
+        "tailGuard": "false",
+        "castIn": "\"vec4<f32>\"",
+        "castOut": "vectorType"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RotaryEmbedding.Quad",
+          "shader": "rotary-embedding-slices.wgsl.jinja",
+          "bindings": ["x", "position_ids", "cos_cache", "sin_cache", "y", "params"],
+          "dispatch": {
+            "x": "min(quadBlocks, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(quadBlocks, DISPATCH_FOLD_WIDTH)",
+            "z": "sliceCount"
+          }
+        }
+      ]
+    },
+    {
+      "id": "pair",
+      "priority": 0,
+      "when": ["posTableOk or posOffsetOk", "pairOk"],
+      "derive": {
+        "rank": "ranks.x",
+        "posOffset": "posOffsetOk",
+        "useVec4": "false",
+        "headUnits": "pairHeadUnits",
+        "headStride": "headSize",
+        "rotStart": "rotaryDim",
+        "rotUnits": "halfRotaryDim",
+        "hasTail": "pairHasTail",
+        "tailGuard": "pairTailGuard",
+        "castIn": "\"f32\"",
+        "castOut": "scalarType"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "RotaryEmbedding.Pair",
+          "shader": "rotary-embedding-slices.wgsl.jinja",
+          "bindings": ["x_scalar", "position_ids", "cos_scalar", "sin_scalar", "y_scalar", "params"],
+          "dispatch": {
+            "x": "min(pairBlocks, DISPATCH_FOLD_WIDTH)",
+            "y": "ceilDiv(pairBlocks, DISPATCH_FOLD_WIDTH)",
+            "z": "sliceCount"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,21 @@

+{
+  "name": "com.microsoft.RotaryEmbedding",
+  "id": "_com_microsoft_rotaryembedding_webgpu_91f411d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "8QUIGpk2EMC9wPHWw/ypeXIzxYOfJx9bfNsg8GpbBtc=",
+      "manifest.json": "6Gx395vXiLXEDd5ThN0DsKtZq+E2ttgZ8FPiGKkpaM4=",
+      "rotary-embedding-slices.wgsl.jinja": "Dl6jsIc9PxPRXaC8pn+mTys7XYeg6THauIZGbj8CmiQ=",
+      "test.json": "0DBcHqmXLgRmFXICRr3pcMCeOFxWKR5e3vQqzcnrv5E="
+    }
+  },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.1",
+    "variants": { "quad": ["rotary-embedding-slices.wgsl.jinja"], "pair": ["rotary-embedding-slices.wgsl.jinja"] }
+  }
+}

build/webgpu/rotary-embedding-slices.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,156 @@

+{{ env.wgsl.resourceDeclarations }}
+// Rotary positional embedding. The dispatch pins one slice on z and folds that slice's own
+// rotation groups across x and y, so every row of a slice shares one base offset and all the
+// head geometry below is a compile-time constant: every division in the index math is by a
+// literal rather than by a uniform.
+//
+// A slice is one batch at rank 3, where a token's heads are adjacent and a workgroup covers a
+// whole token, so it reads each cos/sin row once. At rank 4 a token's heads are a sequence
+// apart, so a slice is instead one batch and ITEMS consecutive heads: the invocation gathers
+// its cos/sin row once and rotates the same rotation group in each of those heads.
+const WG: u32 = {{ workgroupSize }}u;
+const ITEMS: u32 = {{ itemsPerLane }}u;
+// Rotation groups per head row, the row's stride, and the number of groups that rotate, all
+// counted in the storage element the bindings use: a vec4 on the four-pair route and a scalar
+// on the single-pair route.
+const HEAD_UNITS: u32 = {{ headUnits }}u;
+const HEAD_STRIDE: u32 = {{ headStride }}u;
+const ROT_UNITS: u32 = {{ rotUnits }}u;
+{% if hasTail %}
+// First element past the rotary window.
+const ROT_END: u32 = {{ rotStart }}u;
+{% endif %}
+const NUM_HEADS: u32 = {{ numHeads }}u;
+{% if rank == 4 %}
+const HEAD_BLOCKS: u32 = {{ headBlocks }}u;
+{% endif %}
+{% if hasTail %}
+{% macro copy_tail(base) %}
+      let tail = ROT_END + (group - ROT_UNITS) * 2u;
+      y[{{ base }} + tail] = x[{{ base }} + tail];
+{% if tailGuard %}
+      // An odd head size leaves the last group with a single element to copy.
+      if (tail + 1u < HEAD_STRIDE) {
+        y[{{ base }} + tail + 1u] = x[{{ base }} + tail + 1u];
+      }
+{% else %}
+      y[{{ base }} + tail + 1u] = x[{{ base }} + tail + 1u];
+{% endif %}
+{% endmacro %}
+{% endif %}
+{% macro rotate(base) %}
+{% if interleaved and useVec4 %}
+      // Four interleaved pairs span two vec4s: (w0.x, w0.y), (w0.z, w0.w), (w1.x, w1.y),
+      // (w1.z, w1.w). Deinterleave, rotate, and reinterleave on the way back out.
+      let pairBase = {{ base }} + group * 2u;
+      let w0 = vec4<f32>(x[pairBase]);
+      let w1 = vec4<f32>(x[pairBase + 1u]);
+      let a = vec4<f32>(w0.x, w0.z, w1.x, w1.z);
+      let b = vec4<f32>(w0.y, w0.w, w1.y, w1.w);
+      let ra = a * cf - b * sf;
+      let rb = a * sf + b * cf;
+      y[pairBase] = {{ castOut }}(vec4<f32>(ra.x, rb.x, ra.y, rb.y));
+      y[pairBase + 1u] = {{ castOut }}(vec4<f32>(ra.z, rb.z, ra.w, rb.w));
+{% elif interleaved %}
+      // Adjacent elements form the rotated pair.
+      let pairBase = {{ base }} + group * 2u;
+      let a = {{ castIn }}(x[pairBase]);
+      let b = {{ castIn }}(x[pairBase + 1u]);
+      y[pairBase] = {{ castOut }}(a * cf - b * sf);
+      y[pairBase + 1u] = {{ castOut }}(a * sf + b * cf);
+{% else %}
+      // Split halves: ROT_UNITS is both the rotating group count and the offset from an
+      // element of the first half to its partner in the second.
+      let a = {{ castIn }}(x[{{ base }} + group]);
+      let b = {{ castIn }}(x[{{ base }} + ROT_UNITS + group]);
+      y[{{ base }} + group] = {{ castOut }}(a * cf - b * sf);
+      y[{{ base }} + ROT_UNITS + group] = {{ castOut }}(a * sf + b * cf);
+{% endif %}
+{% endmacro %}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+{% if rank == 3 %}
+  // wg.y carries the high bits of the group index past the per-axis dispatch fold width.
+  let slice = wg.z;
+  let sliceUnits = params.sliceRows * HEAD_UNITS;
+  let sliceBase = slice * params.sliceRows * HEAD_STRIDE;
+  let laneUnit = (wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u) * (WG * ITEMS) + lid.x;
+{% for item in range(itemsPerLane) %}
+  {
+    let unit = laneUnit + {{ item }}u * WG;
+    if (unit < sliceUnits) {
+      let group = unit % HEAD_UNITS;
+      let row = unit / HEAD_UNITS;
+      let base = sliceBase + row * HEAD_STRIDE;
+      let token = row / NUM_HEADS;
+{% if hasTail %}
+      if (group >= ROT_UNITS) {
+{{ copy_tail("base") }}
+      } else {
+{% endif %}
+{% if posOffset %}
+      // One base offset for the whole request: token s reads cache row p[0] + s.
+      let position = position_ids[0] + token;
+{% else %}
+      // The (batch_size, sequence_length) table. params.sliceRows / NUM_HEADS is the sequence
+      // length, so this slice's table row is slice * sequence_length + token.
+      let position = position_ids[slice * (params.sliceRows / NUM_HEADS) + token];
+{% endif %}
+      let cache = position * ROT_UNITS + group;
+      let cf = {{ castIn }}(cos_cache[cache]);
+      let sf = {{ castIn }}(sin_cache[cache]);
+{{ rotate("base") }}
+{% if hasTail %}
+      }
+{% endif %}
+    }
+  }
+{% endfor %}
+{% else %}
+  // A rank-4 slice is one batch and ITEMS consecutive heads; wg.y carries the high bits of
+  // the group index past the per-axis dispatch fold width.
+  let slice = wg.z;
+  let headBase = slice % HEAD_BLOCKS * ITEMS;
+  let batch = slice / HEAD_BLOCKS;
+  let sliceUnits = params.sliceRows * HEAD_UNITS;
+  let unit = (wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u) * WG + lid.x;
+  if (unit >= sliceUnits) {
+    return;
+  }
+  let group = unit % HEAD_UNITS;
+  let token = unit / HEAD_UNITS;
+  let rowBase = (batch * NUM_HEADS + headBase) * params.sliceRows + token;
+{% if hasTail %}
+  if (group >= ROT_UNITS) {
+{% for item in range(itemsPerLane) %}
+    {
+      let base = (rowBase + {{ item }}u * params.sliceRows) * HEAD_STRIDE;
+{{ copy_tail("base") }}
+    }
+{% endfor %}
+    return;
+  }
+{% endif %}
+{% if posOffset %}
+  // One base offset for the whole request: token s reads cache row p[0] + s.
+  let position = position_ids[0] + token;
+{% else %}
+  // The (batch_size, sequence_length) table; params.sliceRows is the sequence length.
+  let position = position_ids[batch * params.sliceRows + token];
+{% endif %}
+  // Gathered once for the whole head block: cos and sin depend on the token and the rotation
+  // group, never on the head.
+  let cache = position * ROT_UNITS + group;
+  let cf = {{ castIn }}(cos_cache[cache]);
+  let sf = {{ castIn }}(sin_cache[cache]);
+{% for item in range(itemsPerLane) %}
+  {
+    let base = (rowBase + {{ item }}u * params.sliceRows) * HEAD_STRIDE;
+{{ rotate("base") }}
+  }
+{% endfor %}
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1181 @@

+{
+  "fixtureArrays": {
+    "ort_contrib_interleaved_small_llama_offset_f32_input_x": [-1.0408, 0.9166, -1.3042, -1.1097, -0.132, -0.2751, -0.235, 0.0937, -1.2188, 1.1676, -1.0574, -0.1188, -0.7396, -1.2425, -0.1752, 0.699, -0.811, 0.6737, -1.1233, -0.0919, -0.6861, 0.7202, 0.1963, 0.6142],
+    "ort_contrib_interleaved_small_llama_offset_f32_input_cos": [1, 1, 0.5403, 0.9999, -0.4161, 0.9998, -0.99, 0.9996, -0.6536, 0.9992, 0.2837, 0.9988, 0.9602, 0.9982, 0.7539, 0.9976],
+    "ort_contrib_interleaved_small_llama_offset_f32_input_sin": [0, 0, 0.8415, 0.01, 0.9093, 0.02, 0.1411, 0.03, -0.7568, 0.04, -0.9589, 0.05, -0.2794, 0.06, 0.657, 0.0699],
+    "ort_contrib_interleaved_small_llama_offset_f32_output_y": [-1.0408, 0.9166, -1.3042, -1.1097, -0.132, -0.2751, -0.235, 0.0937, -1.6411, -0.3948, -1.0561, -0.1294, 0.646, -1.2937, -0.1822, 0.6972, -0.2751, -1.0178, -1.1212, -0.1143, -0.3694, -0.9235, 0.184, 0.618],
+    "ort_contrib_not_interleaved_small_llama_table_f32_input_x": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, 1.0076, -0.7529, -0.225, -0.4327, -1.5071, -0.4586, -0.8663, -0.2656, 0.1665, 0.7911, -0.932, -0.8579, -1.0574, -0.1188, -0.9078, 0.3452, -0.5713, -0.2351, -0.848, 0.5266, -1.2944, -0.0243, -0.2354, -0.7087, -0.9647, -0.0991, -0.2994, -0.065, -1.572, -1.3211],
+    "ort_contrib_not_interleaved_small_llama_table_f32_output_y": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, 1.0076, -0.7529, -0.225, -0.4327, -1.5071, -0.4586, -0.8663, -0.2656, 0.1665, 0.7911, -0.932, -0.8579, -0.8618, -0.0922, -0.9073, -0.7032, -0.5762, -0.2371, -0.4377, 0.537, -1.2929, -0.7267, -0.2107, -0.7115, -0.4666, -0.0261, -0.2965, -0.8469, -1.5749, -1.3217],
+    "ort_contrib_interleaved_large_llama_offset_f32_input_x": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, -1.019, 0.3157, -1.6036, 1.8493, 0.0447, 1.5853, 0.1036, -0.3514, 0.2421, 0.6463, 0.873, -0.9276, 1.0311, -1.9557, -0.1482, 1.7376, 2.2039, -0.6589, -1.0574, -0.1188, -0.9078, 0.3452, -0.5713, -0.2351, -0.5912, 1.1312, 0.7562, -1.2023, -0.5833, -0.4407, 0.1766, 1.0224, -0.4826, -0.5421, -0.5342, -0.6413, 1.3314, -0.4498, 0.5493, 0.0539, 0.2601, 0.857, 1.0076, -0.7529, -0.225, -0.4327, -1.5071, -0.4586, -1.9791, 0.7787, -0.7749, -0.1398, 1.1414, -0.6354, 0.0352, -0.4765, -0.0409, 1.1993, 0.5374, -0.193, 2.5211, -0.0452, -0.3105, -0.9407, -0.0034, 1.5199, -0.848, 0.5266, 0.0299, -0.0498, 1.0651, 0.886, -1.4702, -0.2134, -0.8707, 1.6159, -0.2356, 0.9444, 0.5937, 0.7203, 0.5061, 1.5192, -0.4897, 0.9231, 0.2654, -0.1441, 0.5407, -1.5476, 0.6455, -1.1382, 0.464, -0.4986, 0.1289, 2.7631, 0.1405, 1.1191, 2.1134, -0.9754, 0.1757, -0.1319, -0.2735, 0.3355, -0.6008, -1.1164, 0.2577, -0.7226, -0.9244, 1.8737, 0.6052, 1.1904, 1.2195, -0.047, -1.0914, 1.0223, 0.3152, 1.7528, -0.765, 1.8299, -0.2784, -0.2719, 0.1885, 2.1432, 0.8527, 0.0965, -0.0625, 0.8269, 1.0122, -1.4482, -0.0644, 0.3215, 0.5908, -1.4197, 0.2113, 0.0306, 0.3604, 0.3166, -0.8975, -0.6393, -1.2944, -0.0243, -0.2354, -0.7087, 1.1566, 0.4296, 0.5599, -0.7776, 0.3339, 0.1759, 2.1108, 1.0702, 0.8279, -0.2969, 0.712, -0.2068, -0.1548, 0.1553, 0.6207, -0.169, -0.5816, 1.2632, 0.0695, 1.1862, -1.1874, -0.7468, -0.932, -0.8579, -0.9647, -0.0991, 0.0195, 1.1213, -1.4873, -0.2043, -1.0466, -1.5772, -0.0489, 0.343, 0.1264, 0.1519, -1.3639, -1.6593, 1.8127, -1.4459, -0.2158, -0.9792, -1.4392, 0.6508, 0.8964, 0.5717, -0.239, 0.6983, -1.3416, 0.2715, -0.2852, 0.6051, 0.2167, -0.2181, -1.6306, 1.4788, 0.2754, -0.0261, -0.4618, -0.5646, -1.0389, 0.5819, 1.3697, 0.0002, 1.5333, -1.0556, -0.1254, 0.1527, -0.5996, -1.0962, 1.6327, 1.3951, 0.8784, 0.3389, 1.2907, 0.3124, 0.7299, 1.422, 0.3375, 0.0438, 1.8698, -0.2635, -2.0799, -0.6313, 0.409, -1.1458, 0.0784, -1.8848, -1.6165, 0.6179, 0.9905, -0.0729, 0.5054, -0.6681, -1.4382, 1.7547, -0.9605, -0.4558, -1.6105, 0.2979, 1.1537, -1.5604, 1.2779, -1.2514, 0.6056, 0.5763, -3.3558, 0.2836, 0.6909, -0.7631, 2.4451, -0.35, 1.3289, -0.6494, 0.3478, 1.0038, -0.2937, 0.9238, -1.2185, 0.4138, 0.5033, 0.9174, 1.8131, 1.4436, -0.4207, 0.022, -0.6807, -1.3306, 1.5646, 0.3338, 0.7105, 0.4683, -0.6179, 0.0818, -0.0488, -0.981, -1.3632, 0.0929, -1.7926, -0.2921, -0.4792, 0.6756, -0.3413, -0.2242, -0.2111, 0.6282, 0.1667, -1.4055, 1.5895, 1.0838, -0.9077, -0.806, 0.7967, -2.9351, 2.4179, -0.4026, 0.6451, 1.6845, -0.0901, 0.6106, 2.3603, 1.3908, -0.7917, -0.6734, -0.1213, -1.1116, -0.7401, -0.7879, 0.0606, -2.3337, -1.2603, -1.7245, -0.3533, -0.9421, -0.1776, 0.3992, -1.7142, -0.5319, -0.8848, 0.6513, 1.0002, -1.4699, -1.4254, 0.7013, 0.2414, 0.2551, -0.7457, 0.3133, -1.0941, -0.3682, -0.0163, -0.0645, -0.8101, 0.1415, 0.0551, 0.5873, -0.5887, -1.4733, -0.8565, 0.74, -0.5033, 0.0553, 0.9265, -0.8652, -0.0288, -0.2209, 0.061, 0.6776, 0.4361, -0.8052, 0.3955, 0.8988, 0.8238, 0.2262, 1.2912, 0.6488, 1.2114, 1.3569, 0.2983, 0.4718, -1.1936, 0.7928, -0.8665, 0.9468, 1.1629, 0.0616, -1.3136, -0.2764, 0.0277, -0.1126, 0.2342, -0.5866, -1.8219, 1.1079, 0.5795, -1.4249],
+    "ort_contrib_interleaved_large_llama_offset_f32_input_cos": [1, 1, 1, 0.5403, 0.9989, 1, -0.4161, 0.9957, 1, -0.99, 0.9903, 1, -0.6536, 0.9828, 1, 0.2837, 0.9732, 0.9999, 0.9602, 0.9615, 0.9999, 0.7539, 0.9477, 0.9999, -0.1455, 0.9318, 0.9999, -0.9111, 0.914, 0.9998, -0.8391, 0.8942, 0.9998, 0.0044, 0.8725, 0.9997, 0.8439, 0.8488, 0.9997, 0.9074, 0.8234, 0.9996, 0.1367, 0.7962, 0.9995, -0.7597, 0.7673, 0.9995],
+    "ort_contrib_interleaved_large_llama_offset_f32_input_sin": [0, 0, 0, 0.8415, 0.0464, 0.0022, 0.9093, 0.0927, 0.0043, 0.1411, 0.1388, 0.0065, -0.7568, 0.1846, 0.0086, -0.9589, 0.23, 0.0108, -0.2794, 0.2749, 0.0129, 0.657, 0.3192, 0.0151, 0.9894, 0.3629, 0.0172, 0.4121, 0.4057, 0.0194, -0.544, 0.4477, 0.0215, -1, 0.4887, 0.0237, -0.5366, 0.5286, 0.0259, 0.4202, 0.5675, 0.028, 0.9906, 0.605, 0.0302, 0.6503, 0.6413, 0.0323]
+  },
+  "cases": [
+    {
+      "name": "ort_contrib_interleaved_small_llama_offset_f32",
+      "attrs": { "interleaved": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_Interleaved_SmallData_LlamaMSFT",
+        "notes": "Format-0 position ids (a one-element base offset) with num_heads omitted, so head_size is 2 * cos_cache width and the head count is inferred from hidden_size."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 8],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_x" }
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_cos" }
+          }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_sin" }
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 3, 8],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_output_y" }
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_contrib_interleaved_small_llama_offset_f16",
+      "attrs": { "interleaved": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_Interleaved_SmallData_LlamaMSFT",
+        "notes": "The float16 arm of the same upstream case; ORT runs every RunTests case at float32 and float16 against the same expected values with a 0.002 absolute tolerance."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 3, 8],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_x" }
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_cos" }
+          }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_sin" }
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 3, 8],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_output_y" }
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "ort_contrib_not_interleaved_small_llama_table_f32",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_NotInterleaved_SmallData_LlamaMSFT",
+        "notes": "Format-1 position ids over three heads of six, with num_heads omitted so the head size comes from the cache width."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 18],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_not_interleaved_small_llama_table_f32_input_x" }
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [4, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 1.0, 1.0, 0.5403, 0.9989, 1.0, -0.4161, 0.9957, 1.0, -0.99, 0.9903, 1.0]
+          }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [4, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, 0.0, 0.8415, 0.0464, 0.0022, 0.9093, 0.0927, 0.0043, 0.1411, 0.1388, 0.0065]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 18],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_not_interleaved_small_llama_table_f32_output_y" }
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_contrib_not_interleaved_small_llama_table_f16",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_NotInterleaved_SmallData_LlamaMSFT",
+        "notes": "The float16 arm of the same upstream case."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 18],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_not_interleaved_small_llama_table_f32_input_x" }
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [4, 3],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 1.0, 1.0, 0.5403, 0.9989, 1.0, -0.4161, 0.9957, 1.0, -0.99, 0.9903, 1.0]
+          }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [4, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, 0.0, 0.8415, 0.0464, 0.0022, 0.9093, 0.0927, 0.0043, 0.1411, 0.1388, 0.0065]
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 2, 18],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_not_interleaved_small_llama_table_f32_output_y" }
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "ort_contrib_interleaved_large_llama_offset_f32",
+      "attrs": { "interleaved": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_Interleaved_LargeData_LlamaMSFT",
+        "notes": "Two batches of eight tokens over four heads of six, gathered through a format-0 base offset of zero."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8, 24],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_x" }
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_cos" }
+          }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_sin" }
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 8, 24],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, -1.019, 0.3157, -1.6036, 1.8493, 0.0447, 1.5853, 0.1036, -0.3514, 0.2421, 0.6463, 0.873, -0.9276, 1.0311, -1.9557, -0.1482, 1.7376, 2.2039, -0.6589, -0.4713, -0.954, -0.9229, 0.3027, -0.5708, -0.2363, -1.2713, 0.1137, 0.8112, -1.1659, -0.5824, -0.4419, -0.7649, 0.7011, -0.4569, -0.5639, -0.5328, -0.6424, 1.0979, 0.8773, 0.5462, 0.0793, 0.2582, 0.8576, 0.2653, 1.2295, -0.1839, -0.4517, -1.5052, -0.4651, 0.1155, -2.1237, -0.7586, -0.211, 1.1441, -0.6304, 0.4186, 0.2303, -0.1519, 1.1903, 0.5382, -0.1906, -1.008, 2.3112, -0.222, -0.9655, -0.0099, 1.5198, 0.7652, -0.641, 0.0365, -0.0452, 1.0593, 0.8929, 1.4856, 0.0038, -1.0865, 1.4794, -0.2417, 0.9428, -0.6894, -0.6293, 0.2904, 1.5747, -0.4956, 0.9199, -0.2424, 0.1801, 0.7503, -1.4576, 0.6529, -1.134, -0.6807, -0.0252, -0.3834, 2.7394, 0.1308, 1.1203, -2.1196, -0.9618, 0.197, -0.0972, -0.2764, 0.3332, -0.4522, 1.1844, 0.3867, -0.6626, -0.9405, 1.8656, 0.5053, -1.2361, 1.2072, 0.1789, -1.1002, 1.0129, 1.7702, 0.1949, -1.1653, 1.6049, -0.2755, -0.2749, 2.1087, 0.4272, 0.8076, 0.29, -0.0714, 0.8261, -1.1016, -1.3814, -0.1366, 0.2981, 0.606, -1.4132, 0.0893, -0.1939, 0.2779, 0.391, -0.8906, -0.6489, -1.2496, 0.3383, -0.0315, -0.7461, 1.151, 0.4445, 0.3203, -0.9031, 0.2727, 0.2609, 2.0968, 1.0974, 0.712, -0.5164, 0.7415, -0.0031, -0.1568, 0.1533, 0.5487, -0.3357, -0.9064, 1.0546, 0.0542, 1.187, -0.4045, -1.3431, -0.6094, -1.1105, -0.9631, -0.1137, -0.7219, 0.8582, -1.3443, -0.6684, -1.0227, -1.5929, -0.2622, 0.2264, 0.0713, 0.1843, -1.3387, -1.6797, 2.3165, 0.1009, 0.1081, -0.9969, -1.4488, 0.6291, 0.8964, 0.5717, -0.239, 0.6983, -1.3416, 0.2715, -0.2852, 0.6051, 0.2167, -0.2181, -1.6306, 1.4788, 0.2754, -0.0261, -0.4618, -0.5646, -1.0389, 0.5819, 1.3697, 0.0002, 1.5333, -1.0556, -0.1254, 0.1527, 0.5985, -1.0968, 1.5662, 1.4693, 0.8776, 0.3408, 0.4345, 1.2549, 0.6631, 1.4543, 0.3374, 0.0445, 1.232, 1.4311, -2.0483, -0.7272, 0.4114, -1.1449, 1.6283, -0.9524, -1.6435, 0.5422, 0.9907, -0.0708, 0.3972, 0.7376, -1.5947, 1.6138, -0.9586, -0.46, 0.3993, -1.5884, 1.2934, -1.4467, 1.2833, -1.2459, -0.776, 0.3108, -3.3677, -0.0287, 0.6942, -0.7601, -0.6993, 2.369, 1.3834, -0.5234, 0.3435, 1.0053, 0.1604, -0.956, -1.2641, 0.2406, 0.4973, 0.9206, -1.9987, -1.1733, -0.4197, -0.0366, -0.672, -1.335, -1.596, -0.1097, 0.6386, 0.5624, -0.6184, 0.0778, 0.1867, 0.9643, -1.3629, -0.0972, -1.7907, -0.3037, 0.8245, -0.0789, -0.294, -0.2833, -0.2165, 0.6264, -1.1726, 0.7926, 1.3621, 1.3586, -0.9007, -0.8138, -2.7421, 1.3155, 2.4507, 0.0507, 0.6305, 1.69, 0.521, -0.3309, 2.063, 1.8026, -0.7859, -0.6802, -1.1003, -0.199, -0.5391, -0.937, 0.0857, -2.333, -2.0112, 0.7193, -0.1272, -0.9981, -0.1818, 0.3973, -0.9963, 1.4929, -1.0109, 0.4304, 1.016, -1.459, 0.2682, 1.5658, 0.1762, 0.3038, -0.7491, 0.3052, -1.1534, -0.0478, 0.0021, -0.0665, -0.8118, 0.131, 0.2171, 0.5485, -0.161, -1.5784, -0.866, 0.7289, -0.4678, 0.1937, 1.1287, -0.5772, -0.0259, -0.2212, 0.2479, 0.6336, 0.6407, -0.6543, 0.3838, 0.9039, 0.4724, 0.7117, 1.0165, 1.027, 1.1908, 1.375, -0.085, 0.5517, -1.3842, 0.3703, -0.8806, 0.9336, 0.8362, 0.8105, -1.1566, -0.6813, 0.0294, -0.1122, 0.562, -0.2884, -2.0803, 0.4684, 0.6009, -1.416]
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_contrib_not_interleaved_large_llama_offset_f32",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_NotInterleaved_LargeData_LlamaMSFT",
+        "notes": "The split-halves pairing over the same two-batch geometry and a format-0 base offset."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8, 24],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_x" }
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_cos" }
+          }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_sin" }
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 8, 24],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, -1.019, 0.3157, -1.6036, 1.8493, 0.0447, 1.5853, 0.1036, -0.3514, 0.2421, 0.6463, 0.873, -0.9276, 1.0311, -1.9557, -0.1482, 1.7376, 2.2039, -0.6589, -0.8618, -0.0922, -0.9073, -0.7032, -0.5762, -0.2371, 0.6923, 1.1571, 0.7572, -1.1471, -0.5302, -0.4391, 0.5516, 1.0461, -0.4812, -0.1443, -0.4862, -0.6423, 0.674, -0.4614, 0.5475, 1.1495, 0.2389, 0.8582, -0.0259, -0.6099, -0.223, 1.0963, -1.5704, -0.4595, 0.9507, 0.6696, -0.7721, -1.7415, 1.2087, -0.6387, -1.1052, -0.5243, -0.04, -0.4671, 0.4909, -0.1931, -0.1937, -0.0447, -0.3171, 2.6839, -0.0076, 1.5185, 0.8465, 0.3737, 0.0242, -0.0703, 1.1279, 0.8862, 1.2275, -0.1786, -0.8767, -1.8072, -0.263, 0.9387, -0.8021, 0.7813, 0.5001, -1.4202, -0.385, 0.9263, -0.0443, -0.2323, 0.548, 1.5696, 0.6193, -1.1346, 1.7878, -0.516, 0.1192, -2.1572, 0.046, 1.1202, -1.4812, -0.9082, 0.1728, -1.5132, -0.4489, 0.337, -0.1541, -0.9266, 0.2416, 0.927, -1.1146, 1.8758, -0.4312, 1.3714, 1.2106, -0.4272, -0.8529, 1.0328, 1.8441, 1.7698, -0.762, 0.2168, 0.1322, -0.2802, 0.146, 2.1002, 0.8437, -0.1534, 0.4321, 0.836, 0.5955, -1.5452, -0.0491, -0.8794, 0.2418, -1.4203, 0.3635, 0.2362, 0.3672, -0.1128, -0.8664, -0.6354, -1.4409, -0.3413, -0.2409, -0.3188, 1.1054, 0.4265, 0.5867, -1.3279, 0.3201, 0.0125, 1.8157, 1.0745, 0.7372, -0.2429, 0.71, -0.4299, -0.2304, 0.1645, 0.9489, -0.1816, -0.5968, 1.0394, 0.0204, 1.1786, -0.3315, -0.3997, -0.9304, -1.4268, -1.1526, -0.1132, 0.149, 1.3967, -1.4634, -0.1412, -0.6339, -1.5995, -0.1366, 0.7604, 0.1514, 0.0824, -1.183, -1.6572, 2.0099, -0.9108, -0.2256, 0.4527, -1.8254, 0.6475, 0.8964, 0.5717, -0.239, 0.6983, -1.3416, 0.2715, -0.2852, 0.6051, 0.2167, -0.2181, -1.6306, 1.4788, 0.2754, -0.0261, -0.4618, -0.5646, -1.0389, 0.5819, 1.3697, 0.0002, 1.5333, -1.0556, -0.1254, 0.1527, -1.4979, -1.1358, 1.632, 0.2493, 0.8266, 0.3424, -0.4992, 0.2964, 0.7298, 1.8544, 0.3516, 0.0454, 1.5415, -0.2822, -2.0774, 1.2323, 0.3963, -1.1503, -0.4775, -1.9287, -1.6164, 0.3998, 0.902, -0.0764, -1.8059, -0.5762, -1.4362, -0.2706, -1.0183, -0.462, 2.0891, 0.1782, 1.1591, -0.8151, 1.3, -1.2464, -0.5099, 0.5098, -3.3525, 0.4326, 0.7414, -0.7775, -0.4271, -0.3807, 1.3245, 2.4936, 0.3139, 1.0095, 0.2323, 0.845, -1.2244, -0.4511, 0.6266, 0.9095, -1.7981, 1.5241, -0.4121, 0.2341, -0.4737, -1.3333, -1.615, 0.4164, 0.71, -0.2429, -0.5656, 0.0863, 0.0352, -0.7227, -1.3613, -0.0988, -1.9114, -0.3009, 0.1435, 0.7029, -0.3467, 0.5092, -0.0828, 0.6253, 0.7113, -1.2138, 1.5964, -0.8346, -1.1515, -0.7923, -0.8254, -3.0038, 2.4033, -0.3398, 0.0922, 1.7053, 1.1114, 0.7462, 2.366, -0.8409, -0.6654, -0.653, -0.7899, -1.0957, -0.7149, -0.1072, -0.1967, -2.3416, -1.2609, -1.6375, -0.3576, 0.9413, -0.5694, 0.3954, 0.1383, -0.7477, -0.8689, 1.8286, 0.851, -1.4793, -0.1597, 0.8541, 0.238, 1.4392, -0.5644, 0.3158, -1.0686, -0.1313, -0.0181, 0.2438, -0.8801, 0.1413, -0.3587, 0.8002, -0.5982, -1.4301, -0.662, 0.7324, -0.725, 0.061, 0.9293, -0.6902, -0.0125, -0.2089, -0.1664, 0.5428, 0.4245, -0.7901, 0.5665, 0.9044, 0.1948, -0.1723, 1.2705, 1.0303, 1.2202, 1.3762, -0.2959, 0.7237, -1.2077, 0.7937, -0.6705, 0.9287, 1.0583, 0.0496, -1.3118, 0.5556, 0.0459, -0.1324, -0.5513, -0.7409, -1.8002, 0.9892, 0.3619, -1.4522]
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_contrib_custom_rotary_dim_phi_table_f32",
+      "attrs": { "interleaved": 0, "rotary_embedding_dim": 4, "num_heads": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_CustomRotaryDim_SmallData_Phi",
+        "notes": "Partial rotation: four of six head elements rotate and the two-element tail is copied unchanged."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 6],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, 1.0076, -0.7529, -0.225, -0.4327, -1.5071, -0.4586]
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [0, 1] } },
+        "cos": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 0.5403] } },
+        "sin": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.8415] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 6],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, 1.0076, -0.0427, -0.225, -0.8673, -1.5071, -0.4586]
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_contrib_custom_rotary_dim_phi_packed_batching_f32",
+      "attrs": { "interleaved": 0, "rotary_embedding_dim": 4, "num_heads": 1, "is_packed_batching": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/rotary_embedding_op_test.cc",
+        "test": "ContribOpRotaryEmbeddingTest.RotaryEmbedding_CustomRotaryDim_SmallData_Phi_Packed_Batching",
+        "notes": "Packed batching: the sequence is longer than the cache, which upstream permits only under is_packed_batching=1. Every gathered row still indexes the cache."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 6],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, -1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, 1.0076, -0.7529, -0.225, -0.4327, -1.5071, -0.4586]
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [0, 0, 1] } },
+        "cos": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 0.5403] } },
+        "sin": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.8415] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 3, 6],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, -1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, 1.0076, -0.0427, -0.225, -0.8673, -1.5071, -0.4586]
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_rank4_interleaved_small_llama_table_f32",
+      "attrs": { "interleaved": 1, "num_heads": 2 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/llm/rotary_embedding_op_test.cc",
+        "test": "RotaryEmbeddingTest.RotaryEmbedding_Interleaved_SmallData_LlamaMSFT_4D_Input",
+        "notes": "Rank-4 BNSH input with a rank-2 cache and a format-1 position table. The opset-23 file is the only upstream source of numerically checked rank-4 rotary data; the two ops share this math exactly."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.2188, 1.1676, -1.0574, -0.1188, -0.811, 0.6737, -1.1233, -0.0919, -0.132, -0.2751, -0.235, 0.0937, -0.7396, -1.2425, -0.1752, 0.699, -0.6861, 0.7202, 0.1963, 0.6142]
+          }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [0, 1, 2] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_cos" }
+          }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_small_llama_offset_f32_input_sin" }
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [-1.0408, 0.9166, -1.3042, -1.1097, -1.6411, -0.3948, -1.0561, -0.1294, -0.2751, -1.0178, -1.1212, -0.1143, -0.132, -0.2751, -0.235, 0.0937, 0.646, -1.2937, -0.1822, 0.6972, -0.3694, -0.9235, 0.184, 0.618]
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "ort_rank4_not_interleaved_large_llama_table_f32",
+      "attrs": { "interleaved": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/llm/rotary_embedding_op_test.cc",
+        "test": "RotaryEmbeddingTest.RotaryEmbedding_NotInterleaved_LargeData_LlamaMSFT_4D_Input",
+        "notes": "Rank-4 BNSH input over two batches, four heads and eight tokens, with num_heads left unset because a rank-4 head count comes from the shape."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 8, 6],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_x" }
+          }
+        },
+        "positionIds": {
+          "dtype": "uint32",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6, 7, 0, 1, 2, 3, 4, 5, 6, 7] }
+        },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_cos" }
+          }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_contrib_interleaved_large_llama_offset_f32_input_sin" }
+          }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4, 8, 6],
+          "data": {
+            "kind": "values",
+            "values": [-1.04079998, 0.916599989, -1.30420005, -1.10969996, -1.21879995, 1.16760004, -2.10675168, 0.313278645, -1.60708773, 0.141688287, 0.0592993088, 1.58177209, -0.630788624, -0.430816084, 0.246088684, -0.174721941, 0.836671352, -0.926558971, -1.26596439, -2.2426312, -0.143917158, -1.57473576, 1.91107118, -0.659863293, 0.952363968, -0.0112946555, -0.90577817, 0.574617565, -0.583404124, -0.242907077, -1.32060885, 1.23504281, 0.760883927, 0.225809187, -0.307491541, -0.432488948, 0.0181085765, 1.12988913, -0.474278957, -0.569866478, -0.232575566, -0.647461414, 0.968330145, -0.509299397, 0.536304355, 0.91536504, 0.102920607, 0.865208685, 1.00759995, -0.752900004, -0.224999994, -0.432700008, -1.50709999, -0.458600014, -0.951666117, 0.724882483, -0.773502111, -1.74094665, 1.17627621, -0.63710475, -1.10517037, -0.524268031, -0.0400700979, -0.467021376, 0.490917653, -0.193175867, -2.36315632, -0.0442896411, -0.320379347, 1.28702021, -0.00964078028, 1.51788175, 0.51656419, 0.320925027, 0.0222803988, 0.674315631, 1.14399064, 0.886257112, 1.13239074, -0.153492898, -0.880812466, 1.86820555, -0.278367937, 0.934901953, 0.994535208, 0.827187002, 0.4941414, 1.2928561, -0.272836059, 0.929536343, 1.21685827, -0.34260717, 0.557832778, -0.992367864, 0.565743625, -1.12992167, 0.463999987, -0.498600006, 0.128900006, 2.76309991, 0.140499994, 1.11909997, 1.25286388, -0.961636603, 0.174961895, 1.70716047, -0.318457693, 0.335886538, 0.907053053, -1.02590752, 0.249643087, -0.245633602, -1.02391529, 1.87480819, -0.592516303, 1.33033943, 1.21285498, 0.13192372, -0.915585876, 1.03022671, 1.17885363, 1.77404451, -0.762661636, -1.43456602, 0.049955368, -0.27847901, 0.146011293, 2.10013723, 0.843684196, -0.153375596, 0.432110995, 0.83602643, 1.06174159, -1.55485511, -0.046079427, 0.0258956552, 0.169944048, -1.42038882, -0.0487071276, 0.315481633, 0.37001735, 0.377508819, -0.840793192, -0.63379401, -1.29439998, -0.0242999997, -0.235400006, -0.708700001, 1.1566, 0.4296, 0.154494107, -0.874685705, 0.331545562, 0.566194594, 2.07239747, 1.07093453, -0.156445935, -0.281273365, 0.711332202, 0.838858962, -0.181656986, 0.158361614, -0.79273057, -0.177007288, -0.589310288, -1.16298723, 0.0453686491, 1.18241966, 0.126825869, -0.555871427, -0.931147695, 1.45934772, -1.08596635, -0.107115202, -0.190371126, 1.33196723, -1.47011745, -0.0766584575, -0.760652125, -1.5931052, -0.0045129247, 0.70473057, 0.147792324, 0.159517035, -1.21709907, -1.65750349, 2.00992894, -0.910886765, -0.225605503, 0.452725053, -1.82546115, 0.647476315, 0.896399975, 0.571699977, -0.238999993, 0.698300004, -1.34159994, 0.271499991, 0.0294375867, 0.680094242, 0.213446647, -0.357835233, -1.6007297, 1.47927678, 0.398796827, 0.0703182518, -0.464302182, 0.485351264, -1.03685224, 0.579914272, -1.20705771, 0.0176035818, 1.53230751, 1.23830867, -0.124155864, 0.162666455, 1.44771028, -1.23949802, 1.62978542, -0.458060056, 0.660933018, 0.352941215, 1.72972739, 0.2264027, 0.729353964, -0.834230781, 0.400307, 0.0516785383, 1.61899674, -0.365789354, -2.06491113, -1.12859631, 0.320817351, -1.17251611, -0.346854568, -2.10239267, -1.61523759, 0.51734364, 0.337068677, -0.0973018557, 0.505400002, -0.668099999, -1.4382, 1.75469995, -0.960500002, -0.455799997, 0.442923486, 0.238277763, 1.15645313, -2.19831991, 1.29031682, -1.24886191, -0.509867668, 0.509775519, -3.35251856, 0.432666153, 0.741352141, -0.777529955, -2.32901859, -0.394879639, 1.3223753, 0.987909675, 0.295846343, 1.01243794, 0.505126178, 0.815001488, -1.22638965, -0.0481875092, 0.665176749, 0.90692091, 0.535472274, 1.56147265, -0.406287462, -1.7323401, -0.330429256, -1.33501053, 1.63317192, 0.490809381, 0.709373713, 0.0125124454, -0.502349257, 0.0909572691, -0.0978256166, -0.357495785, -1.35865283, 0.0379757136, -2.0119822, -0.312655121, -0.479200006, 0.675599992, -0.341300011, -0.224199995, -0.211099997, 0.628199995, -0.821949601, -1.36183679, 1.59127319, 0.725855172, -0.971916676, -0.802503109, 0.0345773101, -2.98228002, 2.41065669, 0.891961157, 0.370242298, 1.69489694, -0.107042886, 0.714565158, 2.36467719, -1.38960505, -0.699269295, -0.658058047, -0.517001033, -1.10366726, -0.720030189, 0.606771231, -0.145643681, -2.34006476, -1.26092672, -1.63743532, -0.357576013, 0.941227913, -0.569475293, 0.395344406, -1.46400166, -0.786376834, -0.865749776, 1.10432577, 0.81547302, -1.48116696, -1.24220979, 0.902649522, 0.236645028, -0.744167924, -0.482844949, 0.316913843, -1.0941, -0.368200004, -0.0163000003, -0.0644999966, -0.810100019, 0.141499996, 1.26955235, 0.626395524, -0.590327978, -0.749657393, -0.828307152, 0.73870486, 0.99614948, 0.0577319711, 0.927449882, -0.097641021, -0.0235498492, -0.216916054, 0.0532237142, 0.616131902, 0.430257797, 0.805755079, 0.485714525, 0.901634693, -0.0474238396, -0.00131507218, 1.27953064, -1.04750752, 1.23232043, 1.36800432, 0.844843626, 0.658450782, -1.20370615, -0.0611225069, -0.734763801, 0.933814406, 1.03939044, 0.0516136698, -1.31201601, -0.590313554, 0.0435673892, -0.12953417, -0.551326931, -0.740897238, -1.80020177, 0.989115238, 0.361949444, -1.45226824]
+          }
+        }
+      },
+      "tolerance": 0.002,
+      "relTolerance": 0.0001
+    },
+    {
+      "name": "quad_rank3_table_head128_heads4_batch2_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Four-pair route at the Llama head size over two batches whose position rows differ, so a batch or token index mistake changes the values."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": {
+          "dtype": "uint32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [5, 6, 7, 1, 9, 3] }
+        },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 512] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank3_table_head128_heads4_batch2_f16",
+      "attrs": {},
+      "provenance": {
+        "notes": "Float16 rotary embedding with the same rank-three table geometry as the corresponding float32 case."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": {
+          "dtype": "uint32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [5, 6, 7, 1, 9, 3] }
+        },
+        "cos": {
+          "dtype": "float16",
+          "shape": [16, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [16, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3, 512] } },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "quad_rank3_table_head16_heads2_interleaved_f32",
+      "attrs": { "interleaved": 1 },
+      "provenance": {
+        "notes": "Interleaved pairing on the four-pair route: head size 16 is exactly two rotation groups per head, so the deinterleave and reinterleave both run."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [2, 0, 5] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 8],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 8],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 32] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank3_offset_head16_heads2_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Format-0 base offset of three: the four tokens read cache rows 3 through 6, which a token index mistake would move."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 8],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 8],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 32] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank3_offset_head8_scalar_position_f32",
+      "attrs": {},
+      "provenance": { "notes": "A rank-0 scalar position id, the other spelling upstream accepts for a base offset." },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 24],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 4],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 4],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 24] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank3_offset_head16_heads2_interleaved_f16",
+      "attrs": { "interleaved": 1 },
+      "provenance": { "notes": "Interleaved float16 with a base offset, the decode-loop spelling." },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [16, 8],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [16, 8],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 4, 32] } },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "quad_rank3_table_head80_rotary32_tail_f32",
+      "attrs": { "rotary_embedding_dim": 32, "num_heads": 2 },
+      "provenance": {
+        "notes": "Partial rotation at the Phi-2 geometry: 32 of 80 elements rotate and the 48-element tail is copied by the remaining rotation groups."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 160],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [1, 4, 6] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 16],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 16],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 160] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank3_table_head24_rotary8_tail_f16",
+      "attrs": { "rotary_embedding_dim": 8, "num_heads": 2 },
+      "provenance": {
+        "notes": "A short rotary window inside a wide head, so most of the four-pair route's groups take the tail copy."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 2, 48],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 3, 6, 2] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [8, 4],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [8, 4],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 48] } },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "quad_rank3_table_head128_heads32_s8_wide_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Enough rotation groups per slice that the workgroups tile the slice several times over, exercising the strided per-lane item loop rather than its guard."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": {
+          "dtype": "uint32",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6, 7] }
+        },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 4096] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank3_table_head16_rotary8_pairgate_f32",
+      "attrs": { "rotary_embedding_dim": 8, "num_heads": 2 },
+      "provenance": {
+        "notes": "Rotary dimension eight and the paired rotary-dimension-four case check small rotary windows at the same head width."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [1, 5] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 4],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 4],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 32] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank4_table_head128_batch2_heads3_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Rank-4 BNSH with two batches and three heads, so the pinned dispatch slice is a (batch, head) pair and every head of a batch must read the same position row."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 5, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": {
+          "dtype": "uint32",
+          "shape": [2, 5],
+          "data": { "kind": "values", "values": [0, 2, 4, 6, 8, 1, 3, 5, 7, 9] }
+        },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 5, 128] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank4_table_head32_interleaved_f16",
+      "attrs": { "interleaved": 1 },
+      "provenance": { "notes": "Interleaved float16 rank-4 with non-monotonic positions." },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 3, 4, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 4], "data": { "kind": "values", "values": [7, 0, 3, 5] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [8, 16],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [8, 16],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 4, 32] } },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "quad_rank4_offset_head8_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Rank-4 with a format-0 base offset: the token index inside the slice is the row index directly."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 4],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 4],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 3, 8] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank4_table_head96_rotary48_tail_f32",
+      "attrs": { "rotary_embedding_dim": 48, "num_heads": 2 },
+      "provenance": {
+        "notes": "The Phi-3-mini head size of 96 with a 48-element rotary window, which clears the four-pair route's rotary gate and leaves a 48-element tail."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [2, 5, 0] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 24],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 24],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 3, 96] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "quad_rank4_offset_head64_batch2_f16",
+      "attrs": {},
+      "provenance": {
+        "notes": "Two batches sharing one base offset, so a slice-to-batch mistake shows up as a wrong head row rather than a wrong cache row."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 2, 2, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [8, 32],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [8, 32],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 2, 64] } },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "pair_rank3_table_head12_heads3_batch2_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Head size 12 is not a multiple of eight, so the single-pair route carries it; two batches with different position rows."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 3, 36],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": {
+          "dtype": "uint32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1, 3, 5, 8, 2, 11] }
+        },
+        "cos": {
+          "dtype": "float32",
+          "shape": [16, 6],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [16, 6],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 36] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "pair_rank3_table_head16_rotary4_pairgate_f32",
+      "attrs": { "rotary_embedding_dim": 4, "num_heads": 2 },
+      "provenance": {
+        "notes": "A rotary dimension of four fails the four-pair route's rotary gate even though the head size clears it, so this is the high side of that boundary."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [1, 5] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 32] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "pair_rank3_table_head7_rotary4_odd_f32",
+      "attrs": { "rotary_embedding_dim": 4, "num_heads": 2 },
+      "provenance": {
+        "notes": "An odd head size with an even rotary dimension, which the contrib CPU kernel supports: the last tail group copies a single element."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 14],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [0, 4, 7] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 14] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "pair_rank3_offset_head6_heads3_interleaved_f16",
+      "attrs": { "interleaved": 1 },
+      "provenance": { "notes": "Interleaved float16 on the single-pair route with a base offset." },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 3, 18],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "cos": {
+          "dtype": "float16",
+          "shape": [8, 3],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float16",
+          "shape": [8, 3],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 18] } },
+      "tolerance": 0.002,
+      "relTolerance": 0.002
+    },
+    {
+      "name": "pair_rank4_offset_head10_f32",
+      "attrs": {},
+      "provenance": { "notes": "Rank-4 head size 10 on the single-pair route with a zero base offset." },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 5],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 5],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 3, 10] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "pair_rank4_table_head9_rotary4_odd_f32",
+      "attrs": { "rotary_embedding_dim": 4, "num_heads": 2 },
+      "provenance": {
+        "notes": "An odd rank-4 head size with a four-element rotary window; num_heads is set because upstream requires it whenever rotary_embedding_dim is, and is otherwise ignored at rank 4."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 9],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [2, 2], "data": { "kind": "values", "values": [6, 1, 0, 4] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 9] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "pair_rank3_table_head6_heads5_hidden30_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Five heads of six with num_heads omitted, so the head count is inferred as hidden_size divided by twice the cache width."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 30],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 4], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [8, 3],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [8, 3],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 30] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "decode_rank3_offset_one_token_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "Single-token decode with a one-element position vector, the shape a decode loop emits."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [17] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 4096] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "decode_rank3_table_one_token_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "The same single token spelled as a 1x1 position table. Upstream reads a rank-2 tensor as the table format however few elements it holds, so this is the other side of the format boundary."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 1], "data": { "kind": "values", "values": [17] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 4096] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    },
+    {
+      "name": "decode_rank3_table_kv_heads8_f32",
+      "attrs": {},
+      "provenance": {
+        "notes": "The grouped-query key projection of the same decode step: eight heads instead of thirty-two."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 1.0 }
+        },
+        "positionIds": { "dtype": "uint32", "shape": [1, 1], "data": { "kind": "values", "values": [21] } },
+        "cos": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "rotaryCos", "thetaStart": 0.1, "thetaStep": 0.2 }
+        },
+        "sin": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "rotarySin", "thetaStart": 0.1, "thetaStep": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1024] } },
+      "tolerance": 0.000001,
+      "relTolerance": 0.00001
+    }
+  ]
+}