sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +75 -0
build/webgpu/bench.json +80 -0
build/webgpu/loop-add-step.wgsl.jinja +21 -0
build/webgpu/manifest.json +128 -0
build/webgpu/metadata.json +18 -0
build/webgpu/test.json +299 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Loop
+`ai.onnx`  ·  internal tensor lowering (non-standard)  ·  reviewed against ONNX opset 25
+## Description
+Support status: the standard ONNX `Loop` control-flow operator is not implemented because standalone kernel packages cannot carry or execute its body graph. This internal lowering performs one fixed recurrence—adding `step` to a rank-1 loop-carried tensor for up to `M` iterations while writing a dense scan tensor; `step` is not an ONNX `Loop` input, and this lowering must not be treated as ONNX `Loop`.
+See the [standard ONNX `Loop` spec](https://onnx.ai/onnx/operators/onnx__Loop.html) for the contract this internal lowering does not implement.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `M` | `m` | `I` | — | — | Required uint32 trip-count limit encoded as either a scalar tensor or a one-element rank-1 tensor. The lowering executes at most `min(M, scan_output.shape[0])` iterations. | required |
+| `cond` | `cond` | `B` | — | — | Required initial condition encoded as either a scalar tensor or a one-element rank-1 tensor. A false value executes zero iterations; a true value permits all iterations selected by `M`. This fixed lowering does not update the condition inside the loop. | required |
+| `v_initial` | `v_initial` | `T` | `1` | — | Initial rank-1 loop-carried state of shape `[dim]`. | required |
+| `step` | `step` | `T` | `1` | — | Implementation-specific rank-1 increment of shape `[dim]`, added elementwise to the state on every executed iteration. This is not a standard ONNX `Loop` input. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `v_final` | `v_final` | `T` | `1` | same as `v_initial` | Final rank-1 state of shape `[dim]` after the executed additions. | required |
+| `scan_output` | `scan_output` | `T` | `2` | — | Dense tensor of shape `[scan_steps, dim]`. Each executed row contains the updated state for that iteration; rows beyond the executed trip count are zero-filled. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32` |
+| `I` | `uint32` |
+| `B` | `uint32`, `bool` |
+## 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
+- [`loop-add-step.wgsl.jinja`](build/webgpu/loop-add-step.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
+The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
+- `scan_output`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+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.Loop", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { v_final, scan_output } = await kernel({
+  m: { data: mData, shape: [1] },
+  cond: { data: condData, shape: [1] },
+  v_initial: { data: v_initialData, shape: [1] },
+  step: { data: stepData, shape: [1] },
+}, {
+  outputs: { scan_output: { shape: [1, 1], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,80 @@

+{
+  "op": "ai.onnx.Loop",
+  "cases": [
+    {
+      "name": "lowered_HEALTHY_baseline_dim128_steps128",
+      "preset": "smoke",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [128], "dist": "normal", "seed": 8101, "scale": 0.1 },
+        "step": { "dtype": "float32", "shape": [128], "dist": "normal", "seed": 8102, "scale": 0.1 }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [128] },
+        "scan_output": { "dtype": "float32", "shape": [128, 128] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
+    },
+    {
+      "name": "lowered_PATHO_singlelane_serial_dim2048_steps512",
+      "preset": "smoke",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [512] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2048], "dist": "normal", "seed": 8103, "scale": 0.1 },
+        "step": { "dtype": "float32", "shape": [2048], "dist": "normal", "seed": 8104, "scale": 0.1 }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2048] },
+        "scan_output": { "dtype": "float32", "shape": [512, 2048] }
+      },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
+    },
+    {
+      "name": "lowered_PATHO_singlelane_serial_dim4096_steps1024",
+      "preset": "smoke",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1024] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [4096], "dist": "normal", "seed": 8105, "scale": 0.1 },
+        "step": { "dtype": "float32", "shape": [4096], "dist": "normal", "seed": 8106, "scale": 0.1 }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [4096] },
+        "scan_output": { "dtype": "float32", "shape": [1024, 4096] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
+    },
+    {
+      "name": "lowered_PATHO_smalldim2_singlelane_steps8192",
+      "preset": "smoke",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8192] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "dist": "normal", "seed": 8107, "scale": 0.1 },
+        "step": { "dtype": "float32", "shape": [2], "dist": "normal", "seed": 8108, "scale": 0.1 }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2] },
+        "scan_output": { "dtype": "float32", "shape": [8192, 2] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
+    },
+    {
+      "name": "lowered_PATHO_nonpow2_unaligned_dim1000_steps1000",
+      "preset": "smoke",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [1000], "dist": "normal", "seed": 8109, "scale": 0.1 },
+        "step": { "dtype": "float32", "shape": [1000], "dist": "normal", "seed": 8110, "scale": 0.1 }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [1000] },
+        "scan_output": { "dtype": "float32", "shape": [1000, 1000] }
+      },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * 4 * numel(shapes.scan_output)" }] }
+    }
+  ]
+}

build/webgpu/loop-add-step.wgsl.jinja ADDED Viewed

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

+{{ env.wgsl.resourceDeclarations }}
+// One invocation owns each loop-carried element and walks its iterations
+// serially. Columns run independently, but the per-column additions and stores
+// stay sequential, preserving the f32 accumulation order.
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= params.dim) { return; }
+  let iterations = select(0u, min(m[0], params.scanSteps), cond[0] != 0u);
+  var value = v_initial[i];
+  for (var t = 0u; t < params.scanSteps; t = t + 1u) {
+    if (t < iterations) {
+      value = value + step[i];
+      scan_output[t * params.dim + i] = value;
+    } else {
+      scan_output[t * params.dim + i] = 0.0;
+    }
+  }
+  v_final[i] = value;
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,128 @@

+{
+  "domain": "ai.onnx",
+  "name": "Loop",
+  "conformance": "internal-lowering",
+  "sinceVersion": 25,
+  "description": "Support status: the standard ONNX `Loop` control-flow operator is not implemented because standalone kernel packages cannot carry or execute its body graph. This internal lowering performs one fixed recurrence—adding `step` to a rank-1 loop-carried tensor for up to `M` iterations while writing a dense scan tensor; `step` is not an ONNX `Loop` input, and this lowering must not be treated as ONNX `Loop`.",
+  "inputs": [
+    {
+      "role": "M",
+      "dtype": "I",
+      "description": "Required uint32 trip-count limit encoded as either a scalar tensor or a one-element rank-1 tensor. The lowering executes at most `min(M, scan_output.shape[0])` iterations."
+    },
+    {
+      "role": "cond",
+      "dtype": "B",
+      "description": "Required initial condition encoded as either a scalar tensor or a one-element rank-1 tensor. A false value executes zero iterations; a true value permits all iterations selected by `M`. This fixed lowering does not update the condition inside the loop."
+    },
+    {
+      "role": "v_initial",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Initial rank-1 loop-carried state of shape `[dim]`."
+    },
+    {
+      "role": "step",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Implementation-specific rank-1 increment of shape `[dim]`, added elementwise to the state on every executed iteration. This is not a standard ONNX `Loop` input."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "v_final",
+      "dtype": "T",
+      "rank": 1,
+      "description": "Final rank-1 state of shape `[dim]` after the executed additions.",
+      "shape": "shapes.v_initial"
+    },
+    {
+      "role": "scan_output",
+      "dtype": "T",
+      "rank": 2,
+      "description": "Dense tensor of shape `[scan_steps, dim]`. Each executed row contains the updated state for that iteration; rows beyond the executed trip count are zero-filled."
+    }
+  ],
+  "typeConstraints": { "T": ["float32"], "I": ["uint32"], "B": ["uint32", "bool"] },
+  "args": {
+    "m": { "kind": "tensor", "semantic": "M", "role": "input" },
+    "cond": { "kind": "tensor", "semantic": "cond", "role": "input" },
+    "v_initial": { "kind": "tensor", "semantic": "v_initial", "role": "input" },
+    "step": { "kind": "tensor", "semantic": "step", "role": "input" },
+    "v_final": { "kind": "tensor", "semantic": "v_final", "role": "output" },
+    "scan_output": { "kind": "tensor", "semantic": "scan_output", "role": "output" }
+  },
+  "variants": [
+    {
+      "id": "lowered_add_step",
+      "when": ["ranks.M <= 1", "numel(shapes.M) == 1", "ranks.cond <= 1", "numel(shapes.cond) == 1", "ranks.v_initial == 1", "ranks.step == 1", "ranks.v_final == 1", "ranks.scan_output == 2", "dim(shapes.step, 0) == dim(shapes.v_initial, 0)", "dim(shapes.v_final, 0) == dim(shapes.v_initial, 0)", "dim(shapes.scan_output, 1) == dim(shapes.v_initial, 0)"],
+      "derive": { "dimBlocks": "ceil(dim(shapes.v_initial, 0) / 256)" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Loop",
+          "shader": "loop-add-step.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "m",
+              "arg": "m",
+              "semantic": "M",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "u32",
+              "length": 1
+            },
+            {
+              "name": "cond",
+              "arg": "cond",
+              "semantic": "cond",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "u32",
+              "length": 1
+            },
+            {
+              "name": "v_initial",
+              "arg": "v_initial",
+              "semantic": "v_initial",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "step",
+              "arg": "step",
+              "semantic": "step",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "f32"
+            },
+            {
+              "name": "v_final",
+              "arg": "v_final",
+              "semantic": "v_final",
+              "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": "dim", "type": "u32", "value": "dim(shapes.v_initial, 0)" },
+                  { "name": "scanSteps", "type": "u32", "value": "dim(shapes.scan_output, 0)" }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "x": "dimBlocks" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Loop",
+  "id": "_ai_onnx_loop_webgpu_fb2e7a5",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "xq5d0FI/A9ESZ8tyzz5qAxP39qxO5jGxCNji2YREKxI=",
+      "loop-add-step.wgsl.jinja": "Sk4601cYmNLKCMIFX47m3mKo2M5oh/fekvOwLOUN1D8=",
+      "manifest.json": "C5Kc57/iLS4XSFrNuO6zRpUHfXkgOI8Hm+LlxTfs/eg=",
+      "test.json": "UG6yKjMh8XfBVE9AZSF2FfMZjBNwXajveageqJpYjWY="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Loop" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,299 @@

+{
+  "op": "ai.onnx.Loop",
+  "cases": [
+    {
+      "name": "lowered_add_step_scan",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -1.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 2.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [4, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_cond_false_zero_iterations",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "v_initial": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } },
+        "step": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_ort_projection_single_iteration_scan_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.Opset11WithNoVariadicInputsAndOutputs",
+        "notes": "Projection onto the framework's lowered add-step loop: one true iteration produces one scan output value."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "step": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_ort_projection_zero_iterations_initial_passthrough",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ZeroIterations",
+        "notes": "Projection onto the framework's lowered add-step loop: trip count zero must leave carried state unchanged and produce no scan elements."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [123.0, -456.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [99.0, 99.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [0, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_ort_projection_exit_due_to_max_iterations",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ExitDueToMaxIterations",
+        "notes": "Projection onto the framework's lowered add-step loop: two iterations advance the carried state and produce two scan values."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } },
+        "step": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [3.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_trip_count_exceeds_scan_capacity",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [6] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [10.0, -10.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -2.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_zero_trip_count_cond_true",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [5.0] } },
+        "step": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [99.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [1], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [0, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_trip_count_less_than_scan_capacity_tail_zero",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -1.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [5, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_zero_step_preserves_lane",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, -2.0, 5.0] } },
+        "step": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 1.5, -2.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_scalar_trip_count_and_cond",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ExitDueToMaxIterations",
+        "notes": "ONNX Loop trip count and initial condition are scalar values. The project-lowered add-step form should accept scalar logical uint32 metadata as well as [1] tensors."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "cond": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -2.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 3.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_bool_scalar_cond_true",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ExitDueToMaxIterations",
+        "notes": "ONNX Loop conditions are bool tensors. This lowered add-step projection verifies scalar logical bool storage enables iteration."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "cond": { "dtype": "bool", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, -3.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [4.0, 0.5] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0 },
+        "scan_output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "lowered_bool_len1_cond_false_zero_iterations",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ZeroIterations",
+        "notes": "Length-1 logical bool false condition must prevent all iterations, leave state unchanged, and zero-fill the bounded scan output."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "cond": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "v_initial": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, -1.0, 10.0] } },
+        "step": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [100.0, 100.0, 100.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [3], "tolerance": 0 },
+        "scan_output": { "dtype": "float32", "shape": [4, 3], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "lowered_noncanonical_cond_true_value_two",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ExitDueToMaxIterations",
+        "notes": "ONNX only requires the loop condition to be a bool; runtimes that store it as an integer may emit any nonzero value. The lowered kernel must treat any nonzero cond (here 2) as true via cond[0] != 0u, matching the reference's !== 0 test, not a literal == 1 comparison."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "v_initial": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, -1.0] } },
+        "step": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 2.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [2], "tolerance": 0 },
+        "scan_output": { "dtype": "float32", "shape": [3, 2], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "lowered_trip_count_equals_capacity_full_fill_wide_dim64",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ExitDueToMaxIterations",
+        "notes": "Realistic carried-state width (dim 64, a typical small recurrent hidden size) with trip count exactly equal to the scan capacity so every scan row is produced by an iteration and the tail-zero branch is never taken. Exercises the min(m, scanSteps) clamp at its equality boundary across many lanes."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": {
+            "kind": "values",
+            "values": [-0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4]
+          }
+        },
+        "step": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": {
+            "kind": "values",
+            "values": [-0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4]
+          }
+        }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [64], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_clamped_trip_count_partial_fill_wide_dim33_unaligned",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/controlflow/loop_test.cc",
+        "test": "Loop.ExitDueToMaxIterations",
+        "notes": "Non-multiple-of-4 carried-state width (dim 33) combined with a trip count larger than the scan capacity. iterations = min(m, scanSteps) clamps to the capacity, the first rows accumulate, and the unaligned final lane must still be written correctly by the scalar per-lane loop with no vec4 over-read at the dim tail."
+      },
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": {
+          "dtype": "float32",
+          "shape": [33],
+          "data": {
+            "kind": "values",
+            "values": [-0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1, -0.5, 0.2, -0.2, 0.5, 0.1, -0.3, 0.4, 0.0, -0.4, 0.3, -0.1]
+          }
+        },
+        "step": {
+          "dtype": "float32",
+          "shape": [33],
+          "data": {
+            "kind": "values",
+            "values": [-0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3, -0.1, 0.4, 0.0, -0.4, 0.1, -0.3, 0.2, -0.2, 0.3]
+          }
+        }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [33], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [6, 33], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "lowered_positive_m_zero_scan_capacity_preserves_initial",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, -2.0, 3.5] } },
+        "step": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [100.0, 100.0, 100.0] } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [3], "tolerance": 0 },
+        "scan_output": { "dtype": "float32", "shape": [0, 3], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "lowered_dim_straddles_workgroup_boundary_dim258",
+      "inputs": {
+        "m": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "cond": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "v_initial": { "dtype": "float32", "shape": [258], "data": { "kind": "constant", "value": 1.0 } },
+        "step": { "dtype": "float32", "shape": [258], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": {
+        "v_final": { "dtype": "float32", "shape": [258], "tolerance": 0.000001 },
+        "scan_output": { "dtype": "float32", "shape": [3, 258], "tolerance": 0.000001 }
+      }
+    }
+  ]
+}