sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +75 -0
build/webgpu/bench.json +126 -0
build/webgpu/datamove-slice-block.wgsl.jinja +135 -0
build/webgpu/manifest.json +409 -0
build/webgpu/metadata.json +20 -0
build/webgpu/slice-rank2-single-axis-x4.wgsl.jinja +63 -0
build/webgpu/slice.wgsl.jinja +89 -0
build/webgpu/test.json +1547 -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.Slice
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Produces a slice of the input tensor along multiple axes, using `starts`, `ends`, `axes`, and `steps` to select a sub-tensor. Negative indices are resolved relative to the dimension size, and out-of-range values are clamped. Omitting `axes` defaults to all axes in order; omitting `steps` defaults to stride 1.
+See the [ONNX `Slice` spec](https://onnx.ai/onnx/operators/onnx__Slice.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `data` | `data` | `T` | — | — | Tensor of data to extract slices from. | required |
+| `starts` | `starts` | `S` | `1` | — | 1-D tensor of starting indices for each axis in `axes`. | required |
+| `ends` | `ends` | `S` | `1` | — | 1-D tensor of ending indices (exclusive) for each axis in `axes`. | required |
+| `axes` | `axes` | `S` | `1` | — | Optional 1-D tensor of axes that `starts` and `ends` apply to; defaults to all axes if omitted. | optional |
+| `steps` | `steps` | `S` | `1` | — | Optional 1-D tensor of step sizes per axis; negative steps slice backward, defaults to 1. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | same as `data` | — | Sliced data tensor. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
+| `S` | `int32` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`datamove-slice-block.wgsl.jinja`](build/webgpu/datamove-slice-block.wgsl.jinja)
+- [`slice-rank2-single-axis-x4.wgsl.jinja`](build/webgpu/slice-rank2-single-axis-x4.wgsl.jinja)
+- [`slice.wgsl.jinja`](build/webgpu/slice.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:
+- `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.Slice", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({
+  data: { data: dataData, shape: [2, 2] },
+  starts: { data: startsData, shape: [1] },
+  ends: { data: endsData, shape: [1] },
+}, {
+  outputs: { output: { shape: [1, 2], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,126 @@

+{
+  "op": "ai.onnx.Slice",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "rank3_sequence_window",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 128, 768] },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 32 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 96 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 64, 768] } }
+    },
+    {
+      "name": "rank3_end_past_dim",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 128, 768] },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 120 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 999 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 768] } }
+    },
+    {
+      "name": "slice-rows-aligned-f32-4096x4096",
+      "preset": "smoke",
+      "vars": { "outRows": 3584, "cols": 4096 },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 4096], "dist": "normal", "seed": 131, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 256 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 3840 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 0 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3584, 4096] } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "args.outRows * args.cols * 4 * 2" }] }
+    },
+    {
+      "name": "slice-inner-misaligned-f32-4096x4100",
+      "preset": "smoke",
+      "vars": { "rows": 4096, "outCols": 4096 },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 4100], "dist": "normal", "seed": 132, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 2 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 4098 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 4096] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.outCols * 4 * 2" }] }
+    },
+    {
+      "name": "slice-rows-aligned-f16-4096x2048",
+      "preset": "edge",
+      "vars": { "outRows": 3072, "cols": 2048, "dtype": "float16" },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [4096, 2048], "dist": "normal", "seed": 133, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 512 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 3584 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 0 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3072, 2048] } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.outRows * args.cols * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "slice-inner-true-misaligned-f32-4096x4097",
+      "preset": "edge",
+      "vars": { "rows": 4096, "outCols": 4097 },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [8192, 4097], "dist": "normal", "seed": 141, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 2048 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 6144 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 0 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 4097] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.outCols * 4 * 2" }] }
+    },
+    {
+      "name": "slice-no-axes-no-steps-f32-4096x4096",
+      "preset": "edge",
+      "vars": { "outRows": 3584, "cols": 4096 },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 4096], "dist": "normal", "seed": 142, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 256 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 3840 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3584, 4096] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.outRows * args.cols * 4 * 2" }] }
+    },
+    {
+      "name": "slice-inner-vec4-unaligned-reverse-f32-4096x4096",
+      "preset": "stress",
+      "vars": { "rows": 4096, "cols": 4096 },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4096, 4096], "dist": "normal", "seed": 151, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": -1 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": -2147483648 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": -1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 4096], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.rows * args.cols * 4 * 2" }] }
+    },
+    {
+      "name": "slice-rows-aligned-f32-2048x2048",
+      "preset": "smoke",
+      "vars": { "outRows": 1792, "cols": 2048 },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2048, 2048], "dist": "normal", "seed": 134, "scale": 0.5 },
+        "starts": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 128 },
+        "ends": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1920 },
+        "axes": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 0 },
+        "steps": { "dtype": "int32", "shape": [1], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1792, 2048] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.outRows * args.cols * 4 * 2" }] }
+    }
+  ]
+}

build/webgpu/datamove-slice-block.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,135 @@

+// Slice with vec4 loads and stores over the innermost axis.
+//
+// The innermost input and output dimensions are multiples of four, and the
+// output is non-empty. Both buffers can therefore bind as vec4<T>, with every
+// output vector contained within one output row.
+//
+// starts, axes, and steps are runtime tensors, so the GPU resolves innermost-
+// axis slicing once per thread. A unit step with a 16-byte-aligned start uses
+// one vec4 load per vec4 store. Strided, reversed, or misaligned slices instead
+// assemble four components through the vec4 binding and still issue one vec4
+// store. Outer-axis offsets use literal strides and remain multiples of four.
+// The normalization and clamping below preserve out-of-bounds starts and
+// negative steps.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ source.count }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const RANK: u32 = {{ source.rank }}u;
+const SLICE_RANK: u32 = {{ source.sliceRank }}u;
+const INNER_AXIS: u32 = {{ source.rank - 1 }}u;
+const INNER_DIM_DATA: i32 = {{ source.dataShape[source.rank - 1] }};
+const INNER_VECS_OUT: u32 = {{ ((source.outputShape[source.rank - 1] / 4) | int) }}u;
+{% if source.rank != 1 %}
+const DATA_SHAPE: array<i32, {{ source.rank }}> = array<i32, {{ source.rank }}>({% for d in source.dataShape %}{{ d }}{% if not loop.last %}, {% endif %}{% endfor %});
+{% endif %}
+fn normalized_axis(raw_axis: i32) -> u32 {
+  var axis = raw_axis;
+  if (axis < 0) {
+    axis = axis + i32(RANK);
+  }
+  return u32(axis);
+}
+fn normalized_start(raw_start: i32, step: i32, dim: i32) -> i32 {
+  var start = raw_start;
+  if (start < 0) {
+    start = start + dim;
+  }
+  if (step < 0) {
+    return clamp(start, -1, dim - 1);
+  }
+  return clamp(start, 0, dim);
+}
+{% if source.rank != 1 %}
+fn data_coord(axis: u32, out_coord: u32) -> u32 {
+{% if source.sliceRank == 1 %}
+  if (normalized_axis(i32(axes[0])) == axis) {
+    let step = i32(steps[0]);
+    let start = normalized_start(i32(starts[0]), step, DATA_SHAPE[axis]);
+    return u32(start + i32(out_coord) * step);
+  }
+  return out_coord;
+{% else %}
+  var coord = i32(out_coord);
+  for (var i = 0u; i < SLICE_RANK; i = i + 1u) {
+    if (normalized_axis(i32(axes[i])) == axis) {
+      let step = i32(steps[i]);
+      let start = normalized_start(i32(starts[i]), step, DATA_SHAPE[axis]);
+      coord = start + i32(out_coord) * step;
+    }
+  }
+  return u32(coord);
+{% endif %}
+}
+{% endif %}
+// Data offset (in scalars) contributed by the outer axes; always % 4 == 0
+// because every outer data stride is a multiple of the innermost data dim.
+fn outer_data_offset({% if source.rank != 1 %}outer_index: u32{% endif %}) -> u32 {
+{% if source.rank == 1 %}
+  return 0u;
+{% else %}
+  var rem = outer_index;
+  var offset = 0u;
+{% for axis in range(source.rank - 1) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank - 1) %}
+{% set out_stride.value = out_stride.value * source.outputShape[j] %}
+{% endfor %}
+{% set safe_out_stride = out_stride.value %}
+{% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank) %}
+{% set data_stride.value = data_stride.value * source.dataShape[j] %}
+{% endfor %}
+  let out_coord{{ axis }} = rem / {{ safe_out_stride }}u;
+  rem = rem % {{ safe_out_stride }}u;
+  offset = offset + data_coord({{ axis }}u, out_coord{{ axis }}) * {{ data_stride.value }}u;
+{% endfor %}
+  return offset;
+{% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  var inner_step: i32 = 1;
+  var inner_start: i32 = 0;
+  for (var s = 0u; s < SLICE_RANK; s = s + 1u) {
+    if (normalized_axis(i32(axes[s])) == INNER_AXIS) {
+      inner_step = i32(steps[s]);
+      inner_start = normalized_start(i32(starts[s]), inner_step, INNER_DIM_DATA);
+    }
+  }
+  let aligned = inner_step == 1 && (inner_start % 4) == 0;
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < COUNT; i += stride) {
+{% if source.rank != 1 %}
+    let outer_idx = i / INNER_VECS_OUT;
+{% endif %}
+    let inner_vec = i % INNER_VECS_OUT;
+    let base = outer_data_offset({% if source.rank != 1 %}outer_idx{% endif %});
+    if (aligned) {
+      output[i] = data[base / 4u + u32(inner_start) / 4u + inner_vec];
+    } else {
+      let c0 = i32(inner_vec) * 4;
+      let d0 = base + u32(inner_start + c0 * inner_step);
+      let d1 = base + u32(inner_start + (c0 + 1) * inner_step);
+      let d2 = base + u32(inner_start + (c0 + 2) * inner_step);
+      let d3 = base + u32(inner_start + (c0 + 3) * inner_step);
+      output[i] = {{ vectorScalar }}(
+        data[d0 / 4u][d0 % 4u],
+        data[d1 / 4u][d1 % 4u],
+        data[d2 / 4u][d2 % 4u],
+        data[d3 / 4u][d3 % 4u]
+      );
+    }
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,409 @@

+{
+  "domain": "ai.onnx",
+  "name": "Slice",
+  "sinceVersion": 13,
+  "description": "Produces a slice of the input tensor along multiple axes, using `starts`, `ends`, `axes`, and `steps` to select a sub-tensor. Negative indices are resolved relative to the dimension size, and out-of-range values are clamped. Omitting `axes` defaults to all axes in order; omitting `steps` defaults to stride 1.",
+  "inputs": [
+    { "role": "data", "dtype": "T", "description": "Tensor of data to extract slices from." },
+    {
+      "role": "starts",
+      "dtype": "S",
+      "description": "1-D tensor of starting indices for each axis in `axes`.",
+      "rank": 1
+    },
+    {
+      "role": "ends",
+      "dtype": "S",
+      "description": "1-D tensor of ending indices (exclusive) for each axis in `axes`.",
+      "rank": 1
+    },
+    {
+      "role": "axes",
+      "dtype": "S",
+      "optional": true,
+      "description": "Optional 1-D tensor of axes that `starts` and `ends` apply to; defaults to all axes if omitted.",
+      "rank": 1
+    },
+    {
+      "role": "steps",
+      "dtype": "S",
+      "optional": true,
+      "description": "Optional 1-D tensor of step sizes per axis; negative steps slice backward, defaults to 1.",
+      "rank": 1
+    }
+  ],
+  "outputs": [{ "role": "output", "dtype": "T", "description": "Sliced data tensor.", "rank": "ranks.data" }],
+  "typeConstraints": {
+    "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"],
+    "S": ["int32"]
+  },
+  "args": {
+    "data": { "kind": "tensor", "semantic": "data", "role": "input" },
+    "starts": { "kind": "tensor", "semantic": "starts", "role": "starts" },
+    "ends": { "kind": "tensor", "semantic": "ends", "role": "ends" },
+    "axes": { "kind": "tensor", "semantic": "axes", "role": "axes", "required": false },
+    "steps": { "kind": "tensor", "semantic": "steps", "role": "steps", "required": false },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "constants": { "indexScalar": "dtypes.S", "usesF16": "dtypes.T == \"f16\"" },
+  "variants": [
+    {
+      "id": "rank2_single_axis_x4",
+      "priority": 14,
+      "when": ["present.axes", "present.steps", "ranks.data == 2", "ranks.output == 2", "ranks.starts == 1", "ranks.ends == 1", "ranks.axes == 1", "ranks.steps == 1", "dim(shapes.starts, 0) == 1", "dim(shapes.ends, 0) == 1", "dim(shapes.axes, 0) == 1", "dim(shapes.steps, 0) == 1", "dim(shapes.data, 1) == dim(shapes.output, 1)", "dim(shapes.data, 0) != dim(shapes.output, 0)", "dim(shapes.data, 1) % 4 != 0", "numel(shapes.output) > 0", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Slice.rank2SingleAxisX4",
+          "source": {
+            "shader": "slice-rank2-single-axis-x4.wgsl.jinja",
+            "inputs": { "dataShape": "shapes.data", "outputShape": "shapes.output", "count": "numel(shapes.output)" }
+          },
+          "bindings": [
+            {
+              "name": "data",
+              "arg": "data",
+              "semantic": "data",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "starts",
+              "arg": "starts",
+              "semantic": "starts",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar",
+              "length": 1
+            },
+            {
+              "name": "axes",
+              "arg": "axes",
+              "semantic": "axes",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar",
+              "length": 1
+            },
+            {
+              "name": "steps",
+              "arg": "steps",
+              "semantic": "steps",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar",
+              "length": 1
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            }
+          ],
+          "dispatch": { "gridStride": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "inner_vec4",
+      "priority": 12,
+      "when": ["present.axes", "present.steps", "ranks.data >= 1", "ranks.output == ranks.data", "ranks.starts == 1", "ranks.ends == 1", "ranks.axes == 1", "ranks.steps == 1", "dim(shapes.starts, 0) == dim(shapes.ends, 0)", "dim(shapes.starts, 0) == dim(shapes.axes, 0)", "dim(shapes.starts, 0) == dim(shapes.steps, 0)", "dim(shapes.starts, 0) <= ranks.data", "numel(shapes.output) > 0", "dim(shapes.data, ranks.data - 1) % 4 == 0", "dim(shapes.output, ranks.data - 1) % 4 == 0", "f16Ok(dtypes.T)"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Slice.innerVec4",
+          "source": {
+            "shader": "datamove-slice-block.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.output",
+              "rank": "ranks.data",
+              "sliceRank": "dim(shapes.starts, 0)",
+              "count": "numel(shapes.output) / 4"
+            }
+          },
+          "bindings": [
+            {
+              "name": "data",
+              "arg": "data",
+              "semantic": "data",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "starts",
+              "arg": "starts",
+              "semantic": "starts",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "axes",
+              "arg": "axes",
+              "semantic": "axes",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "steps",
+              "arg": "steps",
+              "semantic": "steps",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "generic_metadata",
+      "when": ["present.axes", "present.steps", "ranks.data >= 1", "ranks.output == ranks.data", "ranks.starts == 1", "ranks.ends == 1", "ranks.axes == 1", "ranks.steps == 1", "dim(shapes.starts, 0) == dim(shapes.ends, 0)", "dim(shapes.starts, 0) == dim(shapes.axes, 0)", "dim(shapes.starts, 0) == dim(shapes.steps, 0)", "dim(shapes.starts, 0) <= ranks.data", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T", "presentAxes": true, "presentSteps": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Slice",
+          "source": {
+            "shader": "slice.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.output",
+              "rank": "ranks.data",
+              "sliceRank": "dim(shapes.starts, 0)"
+            }
+          },
+          "bindings": [
+            {
+              "name": "data",
+              "arg": "data",
+              "semantic": "data",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "starts",
+              "arg": "starts",
+              "semantic": "starts",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "axes",
+              "arg": "axes",
+              "semantic": "axes",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "steps",
+              "arg": "steps",
+              "semantic": "steps",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
+              }
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "generic_metadata_no_steps",
+      "when": ["present.axes", "not present.steps", "ranks.data >= 1", "ranks.output == ranks.data", "ranks.starts == 1", "ranks.ends == 1", "ranks.axes == 1", "dim(shapes.starts, 0) == dim(shapes.ends, 0)", "dim(shapes.starts, 0) == dim(shapes.axes, 0)", "dim(shapes.starts, 0) <= ranks.data", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T", "presentAxes": true, "presentSteps": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Slice",
+          "source": {
+            "shader": "slice.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.output",
+              "rank": "ranks.data",
+              "sliceRank": "dim(shapes.starts, 0)"
+            }
+          },
+          "bindings": [
+            {
+              "name": "data",
+              "arg": "data",
+              "semantic": "data",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "starts",
+              "arg": "starts",
+              "semantic": "starts",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "axes",
+              "arg": "axes",
+              "semantic": "axes",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
+              }
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "generic_metadata_no_axes",
+      "when": ["not present.axes", "not present.steps", "ranks.data >= 1", "ranks.output == ranks.data", "ranks.starts == 1", "ranks.ends == 1", "dim(shapes.starts, 0) == dim(shapes.ends, 0)", "dim(shapes.starts, 0) <= ranks.data", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T", "presentAxes": false, "presentSteps": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Slice",
+          "source": {
+            "shader": "slice.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.output",
+              "rank": "ranks.data",
+              "sliceRank": "dim(shapes.starts, 0)"
+            }
+          },
+          "bindings": [
+            {
+              "name": "data",
+              "arg": "data",
+              "semantic": "data",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "starts",
+              "arg": "starts",
+              "semantic": "starts",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
+              }
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "generic_metadata_no_axes_steps",
+      "when": ["not present.axes", "present.steps", "ranks.data >= 1", "ranks.output == ranks.data", "ranks.starts == 1", "ranks.ends == 1", "ranks.steps == 1", "dim(shapes.starts, 0) == dim(shapes.ends, 0)", "dim(shapes.starts, 0) == dim(shapes.steps, 0)", "dim(shapes.starts, 0) <= ranks.data", "f16Ok(dtypes.T)"],
+      "constants": { "scalar": "dtypes.T", "presentAxes": false, "presentSteps": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Slice",
+          "source": {
+            "shader": "slice.wgsl.jinja",
+            "inputs": {
+              "dataShape": "shapes.data",
+              "outputShape": "shapes.output",
+              "rank": "ranks.data",
+              "sliceRank": "dim(shapes.starts, 0)"
+            }
+          },
+          "bindings": [
+            {
+              "name": "data",
+              "arg": "data",
+              "semantic": "data",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "starts",
+              "arg": "starts",
+              "semantic": "starts",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "steps",
+              "arg": "steps",
+              "semantic": "steps",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output)" }]
+              }
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Slice",
+  "id": "_ai_onnx_slice_webgpu_c87b8e2",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "x9yDnzXstMYgGKR+oNwCljthYvd8NHQrJare/so3wFg=",
+      "datamove-slice-block.wgsl.jinja": "xBxRnx+M6qii0SdRIaFx6S1aOpIzUG75wWmlsxxOtRs=",
+      "manifest.json": "E5q+nglUrb3bELWEV12AAnrOKoketNHOHOsiBMV97rs=",
+      "slice-rank2-single-axis-x4.wgsl.jinja": "ykb8cr+AYZNBrLTdyR59y8vyg2lJzB//Kzss0/73744=",
+      "slice.wgsl.jinja": "EPEvvrs04aKI7ZhxJSAie3FC3CFl3iUWBNrNYTd+6YI=",
+      "test.json": "xE6Be8yfOVdyB6LcHzbHD4cUYMc6g45IxpYhTUJWJVY="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Slice" }
+}

build/webgpu/slice-rank2-single-axis-x4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,63 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ source.count }}u;
+const DATA_ROWS: u32 = {{ source.dataShape[0] }}u;
+const DATA_COLS: u32 = {{ source.dataShape[1] }}u;
+const OUT_COLS: u32 = {{ source.outputShape[1] }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+fn normalize_start(raw: i32, step: i32, dim: i32) -> i32 {
+  var start = raw;
+  if (start < 0) { start += dim; }
+  if (step < 0) { return clamp(start, -1, dim - 1); }
+  return clamp(start, 0, dim);
+}
+fn source_offset(out_index: u32, axis: u32, start: i32, step: i32) -> u32 {
+  let row = out_index / OUT_COLS;
+  let col = out_index - row * OUT_COLS;
+  if (axis == 0u) {
+    return u32(start + i32(row) * step) * DATA_COLS + col;
+  }
+  return row * DATA_COLS + u32(start + i32(col) * step);
+}
+// Four adjacent scalar elements per lane. For the common row-slice case, almost
+// every block is one contiguous input/output transaction; only the three values
+// around an odd-width row boundary take the fully generic per-element mapping.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let stride = nwg.x * WG;
+  var block = gid.x;
+  while (block * 4u < COUNT) {
+    let out0 = block * 4u;
+    let axisRaw = i32(axes[0]);
+    let axis = u32(select(axisRaw, axisRaw + 2i, axisRaw < 0));
+    let step = i32(steps[0]);
+    let dim = select(i32(DATA_COLS), i32(DATA_ROWS), axis == 0u);
+    let start = normalize_start(i32(starts[0]), step, dim);
+    let col = out0 % OUT_COLS;
+    if (axis == 0u && step == 1i && col + 3u < OUT_COLS && out0 + 3u < COUNT) {
+      let src0 = source_offset(out0, axis, start, step);
+      output[out0] = data[src0];
+      output[out0 + 1u] = data[src0 + 1u];
+      output[out0 + 2u] = data[src0 + 2u];
+      output[out0 + 3u] = data[src0 + 3u];
+    } else {
+      for (var lane = 0u; lane < 4u; lane++) {
+        let out_index = out0 + lane;
+        if (out_index < COUNT) {
+          output[out_index] = data[source_offset(out_index, axis, start, step)];
+        }
+      }
+    }
+    block += stride;
+  }
+}

build/webgpu/slice.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,89 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if presentAxes %}
+const RANK: u32 = {{ source.rank }}u;
+{% endif %}
+const SLICE_RANK: u32 = {{ source.sliceRank }}u;
+const DATA_SHAPE: array<i32, {{ source.rank }}> = array<i32, {{ source.rank }}>({% for d in source.dataShape %}{{ d }}{% if not loop.last %}, {% endif %}{% endfor %});
+{% if presentAxes %}
+fn normalized_axis(raw_axis: i32) -> u32 {
+  var axis = raw_axis;
+  if (axis < 0) {
+    axis = axis + i32(RANK);
+  }
+  return u32(axis);
+}
+{% endif %}
+fn normalized_start(raw_start: i32, step: i32, dim: i32) -> i32 {
+  var start = raw_start;
+  if (start < 0) {
+    start = start + dim;
+  }
+  if (step < 0) {
+    return clamp(start, -1, dim - 1);
+  }
+  return clamp(start, 0, dim);
+}
+fn data_coord(axis: u32, out_coord: u32) -> u32 {
+  var coord = i32(out_coord);
+  for (var i = 0u; i < SLICE_RANK; i = i + 1u) {
+{% if presentAxes %}
+    let slice_axis = normalized_axis(i32(axes[i]));
+{% else %}
+    // axes omitted -> defaults to [0, 1, ..., sliceRank-1].
+    let slice_axis = i;
+{% endif %}
+    if (slice_axis == axis) {
+{% if presentSteps %}
+      let step = i32(steps[i]);
+{% else %}
+      // steps omitted -> defaults to 1.
+      let step = 1i;
+{% endif %}
+      let start = normalized_start(i32(starts[i]), step, DATA_SHAPE[axis]);
+      coord = start + i32(out_coord) * step;
+    }
+  }
+  return u32(coord);
+}
+fn data_offset(out_index: u32) -> u32 {
+  var rem = out_index;
+  var offset = 0u;
+{% for axis in range(source.rank) %}
+{% set out_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank) %}
+{% set out_stride.value = out_stride.value * source.outputShape[j] %}
+{% endfor %}
+{% set safe_out_stride = 1 if out_stride.value == 0 else out_stride.value %}
+{% set data_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.rank) %}
+{% set data_stride.value = data_stride.value * source.dataShape[j] %}
+{% endfor %}
+  let out_coord{{ axis }} = rem / {{ safe_out_stride }}u;
+  rem = rem % {{ safe_out_stride }}u;
+  offset = offset + data_coord({{ axis }}u, out_coord{{ axis }}) * {{ data_stride.value }}u;
+{% endfor %}
+  return offset;
+}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+// Grid-stride over a dispatch clamped to the device limit, so large outputs
+// remain covered without an invalid one-thread-per-element launch.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < params.count; i += stride) {
+    output[i] = data[data_offset(i)];
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1547 @@

+{
+  "op": "ai.onnx.Slice",
+  "fixtureArrays": {
+    "onnx_backend_slice_input_data": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902, -0.6724604368209839, -0.35955315828323364, -0.8131462931632996, -1.7262825965881348, 0.17742614448070526, -0.4017809331417084, -1.630198359489441, 0.46278226375579834, -0.9072983860969543, 0.05194539576768875, 0.7290905714035034, 0.12898291647434235, 1.1394007205963135, -1.234825849533081, 0.4023416340351105, -0.6848101019859314, -0.8707971572875977, -0.5788496732711792, -0.3115525245666504, 0.056165341287851334, -1.1651498079299927, 0.9008265137672424, 0.4656624495983124, -1.5362436771392822, 1.4882521629333496, 1.895889163017273, 1.1787796020507812, -0.1799248307943344, -1.0707526206970215, 1.0544517040252686, -0.4031769335269928, 1.222445011138916, 0.2082749754190445, 0.9766390323638916, 0.3563663959503174, 0.7065731883049011, 0.01050002034753561, 1.7858705520629883, 0.12691208720207214, 0.4019893705844879, 1.8831506967544556, -1.3477590084075928, -1.2704850435256958, 0.969396710395813, -1.1731233596801758, 1.9436211585998535, -0.4136189818382263, -0.747454822063446, 1.922942042350769, 1.4805147647857666, 1.8675589561462402, 0.9060446619987488, -0.8612256646156311, 1.910064935684204, -0.26800337433815, 0.8024563789367676, 0.9472519755363464, -0.15501008927822113, 0.6140793561935425, 0.922206699848175, 0.37642553448677063, -1.0994007587432861, 0.29823818802833557, 1.3263858556747437, -0.694567859172821, -0.14963454008102417, -0.4351535439491272, 1.8492637872695923, 0.6722947359085083, 0.40746182203292847, -0.7699160575866699, 0.5392491817474365, -0.6743326783180237, 0.0318305566906929, -0.6358460783958435, 0.676433265209198, 0.5765908360481262, -0.20829875767230988, 0.39600670337677, -1.0930615663528442, -1.4912575483322144, 0.43939170241355896, 0.16667349636554718, 0.6350314617156982, 2.3831448554992676, 0.9444794654846191, -0.9128222465515137, 1.117016315460205, -1.31590735912323, -0.46158459782600403, -0.06824160367250443, 1.7133426666259766, -0.7447548508644104, -0.8264385461807251, -0.09845252335071564, -0.6634783148765564, 1.1266359090805054, -1.0799314975738525, -1.1474686861038208, -0.43782004714012146, -0.49803245067596436, 1.9295320510864258, 0.9494208097457886, 0.08755124360322952, -1.225435495376587, 0.8443629741668701, -1.0002152919769287, -1.5447710752487183, 1.1880297660827637, 0.31694260239601135, 0.9208588004112244, 0.31872764229774475, 0.8568305969238281, -0.6510255932807922, -1.034242868423462, 0.6815944910049438, -0.8034096360206604, -0.6895498037338257, -0.4555324912071228, 0.01747915893793106, -0.3539939224720001, -1.3749512434005737, -0.6436184048652649, -2.223403215408325, 0.6252314448356628, -1.602057695388794, -1.1043833494186401, 0.05216507986187935, -0.73956298828125, 1.543014645576477, -1.2928569316864014, 0.2670508623123169, -0.039282817393541336, -1.1680934429168701, 0.523276686668396, -0.1715463250875473, 0.7717905640602112, 0.8235041499137878, 2.163235902786255, 1.336527943611145, -0.3691818416118622, -0.2393791824579239, 1.0996595621109009, 0.6552637219429016, 0.6401315331459045, -1.6169559955596924, -0.024326125159859657, -0.7380309104919434, 0.279924601316452, -0.09815038740634918, 0.9101788997650146, 0.31721821427345276, 0.7863279581069946, -0.4664191007614136, -0.9444462656974792, -0.410049706697464, -0.017020413652062416, 0.37915173172950745, 2.2593090534210205, -0.0422571524977684, -0.9559450149536133, -0.34598177671432495, -0.463595986366272, 0.4814814627170563, -1.5407969951629639, 0.06326199322938919, 0.15650653839111328, 0.23218104243278503, -0.5973160862922668, -0.23792172968387604, -1.4240609407424927, -0.49331986904144287, -0.5428614616394043, 0.4160500466823578, -1.1561824083328247, 0.7811980843544006, 1.494484543800354, -2.0699849128723145, 0.42625874280929565, 0.676908016204834, -0.6374370455741882, -0.3972718119621277, -0.1328805834054947, -0.29779088497161865, -0.3090129792690277, -1.6760038137435913, 1.1523315906524658, 1.0796185731887817, -0.8133642673492432, -1.4664243459701538, 0.5210648775100708, -0.5757879614830017, 0.14195317029953003, -0.3193284273147583, 0.6915387511253357, 0.694749116897583, -0.7255973815917969, -1.383363962173462, -1.5829384326934814, 0.6103793978691101, -1.188859224319458, -0.5068163275718689, -0.596314013004303, -0.05256729573011398, -1.9362797737121582, 0.1887785941362381, 0.523891031742096, 0.08842208981513977, -0.3108861744403839, 0.09740016609430313, 0.3990463316440582, -2.772592782974243, 1.9559123516082764, 0.3900933265686035, -0.6524085998535156, -0.3909533619880676, 0.4937417805194855, -0.11610393971204758, -2.030684471130371, 2.06449294090271, -0.11054065823554993, 1.0201727151870728, -0.6920498609542847, 1.5363770723342896, 0.28634369373321533, 0.6088438630104065, -1.0452533960342407, 1.211145281791687, 0.6898181438446045, 1.3018462657928467, -0.62808758020401, -0.48102712631225586, 2.3039166927337646, -1.0600157976150513, -0.13594970107078552, 1.1368913650512695, 0.09772496670484543, 0.582953691482544, -0.39944902062416077, 0.3700558841228485, -1.3065268993377686, 1.6581306457519531, -0.1181640475988388, -0.6801782250404358, 0.6663830876350403, -0.4607197940349579, -1.3342584371566772, -1.3467174768447876, 0.6937731504440308, -0.15957343578338623, -0.13370156288146973, 1.0777437686920166, -1.1268258094787598, -0.7306777238845825, -0.38487979769706726, 0.09435158967971802, -0.042171452194452286, -0.28688719868659973, -0.06162640079855919, -0.10730527341365814, -0.7196043729782104, -0.8129929900169373, 0.27451634407043457, -0.8909150958061218, -1.1573553085327148, -0.3122922480106354, -0.15766701102256775, 2.256723403930664, -0.7047002911567688, 0.9432607293128967, 0.7471883296966553, -1.188944935798645, 0.7732529640197754, -1.1838806867599487, -2.659172296524048, 0.6063195466995239, -1.7558906078338623, 0.4509344696998596, -0.684010922908783, 1.6595507860183716, 1.0685093402862549, -0.4533858001232147, -0.6878376007080078, -1.214077353477478, -0.4409226179122925, -0.2803554832935333, -0.3646935522556305, 0.1567038595676422, 0.5785214900970459, 0.3496544659137726, -0.7641439437866211, -1.4377914667129517, 1.3645318746566772, -0.6894491910934448, -0.6522936224937439, -0.5211893320083618, -1.8430695533752441, -0.477973997592926, -0.47965580224990845, 0.6203582882881165, 0.6984571218490601, 0.003770889015868306, 0.9318483471870422, 0.33996498584747314, -0.01568211242556572, 0.1609281748533249, -0.1906534880399704, -0.3948495090007782, -0.26773354411125183, -1.1280113458633423, 0.2804417014122009, -0.9931235909461975, 0.841631293296814, -0.24945858120918274, 0.04949498176574707, 0.49383679032325745, 0.6433144807815552, -1.5706233978271484, -0.20690368115901947, 0.8801789283752441, -1.698105812072754, 0.3872804641723633, -2.2555642127990723, -1.022506833076477, 0.03863055258989334, -1.6567151546478271, -0.9855107665061951, -1.4718350172042847, 1.6481349468231201, 0.16422775387763977, 0.5672903060913086, -0.22267509996891022, -0.353431761264801, -1.6164741516113281, -0.2918373644351959, -0.7614921927452087, 0.8579239249229431, 1.1411018371582031, 1.466578722000122, 0.8525519371032715, -0.5986539125442505, -1.1158969402313232, 0.7666631937026978, 0.35629281401634216, -1.768538475036621, 0.3554818034172058, 0.8145198225975037, 0.05892558768391609, -0.18505367636680603, -0.8076484799385071, -1.4465347528457642, 0.8002979755401611, -0.3091144561767578, -0.23346665501594543, 1.732721209526062, 0.6845011115074158, 0.37082499265670776, 0.14206179976463318, 1.519994854927063, 1.719589352607727, 0.9295051097869873, 0.5822246074676514, -2.0946030616760254, 0.12372191250324249, -0.13010695576667786, 0.09395322948694229, 0.9430460929870605, -2.7396771907806396, -0.5693120360374451, 0.26990434527397156, -0.4668455421924591, -1.4169061183929443, 0.8689634799957275, 0.276871919631958, -0.9711045622825623, 0.3148171901702881, 0.8215857148170471, 0.005292646121233702, 0.8005648255348206, 0.07826017588376999, -0.3952289819717407, -1.1594204902648926, -0.08593076467514038, 0.19429293274879456, 0.875832736492157, -0.1151074692606926, 0.4574156105518341, -0.9646120071411133, -0.782629132270813, -0.11038929969072342, -1.0546284914016724, 0.8202478289604187, 0.46313032507896423, 0.27909576892852783, 0.33890411257743835, 2.021043539047241, -0.46886420249938965, -2.2014412879943848, 0.19930019974708557, -0.05060354247689247, -0.5175190567970276, -0.9788298606872559, -0.43918952345848083, 0.18133842945098877, -0.5028166770935059, 2.4124536514282227, -0.9605043530464172, -0.793117344379425, -2.2886199951171875, 0.2514844238758087, -2.016406536102295, -0.5394546389579773, -0.27567052841186523, -0.709727942943573, 1.7388726472854614, 0.9943943619728088, 1.3191368579864502, -0.8824188113212585, 1.1285940408706665, 0.4960009455680847, 0.7714059352874756, 1.029438853263855, -0.9087632298469543, -0.4243176281452179, 0.8625960350036621, -2.6556191444396973, 1.5133280754089355, 0.5531320571899414, -0.045703962445259094, 0.2205076515674591, -1.0299352407455444, -0.3499433696269989, 1.1002843379974365, 1.2980220317840576, 2.6962239742279053, -0.07392466813325882, -0.6585529446601868, -0.5142339468002319, -1.0180418491363525, -0.07785475254058838, 0.38273242115974426, -0.03424227982759476, 1.0963468551635742, -0.23421579599380493, -0.3474506437778473, -0.5812684893608093, -1.6326345205307007, -1.5677677392959595, -1.1791579723358154, 1.3014280796051025, 0.8952602744102478, 1.3749641180038452, -1.3322116136550903, -1.9686247110366821, -0.6600562930107117, 0.17581894993782043, 0.49869027733802795, 1.0479722023010254, 0.2842796742916107, 1.7426687479019165, -0.22260567545890808, -0.9130792021751404, -1.6812182664871216, -0.8889713287353516, 0.2421179562807083, -0.8887202739715576, 0.9367424845695496, 1.4123276472091675, -2.369586944580078, 0.8640522956848145, -2.2396039962768555, 0.40149906277656555, 1.2248705625534058, 0.06485610455274582, -1.2796891927719116, -0.5854312181472778, -0.26164543628692627, -0.18224477767944336, -0.2028968334197998, -0.10988277941942215, 0.21348005533218384, -1.2085736989974976, -0.24201983213424683, 1.518261194229126, -0.3846454322338104, -0.4438360929489136, 1.0781973600387573, -2.5591845512390137, 1.1813786029815674, -0.6319037675857544, 0.1639285683631897, 0.09632135927677155, 0.9424681067466736, -0.26759475469589233, -0.6780257821083069, 1.2978458404541016, -2.3641738891601562, 0.020334182307124138, -1.3479254245758057, -0.7615733742713928, 2.011256694793701, -0.04459542781114578, 0.19506970047950745, -1.7815628051757812, -0.7290446758270264, 0.1965574026107788, 0.354757696390152, 0.616886556148529, 0.00862789899110794, 0.5270041823387146, 0.4537819027900696, -1.8297404050827026, 0.037005722522735596, 0.7679024338722229, 0.5898798108100891, -0.36385881900787354, -0.8056265115737915, -1.118311882019043, -0.13105401396751404, 1.1330798864364624, -1.9518041610717773, -0.6598917245864868, -1.1398024559020996, 0.7849575281143188, -0.554309606552124, -0.47063764929771423, -0.21694956719875336, 0.44539326429367065, -0.3923889994621277, -3.046143054962158, 0.5433118939399719, 0.43904295563697815, -0.2195410281419754, -1.0840365886688232, 0.35178011655807495, 0.37923553586006165, -0.47003287076950073, -0.2167314738035202, -0.9301565289497375, -0.1785890907049179, -1.550429344177246, 0.41731882095336914, -0.9443684816360474, 0.23810315132141113, -1.4059629440307617, -0.5900576710700989, -0.11048940569162369, -1.6606998443603516, 0.11514787375926971, -0.37914755940437317, -1.7423561811447144, -1.303242802619934, 0.605120062828064, 0.8955559730529785, -0.13190864026546478, 0.40476182103157043, 0.2238435596227646, 0.3296229839324951, 1.2859840393066406, -1.5069984197616577, 0.6764607429504395, -0.38200896978378296, -0.2242589294910431, -0.3022497296333313, -0.37514710426330566, -1.2261961698532104, 0.1833391934633255, 1.670943021774292, -0.05613302066922188, -0.0013850427931174636, -0.6872990131378174, -0.11747454851865768, 0.4661664366722107, -0.3702424466609955, -0.45380404591560364, 0.4032645523548126, -0.9180047512054443, 0.2524966299533844, 0.820321798324585, 1.3599485158920288, -0.09038200974464417, 1.367597222328186, 1.0344098806381226, -0.9962126612663269, -1.2179385423660278, -0.3049636483192444, 1.0289355516433716, -0.07228700816631317, -0.60065758228302, 1.5522432327270508, 0.2869044840335846, -2.320594310760498, 0.31716063618659973, 0.5200406312942505, 0.22560866177082062, 0.4497120976448059, -0.06727560609579086, -1.3183958530426025, -0.3707039952278137, -0.9456157684326172, -0.9327409267425537, -1.2630683183670044, 0.45248907804489136, 0.09789614379405975, -0.4481653571128845, -0.6493379473686218, -0.023423105478286743, 1.0791947841644287, -2.004215717315674, 0.3768765330314636, -0.5457119941711426, -1.8845858573913574, -1.9457030296325684, -0.9127835035324097, 0.219509556889534, 0.39306291937828064, -0.9389815926551819, 1.0170209407806396, 1.4229835271835327, 0.39608657360076904, -0.5914026498794556, 1.1244192123413086, 0.7553957104682922, 0.8674074411392212, -0.6564636826515198, -2.834554433822632, 2.116791009902954, -1.610878348350525, -0.03576807305216789, 2.3807454109191895, 0.33057674765586853, 0.9492464661598206, -1.502396583557129, -1.7776669263839722, -0.5327028036117554, 1.090749740600586, -0.34624946117401123, -0.7946363091468811, 0.1979672908782959, 1.081935167312622, -1.444940209388733, -1.2105430364608765, -0.788669228553772, 1.0946383476257324, 0.23482152819633484, 2.1321535110473633, 0.9364457130432129, -0.035095177590847015, 1.265077829360962, 0.2114970088005066, -0.7049213647842407, 0.6799748539924622, -0.6963266730308533, -0.29039710760116577, 1.3277827501296997, -0.10128148645162582, -0.8031414151191711, -0.46433767676353455, 1.021790623664856, -0.55254065990448, -0.3868708610534668, -0.5102927684783936, 0.18392549455165863, -0.3854897618293762, -1.601836085319519, -0.8871809244155884, -0.9327890276908875, 1.2433193922042847, 0.8126740455627441, 0.5872593522071838, -0.5053583383560181, -0.8157915472984314, -0.5075175762176514, -1.051880121231079, 2.4972004890441895, -2.245321750640869, 0.5640085339546204, -1.2845523357391357, -0.1043434888124466, -0.9880019426345825, -1.177628993988037, -1.1401963233947754, 1.7549861669540405, -0.13298842310905457, -0.7657021880149841, 0.5557869672775269, 0.010349314659833908, 0.7200337648391724, -1.8242566585540771, 0.3036039173603058, 0.7726948261260986, -1.6615983247756958, 0.44819527864456177, 1.6961815357208252, -0.014857703819870949, 0.8214059472084045, 0.670570433139801, -0.7075057029724121, 0.03976673632860184, -1.5669946670532227, -0.45130303502082825, 0.2656879723072052, 0.723100483417511, 0.024612125009298325, 0.7199837565422058, -1.1029062271118164, -0.10169727355241776, 0.01927938498556614, 1.8495912551879883, -0.21416665613651276, -0.4990166425704956, 0.021351223811507225, -0.9191134572029114, 0.1927538514137268, -0.3650552034378052, -1.7913275957107544, -0.05858655273914337, -0.317543089389801, -1.6324232816696167, -0.06713415682315826, 1.4893559217453003, 0.5213037729263306, 0.6119272112846375, -1.3414967060089111, 0.4768983721733093, 0.1484495848417282, 0.5290452241897583, 0.42262861132621765, -1.3597806692123413, -0.041400812566280365, -0.7578708529472351, -0.05008409544825554, -0.8974009156227112, 1.3124703168869019, -0.8589723706245422, -0.8989421725273132, 0.0745864063501358, -1.077099084854126, -0.4246633052825928, -0.8299645781517029, 1.4111720323562622, 0.7858038544654846, -0.05746951699256897, -0.39121705293655396, 0.9409176111221313, 0.40520408749580383, 0.4980524182319641, -0.026192236691713333, -1.688230037689209, -0.1124659851193428, -0.5324898958206177, 0.6450552940368652, 1.0118424892425537, -0.6579510569572449, 0.4683852195739746, 1.7358789443969727, -0.6677127480506897, 1.6819217205047607, -0.8525858521461487, 0.0229597557336092, -0.011145612224936485, 0.011498900130391121, -0.8376780152320862, -0.5911831259727478, -0.6677202582359314, 0.32696259021759033, 0.33003512024879456, 2.2259442806243896, 1.3709889650344849, -0.5098432302474976, 0.3248696029186249, 0.9971179962158203, 0.030601823702454567, -0.06964157521724701, 0.0515749417245388, 0.8672766089439392, -0.8483205437660217, -0.32566946744918823, 0.47043314576148987, 0.3114470839500427, 0.23958276212215424, -0.3698011636734009, 0.9725357890129089, 2.1338682174682617, 0.40641549229621887, -0.19317670166492462, 0.7557402849197388, -0.5391326546669006, -0.7496903538703918, 0.03280874714255333, -2.582796573638916, -1.1539503335952759, -0.3479618430137634, -1.3533889055252075, -1.0326430797576904, -0.43674832582473755, -1.642965316772461, -0.4060717821121216, -0.5352701544761658, 0.02540520764887333, 1.1541839838027954, 0.17250441014766693, 0.02106202207505703, 0.09945445507764816, 0.22739277780056, -1.0167386531829834, -0.11477532237768173, 0.30875125527381897, -1.3707599639892578, 0.8656529188156128, 1.081376075744629, -0.6313759684562683, -0.24133779108524323, -0.8781903386116028, 0.6993804574012756, -1.0612223148345947, -0.2224770039319992, -0.8589199185371399, 0.05095427855849266, -1.79422926902771, 1.326461672782898, -0.9646064043045044, 0.059894684702157974, -0.21252304315567017, -0.7621145248413086, -0.8877801299095154, 0.9363985657691956, -0.525640606880188, 0.2711701989173889, -0.8014968633651733, -0.6471814513206482, 0.4722471535205841, 0.9304084777832031, -0.17531640827655792, -1.421919822692871, 1.9979560375213623, -0.8565493226051331, -1.5415873527526855, 2.5944244861602783, -0.40403228998184204, -1.4617327451705933, -0.6834397912025452, 0.36754488945007324, 0.19031155109405518, -0.8517292141914368, 1.822723627090454, -0.5215796828269958, -1.184686541557312, 0.9606934189796448, 1.329062819480896, -0.8174930810928345, -1.401347279548645, 1.0304383039474487, -2.047323703765869, -1.2266216278076172, 0.9674461483955383, -0.05535254627466202, -0.2639373540878296, 0.3528166115283966, -0.15277442336082458, -1.2986867427825928, 1.2760753631591797, 1.3250139951705933, 0.20533256232738495, 0.04513401538133621, 2.3396248817443848, -0.2764328420162201, -0.2595769762992859, 0.36448124051094055, 1.471321940422058, 1.5927706956863403, -0.25857263803482056, 0.3083312511444092, -1.3780834674835205, -0.3119761049747467, -0.840290367603302, -1.0068317651748657, 1.6815767288208008, -0.7922866344451904, -0.5316058993339539, 0.3658487796783447, 1.2978252172470093, 0.4811151325702667, 2.759355068206787, -0.07466797530651093, 0.2587164342403412, 0.2756006717681885, 1.4350494146347046, 0.5072389245033264, -0.11622969806194305, -0.9474886059761047, 0.24444346129894257, 1.4013447761535645, -0.4103817939758301, 0.5289435982704163, 0.2461477816104889, 0.8635196685791016, -0.8047537207603455, 2.346647024154663, -1.2791610956192017, -0.36555108428001404, 0.9380925297737122, 0.29673317074775696, 0.8299861550331116, -0.49610233306884766, -0.07480498403310776, 0.012231983244419098, 1.5692596435546875, 0.6904290318489075, 0.7966721057891846, -0.657926082611084, 0.9688826203346252, 0.22558166086673737, 1.3891453742980957, 2.0140602588653564, -0.30676576495170593, -0.40630313754081726, -0.8640449643135071, -0.14357951283454895, -0.38202545046806335, 0.3595044016838074, -0.14456681907176971, -0.36159926652908325, 1.0645850896835327, -0.9378802180290222, 0.4331079423427582, -0.4059417247772217, 0.7243685126304626, 1.3852615356445312, -0.30309826135635376, 0.44103291630744934, 0.17879286408424377, -0.7994223833084106, 0.24078750610351562, 0.28912049531936646, 0.4128708243370056, -0.1983989030122757, 0.09419230371713638, -1.1476109027862549, -0.3581140637397766],
+    "ort_2d_two_axes_negative_end_input_data": [0, 1, 2, 3, 10, 11, 12, 13, 20, 21, 22, 23, 30, 31, 32, 33, 40, 41, 42, 43, 50, 51, 52, 53],
+    "ort_3d_positive_and_negative_steps_subset_axes_input_data": [27, 20, 2, 4, 26, 11, 26, 5, 17, 0, 21, 6, 22, 13, 29, 19, 17, 27, 4, 20, 12, 3, 9, 24, 17, 6, 24],
+    "ort_5d_subset_axes_offset_input_data": [1, 2, 3, 4, 5, 6, 7, 8, -1, -2, -3, -4, -5, -6, -7, -8],
+    "inner_vec4_uint8_aligned_2d_input_data": [0, 1, 2, 3, 4, 5, 6, 7, 248, 249, 250, 251, 252, 253, 254, 255],
+    "inner_vec4_bool_aligned_2d_input_data": [1, 0, 1, 0, 1, 0, 1, 0, 0, 1, 0, 1, 0, 1, 0, 1]
+  },
+  "cases": [
+    {
+      "name": "int16_step2_boundaries",
+      "inputs": {
+        "data": {
+          "dtype": "int16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-32768, -1, 0, 1, 32767, -32767] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [6] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int16",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1, 1, -32767] }
+        }
+      }
+    },
+    {
+      "name": "rank2_axis1",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 10.0, 11.0, 12.0, 13.0, 14.0, 20.0, 21.0, 22.0, 23.0, 24.0]
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3] } }
+    },
+    {
+      "name": "rank2_axis0_odd_width_x4",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [6, 5], "data": { "kind": "linspace", "start": -1.0, "end": 1.0 } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "rank2_axis0_odd_width_reverse_x4",
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [6, 5], "data": { "kind": "linspace", "start": -1.0, "end": 1.0 } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "end_past_dim_clamped",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 5],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 10.0, 11.0, 12.0, 13.0, 14.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [999] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2] } }
+    },
+    {
+      "name": "empty_output_start_equals_end",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0] } }
+    },
+    {
+      "name": "rank3_axis_step_f16",
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [2, 3, 4] },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 1, 2] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "rank5_subset_axes_offset",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "starts": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 1, 0, 0] } },
+        "ends": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, 2, 2, 2] } },
+        "axes": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "steps": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, 1, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_last_axis_step",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_1d_start_greater_than_end_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_InvalidStartEndRange"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_1d_start_end_out_of_bounds_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_StartAndEndOutOfBounds"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1001] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_start_end_out_of_bounds_empty_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_StartAndEndOutOfBounds"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1000] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [10, 1000] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 0], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_1d_negative_step_regular",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithNegativeSteps_Regular"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_1d_negative_step_start_out_of_bounds",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithNegativeSteps_StartOutOfBounds"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [7] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_positive_steps_subset_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_WithPositiveSteps_1"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 0] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_two_axes_negative_end",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_TwoAxes"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_2d_two_axes_negative_end_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1000, -1] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 0] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_reverse_all_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_ReverseAllAxes"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [-1, -1] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [-2147483648, -2147483648] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [-1, -1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_reverse_negative_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_ReverseSubsetOfNegAxes_1"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_mixed_negative_steps_single_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_WithNegativeSteps_2"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [-1, -2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_3d_positive_and_negative_steps_subset_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice3D_WithPositiveAndNegativeSteps_SubsetOfAxes_1"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_3d_positive_and_negative_steps_subset_axes_input_data" }
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 4] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1000, 1] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [3, -2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [3, 1, 1],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [11, 29, 24] }
+        }
+      }
+    },
+    {
+      "name": "ort_3d_positive_and_negative_steps_subset_axes_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice3D_WithPositiveAndNegativeSteps_SubsetOfAxes_2"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_3d_positive_and_negative_steps_subset_axes_input_data" }
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 4] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1000, 2] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [3, -2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 1, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_5d_subset_axes_offset",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice5D_SubsetOfAxes_Flatten2Dims_OffsetInput"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 1, 1, 0] } },
+        "ends": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, 2, 2147483647, 6] } },
+        "axes": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "steps": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, 1, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 1, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_5d_large_step",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice5D_LargeStep"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2147483647] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_3d_flatten_innermost_dims_copy",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice3D_FlattenInnermostDimsIncopy"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_3d_positive_and_negative_steps_subset_axes_input_data" }
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 0, 1] } },
+        "ends": {
+          "dtype": "int32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [2147483647, 2147483647, 2147483647] }
+        },
+        "axes": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [2, 1, 0] } },
+        "steps": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [1, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_5d_copy_axis2_large_block",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice5D_CopyAxis2LargeBlock"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 3, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0]
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_coalesce_dims_negative_and_positive_steps",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.CoalesceDims / negative and positive steps"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 2] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [-1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_coalesce_dims_middle_axis_full_extent",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.CoalesceDims / middle axis full extent"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2147483647] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_coalesce_dims_two_subset_axes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.CoalesceDims / two subset axes"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2147483647, 2147483647] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 2, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_reverse_last_axis",
+      "inputs": {
+        "data": {
+          "dtype": "float16",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_empty_dim_positive_step",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc", "test": "SliceTest.EmptyDim" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [0, 6], "data": { "kind": "values", "values": [] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_empty_dim_negative_step",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc", "test": "SliceTest.EmptyDim" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [0, 6], "data": { "kind": "values", "values": [] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_1d_regular_int32_data",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_Int32"
+      },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_slice",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [3, 10] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 10, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice",
+        "notes": "ONNX int64 index values are adapted to supported int32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_slice_default_axes",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 0, 3] } },
+        "ends": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [20, 10, 4] } },
+        "axes": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 1, 2] } },
+        "steps": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [20, 10, 1] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice_default_axes",
+        "notes": "ONNX int64 index values use supported int32 storage where representable. Omitted axes and steps use their standard defaults; this also subsumes test_slice_default_steps."
+      }
+    },
+    {
+      "name": "onnx_backend_slice_end_out_of_bounds",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [20, 9, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice_end_out_of_bounds",
+        "notes": "ONNX int64 index values are adapted to supported int32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_slice_neg",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [20, 9, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice_neg",
+        "notes": "ONNX int64 index values are adapted to supported int32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_slice_neg_steps",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [20, 10, 4] } },
+        "ends": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 0, 1] } },
+        "axes": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 1, 2] } },
+        "steps": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-1, -3, -2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [19, 3, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice_neg_steps",
+        "notes": "ONNX int64 index values are adapted to supported int32 where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_slice_negative_axes",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 0, 3] } },
+        "ends": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [20, 10, 4] } },
+        "axes": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, -2, -1] } },
+        "steps": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [20, 10, 1] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice_negative_axes",
+        "notes": "ONNX int64 index values are adapted to supported int32 where representable. ONNX omitted steps input represented with default unit steps."
+      }
+    },
+    {
+      "name": "ort_negative_step_regular_1d",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithNegativeSteps_Regular"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_step_end_out_of_bounds_1d",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithNegativeSteps_EndOutOfBounds_2"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-10] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_step_valid_range_1d",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithNegativeSteps_ValidStartEndRange"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reverse_subset_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_ReverseSubsetOfAxes_1"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_positive_negative_steps_subset_axes_int32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice3D_WithPositiveAndNegativeSteps_SubsetOfAxes_1"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_3d_positive_and_negative_steps_subset_axes_input_data" }
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 4] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1000, 1] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [3, -2] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_5d_subset_axes_flatten_offset",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice5D_SubsetOfAxes_Flatten2Dims_OffsetInput"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 1, 1, 0] } },
+        "ends": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, 2, 2147483647, 6] } },
+        "axes": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "steps": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, 1, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_5d_large_step_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice5D_LargeStep"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_5d_subset_axes_offset_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2147483647] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_slice_start_out_of_bounds",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [20, 10, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_slice_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1000] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [20, 0, 5] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_slice_start_out_of_bounds",
+        "notes": "ONNX int64 index values are adapted to supported int32 where representable."
+      }
+    },
+    {
+      "name": "ort_1d_positive_step_every_other",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithPositiveSteps"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [6] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 2.0, 4.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_1d_reverse_all_axes_end_int32_max",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_ReverseAllAxes_1",
+        "notes": "Uses int32 max as the framework-representable sentinel for the ONNX int64 max end."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2147483647] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [4.0, 3.0, 2.0, 1.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_1d_reverse_all_axes_end_int32_min",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_ReverseAllAxes_2",
+        "notes": "Uses int32 min as the framework-representable sentinel for the ONNX int64 min end."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [4.0, 3.0, 2.0, 1.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_1d_reverse_all_axes_end_before_negative_dim",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_ReverseAllAxes_3"
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-5] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [4.0, 3.0, 2.0, 1.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_2d_reverse_subset_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_ReverseSubsetOfAxes_2",
+        "notes": "Uses int32 min as the framework-representable sentinel for the ONNX int64 min end."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [3.0, 4.0, 1.0, 2.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_2d_implicit_full_axis_copy",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_ImplicitCopyBySlicingADimensionFully",
+        "notes": "Uses int32 max as the framework-representable oversized positive end."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2147483647] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_int8_reverse_all_axes_end_int32_min",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_ReverseAllAxes_2",
+        "notes": "Reverse-slice sentinel case with logical int8 edge values."
+      },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, -1, 0, 127] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [127, 0, -1, -128] }
+        }
+      }
+    },
+    {
+      "name": "uint8_positive_step_every_other_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_WithPositiveSteps",
+        "notes": "Positive-step slice with logical uint8 edge values."
+      },
+      "inputs": {
+        "data": { "dtype": "uint8", "shape": [6], "data": { "kind": "values", "values": [0, 1, 127, 128, 254, 255] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [6] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 127, 254] }
+        }
+      }
+    },
+    {
+      "name": "inner_vec4_step1_aligned_2d_f32",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [6, 16] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 12], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_step1_misaligned_start_2d_f32",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_reverse_inner_2d_f32",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.11 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_outer_slice_rank3_f16",
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [3, 5, 8] },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 4, 8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "inner_vec4_inner_step2_f16",
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [2, 16] },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [16] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "inner_vec4_outer_slice_u32",
+      "inputs": {
+        "data": {
+          "dtype": "uint32",
+          "shape": [8, 4],
+          "data": {
+            "kind": "values",
+            "values": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31]
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint32",
+          "shape": [7, 4],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31]
+          }
+        }
+      }
+    },
+    {
+      "name": "inner_vec4_rank1_tail_f32",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [4.0, 5.0, 6.0, 7.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_1d_regular",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_Regular"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [2.0, 3.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_1d_end_out_of_bounds_full_copy",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice1D_EndOutOfBounds"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [6],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_2d_one_axis_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_OneAxis"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_2d_two_axes_negative_end_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [10.0, 11.0, 12.0, 13.0, 20.0, 21.0, 22.0, 23.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_2d_one_axis_default_steps",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_OneAxis",
+        "notes": "ORT omits the optional steps input here; ONNX defaults it to unit steps."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_2d_two_axes_negative_end_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [10.0, 11.0, 12.0, 13.0, 20.0, 21.0, 22.0, 23.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_2d_default_axes_and_steps",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_DefaultAxes",
+        "notes": "ORT omits optional axes and steps; ONNX defaults axes to [0] and steps to [1]."
+      },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1.0, 2.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_3d_clamped_multi_axis",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc", "test": "SliceTest.Slice3D" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [111.0, 112.0, 113.0, 121.0, 122.0, 123.0, 131.0, 132.0, 133.0, 211.0, 212.0, 213.0, 221.0, 222.0, 223.0, 231.0, 232.0, 233.0, 311.0, 312.0, 313.0, 321.0, 322.0, 323.0, 331.0, 332.0, 333.0]
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 1, 1] } },
+        "ends": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1000, 1000, 1000] } },
+        "axes": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, 1, 2] } },
+        "steps": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 1, 1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3, 2, 2],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [122.0, 123.0, 132.0, 133.0, 222.0, 223.0, 232.0, 233.0, 322.0, 323.0, 332.0, 333.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_bool_reverse_subset_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/slice_op.test.cc",
+        "test": "SliceTest.Slice2D_ReverseSubsetOfAxes_1",
+        "notes": "Same reverse-subset slice behavior with a bool payload."
+      },
+      "inputs": {
+        "data": { "dtype": "bool", "shape": [2, 2], "data": { "kind": "values", "values": [1, 0, 0, 1] } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2147483648] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 1, 1, 0] }
+        }
+      }
+    },
+    {
+      "name": "ort_caseB_empty",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 10.0, 11.0, 12.0, 13.0, 14.0, 20.0, 21.0, 22.0, 23.0, 24.0]
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "ends": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "axes": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "steps": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 10.0, 11.0, 12.0, 13.0, 14.0, 20.0, 21.0, 22.0, 23.0, 24.0]
+          },
+          "tolerance": 0.001
+        }
+      }
+    },
+    {
+      "name": "inner_vec4_int8_aligned_2d",
+      "provenance": {
+        "notes": "int8 aligned inner_vec4 store (step=1, inner_start=0 %4==0, data/out last dim 8 %4==0). int8 maps to i32 storage so binding is vec4<i32>. Previously only an int8 reverse case existed, which hits the per-component branch, leaving the aligned vec4<i8> store path GPU-unvalidated."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int8",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": [-128, -1, 0, 1, 2, 3, 4, 127, 10, 20, 30, 40, 50, 60, 70, 80] }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 0] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 8] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int8",
+          "shape": [1, 8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80] }
+        }
+      }
+    },
+    {
+      "name": "inner_vec4_uint8_aligned_2d",
+      "provenance": {
+        "notes": "uint8 aligned inner_vec4 full-copy store (step=1, start=0, last dim 8 %4==0). uint8 maps to u32 storage -> binding vec4<u32>. Edge values 0..255 to catch any narrowing."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "uint8",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/inner_vec4_uint8_aligned_2d_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 8] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [2, 8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/inner_vec4_uint8_aligned_2d_input_data" } }
+        }
+      }
+    },
+    {
+      "name": "inner_vec4_bool_aligned_2d",
+      "provenance": {
+        "notes": "bool aligned inner_vec4 store (step=1, start=0, last dim 8 %4==0). bool maps to u32 storage -> binding vec4<u32>. ORT computes bool Slice. The only prior bool case is _gpu_gap-named so render coverage skips it, leaving bool Slice GPU-unvalidated; this case exercises bool through the aligned vec4 store on a non-skipped name."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "bool",
+          "shape": [2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/inner_vec4_bool_aligned_2d_input_data" } }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 0] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 8] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [2, 8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/inner_vec4_bool_aligned_2d_input_data" } }
+        }
+      }
+    },
+    {
+      "name": "inner_vec4_reversed_axes_order_rank2",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [4, 1] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [16, 5] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 0] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 12], "tolerance": 0 } }
+    },
+    {
+      "name": "inner_vec4_negative_inner_axis_step2_rank3",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 5, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.17 }
+        },
+        "starts": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 0] } },
+        "ends": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [3, 16] } },
+        "axes": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, -1] } },
+        "steps": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 2, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "generic_metadata_negative_axis_rank4",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "generic_metadata_no_axes_rank3_negative_start",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 6, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 6, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "steps_without_axes_reverse_stride2",
+      "provenance": {
+        "notes": "Pins the independent optional-input combination in which steps is supplied while axes is omitted. Implicit axis 0 with a negative stride of magnitude two must select rows 4 and 2, so neither silently defaulting steps nor requiring axes can pass."
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [5, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 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]
+          }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [12.0, 13.0, 14.0, 6.0, 7.0, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "rank7_last_axis",
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [6] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1, 2, 1, 2, 1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "rank8_last_axis",
+      "attrs": {},
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 2, 1, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [7] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1, 2, 1, 2, 1, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "rank2_axis0_odd_width_x4_f16",
+      "provenance": {
+        "notes": "A float16 odd-width x4 row slice checks half-precision storage and its guarded final group. Slice is a copy, so the tolerance remains zero."
+      },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [6, 5], "data": { "kind": "linspace", "start": -1.0, "end": 1.0 } },
+        "starts": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "ends": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "axes": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "steps": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 5], "tolerance": 0 } }
+    }
+  ]
+}