sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +14 -10
build/webgpu/bench.json +0 -1
build/webgpu/depth-to-space-nchw-coarsened4.wgsl.jinja +14 -17
build/webgpu/manifest.json +69 -94
build/webgpu/metadata.json +15 -8
build/webgpu/space-depth-permute.wgsl.jinja +27 -30
build/webgpu/test.json +6 -7

README.md CHANGED Viewed

@@ -18,15 +18,15 @@ See the [ONNX `DepthToSpace` spec](https://onnx.ai/onnx/operators/onnx__DepthToS
 ## 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
@@ -34,8 +34,8 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
-| `mode` | `"DCR"` | Element ordering within each block: `DCR`, the default depth-column-row order, or `CRD`, the column-row-depth order. |
 | `blocksize` | — | Side length of the spatial blocks; input channels are divided by `blocksize * blocksize`. |
 ## Type constraints
@@ -45,7 +45,7 @@ Attributes and default values (overridable per request):
 ## 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
@@ -54,10 +54,14 @@ Attributes and default values (overridable per request):
 ## 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.

 ## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `input` | `T` | `4` | — | 4-D input tensor of shape `[N, C, H, W]`. | required |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `output` | `T` | `4` | derived | 4-D output tensor of shape `[N, C/(blocksize*blocksize), H*blocksize, W*blocksize]`. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `blocksize` | — | Side length of the spatial blocks; input channels are divided by `blocksize * blocksize`. |
+| `mode` | `"DCR"` | Element ordering within each block: `DCR`, the default depth-column-row order, or `CRD`, the column-row-depth order. |
 ## Type constraints
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.DepthToSpace",
   "cases": [
     {
       "name": "nchw_1x256x64x64",

 {
   "cases": [
     {
       "name": "nchw_1x256x64x64",

build/webgpu/depth-to-space-nchw-coarsened4.wgsl.jinja CHANGED Viewed

@@ -1,26 +1,23 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // Four adjacent output columns share N/C/H decoding and the block-row lookup.
 // Keeping scalar bindings makes this valid for odd output widths; only the last
-// group in each row is predicated. Static dimensions let the shader compiler
-// eliminate all division and modulo operations.
-const GROUPS: u32 = {{ source.groups }}u;
-const OUT_W_GROUPS: u32 = max(1u, {{ source.outWGroups }}u);
-const OUT_W: u32 = max(1u, {{ source.outW }}u);
-const OUT_H: u32 = max(1u, {{ source.outH }}u);
-const OUT_C: u32 = max(1u, {{ source.outC }}u);
-const IN_C: u32 = {{ source.inC }}u;
-const IN_H: u32 = {{ source.inH }}u;
-const IN_W: u32 = {{ source.inW }}u;
-const BS: u32 = {{ source.blocksize }}u;
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
-  let q = gid.x + gid.y * nwg.x * WG;
   if (q >= GROUPS) { return; }
   let owBase = (q % OUT_W_GROUPS) * 4u;
@@ -38,7 +35,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
     if (ow < OUT_W) {
       let bw = ow % BS;
       let iw = ow / BS;
-{% if source.mode == "CRD" %}
       let ic = (oc * BS + bh) * BS + bw;
 {% else %}
       let ic = (bh * BS + bw) * OUT_C + oc;

 {{ env.wgsl.resourceDeclarations }}
 // Four adjacent output columns share N/C/H decoding and the block-row lookup.
 // Keeping scalar bindings makes this valid for odd output widths; only the last
+// group in each row is predicated. Every division and modulo operation uses a
+// static dimension as its divisor.
+const GROUPS: u32 = {{ groups }}u;
+const OUT_W_GROUPS: u32 = max(1u, {{ outWGroups }}u);
+const OUT_W: u32 = max(1u, {{ outW }}u);
+const OUT_H: u32 = max(1u, {{ outH }}u);
+const OUT_C: u32 = max(1u, {{ outC }}u);
+const IN_C: u32 = {{ inC }}u;
+const IN_H: u32 = {{ inH }}u;
+const IN_W: u32 = {{ inW }}u;
+const BS: u32 = {{ blocksize }}u;
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let q = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   if (q >= GROUPS) { return; }
   let owBase = (q % OUT_W_GROUPS) * 4u;
     if (ow < OUT_W) {
       let bw = ow % BS;
       let iw = ow / BS;
+{% if modeSpec == "CRD" %}
       let ic = (oc * BS + bh) * BS + bw;
 {% else %}
       let ic = (bh * BS + bw) * OUT_C + oc;

build/webgpu/manifest.json CHANGED Viewed

@@ -2,141 +2,116 @@
   "domain": "ai.onnx",
   "name": "DepthToSpace",
   "sinceVersion": 13,
-  "description": "Rearranges data from the depth dimension into spatial blocks, expanding height and width by `blocksize` while reducing channels by `blocksize * blocksize`. The inverse of SpaceToDepth; supports `DCR` (depth-column-row) and `CRD` (column-row-depth) element orderings.",
-  "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": { "mode": "DCR" },
-  "attributeDescriptions": {
-    "blocksize": "Side length of the spatial blocks; input channels are divided by `blocksize * blocksize`.",
-    "mode": "Element ordering within each block: `DCR`, the default depth-column-row order, or `CRD`, the column-row-depth order."
   },
   "attributeConstraints": { "blocksize": { "required": true }, "mode": { "values": ["DCR", "CRD"] } },
   "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) * attrs.blocksize * attrs.blocksize == dim(shapes.input, 1) and dim(shapes.output, 2) == dim(shapes.input, 2) * attrs.blocksize and dim(shapes.output, 3) == dim(shapes.input, 3) * attrs.blocksize and f16Ok(dtypes.T)"
-  },
-  "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"],
-      "constants": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "DepthToSpace",
-          "source": {
-            "shader": "space-depth-permute.wgsl.jinja",
-            "inputs": {
-              "direction": "\"depthToSpace\"",
-              "mode": "attrs.mode",
-              "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": "DepthToSpace",
-          "source": {
-            "shader": "space-depth-permute.wgsl.jinja",
-            "inputs": {
-              "direction": "\"depthToSpace\"",
-              "mode": "attrs.mode",
-              "vectorized": false,
-              "blocksize": "attrs.blocksize",
-              "count": "numel(shapes.output)",
-              "outWUnits": "dim(shapes.output, 3)",
-              "outC": "dim(shapes.output, 1)",
-              "outH": "dim(shapes.output, 2)",
-              "outW": "dim(shapes.output, 3)",
-              "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" }
         }
       ]
     },
     {
       "id": "nchw_coarsened4",
       "priority": 15,
-      "when": ["shapeContract", "dim(shapes.output, 3) % 4 != 0", "numel(shapes.output) >= 4096"],
-      "constants": { "outputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "DepthToSpace.Coarsened4",
-          "source": {
-            "shader": "depth-to-space-nchw-coarsened4.wgsl.jinja",
-            "inputs": {
-              "mode": "attrs.mode",
-              "blocksize": "attrs.blocksize",
-              "groups": "dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2) * ceil(dim(shapes.output, 3) / 4)",
-              "outWGroups": "ceil(dim(shapes.output, 3) / 4)",
-              "outC": "dim(shapes.output, 1)",
-              "outH": "dim(shapes.output, 2)",
-              "outW": "dim(shapes.output, 3)",
-              "inC": "dim(shapes.input, 1)",
-              "inH": "dim(shapes.input, 2)",
-              "inW": "dim(shapes.input, 3)"
-            }
           },
-          "bindings": "io",
           "dispatch": {
-            "threads": "dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2) * ceil(dim(shapes.output, 3) / 4)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]

   "domain": "ai.onnx",
   "name": "DepthToSpace",
   "sinceVersion": 13,
+  "inputs": { "input": { "dtype": "T", "rank": 4 } },
+  "outputs": {
+    "output": {
       "dtype": "T",
       "rank": 4,
       "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": { "mode": { "default": "DCR" }, "blocksize": {} },
   "attributeConstraints": { "blocksize": { "required": true }, "mode": { "values": ["DCR", "CRD"] } },
   "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "int8", "uint32", "uint8", "bool"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 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) * attrs.blocksize * attrs.blocksize == dim(shapes.input, 1) and dim(shapes.output, 2) == dim(shapes.input, 2) * attrs.blocksize and dim(shapes.output, 3) == dim(shapes.input, 3) * attrs.blocksize and f16Ok(dtypes.T)",
+    "scalar": "dtypes.T"
   },
+  "when": ["shapeContract"],
+  "bindings": { "output": { "buffer": "storage", "elementType": "$outputElement" } },
   "variants": [
     {
       "id": "nchw_vec4",
       "priority": 20,
+      "when": ["dim(shapes.output, 3) % 4 == 0"],
+      "derive": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "DepthToSpace",
+          "shader": "space-depth-permute.wgsl.jinja",
+          "derive": {
+            "direction": "\"depthToSpace\"",
+            "modeSpec": "attrs.mode",
+            "vectorizedSpec": 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": ["input", "output"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "nchw_const",
       "priority": 10,
+      "derive": { "outputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "DepthToSpace",
+          "shader": "space-depth-permute.wgsl.jinja",
+          "derive": {
+            "direction": "\"depthToSpace\"",
+            "modeSpec": "attrs.mode",
+            "vectorizedSpec": false,
+            "blocksize": "attrs.blocksize",
+            "count": "numel(shapes.output)",
+            "outWUnits": "dim(shapes.output, 3)",
+            "outC": "dim(shapes.output, 1)",
+            "outH": "dim(shapes.output, 2)",
+            "outW": "dim(shapes.output, 3)",
+            "inC": "dim(shapes.input, 1)",
+            "inH": "dim(shapes.input, 2)",
+            "inW": "dim(shapes.input, 3)"
           },
+          "bindings": ["input", "output"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.output)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.output)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "nchw_coarsened4",
       "priority": 15,
+      "when": ["dim(shapes.output, 3) % 4 != 0", "numel(shapes.output) >= 4096"],
+      "derive": { "outputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "DepthToSpace.Coarsened4",
+          "shader": "depth-to-space-nchw-coarsened4.wgsl.jinja",
+          "derive": {
+            "modeSpec": "attrs.mode",
+            "blocksize": "attrs.blocksize",
+            "groups": "dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2) * ceil(dim(shapes.output, 3) / 4)",
+            "outWGroups": "ceil(dim(shapes.output, 3) / 4)",
+            "outC": "dim(shapes.output, 1)",
+            "outH": "dim(shapes.output, 2)",
+            "outW": "dim(shapes.output, 3)",
+            "inC": "dim(shapes.input, 1)",
+            "inH": "dim(shapes.input, 2)",
+            "inW": "dim(shapes.input, 3)"
           },
+          "bindings": ["input", "output"],
           "dispatch": {
+            "x": "min(ceilDiv((dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2) * ceil(dim(shapes.output, 3) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2) * ceil(dim(shapes.output, 3) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,26 @@
 {
   "name": "ai.onnx.DepthToSpace",
-  "id": "_ai_onnx_depthtospace_webgpu_ea6e538",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "437M7+npzBe6ifdJlYIgSUZqHw1gO0A7/aRUKJ2NJVY=",
-      "depth-to-space-nchw-coarsened4.wgsl.jinja": "onMG0dd/K41+1hEnl7TddZRLG3CBjjy40MOlZPnrkV0=",
-      "manifest.json": "n4i4VOtp2TXrMC5FwnUM2I16OunwYl0aBPPsFYhDsgU=",
-      "space-depth-permute.wgsl.jinja": "TNRb9rw4b8VDiWMXKSCAaZQUOI+jx4LngtZHiqG3Dg8=",
-      "test.json": "MmJ58u1S776RqIS1psfK6ogGNMRWfjVOghdVgodij3c="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.DepthToSpace" }
 }

 {
   "name": "ai.onnx.DepthToSpace",
+  "id": "_ai_onnx_depthtospace_webgpu_af18dda",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "JU/Zd5udt4ll47K2EnVjXxjt0sEh9KmpBf5fZnZrXFY=",
+      "depth-to-space-nchw-coarsened4.wgsl.jinja": "4xGeu/kV7B8e2wY47sCv1CFLAhW7NbhUkNFAO8N01jw=",
+      "manifest.json": "4syDCQ7+CYWnOP8y3nQeAL0uF3R0PuXa2WGfqCUwI5I=",
+      "space-depth-permute.wgsl.jinja": "AdqL6I5mdEY/Q+Pd8AMR7O4/IasCj9BhRd+8QDccF5A=",
+      "test.json": "hs7y38XH3BuXDsaABYqenTgiC4pdBPPan6x4uUc0EU0="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "nchw_vec4": ["space-depth-permute.wgsl.jinja"],
+      "nchw_const": ["space-depth-permute.wgsl.jinja"],
+      "nchw_coarsened4": ["depth-to-space-nchw-coarsened4.wgsl.jinja"]
+    }
+  }
 }

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

@@ -1,55 +1,52 @@
 {% 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;
@@ -58,14 +55,14 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   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;
-{% for lane in range(4 if source.vectorized else 1) %}
   let ow{{ lane }} = ow_base + {{ lane }}u;
   let bh{{ lane }} = oh % BLOCK_SIZE;
   let bw{{ lane }} = ow{{ lane }} % BLOCK_SIZE;
   let ih{{ lane }} = oh / BLOCK_SIZE;
   let iw{{ lane }} = ow{{ lane }} / BLOCK_SIZE;
-{% if source.mode == "CRD" %}
   let ic{{ lane }} = (oc * BLOCK_SIZE + bh{{ lane }}) * BLOCK_SIZE + bw{{ lane }};
 {% else %}
   let ic{{ lane }} = (bh{{ lane }} * BLOCK_SIZE + bw{{ lane }}) * OUT_C + oc;
@@ -73,7 +70,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let value{{ lane }} = input[((n * IN_C + ic{{ lane }}) * IN_H + ih{{ lane }}) * IN_W + iw{{ lane }}];
 {% endfor %}
-{% if source.vectorized %}
   output[out_unit] = {{ vectorScalar }}(value0, value1, value2, value3);
 {% else %}
   output[out_unit] = value0;

 {% 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 dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; 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 dispatch's
+  // per-axis workgroup fold width (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
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {%- endif %}
 {% endmacro %}
 {{ 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 = {{ count }}u;
+const OUT_W_UNITS: u32 = max(1u, {{ outWUnits }}u);
+const OUT_H: u32 = max(1u, {{ outH }}u);
+const OUT_C: u32 = max(1u, {{ outC }}u);
+const IN_C: u32 = max(1u, {{ inC }}u);
+const IN_H: u32 = {{ inH }}u;
+const IN_W: u32 = {{ inW }}u;
+const BLOCK_SIZE: u32 = {{ blocksize }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d("out_unit", "COUNT", note="") }}
   let ow_unit = out_unit % OUT_W_UNITS;
   remaining /= OUT_H;
   let oc = remaining % OUT_C;
   let n = remaining / OUT_C;
+  let ow_base = ow_unit * {{ 4 if vectorizedSpec else 1 }}u;
+{% for lane in range(4 if vectorizedSpec else 1) %}
   let ow{{ lane }} = ow_base + {{ lane }}u;
   let bh{{ lane }} = oh % BLOCK_SIZE;
   let bw{{ lane }} = ow{{ lane }} % BLOCK_SIZE;
   let ih{{ lane }} = oh / BLOCK_SIZE;
   let iw{{ lane }} = ow{{ lane }} / BLOCK_SIZE;
+{% if modeSpec == "CRD" %}
   let ic{{ lane }} = (oc * BLOCK_SIZE + bh{{ lane }}) * BLOCK_SIZE + bw{{ lane }};
 {% else %}
   let ic{{ lane }} = (bh{{ lane }} * BLOCK_SIZE + bw{{ lane }}) * OUT_C + oc;
   let value{{ lane }} = input[((n * IN_C + ic{{ lane }}) * IN_H + ih{{ lane }}) * IN_W + iw{{ lane }}];
 {% endfor %}
+{% if vectorizedSpec %}
   output[out_unit] = {{ vectorScalar }}(value0, value1, value2, value3);
 {% else %}
   output[out_unit] = value0;

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.DepthToSpace",
   "fixtureArrays": {
     "dcr_input_input": [0, 1, 2, 3, 10, 11, 12, 13, 20, 21, 22, 23, 30, 31, 32, 33],
     "onnx_backend_dcr_1x8x2x3_input_input": [0, 1, 2, 3, 4, 5, 9, 10, 11, 12, 13, 14, 18, 19, 20, 21, 22, 23, 27, 28, 29, 30, 31, 32, 36, 37, 38, 39, 40, 41, 45, 46, 47, 48, 49, 50, 54, 55, 56, 57, 58, 59, 63, 64, 65, 66, 67, 68],
@@ -69,7 +68,7 @@
     {
       "name": "dcr_high_channel_vec4_1x256x4x4",
       "provenance": {
-        "notes": "Compact high-channel DCR sibling for the 1x256x64x64 DepthToSpace benchmark; preserves the vec4/channel-heavy shape family without benchmark-scale spatial size."
       },
       "attrs": { "blocksize": 2, "mode": "DCR" },
       "inputs": {
@@ -84,7 +83,7 @@
     {
       "name": "crd_high_channel_vec4_1x256x4x4",
       "provenance": {
-        "notes": "Compact high-channel CRD sibling for the 1x256x64x64 DepthToSpace benchmark; preserves the mode and channel-heavy vec4 regime without benchmark-scale spatial size."
       },
       "attrs": { "blocksize": 2, "mode": "CRD" },
       "inputs": {
@@ -195,7 +194,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
         "test": "TensorOpTest.DepthToSpaceTest_3<uint8_t>",
-        "notes": "ORT typed coverage includes uint8 DepthToSpace; this compact DCR/default-mode case keeps logical uint8 coverage small."
       },
       "inputs": {
         "input": {
@@ -353,7 +352,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
         "test": "TensorOpTest.DepthToSpaceTest_4",
-        "notes": "Typed ORT coverage includes uint8; this framework stores logical uint8 values in 32-bit slots."
       },
       "inputs": {
         "input": {
@@ -377,7 +376,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
         "test": "TensorOpTest.DepthToSpaceTest_5",
-        "notes": "Typed ORT coverage includes uint8; this framework stores logical uint8 values in 32-bit slots."
       },
       "inputs": {
         "input": {
@@ -497,7 +496,7 @@
       "name": "ort_standard_int16_dcr_extremes",
       "provenance": {
         "source": "onnx/docs/Operators.md#DepthToSpace",
-        "notes": "ONNX DepthToSpace-13 permits int16; both signed extrema guard the runtime's widened i32 storage route."
       },
       "attrs": { "blocksize": 2, "mode": "DCR" },
       "inputs": {

 {
   "fixtureArrays": {
     "dcr_input_input": [0, 1, 2, 3, 10, 11, 12, 13, 20, 21, 22, 23, 30, 31, 32, 33],
     "onnx_backend_dcr_1x8x2x3_input_input": [0, 1, 2, 3, 4, 5, 9, 10, 11, 12, 13, 14, 18, 19, 20, 21, 22, 23, 27, 28, 29, 30, 31, 32, 36, 37, 38, 39, 40, 41, 45, 46, 47, 48, 49, 50, 54, 55, 56, 57, 58, 59, 63, 64, 65, 66, 67, 68],
     {
       "name": "dcr_high_channel_vec4_1x256x4x4",
       "provenance": {
+        "notes": "A compact high-channel DCR input exercises the vec4 channel-heavy mapping without benchmark-scale spatial dimensions."
       },
       "attrs": { "blocksize": 2, "mode": "DCR" },
       "inputs": {
     {
       "name": "crd_high_channel_vec4_1x256x4x4",
       "provenance": {
+        "notes": "A compact high-channel CRD input exercises the vec4 channel-heavy mapping without benchmark-scale spatial dimensions."
       },
       "attrs": { "blocksize": 2, "mode": "CRD" },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
         "test": "TensorOpTest.DepthToSpaceTest_3<uint8_t>",
+        "notes": "A compact logical uint8 input exercises the default DCR channel-to-space mapping."
       },
       "inputs": {
         "input": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
         "test": "TensorOpTest.DepthToSpaceTest_4",
+        "notes": "Logical uint8 values are stored in 32-bit slots while preserving the typed source case."
       },
       "inputs": {
         "input": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc",
         "test": "TensorOpTest.DepthToSpaceTest_5",
+        "notes": "Logical uint8 values are stored in 32-bit slots while preserving the typed source case."
       },
       "inputs": {
         "input": {
       "name": "ort_standard_int16_dcr_extremes",
       "provenance": {
         "source": "onnx/docs/Operators.md#DepthToSpace",
+        "notes": "Both int16 extrema exercise the package's widened int32 storage while preserving logical signed values."
       },
       "attrs": { "blocksize": 2, "mode": "DCR" },
       "inputs": {