sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +84 -0
build/webgpu/bench.json +43 -0
build/webgpu/manifest.json +296 -0
build/webgpu/metadata.json +18 -0
build/webgpu/split-to-sequence.wgsl.jinja +62 -0
build/webgpu/test.json +642 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.SplitToSequence
+`ai.onnx`  ·  internal tensor lowering (non-standard)  ·  reviewed against ONNX opset 11
+## Description
+Support status: the standard ONNX `SplitToSequence` operator is not implemented because this kernel ABI exposes tensors rather than an ONNX sequence value. This internal lowering copies an input into two, three, four, or six caller-shaped tensor outputs and therefore does not implement the standard `output_sequence` contract; it must not be treated as ONNX `SplitToSequence`.
+See the [standard ONNX `SplitToSequence` spec](https://onnx.ai/onnx/operators/onnx__SplitToSequence.html) for the contract this internal lowering does not implement.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | — | — | The tensor to split. | required |
+| `split` | `split` | `S` | — | — | Length of each output slice: a scalar for uniform chunks or a 1-D tensor of per-output lengths. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y0` | `y0` | `T` | derived | — | First output tensor slice. | required |
+| `Y1` | `y1` | `T` | derived | — | Second output tensor slice. | required |
+| `Y2` | `y2` | `T` | derived | — | Third output tensor slice. | optional |
+| `Y3` | `y3` | `T` | derived | — | Fourth output tensor slice. | optional |
+| `Y4` | `y4` | `T` | derived | — | Fifth output tensor slice. | optional |
+| `Y5` | `y5` | `T` | derived | — | Sixth output tensor slice. | optional |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `0` | Axis along which to split; negative values count from the back. Accepted range is `[-rank, rank-1]`. |
+| `keepdims` | `1` | Whether to keep the split dimension in the output (default `1`). Ignored when `split` is provided. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `bool` |
+| `S` | `uint32` |
+## 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
+- [`split-to-sequence.wgsl.jinja`](build/webgpu/split-to-sequence.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:
+- `y0`
+- `y1`
+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.SplitToSequence", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { y0, y1 } = await kernel({ input: { data: inputData, shape: [3, 2] } }, {
+  outputs: {
+    y0: { shape: [1, 2], dtype: "float32" },
+    y1: { shape: [1, 2], dtype: "float32" },
+  },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,43 @@

+{
+  "op": "ai.onnx.SplitToSequence",
+  "cases": [
+    {
+      "name": "split3-1536x256-f32",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1536, 256] },
+        "split": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [512, 512, 512] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [512, 256] },
+        "y1": { "dtype": "float32", "shape": [512, 256] },
+        "y2": { "dtype": "float32", "shape": [512, 256] }
+      },
+      "preset": "smoke"
+    },
+    {
+      "name": "rank6_squeeze_six_outputs_addr_arith_f32_stress",
+      "preset": "stress",
+      "attrs": { "axis": 1, "keepdims": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [8, 6, 4, 4, 16, 64], "dist": "normal", "seed": 1234, "scale": 2 },
+        "split": { "dtype": "uint32", "shape": [], "dist": "constant", "value": 1 }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [8, 4, 4, 16, 64], "dist": "empty" },
+        "y1": { "dtype": "float32", "shape": [8, 4, 4, 16, 64], "dist": "empty" },
+        "y2": { "dtype": "float32", "shape": [8, 4, 4, 16, 64], "dist": "empty" },
+        "y3": { "dtype": "float32", "shape": [8, 4, 4, 16, 64], "dist": "empty" },
+        "y4": { "dtype": "float32", "shape": [8, 4, 4, 16, 64], "dist": "empty" },
+        "y5": { "dtype": "float32", "shape": [8, 4, 4, 16, 64], "dist": "empty" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "elements", "value": "numel(shapes.input)" },
+          { "type": "bandwidth", "value": "numel(shapes.input)*4*2", "unit": "GB/s" }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,296 @@

+{
+  "domain": "ai.onnx",
+  "name": "SplitToSequence",
+  "conformance": "internal-lowering",
+  "sinceVersion": 11,
+  "description": "Support status: the standard ONNX `SplitToSequence` operator is not implemented because this kernel ABI exposes tensors rather than an ONNX sequence value. This internal lowering copies an input into two, three, four, or six caller-shaped tensor outputs and therefore does not implement the standard `output_sequence` contract; it must not be treated as ONNX `SplitToSequence`.",
+  "inputs": [
+    { "role": "input", "dtype": "T", "description": "The tensor to split." },
+    {
+      "role": "split",
+      "dtype": "S",
+      "optional": true,
+      "description": "Length of each output slice: a scalar for uniform chunks or a 1-D tensor of per-output lengths."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y0",
+      "dtype": "T",
+      "description": "First output tensor slice.",
+      "rank": "ranks.input if attrs.keepdims == 1 else ranks.input - 1"
+    },
+    {
+      "role": "Y1",
+      "dtype": "T",
+      "description": "Second output tensor slice.",
+      "rank": "ranks.input if attrs.keepdims == 1 else ranks.input - 1"
+    },
+    {
+      "role": "Y2",
+      "dtype": "T",
+      "optional": true,
+      "description": "Third output tensor slice.",
+      "rank": "ranks.input if attrs.keepdims == 1 else ranks.input - 1"
+    },
+    {
+      "role": "Y3",
+      "dtype": "T",
+      "optional": true,
+      "description": "Fourth output tensor slice.",
+      "rank": "ranks.input if attrs.keepdims == 1 else ranks.input - 1"
+    },
+    {
+      "role": "Y4",
+      "dtype": "T",
+      "optional": true,
+      "description": "Fifth output tensor slice.",
+      "rank": "ranks.input if attrs.keepdims == 1 else ranks.input - 1"
+    },
+    {
+      "role": "Y5",
+      "dtype": "T",
+      "optional": true,
+      "description": "Sixth output tensor slice.",
+      "rank": "ranks.input if attrs.keepdims == 1 else ranks.input - 1"
+    }
+  ],
+  "attributes": { "axis": 0, "keepdims": 1 },
+  "attributeDescriptions": {
+    "axis": "Axis along which to split; negative values count from the back. Accepted range is `[-rank, rank-1]`.",
+    "keepdims": "Whether to keep the split dimension in the output (default `1`). Ignored when `split` is provided."
+  },
+  "attributeConstraints": { "keepdims": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32", "float16", "bool"], "S": ["uint32"] },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "split": { "kind": "tensor", "semantic": "split", "role": "split", "required": false },
+    "y0": { "kind": "tensor", "semantic": "Y0", "role": "output0" },
+    "y1": { "kind": "tensor", "semantic": "Y1", "role": "output1" },
+    "y2": { "kind": "tensor", "semantic": "Y2", "role": "output2", "required": false },
+    "y3": { "kind": "tensor", "semantic": "Y3", "role": "output3", "required": false },
+    "y4": { "kind": "tensor", "semantic": "Y4", "role": "output4", "required": false },
+    "y5": { "kind": "tensor", "semantic": "Y5", "role": "output5", "required": false }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "baseContract": "ranks.input >= 1 and (attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.input) >= 0 and (attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.input) < ranks.input and f16Ok(dtypes.T)",
+    "outputRankContract": "ranks.Y0 == ranks.input if present.split else (ranks.Y0 == ranks.input if attrs.keepdims == 1 else ranks.Y0 == ranks.input - 1)",
+    "twoOutputContract": "baseContract and not present.y2 and ranks.Y0 == ranks.Y1 and outputRankContract",
+    "threeOutputContract": "baseContract and present.y2 and not present.y3 and ranks.Y0 == ranks.Y1 and ranks.Y0 == ranks.Y2 and outputRankContract",
+    "fourOutputContract": "baseContract and present.y2 and present.y3 and not present.y4 and ranks.Y0 == ranks.Y1 and ranks.Y0 == ranks.Y2 and ranks.Y0 == ranks.Y3 and outputRankContract",
+    "sixOutputContract": "baseContract and present.y2 and present.y3 and present.y4 and present.y5 and ranks.Y0 == ranks.Y1 and ranks.Y0 == ranks.Y2 and ranks.Y0 == ranks.Y3 and ranks.Y0 == ranks.Y4 and ranks.Y0 == ranks.Y5 and outputRankContract"
+  },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "variants": [
+    {
+      "id": "two_outputs",
+      "when": ["twoOutputContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "SplitToSequence2",
+          "source": {
+            "shader": "split-to-sequence.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.input",
+              "rank": "ranks.input",
+              "axis": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.input",
+              "squeeze": "ranks.Y0 < ranks.input",
+              "numOutputs": "2",
+              "outShapes": ["shapes.Y0", "shapes.Y1"]
+            }
+          },
+          "bindings": [
+            {
+              "name": "input",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "y0", "arg": "y0", "semantic": "Y0", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y1", "arg": "y1", "semantic": "Y1", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "y0Count", "type": "u32", "value": "numel(shapes.Y0)" },
+                  { "name": "y1Count", "type": "u32", "value": "numel(shapes.Y1)" }
+                ]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "max(numel(shapes.Y0), numel(shapes.Y1))",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "three_outputs",
+      "priority": 20,
+      "when": ["threeOutputContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "SplitToSequence3",
+          "source": {
+            "shader": "split-to-sequence.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.input",
+              "rank": "ranks.input",
+              "axis": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.input",
+              "squeeze": "ranks.Y0 < ranks.input",
+              "numOutputs": "3",
+              "outShapes": ["shapes.Y0", "shapes.Y1", "shapes.Y2"]
+            }
+          },
+          "bindings": [
+            {
+              "name": "input",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "y0", "arg": "y0", "semantic": "Y0", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y1", "arg": "y1", "semantic": "Y1", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y2", "arg": "y2", "semantic": "Y2", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "y0Count", "type": "u32", "value": "numel(shapes.Y0)" },
+                  { "name": "y1Count", "type": "u32", "value": "numel(shapes.Y1)" },
+                  { "name": "y2Count", "type": "u32", "value": "numel(shapes.Y2)" }
+                ]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "max(max(numel(shapes.Y0), numel(shapes.Y1)), numel(shapes.Y2))",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "four_outputs",
+      "priority": 30,
+      "when": ["fourOutputContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "SplitToSequence4",
+          "source": {
+            "shader": "split-to-sequence.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.input",
+              "rank": "ranks.input",
+              "axis": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.input",
+              "squeeze": "ranks.Y0 < ranks.input",
+              "numOutputs": "4",
+              "outShapes": ["shapes.Y0", "shapes.Y1", "shapes.Y2", "shapes.Y3"]
+            }
+          },
+          "bindings": [
+            {
+              "name": "input",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "y0", "arg": "y0", "semantic": "Y0", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y1", "arg": "y1", "semantic": "Y1", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y2", "arg": "y2", "semantic": "Y2", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y3", "arg": "y3", "semantic": "Y3", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "y0Count", "type": "u32", "value": "numel(shapes.Y0)" },
+                  { "name": "y1Count", "type": "u32", "value": "numel(shapes.Y1)" },
+                  { "name": "y2Count", "type": "u32", "value": "numel(shapes.Y2)" },
+                  { "name": "y3Count", "type": "u32", "value": "numel(shapes.Y3)" }
+                ]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "max(max(max(numel(shapes.Y0), numel(shapes.Y1)), numel(shapes.Y2)), numel(shapes.Y3))",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "six_outputs",
+      "priority": 40,
+      "when": ["sixOutputContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "SplitToSequence6",
+          "source": {
+            "shader": "split-to-sequence.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.input",
+              "rank": "ranks.input",
+              "axis": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.input",
+              "squeeze": "ranks.Y0 < ranks.input",
+              "numOutputs": "6",
+              "outShapes": ["shapes.Y0", "shapes.Y1", "shapes.Y2", "shapes.Y3", "shapes.Y4", "shapes.Y5"]
+            }
+          },
+          "bindings": [
+            {
+              "name": "input",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "y0", "arg": "y0", "semantic": "Y0", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y1", "arg": "y1", "semantic": "Y1", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y2", "arg": "y2", "semantic": "Y2", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y3", "arg": "y3", "semantic": "Y3", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y4", "arg": "y4", "semantic": "Y4", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            { "name": "y5", "arg": "y5", "semantic": "Y5", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "y0Count", "type": "u32", "value": "numel(shapes.Y0)" },
+                  { "name": "y1Count", "type": "u32", "value": "numel(shapes.Y1)" },
+                  { "name": "y2Count", "type": "u32", "value": "numel(shapes.Y2)" },
+                  { "name": "y3Count", "type": "u32", "value": "numel(shapes.Y3)" },
+                  { "name": "y4Count", "type": "u32", "value": "numel(shapes.Y4)" },
+                  { "name": "y5Count", "type": "u32", "value": "numel(shapes.Y5)" }
+                ]
+              }
+            }
+          ],
+          "dispatch": {
+            "threads": "max(max(max(max(max(numel(shapes.Y0), numel(shapes.Y1)), numel(shapes.Y2)), numel(shapes.Y3)), numel(shapes.Y4)), numel(shapes.Y5))",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.SplitToSequence",
+  "id": "_ai_onnx_splittosequence_webgpu_e21e7a0",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "EENkbgxvL5j8em9SHo+NktW1UwzMESrX+8p/FKNCoQc=",
+      "manifest.json": "bKxHYk9Csq9b7FX6sH6WaiOZyBawG9Om33PdtCe3kKs=",
+      "split-to-sequence.wgsl.jinja": "LTXzbRgOGFK2BVTVM42WPJ9DWr+mx8PcoY9jqf4KBUc=",
+      "test.json": "r6D3suRTZe+ljNzJvTQu6Le3qDRRZ6RR68+ObgTAIMY="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.SplitToSequence" }
+}

build/webgpu/split-to-sequence.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,62 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// SplitToSequence moves contiguous slices along the split axis; the output
+// tensors form the sequence. For size-one slices with
+// keepdims=0, each output drops the axis dimension. The copy reconstructs the
+// full input coordinate by reinserting that axis position.
+{% set R = source.rank %}
+{% set ax = source.axis %}
+{% set acc = namespace(offset=0) %}
+{% for k in range(source.numOutputs) %}
+{% set oShape = source.outShapes[k] %}
+fn input_offset_y{{ k }}(out_index: u32) -> u32 {
+  var rem = out_index;
+  var offset = 0u;
+{% if source.squeeze %}
+{% for d in range(R - 1) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(d + 1, R - 1) %}{% set out_stride.value = out_stride.value * oShape[j] %}{% endfor %}
+{% set in_dim = d if d < ax else d + 1 %}
+{% set in_stride = namespace(value=1) %}
+{% for j in range(in_dim + 1, R) %}{% set in_stride.value = in_stride.value * source.inputShape[j] %}{% endfor %}
+  let c{{ d }} = rem / {{ out_stride.value if out_stride.value > 0 else 1 }}u;
+  rem = rem % {{ out_stride.value if out_stride.value > 0 else 1 }}u;
+  offset = offset + c{{ d }} * {{ in_stride.value }}u;
+{% endfor %}
+{% set axis_stride = namespace(value=1) %}
+{% for j in range(ax + 1, R) %}{% set axis_stride.value = axis_stride.value * source.inputShape[j] %}{% endfor %}
+  offset = offset + {{ acc.offset }}u * {{ axis_stride.value }}u;
+{% set acc.offset = acc.offset + 1 %}
+{% else %}
+{% for d in range(R) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(d + 1, R) %}{% set out_stride.value = out_stride.value * oShape[j] %}{% endfor %}
+{% set in_stride = namespace(value=1) %}
+{% for j in range(d + 1, R) %}{% set in_stride.value = in_stride.value * source.inputShape[j] %}{% endfor %}
+  var c{{ d }} = rem / {{ out_stride.value if out_stride.value > 0 else 1 }}u;
+  rem = rem % {{ out_stride.value if out_stride.value > 0 else 1 }}u;
+{% if d == ax %}
+  c{{ d }} = c{{ d }} + {{ acc.offset }}u;
+{% endif %}
+  offset = offset + c{{ d }} * {{ in_stride.value }}u;
+{% endfor %}
+{% set acc.offset = acc.offset + oShape[ax] %}
+{% endif %}
+  return offset;
+}
+{% endfor %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{% for k in range(source.numOutputs) %}
+  if (i < params.y{{ k }}Count) {
+    y{{ k }}[i] = input[input_offset_y{{ k }}(i)];
+  }
+{% endfor %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,642 @@

+{
+  "op": "ai.onnx.SplitToSequence",
+  "fixtureArrays": {
+    "ort_positive_axis_scalar_split_rank3_input_input": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24],
+    "onnx_backend_split_scalar_axis1_three_outputs_input_input": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_two_outputs_axis0_f32",
+      "attrs": { "axis": 0, "keepdims": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [16777217, 1],
+          "data": { "kind": "linspace", "start": 0.0, "end": 1.0 }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [16777216, 1] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [16777216, 1], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [1, 1], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "split_1d_three_chunks_f32",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.37 }
+        },
+        "split": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "split_1d_unequal_two_chunks_f32",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.29 }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "scalar_split_two_chunks_f32",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.33, "cosStep": 0.17 }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [1, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "no_split_keepdims1_three_f32",
+      "attrs": { "axis": 0, "keepdims": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [1, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [1, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [1, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "no_split_keepdims0_three_f32",
+      "attrs": { "axis": 0, "keepdims": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "axis1_two_chunks_f32",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.31 }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "split_1d_three_chunks_f16",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [6, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.37 }
+        },
+        "split": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float16", "shape": [2, 2], "tolerance": 0.002 },
+        "y1": { "dtype": "float16", "shape": [2, 2], "tolerance": 0.002 },
+        "y2": { "dtype": "float16", "shape": [2, 2], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "ort_default_axis0_equal_split_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_DefaultAxis0EqualSplitFloat"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "split": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_default_axis0_equal_split_float16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_DefaultAxis0EqualSplitMLFloat16"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "split": { "dtype": "uint32", "shape": [1, 2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float16", "shape": [2, 2], "tolerance": 0.002 },
+        "y1": { "dtype": "float16", "shape": [2, 2], "tolerance": 0.002 }
+      }
+    },
+    {
+      "name": "ort_default_axis0_scalar_split_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_DefaultAxis0EqualSplitFloatScalarSplit"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_default_axis0_uneven_scalar_split_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_DefaultAxis0UnevenSplitFloat"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0] }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [1, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_positive_axis_scalar_split_rank3",
+      "attrs": { "axis": 2 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_PositiveAxisScalarSplit"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 6],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_axis_scalar_split_rank3_input_input" }
+          }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_positive_axis_uneven_scalar_split",
+      "attrs": { "axis": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_PositiveAxisUnevenSplit"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 5],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0] }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [2, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_axis0_keepdims0_rank3",
+      "attrs": { "axis": 0, "keepdims": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_Axis0DefaultSplitFloatSetAxisExplicitlyDontKeepDims3Dim"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_axis_scalar_split_rank3_input_input" }
+          }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_axis0_keepdims0_rank2",
+      "attrs": { "axis": 0, "keepdims": 0 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_Axis0DefaultSplitFloatSetAxisExplicitlyDontKeepDims2Dim"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_default_axis0_four_outputs",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_Axis0DefaultSplitFloatSetAxisExplicitly",
+        "notes": "Valid ONNX sequence output modelled as four fixed tensors."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": {
+        "y0": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [1.0, 2.0] }
+        },
+        "y1": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [3.0, 4.0] }
+        },
+        "y2": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [5.0, 6.0] }
+        },
+        "y3": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [7.0, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_bool_split_axis0_four_outputs",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_BoolSplit",
+        "notes": "Valid ONNX bool sequence output modelled as four fixed tensors."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [4, 2], "data": { "kind": "values", "values": [1, 1, 1, 1, 0, 0, 0, 0] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "bool", "shape": [1, 2], "tolerance": 0, "data": { "kind": "values", "values": [1, 1] } },
+        "y1": { "dtype": "bool", "shape": [1, 2], "tolerance": 0, "data": { "kind": "values", "values": [1, 1] } },
+        "y2": { "dtype": "bool", "shape": [1, 2], "tolerance": 0, "data": { "kind": "values", "values": [0, 0] } },
+        "y3": { "dtype": "bool", "shape": [1, 2], "tolerance": 0, "data": { "kind": "values", "values": [0, 0] } }
+      }
+    },
+    {
+      "name": "ort_positive_axis_keepdims0_four_outputs",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_PositiveAxisDontKeepDims",
+        "notes": "Valid ONNX keepdims=0 split along a non-leading axis, modelled as four fixed tensors."
+      },
+      "attrs": { "axis": 2, "keepdims": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_axis_scalar_split_rank3_input_input" }
+          }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 },
+        "y3": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_default_axis0_uneven_six_outputs",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/sequence/sequence_ops_test.cc",
+        "test": "SequenceOpsTest.SplitToSequence_DefaultAxis0UnevenSplitFloat2",
+        "notes": "Valid ONNX uneven scalar split producing six sequence elements."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [17, 2],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0]
+          }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y3": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y4": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y5": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "onnx_backend_split_scalar_axis1_three_outputs",
+      "attrs": { "axis": 1 },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_split_to_sequence_1",
+        "notes": "Official sequence output is represented as fixed tensor outputs y0/y1/y2; int64 split scalar is stored as uint32."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 6],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_split_scalar_axis1_three_outputs_input_input" }
+          }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 },
+        "y2": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "onnx_backend_split_vector_axis0_two_outputs",
+      "attrs": { "axis": 0 },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_split_to_sequence_2",
+        "notes": "Official sequence output is represented as fixed tensor outputs y0/y1; int64 split vector is stored as uint32."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 6],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_split_scalar_axis1_three_outputs_input_input" }
+          }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [1, 6], "tolerance": 0.000001 },
+        "y1": { "dtype": "float32", "shape": [2, 6], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [6, 0], "data": { "kind": "values", "values": [] } },
+        "split": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 },
+        "y2": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "two_outputs_zero_dim",
+      "provenance": {
+        "notes": "Two-chunk split of a tensor with a zero-width trailing axis. The two-output kernel floors each chunk stride at one so a zero-width output cannot divide by zero; the existing zero-dim case splits into three chunks, which is a different kernel."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4, 0], "data": { "kind": "values", "values": [] } },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "two_outputs_zero_dim_squeeze",
+      "provenance": {
+        "notes": "Squeezed twin of the case above: keepdims=0 drops the split axis, which is a separate coordinate walk with its own chunk-stride floor. The zero-width trailing axis makes that floor fire."
+      },
+      "attrs": { "axis": 0, "keepdims": 0 },
+      "inputs": { "input": { "dtype": "float32", "shape": [2, 3, 0], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [3, 0], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [3, 0], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "ort_caseB_empty",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.15191493928432465, 0.14055313169956207, 0.12134978920221329, 0.10187792778015137, 0.08903708308935165, 0.08793208003044128, 0.1010342463850975, 0.1277448534965515, 0.1644245833158493, 0.20488715171813965, 0.24129025638103485, 0.2653009593486786]
+          }
+        },
+        "split": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "y0": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "data": { "kind": "values", "values": [0.15191493928432465, 0.14055313169956207] },
+          "tolerance": 0.001
+        },
+        "y1": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "data": { "kind": "values", "values": [0.12134978920221329, 0.10187792778015137] },
+          "tolerance": 0.001
+        },
+        "y2": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "data": { "kind": "values", "values": [0.08903708308935165, 0.08793208003044128] },
+          "tolerance": 0.001
+        }
+      }
+    },
+    {
+      "name": "keepdims0_axis2_three_outputs_squeeze_f32",
+      "attrs": { "axis": 2, "keepdims": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_split_scalar_axis1_three_outputs_input_input" }
+          }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 3], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [2, 3], "tolerance": 0 },
+        "y2": { "dtype": "float32", "shape": [2, 3], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "negative_axis_neg1_scalar_split_rank3_f32",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 6],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_axis_scalar_split_rank3_input_input" }
+          }
+        },
+        "split": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0 },
+        "y2": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "negative_axis_neg2_keepdims0_three_outputs_squeeze_f32",
+      "attrs": { "axis": -2, "keepdims": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_positive_axis_scalar_split_rank3_input_input" }
+          }
+        }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 },
+        "y2": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rank6_middle_axis_two_outputs_deep_strides_f32",
+      "attrs": { "axis": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 4, 2, 2, 3],
+          "data": { "kind": "linspace", "start": 0.0, "end": 1.0 }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [2, 2, 3, 2, 2, 3], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [2, 2, 1, 2, 2, 3], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rank7_two_outputs",
+      "attrs": { "axis": 5, "keepdims": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 2, 3],
+          "data": { "kind": "linspace", "start": 0.0, "end": 5.0 }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 1, 3], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 1, 3], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rank8_two_outputs",
+      "attrs": { "axis": 6, "keepdims": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 1, 2, 3],
+          "data": { "kind": "linspace", "start": 1.0, "end": 6.0 }
+        },
+        "split": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": {
+        "y0": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 1, 1, 3], "tolerance": 0 },
+        "y1": { "dtype": "float32", "shape": [1, 1, 1, 1, 1, 1, 1, 3], "tolerance": 0 }
+      }
+    }
+  ]
+}