sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +66 -0
build/webgpu/bench.json +49 -0
build/webgpu/manifest.json +114 -0
build/webgpu/metadata.json +18 -0
build/webgpu/space-depth-permute.wgsl.jinja +80 -0
build/webgpu/test.json +361 -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.SpaceToDepth
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Rearranges blocks of spatial data into depth by moving values from the height and width dimensions into the channel dimension. An NCHW input of shape `[N, C, H, W]` produces an output of shape `[N, C * blocksize * blocksize, H / blocksize, W / blocksize]`.
+See the [ONNX `SpaceToDepth` spec](https://onnx.ai/onnx/operators/onnx__SpaceToDepth.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | `4` | — | 4-D input tensor of shape `[N, C, H, W]`. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | `4` | derived; see description | 4-D output tensor of shape `[N, C * blocksize * blocksize, H / blocksize, W / blocksize]`. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `blocksize` | — | Size of the spatial block to collapse into depth; each `blocksize x blocksize` patch of pixels becomes additional channels. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `int16`, `int8`, `uint32`, `uint8`, `bool` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`space-depth-permute.wgsl.jinja`](build/webgpu/space-depth-permute.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.SpaceToDepth", { version: 1 });
+const { output } = await kernel({ input: { data: inputData, shape: [1, 1, 2, 4] } }, {
+  attrs: { blocksize: 2 },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,49 @@

+{
+  "op": "ai.onnx.SpaceToDepth",
+  "cases": [
+    {
+      "name": "nchw_1x64x128x128",
+      "attrs": { "blocksize": 2 },
+      "inputs": { "input": { "dtype": "float32", "shape": [1, 64, 128, 128] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 256, 64, 64] } }
+    },
+    {
+      "name": "nchw_1x32x192x192_block3",
+      "attrs": { "blocksize": 3 },
+      "inputs": { "input": { "dtype": "float32", "shape": [1, 32, 192, 192] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 288, 64, 64] } }
+    },
+    {
+      "name": "nchw_f16_1x64x128x128",
+      "preset": "smoke",
+      "attrs": { "blocksize": 2 },
+      "vars": { "dtype": "float16", "count": 1048576 },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 4471, "scale": 1 }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 256, 64, 64] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "nchw_dispatch_cliff_1x4194241x2x2",
+      "attrs": { "blocksize": 2 },
+      "vars": { "dtype": "float32", "count": 16776964 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 4194241, 2, 2], "dist": "normal", "seed": 4472, "scale": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16776964, 1, 1] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "nchw-const-fallback-odd-width-block2-f32-1x128x2048x6",
+      "preset": "stress",
+      "attrs": { "blocksize": 2 },
+      "vars": { "dtype": "float32", "count": 1572864 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 128, 2048, 6], "dist": "normal", "seed": 5201, "scale": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 512, 1024, 3] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,114 @@

+{
+  "domain": "ai.onnx",
+  "name": "SpaceToDepth",
+  "sinceVersion": 13,
+  "description": "Rearranges blocks of spatial data into depth by moving values from the height and width dimensions into the channel dimension. An NCHW input of shape `[N, C, H, W]` produces an output of shape `[N, C * blocksize * blocksize, H / blocksize, W / blocksize]`.",
+  "inputs": [{ "role": "input", "dtype": "T", "rank": 4, "description": "4-D input tensor of shape `[N, C, H, W]`." }],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": 4,
+      "description": "4-D output tensor of shape `[N, C * blocksize * blocksize, H / blocksize, W / blocksize]`.",
+      "shape": ["dim(shapes.input, 0)", "dim(shapes.input, 1) * attrs.blocksize * attrs.blocksize", "dim(shapes.input, 2) / attrs.blocksize", "dim(shapes.input, 3) / attrs.blocksize"]
+    }
+  ],
+  "attributes": {},
+  "attributeDescriptions": {
+    "blocksize": "Size of the spatial block to collapse into depth; each `blocksize x blocksize` patch of pixels becomes additional channels."
+  },
+  "attributeConstraints": { "blocksize": { "required": true } },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "int8", "uint32", "uint8", "bool"] },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "shapeContract": "attrs.blocksize > 0 and ranks.input == 4 and ranks.output == 4 and dim(shapes.output, 0) == dim(shapes.input, 0) and dim(shapes.output, 1) == dim(shapes.input, 1) * attrs.blocksize * attrs.blocksize and dim(shapes.output, 2) * attrs.blocksize == dim(shapes.input, 2) and dim(shapes.output, 3) * attrs.blocksize == dim(shapes.input, 3) and f16Ok(dtypes.T)",
+    "reportedWideExecution": "has(device.adapterInfo, \"subgroupMinSize\") and device.adapterInfo.subgroupMinSize >= 32",
+    "f16VectorPermutationPreferred": "dtypes.T != \"f16\" or device.features.has(\"subgroups\") or reportedWideExecution"
+  },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "bindingSets": {
+    "io": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scalar"
+      },
+      {
+        "name": "output",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputElement"
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "nchw_vec4",
+      "priority": 20,
+      "when": ["shapeContract", "dim(shapes.output, 3) % 4 == 0", "f16VectorPermutationPreferred"],
+      "constants": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SpaceToDepth",
+          "source": {
+            "shader": "space-depth-permute.wgsl.jinja",
+            "inputs": {
+              "direction": "\"spaceToDepth\"",
+              "mode": "\"DCR\"",
+              "vectorized": true,
+              "blocksize": "attrs.blocksize",
+              "count": "numel(shapes.output) / 4",
+              "outWUnits": "dim(shapes.output, 3) / 4",
+              "outH": "dim(shapes.output, 2)",
+              "outC": "dim(shapes.output, 1)",
+              "inC": "dim(shapes.input, 1)",
+              "inH": "dim(shapes.input, 2)",
+              "inW": "dim(shapes.input, 3)"
+            }
+          },
+          "bindings": "io",
+          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw_const",
+      "priority": 10,
+      "when": ["shapeContract"],
+      "constants": { "outputElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "SpaceToDepth",
+          "source": {
+            "shader": "space-depth-permute.wgsl.jinja",
+            "inputs": {
+              "direction": "\"spaceToDepth\"",
+              "mode": "\"DCR\"",
+              "vectorized": false,
+              "blocksize": "attrs.blocksize",
+              "count": "numel(shapes.output)",
+              "outWUnits": "dim(shapes.output, 3)",
+              "outW": "dim(shapes.output, 3)",
+              "outH": "dim(shapes.output, 2)",
+              "outC": "dim(shapes.output, 1)",
+              "inC": "dim(shapes.input, 1)",
+              "inH": "dim(shapes.input, 2)",
+              "inW": "dim(shapes.input, 3)"
+            }
+          },
+          "bindings": "io",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.SpaceToDepth",
+  "id": "_ai_onnx_spacetodepth_webgpu_cbf603a",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "oDNgQ+aqxMurIw6vhcAUnPq8eIWuPYweL5P+hSxQqa8=",
+      "manifest.json": "XHyEI7bli7+h03YIu3a1rT9KFb1RBsgXhYWECC76yO4=",
+      "space-depth-permute.wgsl.jinja": "KbYVOfLYDd2L7BSJvB0dQYoKUYbQfHqvrzN+VJbkqew=",
+      "test.json": "IvLVrset+vkCY0gRqAbD44jO6hGRKbFHvti1sRY8iKo="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.SpaceToDepth" }
+}

build/webgpu/space-depth-permute.wgsl.jinja ADDED Viewed

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

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Specializing all geometry lets division and modulo lower to constant
+// multiply-high/shift sequences. max(1) keeps zero-sized tensors compilable;
+// they dispatch no work.
+const COUNT: u32 = {{ source.count }}u;
+const OUT_W_UNITS: u32 = max(1u, {{ source.outWUnits }}u);
+const OUT_H: u32 = max(1u, {{ source.outH }}u);
+const OUT_C: u32 = max(1u, {{ source.outC }}u);
+const IN_C: u32 = max(1u, {{ source.inC }}u);
+const IN_H: u32 = {{ source.inH }}u;
+const IN_W: u32 = {{ source.inW }}u;
+const BLOCK_SIZE: u32 = {{ source.blocksize }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("out_unit", "COUNT", note="") }}
+  let ow_unit = out_unit % OUT_W_UNITS;
+  var remaining = out_unit / OUT_W_UNITS;
+  let oh = remaining % OUT_H;
+  remaining /= OUT_H;
+  let oc = remaining % OUT_C;
+  let n = remaining / OUT_C;
+  let ow_base = ow_unit * {{ 4 if source.vectorized else 1 }}u;
+  let c = oc % IN_C;
+  let block = oc / IN_C;
+  let bh = block / BLOCK_SIZE;
+  let bw = block % BLOCK_SIZE;
+  let ih = oh * BLOCK_SIZE + bh;
+  let input_base = ((n * IN_C + c) * IN_H + ih) * IN_W + bw;
+{% for lane in range(4 if source.vectorized else 1) %}
+  let ow{{ lane }} = ow_base + {{ lane }}u;
+  let value{{ lane }} = input[input_base + ow{{ lane }} * BLOCK_SIZE];
+{% endfor %}
+{% if source.vectorized %}
+  output[out_unit] = {{ vectorScalar }}(value0, value1, value2, value3);
+{% else %}
+  output[out_unit] = value0;
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,361 @@

+{
+  "op": "ai.onnx.SpaceToDepth",
+  "fixtureArrays": {
+    "onnx_backend_example_1x1x4x6_input_input": [0, 6, 1, 7, 2, 8, 12, 18, 13, 19, 14, 20, 3, 9, 4, 10, 5, 11, 15, 21, 16, 22, 17, 23]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_nchw_channel_heavy",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4194241, 2, 2],
+          "data": { "kind": "linspace", "start": 0.0, "end": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16776964, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 10.0, 1.0, 11.0, 20.0, 30.0, 21.0, 31.0, 2.0, 12.0, 3.0, 13.0, 22.0, 32.0, 23.0, 33.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2, 2] } }
+    },
+    {
+      "name": "nchw_block3_multichannel",
+      "attrs": { "blocksize": 3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 6, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 18, 2, 1] } }
+    },
+    {
+      "name": "nchw_f16_batch2_block2",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [2, 1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, -1.0, -2.0, -3.0, -4.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 4, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_height_block2",
+      "attrs": { "blocksize": 2 },
+      "inputs": { "input": { "dtype": "float32", "shape": [1, 1, 0, 4] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 0, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_example_1x1x4x6",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 6],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_example_1x1x4x6_input_input" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_block2_two_channels_wide",
+      "attrs": { "blocksize": 2 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
+        "test": "TensorOpTest.SpaceToDepthTest_1"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.1, 0.2, 0.3, 1.0, 1.1, 1.2, 1.3, 2.0, 2.1, 2.2, 2.3, 3.0, 3.1, 3.2, 3.3]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_block2_two_channels_tall_and_wide",
+      "attrs": { "blocksize": 2 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
+        "test": "TensorOpTest.SpaceToDepthTest_3"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 4, 8],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 1.0, 1.1, 1.2, 1.3, 1.4, 1.5, 1.6, 1.7, 2.0, 2.1, 2.2, 2.3, 2.4, 2.5, 2.6, 2.7, 3.0, 3.1, 3.2, 3.3, 3.4, 3.5, 3.6, 3.7, 4.0, 4.1, 4.2, 4.3, 4.4, 4.5, 4.6, 4.7, 5.0, 5.1, 5.2, 5.3, 5.4, 5.5, 5.6, 5.7, 6.0, 6.1, 6.2, 6.3, 6.4, 6.5, 6.6, 6.7, 7.0, 7.1, 7.2, 7.3, 7.4, 7.5, 7.6, 7.7]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f16_block3_single_pixel_output",
+      "attrs": { "blocksize": 3 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 9, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_block3_batch2_channels3",
+      "attrs": { "blocksize": 3 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
+        "test": "TensorOpTest.SpaceToDepthTest_2",
+        "notes": "Uses the ORT shape/blocksize; fixture values are generated deterministically."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.01, "cosStep": 0.0, "scale": 100.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 27, 1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_spacetodepth",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 6, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.54881352186203, 0.7151893377304077, 0.6027633547782898, 0.5448831915855408, 0.42365479469299316, 0.6458941102027893, 0.4375872015953064, 0.891772985458374, 0.9636627435684204, 0.3834415078163147, 0.7917250394821167, 0.5288949012756348, 0.5680445432662964, 0.9255966544151306, 0.07103605568408966, 0.08712930232286453, 0.020218396559357643, 0.832619845867157, 0.7781567573547363, 0.8700121641159058, 0.978618323802948, 0.7991585731506348, 0.4614793658256531, 0.7805292010307312, 0.11827442795038223, 0.6399210095405579, 0.14335328340530396, 0.9446688890457153, 0.5218483209609985, 0.4146619439125061, 0.26455560326576233, 0.7742336988449097, 0.4561503231525421, 0.568433940410614, 0.018789799883961678, 0.6176354885101318, 0.6120957136154175, 0.6169340014457703, 0.9437480568885803, 0.681820273399353, 0.35950788855552673, 0.43703195452690125, 0.6976311802864075, 0.0602254718542099, 0.6667667031288147, 0.670637845993042, 0.21038256585597992, 0.12892629206180573, 0.31542834639549255, 0.36371076107025146, 0.5701967477798462, 0.4386015236377716, 0.9883738160133362, 0.10204481333494186, 0.20887675881385803, 0.16130951046943665, 0.6531082987785339, 0.25329160690307617, 0.4663107693195343, 0.24442559480667114, 0.15896958112716675, 0.11037514358758926, 0.6563295722007751, 0.13818295300006866, 0.1965823620557785, 0.3687251806259155, 0.8209932446479797, 0.09710127860307693, 0.8379449248313904, 0.0960984081029892, 0.9764594435691833, 0.4686512053012848, 0.9767611026763916, 0.6048455238342285, 0.7392635941505432, 0.03918779268860817, 0.28280696272850037, 0.12019655853509903, 0.296140193939209, 0.11872772127389908, 0.3179831802845001, 0.414262980222702, 0.06414749473333359, 0.6924721002578735, 0.5666014552116394, 0.26538950204849243, 0.5232480764389038, 0.09394051134586334, 0.5759465098381042, 0.9292961955070496, 0.3185689449310303, 0.6674103736877441, 0.13179786503314972, 0.7163271903991699, 0.28940609097480774, 0.18319135904312134, 0.5865129232406616, 0.02010754682123661, 0.8289400339126587, 0.004695476032793522, 0.6778165102005005, 0.2700079679489136, 0.7351940274238586, 0.9621885418891907, 0.2487531453371048, 0.5761573314666748, 0.5920419096946716, 0.5722519159317017, 0.22308163344860077, 0.9527490139007568, 0.4471253752708435, 0.8464086651802063, 0.6994792819023132, 0.2974369525909424, 0.8137978315353394, 0.396505743265152, 0.8811032176017761, 0.5812729001045227, 0.8817353844642639, 0.6925315856933594, 0.7252542972564697, 0.5013243556022644, 0.9560836553573608, 0.6439902186393738, 0.4238550364971161, 0.6063932180404663, 0.019193198531866074, 0.30157482624053955, 0.6601735353469849, 0.2900775969028473, 0.6180154085159302, 0.42876869440078735, 0.1354740709066391, 0.29828232526779175, 0.5699648857116699, 0.5908727645874023, 0.5743252635002136, 0.6532008051872253, 0.6521032452583313, 0.43141844868659973, 0.8965466022491455, 0.36756187677383423, 0.4358649253845215, 0.8919233679771423]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 8, 3, 3] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_spacetodepth" }
+    },
+    {
+      "name": "onnx_backend_spacetodepth_example",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 6],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_example_1x1x4x6_input_input" } }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2, 3] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_spacetodepth_example" }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "blocksize": 2 },
+      "inputs": { "input": { "dtype": "float32", "shape": [0, 1, 4, 4], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 4, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_dim_f16",
+      "attrs": { "blocksize": 2 },
+      "inputs": { "input": { "dtype": "float16", "shape": [0, 1, 4, 4], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [0, 4, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw_block4_multichannel",
+      "attrs": { "blocksize": 4 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 3, 8, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.07 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 48, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw_block4_f16",
+      "attrs": { "blocksize": 4 },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 2, 4, 16],
+          "data": {
+            "kind": "values",
+            "values": [-8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, -8.0, -7.0, -6.0, -5.0, -4.0, -3.0, -2.0, -1.0, 0.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 32, 1, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw_block1_identity",
+      "attrs": { "blocksize": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [-5.0, -3.5, -2.0, -0.5, 1.0, 2.5, 4.0, 5.5, 7.0, 8.5, 10.0, 11.5, 13.0, 14.5, 16.0, 17.5]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw_const_odd_width_block2_addressing",
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 4, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw_const_block3_odd_width_multichannel",
+      "attrs": { "blocksize": 3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 6, 15],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.17 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 18, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "nchw_vec4_block3_width_mul4_multichannel",
+      "attrs": { "blocksize": 3 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 6, 12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 18, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_standard_int32_block2_extremes",
+      "provenance": {
+        "source": "onnx/docs/Operators.md#SpaceToDepth",
+        "notes": "ONNX SpaceToDepth-13 permits int32; signed extrema prove the widened i32 route is a bit-exact permutation."
+      },
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 4],
+          "data": { "kind": "values", "values": [-2147483648, 2147483647, -16777217, 16777217, -1, 0, 42, -42] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_standard_int16_block2_extremes",
+      "provenance": {
+        "source": "onnx/docs/Operators.md#SpaceToDepth",
+        "notes": "ONNX SpaceToDepth-13 permits int16; both signed extrema guard the runtime's widened i32 storage route."
+      },
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "int16",
+          "shape": [1, 1, 2, 4],
+          "data": { "kind": "values", "values": [-32768, 32767, -12345, 12345, -1, 0, 42, -42] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int16", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_standard_int8_block2_extremes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
+        "test": "TensorOpTest.SpaceToDepthTest_int_opset13",
+        "notes": "Uses both int8 extrema so the widened i32 storage route cannot hide narrowing or sign loss."
+      },
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "int8",
+          "shape": [1, 1, 2, 4],
+          "data": { "kind": "values", "values": [-128, 127, -64, 63, -1, 0, 42, -42] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_standard_uint32_block2_extremes",
+      "provenance": {
+        "source": "onnx/docs/Operators.md#SpaceToDepth",
+        "notes": "ONNX SpaceToDepth-13 permits uint32; values across the signed boundary and UINT32_MAX guard exact u32 payload movement."
+      },
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [1, 1, 2, 4],
+          "data": { "kind": "values", "values": [0, 4294967295, 16777217, 2147483648, 1, 42, 4000000000, 255] }
+        }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_standard_uint8_block2_extremes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
+        "test": "TensorOpTest.SpaceToDepthTest_int_opset13",
+        "notes": "Includes UINT8_MAX and values around the signed boundary on the widened u32 storage route."
+      },
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 4],
+          "data": { "kind": "values", "values": [0, 255, 127, 128, 1, 42, 254, 64] }
+        }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_standard_bool_block2_pattern",
+      "provenance": {
+        "source": "onnx/docs/Operators.md#SpaceToDepth",
+        "notes": "ONNX SpaceToDepth-13 permits bool; an asymmetric pattern checks that widened u32 boolean payloads are reordered without conversion."
+      },
+      "attrs": { "blocksize": 2 },
+      "inputs": {
+        "input": {
+          "dtype": "bool",
+          "shape": [1, 1, 2, 4],
+          "data": { "kind": "values", "values": [0, 1, 1, 0, 1, 1, 0, 0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "bool", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    }
+  ]
+}