sync 2e7068faf55e

Browse files

Files changed (10) hide show

README.md +73 -0
build/webgpu/bench.json +146 -0
build/webgpu/manifest.json +266 -0
build/webgpu/metadata.json +22 -0
build/webgpu/scan-coop-channel.wgsl.jinja +121 -0
build/webgpu/scan-multichunk-apply.wgsl.jinja +16 -0
build/webgpu/scan-multichunk-carries.wgsl.jinja +17 -0
build/webgpu/scan-multichunk-local.wgsl.jinja +81 -0
build/webgpu/scan-prefix-sum.wgsl.jinja +18 -0
build/webgpu/test.json +360 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Scan
+`ai.onnx`  ·  internal tensor lowering (non-standard)  ·  reviewed against ONNX opset 25
+## Description
+Support status: the standard variadic ONNX `Scan` control-flow operator is not implemented because standalone kernel packages cannot carry or execute its body graph. This internal lowering computes a forward or reverse additive prefix scan for one rank-1 state and one rank-2 input and must not be treated as ONNX `Scan`.
+See the [standard ONNX `Scan` spec](https://onnx.ai/onnx/operators/onnx__Scan.html) for the contract this internal lowering does not implement.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `initial_state` | `initial_state` | `T` | `1` | — | Initial rank-1 additive state of shape `[dim]`. | required |
+| `scan_input` | `scan_input` | `T` | `2` | — | Rank-2 input of shape `[steps, dim]`. Each row is added to the running state in forward or reverse traversal order. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `final_state` | `final_state` | `T` | `1` | same as `initial_state` | Final rank-1 state of shape `[dim]` after all input rows have been accumulated. | required |
+| `scan_output` | `scan_output` | `T` | `2` | same as `scan_input` | Inclusive additive prefix results with the same `[steps, dim]` shape as `scan_input`. Reverse traversal still writes each result at its corresponding input row. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `reverse` | `0` | When non-zero, the scan input sequence is traversed in reverse order (equivalent to `scan_input_directions=1`); default `0` scans forward. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`scan-coop-channel.wgsl.jinja`](build/webgpu/scan-coop-channel.wgsl.jinja)
+- [`scan-multichunk-apply.wgsl.jinja`](build/webgpu/scan-multichunk-apply.wgsl.jinja)
+- [`scan-multichunk-carries.wgsl.jinja`](build/webgpu/scan-multichunk-carries.wgsl.jinja)
+- [`scan-multichunk-local.wgsl.jinja`](build/webgpu/scan-multichunk-local.wgsl.jinja)
+- [`scan-prefix-sum.wgsl.jinja`](build/webgpu/scan-prefix-sum.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Scan", { version: 1 });
+const { final_state, scan_output } = await kernel({
+  initial_state: { data: initial_stateData, shape: [1] },
+  scan_input: { data: scan_inputData, shape: [3, 1] },
+});
+```

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "ai.onnx.Scan",
+  "cases": [
+    {
+      "name": "lowered_prefix_sum_256x512",
+      "preset": "smoke",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 741, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [256, 512], "dist": "normal", "seed": 742, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [512] },
+        "scan_output": { "dtype": "float32", "shape": [256, 512] }
+      },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "256 * 512 * 4 * 2" }] }
+    },
+    {
+      "name": "lowered_prefix_sum_2048x512",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 743, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [2048, 512], "dist": "normal", "seed": 744, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [512] },
+        "scan_output": { "dtype": "float32", "shape": [2048, 512] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2048 * 512 * 4 * 2" }] }
+    },
+    {
+      "name": "lowered_prefix_sum_256x128",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [128] },
+        "scan_input": { "dtype": "float32", "shape": [256, 128] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [128] },
+        "scan_output": { "dtype": "float32", "shape": [256, 128] }
+      }
+    },
+    {
+      "name": "lowered_fewlane_serial_dim32_steps4096",
+      "preset": "smoke",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [32], "dist": "normal", "seed": 751, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [4096, 32], "dist": "normal", "seed": 752, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [32] },
+        "scan_output": { "dtype": "float32", "shape": [4096, 32] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "coop_healthy_dim512_steps4096",
+      "preset": "smoke",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 753, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [4096, 512], "dist": "normal", "seed": 754, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [512] },
+        "scan_output": { "dtype": "float32", "shape": [4096, 512] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "lowered_largedim_fallback_dim4096_steps1024",
+      "preset": "smoke",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [4096], "dist": "normal", "seed": 755, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [1024, 4096], "dist": "normal", "seed": 756, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [4096] },
+        "scan_output": { "dtype": "float32", "shape": [1024, 4096] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "lowered_smalldim2_singlelane_steps8192",
+      "preset": "smoke",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "dist": "normal", "seed": 757, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [8192, 2], "dist": "normal", "seed": 758, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2] },
+        "scan_output": { "dtype": "float32", "shape": [8192, 2] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "coop_reverse_partialchunk_dim512_steps300",
+      "preset": "smoke",
+      "attrs": { "reverse": 1 },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 759, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [300, 512], "dist": "normal", "seed": 760, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [512] },
+        "scan_output": { "dtype": "float32", "shape": [300, 512] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "lowered_singlelane_dim1_steps65536",
+      "preset": "stress",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1], "dist": "normal", "seed": 811, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [65536, 1], "dist": "normal", "seed": 812, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1] },
+        "scan_output": { "dtype": "float32", "shape": [65536, 1] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "lowered_dim63_steps32768_one_workgroup_launchbound",
+      "preset": "stress",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [63], "dist": "normal", "seed": 813, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [32768, 63], "dist": "normal", "seed": 814, "scale": 0.1 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [63] },
+        "scan_output": { "dtype": "float32", "shape": [32768, 63] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    },
+    {
+      "name": "streaming_recurrence_dim8_steps131072_multichunk_pathology",
+      "preset": "stress",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [8], "dist": "normal", "seed": 815, "scale": 0.1 },
+        "scan_input": { "dtype": "float32", "shape": [131072, 8], "dist": "normal", "seed": 816, "scale": 0.01 }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [8] },
+        "scan_output": { "dtype": "float32", "shape": [131072, 8] }
+      },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_input)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,266 @@

+{
+  "domain": "ai.onnx",
+  "name": "Scan",
+  "conformance": "internal-lowering",
+  "sinceVersion": 25,
+  "description": "Support status: the standard variadic ONNX `Scan` control-flow operator is not implemented because standalone kernel packages cannot carry or execute its body graph. This internal lowering computes a forward or reverse additive prefix scan for one rank-1 state and one rank-2 input and must not be treated as ONNX `Scan`.",
+  "inputs": [
+    {
+      "role": "initial_state",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Initial rank-1 additive state of shape `[dim]`."
+    },
+    {
+      "role": "scan_input",
+      "dtype": "T",
+      "rank": 2,
+      "description": "Rank-2 input of shape `[steps, dim]`. Each row is added to the running state in forward or reverse traversal order."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "final_state",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Final rank-1 state of shape `[dim]` after all input rows have been accumulated.",
+      "shape": "shapes.initial_state"
+    },
+    {
+      "role": "scan_output",
+      "dtype": "T",
+      "rank": 2,
+      "description": "Inclusive additive prefix results with the same `[steps, dim]` shape as `scan_input`. Reverse traversal still writes each result at its corresponding input row.",
+      "shape": "shapes.scan_input"
+    }
+  ],
+  "attributes": { "reverse": 0 },
+  "attributeDescriptions": {
+    "reverse": "When non-zero, the scan input sequence is traversed in reverse order (equivalent to `scan_input_directions=1`); default `0` scans forward."
+  },
+  "attributeConstraints": { "reverse": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32"] },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
+    "subgroupsWave32": "device.features.has(\"subgroups\") and wave32Adapter",
+    "variableNarrowSubgroups": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < device.adapterInfo.subgroupMaxSize and device.adapterInfo.subgroupMaxSize <= 16",
+    "parallelSubgroupScan": "subgroupsWave32 or variableNarrowSubgroups",
+    "scanUseSubgroups": "parallelSubgroupScan"
+  },
+  "args": {
+    "initial_state": { "kind": "tensor", "semantic": "initial_state", "role": "input" },
+    "scan_input": { "kind": "tensor", "semantic": "scan_input", "role": "input" },
+    "final_state": { "kind": "tensor", "semantic": "final_state", "role": "output" },
+    "scan_output": { "kind": "tensor", "semantic": "scan_output", "role": "output" }
+  },
+  "bindingSets": {
+    "scanState": [
+      {
+        "name": "initial_state",
+        "arg": "initial_state",
+        "semantic": "initial_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scan_input",
+        "arg": "scan_input",
+        "semantic": "scan_input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "final_state",
+        "arg": "final_state",
+        "semantic": "final_state",
+        "buffer": { "type": "storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "scan_output",
+        "arg": "scan_output",
+        "semantic": "scan_output",
+        "buffer": { "type": "storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "steps", "type": "u32", "value": "dim(shapes.scan_input, 0)" },
+            { "name": "dim", "type": "u32", "value": "dim(shapes.scan_input, 1)" },
+            { "name": "reverse", "type": "u32", "value": "attrs.reverse" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "multichunk_small_state",
+      "priority": 30,
+      "when": ["ranks.initial_state == 1", "ranks.scan_input == 2", "ranks.final_state == 1", "ranks.scan_output == 2", "dim(shapes.initial_state, 0) == dim(shapes.scan_input, 1)", "dim(shapes.final_state, 0) == dim(shapes.initial_state, 0)", "dim(shapes.scan_output, 0) == dim(shapes.scan_input, 0)", "dim(shapes.scan_output, 1) == dim(shapes.scan_input, 1)", "dim(shapes.scan_input, 1) >= 1", "dim(shapes.scan_input, 1) <= 32", "dim(shapes.scan_input, 0) >= 4096", "dim(shapes.scan_input, 1) >= 1"],
+      "derive": {
+        "chunks": "ceil(dim(shapes.scan_input, 0) / tunables.WORKGROUP_SIZE)",
+        "totalChunks": "dim(shapes.scan_input, 1) * chunks"
+      },
+      "intermediates": [
+        { "id": "chunkTotals", "dtype": "float32", "shape": "[totalChunks]" },
+        { "id": "chunkCarries", "dtype": "float32", "shape": "[totalChunks]" }
+      ],
+      "passes": [
+        {
+          "id": "local",
+          "name": "Scan.MultichunkLocal",
+          "source": {
+            "shader": "scan-multichunk-local.wgsl.jinja",
+            "inputs": { "chunks": "chunks", "totalChunks": "totalChunks", "useSubgroups": "scanUseSubgroups" }
+          },
+          "bindings": [
+            {
+              "name": "scan_input",
+              "arg": "scan_input",
+              "semantic": "scan_input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "scan_output",
+              "arg": "scan_output",
+              "semantic": "scan_output",
+              "buffer": { "type": "storage" },
+              "elementType": "f32"
+            },
+            { "name": "chunkTotals", "semantic": "chunkTotals", "buffer": { "type": "storage" }, "elementType": "f32" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "steps", "type": "u32", "value": "dim(shapes.scan_input, 0)" },
+                  { "name": "dim", "type": "u32", "value": "dim(shapes.scan_input, 1)" },
+                  { "name": "reverse", "type": "u32", "value": "attrs.reverse" }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "workgroups": "totalChunks" }
+        },
+        {
+          "id": "carries",
+          "name": "Scan.MultichunkCarries",
+          "source": { "shader": "scan-multichunk-carries.wgsl.jinja", "inputs": { "chunks": "chunks" } },
+          "bindings": [
+            {
+              "name": "initial_state",
+              "arg": "initial_state",
+              "semantic": "initial_state",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "chunkTotals",
+              "semantic": "chunkTotals",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "chunkCarries",
+              "semantic": "chunkCarries",
+              "buffer": { "type": "storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "final_state",
+              "arg": "final_state",
+              "semantic": "final_state",
+              "buffer": { "type": "storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "dim", "type": "u32", "value": "dim(shapes.scan_input, 1)" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "dim(shapes.scan_input, 1)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "apply",
+          "name": "Scan.MultichunkApply",
+          "source": { "shader": "scan-multichunk-apply.wgsl.jinja", "inputs": { "chunks": "chunks" } },
+          "bindings": [
+            {
+              "name": "chunkCarries",
+              "semantic": "chunkCarries",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "scan_output",
+              "arg": "scan_output",
+              "semantic": "scan_output",
+              "buffer": { "type": "storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "count", "type": "u32", "value": "numel(shapes.scan_output)" },
+                  { "name": "steps", "type": "u32", "value": "dim(shapes.scan_input, 0)" },
+                  { "name": "dim", "type": "u32", "value": "dim(shapes.scan_input, 1)" },
+                  { "name": "reverse", "type": "u32", "value": "attrs.reverse" }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.scan_output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "coop_channel_prefix_sum",
+      "priority": 10,
+      "when": ["ranks.initial_state == 1", "ranks.scan_input == 2", "ranks.final_state == 1", "ranks.scan_output == 2", "dim(shapes.initial_state, 0) == dim(shapes.scan_input, 1)", "dim(shapes.final_state, 0) == dim(shapes.initial_state, 0)", "dim(shapes.scan_output, 0) == dim(shapes.scan_input, 0)", "dim(shapes.scan_output, 1) == dim(shapes.scan_input, 1)", "dim(shapes.scan_input, 1) >= 1", "dim(shapes.scan_input, 1) <= device.limits.maxComputeWorkgroupsPerDimension", "dim(shapes.scan_input, 0) >= 64", "dim(shapes.scan_input, 1) >= 1"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Scan.CoopChannel",
+          "source": {
+            "shader": "scan-coop-channel.wgsl.jinja",
+            "inputs": { "reverse": "attrs.reverse != 0", "useSubgroups": "scanUseSubgroups" }
+          },
+          "bindings": "scanState",
+          "dispatch": { "workgroups": "dim(shapes.scan_input, 1)" }
+        }
+      ]
+    },
+    {
+      "id": "lowered_prefix_sum",
+      "when": ["ranks.initial_state == 1", "ranks.scan_input == 2", "ranks.final_state == 1", "ranks.scan_output == 2", "dim(shapes.initial_state, 0) == dim(shapes.scan_input, 1)", "dim(shapes.final_state, 0) == dim(shapes.initial_state, 0)", "dim(shapes.scan_output, 0) == dim(shapes.scan_input, 0)", "dim(shapes.scan_output, 1) == dim(shapes.scan_input, 1)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Scan",
+          "shader": "scan-prefix-sum.wgsl.jinja",
+          "bindings": "scanState",
+          "dispatch": { "threads": "dim(shapes.scan_input, 1)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Scan",
+  "id": "_ai_onnx_scan_webgpu_b9f0a4d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "0+C2HUd1flyDUAL11YlBp9o6XTgmDufVT4+as0ZGYZs=",
+      "manifest.json": "sOeqaF0C7CiKDFazgO3dRk9xClpZmysMuCawCo1rO7M=",
+      "scan-coop-channel.wgsl.jinja": "oxDxT4gqbqgj7CWwFwAm+RSjXJoNzaio+2Kn7ST6qwA=",
+      "scan-multichunk-apply.wgsl.jinja": "zrBQlGJuesvSZthECsjSYiSaULCb0AiLsIs1j0qCd0U=",
+      "scan-multichunk-carries.wgsl.jinja": "9mRyBR64N4aMcggHqIfg0GdfN7WSe2dxfiqiXb5hmgg=",
+      "scan-multichunk-local.wgsl.jinja": "BIVowPnPqSNj7WQwZFBnpmwrwmmCzJsYiuXSUcPt/9I=",
+      "scan-prefix-sum.wgsl.jinja": "/gbGyAZ3J8GXqcPXK86J1BKxNHjezV3/nvXsBTi5OHM=",
+      "test.json": "DCu1yx2q9qHCxbqB0yyT2mZN656tT+ja1Mrf9knXxKM="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Scan" }
+}

build/webgpu/scan-coop-channel.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,121 @@

+// Workgroup-per-channel prefix scan for small channel counts. Each workgroup
+// scans the sequential steps axis in workgroup-sized chunks; channel i at step
+// t is stored at t*dim + i. Reverse mode changes traversal direction.
+//
+{% if source.useSubgroups %}
+// Each chunk combines an inclusive subgroup scan with a cross-subgroup carry.
+{% else %}
+// Each chunk uses a workgroup-wide Hillis-Steele scan.
+{% endif %}
+// Chunk-local sums reassociate the additions, while the cross-chunk carry
+// preserves prefix dependencies. Float results therefore depend on this tree.
+{% if source.useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const IDENTITY: f32 = 0.0;
+{% if source.useSubgroups %}
+// One partial per subgroup; subgroup size bounds the count to WG/minSubgroup.
+var<workgroup> sgTotals: array<f32, WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(subgroup_invocation_id) sgLane: u32,
+        @builtin(subgroup_size) sgSize: u32) {
+  // wg.y carries channel-index bits past the per-dimension dispatch limit.
+  let i = wg.x + wg.y * nwg.x;
+  let tid = lid.x;
+  if (i >= params.dim) {
+    return;
+  }
+  let safeSg = max(sgSize, 1u);
+  let sgId = tid / safeSg;
+  let numSub = (WG + safeSg - 1u) / safeSg;
+  var carry: f32 = initial_state[i];
+  let chunks = (params.steps + WG - 1u) / WG;
+  for (var c = 0u; c < chunks; c = c + 1u) {
+    let j = c * WG + tid;
+    let valid = j < params.steps;
+    let t = select(j, params.steps - 1u - j, params.reverse != 0u);
+    let addr = t * params.dim + i;
+    var v: f32 = IDENTITY;
+    if (valid) {
+      v = scan_input[addr];
+    }
+    let incSg = subgroupInclusiveAdd(v);
+    let totSg = subgroupAdd(v);
+    if (sgLane == 0u) {
+      sgTotals[min(sgId, WG - 1u)] = totSg;
+    }
+    workgroupBarrier();
+    var sgOffset: f32 = IDENTITY;
+    for (var s = 0u; s < sgId; s = s + 1u) {
+      sgOffset = sgOffset + sgTotals[s];
+    }
+    var chunkTotal: f32 = IDENTITY;
+    for (var s = 0u; s < numSub; s = s + 1u) {
+      chunkTotal = chunkTotal + sgTotals[s];
+    }
+    let value = carry + sgOffset + incSg;
+    if (valid) {
+      scan_output[addr] = value;
+    }
+    carry = carry + chunkTotal;
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    final_state[i] = carry;
+  }
+}
+{% else %}
+var<workgroup> wgScan: array<f32, WG>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  let i = wg.x + wg.y * nwg.x;
+  let tid = lid.x;
+  if (i >= params.dim) {
+    return;
+  }
+  var carry: f32 = initial_state[i];
+  let chunks = (params.steps + WG - 1u) / WG;
+  for (var c = 0u; c < chunks; c = c + 1u) {
+    let j = c * WG + tid;
+    let valid = j < params.steps;
+    let t = select(j, params.steps - 1u - j, params.reverse != 0u);
+    let addr = t * params.dim + i;
+    var v: f32 = IDENTITY;
+    if (valid) {
+      v = scan_input[addr];
+    }
+    workgroupBarrier();
+    wgScan[tid] = v;
+    for (var step = 1u; step < WG; step = step << 1u) {
+      workgroupBarrier();
+      var prev: f32 = IDENTITY;
+      if (tid >= step) {
+        prev = wgScan[tid - step];
+      }
+      workgroupBarrier();
+      wgScan[tid] = prev + wgScan[tid];
+    }
+    workgroupBarrier();
+    let value = carry + wgScan[tid];
+    if (valid) {
+      scan_output[addr] = value;
+    }
+    carry = carry + wgScan[WG - 1u];
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    final_state[i] = carry;
+  }
+}
+{% endif %}

build/webgpu/scan-multichunk-apply.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,16 @@

+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const CHUNKS: u32 = {{ source.chunks }}u;
+@compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let i = gid.x + gid.y * nwg.x * WG;
+  if (i >= params.count) { return; }
+  let t = i / params.dim;
+  let channel = i % params.dim;
+  let iter = select(t, params.steps - 1u - t, params.reverse != 0u);
+  let chunk = iter / WG;
+  scan_output[i] += chunkCarries[channel * CHUNKS + chunk];
+}

build/webgpu/scan-multichunk-carries.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,17 @@

+{{ env.wgsl.resourceDeclarations }}
+const CHUNKS: u32 = {{ source.chunks }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let channel = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (channel >= params.dim) { return; }
+  var carry = initial_state[channel];
+  let base = channel * CHUNKS;
+  for (var chunk = 0u; chunk < CHUNKS; chunk += 1u) {
+    chunkCarries[base + chunk] = carry;
+    carry += chunkTotals[base + chunk];
+  }
+  final_state[channel] = carry;
+}

build/webgpu/scan-multichunk-local.wgsl.jinja ADDED Viewed

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

+// Workgroup-parallel local prefix chunks for small-state, long Scan streams.
+{% if source.useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const CHUNKS: u32 = {{ source.chunks }}u;
+const TOTAL_CHUNKS: u32 = {{ source.totalChunks }}u;
+{% if source.useSubgroups %}
+var<workgroup> sgTotals: array<f32, WG>;
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(subgroup_invocation_id) sgLane: u32,
+        @builtin(subgroup_size) sgSize: u32) {
+  let group = wg.x + wg.y * nwg.x;
+  if (group >= TOTAL_CHUNKS) { return; }
+  let channel = group / CHUNKS;
+  let chunk = group % CHUNKS;
+  let j = chunk * WG + lid.x;
+  let valid = j < params.steps;
+  let t = select(j, params.steps - 1u - j, params.reverse != 0u);
+  let addr = t * params.dim + channel;
+  var v = 0.0;
+  if (valid) { v = scan_input[addr]; }
+  let inclusive = subgroupInclusiveAdd(v);
+  let subgroupTotal = subgroupAdd(v);
+  let safeSg = max(sgSize, 1u);
+  let subgroupId = lid.x / safeSg;
+  let subgroupCount = (WG + safeSg - 1u) / safeSg;
+  if (sgLane == 0u) { sgTotals[subgroupId] = subgroupTotal; }
+  workgroupBarrier();
+  var offset = 0.0;
+  for (var s = 0u; s < subgroupId; s += 1u) {
+    offset += sgTotals[s];
+  }
+  if (valid) { scan_output[addr] = offset + inclusive; }
+  if (lid.x == 0u) {
+    var total = 0.0;
+    for (var s = 0u; s < subgroupCount; s += 1u) {
+      total += sgTotals[s];
+    }
+    chunkTotals[group] = total;
+  }
+}
+{% else %}
+var<workgroup> values: array<f32, WG>;
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  let group = wg.x + wg.y * nwg.x;
+  if (group >= TOTAL_CHUNKS) { return; }
+  let channel = group / CHUNKS;
+  let chunk = group % CHUNKS;
+  let j = chunk * WG + lid.x;
+  let valid = j < params.steps;
+  let t = select(j, params.steps - 1u - j, params.reverse != 0u);
+  let addr = t * params.dim + channel;
+  var v = 0.0;
+  if (valid) { v = scan_input[addr]; }
+  values[lid.x] = v;
+  for (var step = 1u; step < WG; step <<= 1u) {
+    workgroupBarrier();
+    var addend = 0.0;
+    if (lid.x >= step) { addend = values[lid.x - step]; }
+    workgroupBarrier();
+    values[lid.x] += addend;
+  }
+  workgroupBarrier();
+  if (valid) { scan_output[addr] = values[lid.x]; }
+  if (lid.x == 0u) { chunkTotals[group] = values[WG - 1u]; }
+}
+{% endif %}

build/webgpu/scan-prefix-sum.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,18 @@

+{{ env.wgsl.resourceDeclarations }}
+// One invocation owns each channel and runs its prefix sum serially. Channels
+// run independently, but the cross-step recurrence and per-channel addition
+// order remain sequential.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // gid.y carries channel-index bits past the per-dimension dispatch limit.
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.dim) { return; }
+  var state = initial_state[i];
+  for (var iter = 0u; iter < params.steps; iter = iter + 1u) {
+    let t = select(iter, params.steps - 1u - iter, params.reverse != 0u);
+    state = state + scan_input[t * params.dim + i];
+    scan_output[t * params.dim + i] = state;
+  }
+  final_state[i] = state;
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,360 @@

+{
+  "op": "ai.onnx.Scan",
+  "cases": [
+    {
+      "name": "lowered_prefix_sum_state",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 10.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_reverse",
+      "attrs": { "reverse": 1 },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [4, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [4, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_ort_projection_short_sequence_two_state_lanes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/scan_test.cc",
+        "test": "Scan8.ShortSequenceTwoInBatchOneLoopStateVar",
+        "notes": "Projection onto the framework's lowered prefix-sum variant using two state lanes from ORT's short sequence data."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 10.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [1.0, -1.0, 4.0, -4.0] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_ort_projection_reverse_two_state_lanes",
+      "attrs": { "reverse": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/scan_test.cc",
+        "test": "Scan8.MixedSequenceLensReverse",
+        "notes": "Projection onto the framework's lowered reverse prefix-sum variant using reverse-direction values from ORT's Scan coverage."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 10.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [1.0, -1.0, 4.0, -4.0] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_ort_projection_scalar_loop_state",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/scan_test.cc",
+        "test": "Scan8.OnnxScalarLoopState",
+        "notes": "Projection onto the framework's lowered prefix-sum variant using a scalar carried state."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "scan_input": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_zero_steps",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, -3.0] } },
+        "scan_input": { "dtype": "float32", "shape": [0, 2], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [0, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_reverse_rank2_signed",
+      "attrs": { "reverse": 1 },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [10.0, -10.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, -1.0, 2.0, -2.0, 3.0, -3.0] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_reverse_zero_steps",
+      "attrs": { "reverse": 1 },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [7.0, -8.0] } },
+        "scan_input": { "dtype": "float32", "shape": [0, 2], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [0, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_zero_delta_lane",
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 10.0, -5.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [0.0, 2.0, -1.0, 0.0, -3.0, 4.0, 0.0, 1.0, -2.0] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_128x64",
+      "provenance": {
+        "notes": "Cross-gate case for the coop_channel_prefix_sum variant (dim>=64 and steps>=64). dim=64 -> 64 channel workgroups; steps=128 spans multiple WG-sized chunks with a cross-chunk carry. Tolerance covers the within-chunk subgroup/Hillis-Steele reassociation vs the strict serial loop."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [64] },
+        "scan_input": { "dtype": "float32", "shape": [128, 64] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [64], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_reverse_300x96",
+      "attrs": { "reverse": 1 },
+      "provenance": {
+        "notes": "Reverse-direction cross-gate case for coop_channel_prefix_sum. steps=300 is NOT a multiple of the 256 workgroup size, exercising the partial-chunk identity padding; dim=96 channel workgroups."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [96] },
+        "scan_input": { "dtype": "float32", "shape": [300, 96] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [96], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [300, 96], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_dim512_steps256_perf_compact",
+      "provenance": {
+        "notes": "Compact correctness sibling for the Scan no-MMA prefix-sum cliff: dim=512 and steps=256 select coop_channel_prefix_sum with one full workgroup-sized scan chunk."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [512] },
+        "scan_input": { "dtype": "float32", "shape": [256, 512] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [512], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [256, 512], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_dim1024_steps256",
+      "provenance": {
+        "notes": "Exercises a wide cooperative dispatch (1024 channel workgroups) and one full workgroup-sized scan chunk. Tolerance covers within-chunk subgroup/Hillis-Steele reassociation."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1024] },
+        "scan_input": { "dtype": "float32", "shape": [256, 1024] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1024], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [256, 1024], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "coop_dim1025_above_old_fixed_gate",
+      "provenance": {
+        "notes": "Regression for the former hard-coded dim<=1024 ceiling. Device-limit-derived cooperative coverage keeps this realistic 1025-channel scan off the serial fallback."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1025] },
+        "scan_input": { "dtype": "float32", "shape": [256, 1025] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1025], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [256, 1025], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_dim64_steps64_min_gate",
+      "provenance": {
+        "notes": "Both lower coop gates at their exact inclusive thresholds: dim=64 (dim>=64) and steps=64 (steps>=64). This is the smallest shape that still selects coop_channel_prefix_sum. An off-by-one tightening of either >=64 gate (to >64) would demote this corner to the lowered fallback. Verifies coop at the minimum supported width and exactly one partial-free chunk boundary (steps==WG/4)."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [64] },
+        "scan_input": { "dtype": "float32", "shape": [64, 64] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [64], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [64, 64], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "lowered_dim64_steps63_below_steps_gate",
+      "provenance": {
+        "notes": "Just below the coop steps gate: dim=64 passes dim>=64 but steps=63 fails steps>=64, so this MUST fall to lowered_prefix_sum. Pins the steps off-by-one: an erroneous steps>=63 (or >63 vs >=64) gate would mis-route this short sequence. Both variants are spec-correct; this checks the lowered path at the exact step below the coop threshold."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [64] },
+        "scan_input": { "dtype": "float32", "shape": [63, 64] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [64], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [63, 64], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_subnormal_residuals_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes denormals to zero in the ALU; this lowered prefix-sum accumulates subnormal residuals that flush on GPU. Permanent FTZ limitation."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/scan_test.cc",
+        "test": "Scan8.ShortSequenceTwoInBatchOneLoopStateVar",
+        "notes": "Projection onto the lowered prefix-sum variant: finite subnormal loop-state increments are valid residuals and must not flush to zero."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 0.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-40, -1e-40, -1e-40, 1e-40, 2e-40, -2e-40] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0 },
+        "scan_output": { "dtype": "float32", "shape": [4, 2], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "lowered_prefix_sum_reverse_subnormal_residuals_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes denormals to zero in the ALU; this lowered prefix-sum accumulates subnormal residuals that flush on GPU. Permanent FTZ limitation."
+      },
+      "attrs": { "reverse": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/scan_test.cc",
+        "test": "Scan8.MixedSequenceLensReverse",
+        "notes": "Reverse companion for subnormal prefix residuals; the scan-output slots are written in reverse traversal order but still contain finite float32 subnormal sums."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [3, 1],
+          "data": { "kind": "values", "values": [1e-40, 2e-40, 3e-40] }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1], "tolerance": 0 },
+        "scan_output": { "dtype": "float32", "shape": [3, 1], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_dim64_steps257_chunk_carry",
+      "provenance": {
+        "notes": "Routes to coop_channel_prefix_sum (dim=64>=64, dim<=1024, steps=257>=64). chunks=ceil(257/256)=2: chunk 0 = steps 0..255, chunk 1 = ONLY step 256 (one valid lane, 255 identity-padded). Isolates cross-chunk carry propagation into a near-empty second chunk and the final_state = carry-after-both-chunks path. An off-by-one chunk bound, stale carry, or padded-lane double-count would corrupt scan_output[256*64+i] and final_state. Verifies BOTH outputs against the TS reference; tolerance 1e-4 covers within-chunk subgroup/Hillis-Steele reassociation."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [64] },
+        "scan_input": { "dtype": "float32", "shape": [257, 64] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [64], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [257, 64], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "coop_channel_prefix_sum_reverse_dim64_steps257_chunk_carry",
+      "attrs": { "reverse": 1 },
+      "provenance": {
+        "notes": "Reverse coop_channel_prefix_sum (dim=64>=64, dim<=1024, steps=257>=64, reverse=1). chunks=2: chunk 0 (j=0..255) -> original steps 256..1, chunk 1 (j=256) -> original step 0. Pins the minimal 1-step residual chunk under reverse: verifies the reversed read/write addresses at the chunk boundary and the scan-order cross-chunk carry land at the correct reversed scan_output slots and final_state. Both outputs checked vs the TS reference; tolerance 1e-4 covers within-chunk reassociation."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [64] },
+        "scan_input": { "dtype": "float32", "shape": [257, 64] }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [64], "tolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [257, 64], "tolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "lowered_dim1_steps2000_singlechannel_serial",
+      "provenance": {
+        "notes": "dim=1 fails coop 'dim>=64' -> lowered_prefix_sum, one active lane (i=0) running the serial 2000-step recurrence with addressing t*1+0=t. Covers the single-channel serial scan over a long sequence (existing dim=1 cases stop at steps<=4). Serial per-channel accumulation is bit-identical to the reference, so tolerance stays at 1e-6; both outputs verified vs the TS reference."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [3.5] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [2000, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "scale": 0.25, "offset": 0.01 }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [2000, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "multichunk_small_state_dim2_steps4096",
+      "provenance": {
+        "notes": "Compact correctness lock for a realistic long-stream recurrence with only two independent state lanes. The optimized path partitions the sequential axis into workgroup-sized chunks, scans chunk totals, then applies carries."
+      },
+      "inputs": {
+        "initial_state": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.5, -2.0] } },
+        "scan_input": {
+          "dtype": "float32",
+          "shape": [4096, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.013, "scale": 0.01, "offset": 0.0001 }
+        }
+      },
+      "outputs": {
+        "final_state": { "dtype": "float32", "shape": [2], "tolerance": 0.001, "relTolerance": 0.0001 },
+        "scan_output": { "dtype": "float32", "shape": [4096, 2], "tolerance": 0.001, "relTolerance": 0.0001 }
+      }
+    }
+  ]
+}