sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +66 -0
build/webgpu/bench.json +18 -0
build/webgpu/manifest.json +65 -0
build/webgpu/metadata.json +18 -0
build/webgpu/shape.wgsl.jinja +13 -0
build/webgpu/test.json +410 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Shape
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 15
+## Description
+Returns a 1-D tensor containing the shape of the input tensor. Optional `start` and `end` attributes select a slice of the shape axes; negative values count from the back, and axes are clamped to `[0, rank]`.
+See the [ONNX `Shape` spec](https://onnx.ai/onnx/operators/onnx__Shape.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `data` | `data` | `T` | — | — | The input tensor whose shape is computed. | required |
+## Outputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `shape` | `shape` | `S` | `uint32` | `1` | derived; see description | Logical int64 1-D tensor of dimension sizes for the selected axes of the input; WebGPU emits bounded uint32 values. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `start` | `0` | First axis (inclusive) of the shape slice; negative values count from the back, default is 0. |
+| `end` | — | Last axis (exclusive) of the shape slice; negative values count from the back; if omitted, all axes through the last are included. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `int16`, `uint32`, `int8`, `uint8`, `bool` |
+| `S` | `int64` |
+## 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
+- [`shape.wgsl.jinja`](build/webgpu/shape.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Shape", { version: 1 });
+const { shape } = await kernel({ data: { data: dataData, shape: [1, 2, 2] } });
+```

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "ai.onnx.Shape",
+  "cases": [
+    {
+      "name": "rank4",
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 16, 128, 128] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [4] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "name": "write-BW", "value": "4 * numel(shapes.shape)" }] },
+      "preset": "smoke"
+    },
+    {
+      "name": "rank7_full_shape_gpu_gap",
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 2, 1, 3, 1, 4, 1] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [7] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "name": "write-BW", "value": "4 * numel(shapes.shape)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,65 @@

+{
+  "domain": "ai.onnx",
+  "name": "Shape",
+  "sinceVersion": 15,
+  "description": "Returns a 1-D tensor containing the shape of the input tensor. Optional `start` and `end` attributes select a slice of the shape axes; negative values count from the back, and axes are clamped to `[0, rank]`.",
+  "inputs": [{ "role": "data", "dtype": "T", "description": "The input tensor whose shape is computed." }],
+  "outputs": [
+    {
+      "role": "shape",
+      "dtype": "S",
+      "rank": 1,
+      "description": "Logical int64 1-D tensor of dimension sizes for the selected axes of the input; WebGPU emits bounded uint32 values.",
+      "shape": ["max(0, endAxis - startAxis)"]
+    }
+  ],
+  "attributes": { "start": 0 },
+  "attributeDescriptions": {
+    "start": "First axis (inclusive) of the shape slice; negative values count from the back, default is 0.",
+    "end": "Last axis (exclusive) of the shape slice; negative values count from the back; if omitted, all axes through the last are included."
+  },
+  "typeConstraints": {
+    "T": ["float32", "float16", "int32", "int16", "uint32", "int8", "uint8", "bool"],
+    "S": ["int64"]
+  },
+  "args": {
+    "data": { "kind": "tensor", "semantic": "data", "role": "input" },
+    "shape": { "kind": "tensor", "semantic": "shape", "role": "output", "dtype": "uint32" }
+  },
+  "derive": {
+    "startAxis": "max(0, min(ranks.data, attrs.start + (ranks.data if attrs.start < 0 else 0)))",
+    "endParam": "attrs.end if has(attrs, \"end\") else ranks.data",
+    "endAxis": "min(ranks.data, endParam + (ranks.data if endParam < 0 else 0))"
+  },
+  "variants": [
+    {
+      "id": "full_shape",
+      "when": ["ranks.shape == 1", "dim(shapes.shape, 0) == max(0, endAxis - startAxis)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Shape",
+          "source": {
+            "shader": "shape.wgsl.jinja",
+            "inputs": {
+              "inputShape": "shapes.data",
+              "rank": "ranks.data",
+              "start": "startAxis",
+              "outputRank": "dim(shapes.shape, 0)"
+            }
+          },
+          "bindings": [
+            {
+              "name": "shape",
+              "arg": "shape",
+              "semantic": "shape",
+              "buffer": { "type": "storage" },
+              "elementType": "u32"
+            }
+          ],
+          "dispatch": { "x": 1 }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Shape",
+  "id": "_ai_onnx_shape_webgpu_353de3d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "kiw8iMiBPBGaaz6TKqd9RaaOn38QOi/xilG0TXGnttQ=",
+      "manifest.json": "Y8kEIxBXPtwm8XOQcctfW9bFg8RGzmi1r7HwvpwrEPY=",
+      "shape.wgsl.jinja": "J1oIMamgxFx04z+NWiHv6zFj8vclOHR99dd4R/GsPzY=",
+      "test.json": "U5HvTdgN4iqjYIRxu51OW8NZCuqqj/oRlYrEEruggIM="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Shape" }
+}

build/webgpu/shape.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,13 @@

+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  if (gid.x != 0u) {
+    return;
+  }
+{% for dim in source.inputShape %}
+{% if loop.index0 >= source.start and loop.index0 < source.start + source.outputRank %}
+  shape[{{ loop.index0 - source.start }}u] = {{ dim }}u;
+{% endif %}
+{% endfor %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,410 @@

+{
+  "op": "ai.onnx.Shape",
+  "cases": [
+    {
+      "name": "rank3_f32",
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 3, 4] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [3] } }
+    },
+    {
+      "name": "rank4_u32",
+      "inputs": { "data": { "dtype": "uint32", "shape": [1, 2, 3, 4] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [4] } }
+    },
+    {
+      "name": "rank0_scalar_input_empty_shape_vector",
+      "inputs": { "data": { "dtype": "float32", "shape": [] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_interior_zero_dim_reports_zero",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_Default",
+        "notes": "Additional node-level zero-dimension edge: Shape must report an interior zero dimension exactly."
+      },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "shape": { "dtype": "uint32", "shape": [3], "tolerance": 0, "data": { "kind": "values", "values": [2, 0, 3] } }
+      }
+    },
+    {
+      "name": "start_end_middle_dims",
+      "attrs": { "start": 1, "end": 3 },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "negative_start_negative_end",
+      "attrs": { "start": -3, "end": -1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_start_equals_end_empty_slice",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_StartAndEndProducingEmptySlice"
+      },
+      "attrs": { "start": 2, "end": 2 },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_shape_3x4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestInt8"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int8",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int16_shape_3x4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestInt16"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int16",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_shape_12x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestUint8"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "uint8",
+          "shape": [12, 1],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_shape_4x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestInt32"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [4, 3],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_default_opset15",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_Default"
+      },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_start_only_int32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_StartOnly"
+      },
+      "attrs": { "start": 1 },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_end_only_int32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_EndOnly"
+      },
+      "attrs": { "end": 2 },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_start_and_end_int32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_StartAndEnd"
+      },
+      "attrs": { "start": 1, "end": 2 },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_start_and_end_int32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeOpset15_StartAndEndNegative"
+      },
+      "attrs": { "start": -2, "end": -1 },
+      "inputs": {
+        "data": { "dtype": "int32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_float_shape_2x6",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestFloat"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0] }
+        }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_exact_shape_2x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestBool"
+      },
+      "inputs": {
+        "data": { "dtype": "bool", "shape": [2, 3], "data": { "kind": "values", "values": [1, 1, 0, 0, 1, 0] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_shape_12x1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestUint32"
+      },
+      "inputs": {
+        "data": {
+          "dtype": "uint32",
+          "shape": [12, 1],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12] }
+        }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_shape_logical_uint32_2x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestBool",
+        "notes": "Logical bool fixture represented with uint32 storage slots."
+      },
+      "inputs": {
+        "data": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 1, 0, 0, 1, 0] } }
+      },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_shape_exact_2x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/shape_op_test.cc",
+        "test": "ShapeOpTest.ShapeTestBool"
+      },
+      "inputs": {
+        "data": { "dtype": "bool", "shape": [2, 3], "data": { "kind": "values", "values": [1, 1, 0, 0, 1, 0] } }
+      },
+      "outputs": {
+        "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0, "data": { "kind": "values", "values": [2, 3] } }
+      }
+    },
+    {
+      "name": "onnx_backend_shape",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_clip_end",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_clip_end",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "end": 10 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_clip_start",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_clip_start",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "start": -10 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_end_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_end_1",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "end": 1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_end_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_end_negative_1",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "end": -1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_example",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_example",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 3] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_start_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_start_1",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "start": 1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_start_1_end_2",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_start_1_end_2",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "end": 2, "start": 1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_start_1_end_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_start_1_end_negative_1",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "end": -1, "start": 1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_start_greater_than_end",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_start_greater_than_end",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "end": 1, "start": 2 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_shape_start_negative_1",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_shape_start_negative_1",
+        "notes": "ONNX shape/size outputs are int64; this framework uses uint32 shape/index storage."
+      },
+      "attrs": { "start": -1 },
+      "inputs": { "data": { "dtype": "float32", "shape": [3, 4, 5] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_caseB_empty",
+      "inputs": { "data": { "dtype": "float32", "shape": [0, 3, 4], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 3, 4] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "rank3_f16",
+      "inputs": { "data": { "dtype": "float16", "shape": [2, 3, 4] } },
+      "outputs": { "shape": { "dtype": "uint32", "shape": [3] } }
+    },
+    {
+      "name": "rank6_negative_start_clip",
+      "attrs": { "start": -7 },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 3, 4, 5, 6, 7] } },
+      "outputs": {
+        "shape": {
+          "dtype": "uint32",
+          "shape": [6],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7] }
+        }
+      }
+    },
+    {
+      "name": "rank6_start_end_middle_negative",
+      "attrs": { "start": -4, "end": -1 },
+      "inputs": { "data": { "dtype": "uint32", "shape": [2, 3, 4, 5, 6, 7] } },
+      "outputs": {
+        "shape": { "dtype": "uint32", "shape": [3], "tolerance": 0, "data": { "kind": "values", "values": [4, 5, 6] } }
+      }
+    },
+    {
+      "name": "rank7_full_shape",
+      "provenance": {
+        "source": "onnx spec: Shape returns the input tensor dimensions",
+        "notes": "Small rank-7 sentinel. ORT CPU returns [1,2,1,3,1,4,1]; full_shape is rank-agnostic, and rank8_full_shape covers the band above it."
+      },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 2, 1, 3, 1, 4, 1] } },
+      "outputs": {
+        "shape": {
+          "dtype": "uint32",
+          "shape": [7],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 2, 1, 3, 1, 4, 1] }
+        }
+      }
+    },
+    {
+      "name": "rank8_full_shape",
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 2, 1, 3, 1, 4, 1, 2] } },
+      "outputs": {
+        "shape": {
+          "dtype": "uint32",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 2, 1, 3, 1, 4, 1, 2] }
+        }
+      }
+    }
+  ]
+}