Xenova HF Staff commited on 7 days ago

Commit

410baf2

verified ·

1 Parent(s): 77b755a

sync 2e7068faf55e

Browse files

Files changed (19) hide show

README.md +93 -0
build/webgpu/bench.json +452 -0
build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja +402 -0
build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja +121 -0
build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja +306 -0
build/webgpu/conv-transpose-empty-input.wgsl.jinja +57 -0
build/webgpu/conv-transpose1d-ncl.wgsl.jinja +102 -0
build/webgpu/conv-transpose1d-phase-tiled.wgsl.jinja +102 -0
build/webgpu/conv-transpose2d-col2im.wgsl.jinja +96 -0
build/webgpu/conv-transpose2d-grouped-stride-phase.wgsl.jinja +70 -0
build/webgpu/conv-transpose2d-nchw.wgsl.jinja +124 -0
build/webgpu/conv-transpose2d-phase-gemm-tiled.wgsl.jinja +155 -0
build/webgpu/conv-transpose2d-weight-reorder.wgsl.jinja +57 -0
build/webgpu/conv-transpose3d-col2im.wgsl.jinja +88 -0
build/webgpu/conv-transpose3d-ncdhw.wgsl.jinja +73 -0
build/webgpu/conv-transpose3d-weight-reorder.wgsl.jinja +56 -0
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +31 -0
build/webgpu/test.json +1856 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.ConvTranspose
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 11
+## Description
+Computes the transpose of a convolution, also known as a fractionally strided convolution or deconvolution, from input tensor `X`, filter weights `W`, and an optional bias `B`. Output spatial dimensions follow the stride, dilation, padding, and optional `output_padding` attributes. Supports grouped convolution through `group`.
+See the [ONNX `ConvTranspose` spec](https://onnx.ai/onnx/operators/onnx__ConvTranspose.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Input data tensor of shape `(N x C x D1 x ... x Dn)`, where `N` is batch size and `C` is the number of input channels. | required |
+| `W` | `w` | `T` | — | — | Filter weight tensor of shape `(C x M/group x k1 x ... x kn)`, where `M` is the number of output feature maps. | required |
+| `B` | `bias` | `T` | `1` | — | Optional 1-D bias of length M added to each output channel. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | derived; see description | Output tensor whose spatial dimensions are computed from the input size, kernel shape, strides, dilations, and padding. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `auto_pad` | `"NOTSET"` | Padding mode: `NOTSET` uses explicit pads; `SAME_UPPER` and `SAME_LOWER` make output spatial size equal input size times stride, with any odd extra padding added at the end or beginning respectively; `VALID` applies no padding. |
+| `group` | `1` | Number of groups that input and output channels are divided into for grouped (depthwise) convolution. |
+| `dilations` | — | Dilation factors for each spatial axis; defaults to one on every axis. |
+| `kernel_shape` | — | Kernel dimensions for each spatial axis. When omitted, they are inferred from the spatial dimensions of `W`. |
+| `output_padding` | — | Additional size on the high-index end of each output spatial axis; each value must be smaller than the corresponding stride or dilation. |
+| `output_shape` | — | Requested output spatial dimensions. When present, it must match the declared spatial shape of `Y`. |
+| `pads` | — | Padding at the beginning of every spatial axis followed by padding at the end of every spatial axis; defaults to zeros. |
+| `strides` | — | Stride factors for each spatial axis; defaults to one on every axis. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Device requirements
+Some implementation variants require `subgroup-matrix`, `shader-f16`, and `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
+## 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
+- [`conv-1x1-gemm-tiled-reg.wgsl.jinja`](build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja)
+- [`conv-1x1-gemm-tiled.wgsl.jinja`](build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja)
+- [`conv-1x1-subgroup-matrix.wgsl.jinja`](build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja)
+- [`conv-transpose-empty-input.wgsl.jinja`](build/webgpu/conv-transpose-empty-input.wgsl.jinja)
+- [`conv-transpose1d-ncl.wgsl.jinja`](build/webgpu/conv-transpose1d-ncl.wgsl.jinja)
+- [`conv-transpose1d-phase-tiled.wgsl.jinja`](build/webgpu/conv-transpose1d-phase-tiled.wgsl.jinja)
+- [`conv-transpose2d-col2im.wgsl.jinja`](build/webgpu/conv-transpose2d-col2im.wgsl.jinja)
+- [`conv-transpose2d-grouped-stride-phase.wgsl.jinja`](build/webgpu/conv-transpose2d-grouped-stride-phase.wgsl.jinja)
+- [`conv-transpose2d-nchw.wgsl.jinja`](build/webgpu/conv-transpose2d-nchw.wgsl.jinja)
+- [`conv-transpose2d-phase-gemm-tiled.wgsl.jinja`](build/webgpu/conv-transpose2d-phase-gemm-tiled.wgsl.jinja)
+- [`conv-transpose2d-weight-reorder.wgsl.jinja`](build/webgpu/conv-transpose2d-weight-reorder.wgsl.jinja)
+- [`conv-transpose3d-col2im.wgsl.jinja`](build/webgpu/conv-transpose3d-col2im.wgsl.jinja)
+- [`conv-transpose3d-ncdhw.wgsl.jinja`](build/webgpu/conv-transpose3d-ncdhw.wgsl.jinja)
+- [`conv-transpose3d-weight-reorder.wgsl.jinja`](build/webgpu/conv-transpose3d-weight-reorder.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.ConvTranspose", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 1, 3] },
+  w: { data: wData, shape: [1, 2, 2] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,452 @@

+{
+  "op": "ai.onnx.ConvTranspose",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "conv-transpose-f32-b1c16m16-16x16-k3s2",
+      "preset": "smoke",
+      "vars": {
+        "batch": 1,
+        "inChannels": 16,
+        "outChannels": 16,
+        "inH": 16,
+        "inW": 16,
+        "kernelH": 3,
+        "kernelW": 3,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 16, 16, 16], "dtype": "float32", "dist": "normal", "seed": 114, "scale": 0.2 },
+        "w": { "shape": [16, 16, 3, 3], "dtype": "float32", "dist": "normal", "seed": 115, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 16, 31, 31], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-non-vector-bias-b1c3m5-32x32-k3s2",
+      "preset": "edge",
+      "vars": {
+        "batch": 1,
+        "inChannels": 3,
+        "outChannels": 5,
+        "inH": 32,
+        "inW": 32,
+        "kernelH": 3,
+        "kernelW": 3,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 3, 32, 32], "dtype": "float32", "dist": "normal", "seed": 116, "scale": 0.2 },
+        "w": { "shape": [3, 5, 3, 3], "dtype": "float32", "dist": "normal", "seed": 117, "scale": 0.1 },
+        "bias": { "shape": [5], "dtype": "float32", "dist": "normal", "seed": 118, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 5, 63, 63], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f16-b1c64m64-32x32-k3s2",
+      "preset": "smoke",
+      "vars": {
+        "batch": 1,
+        "inChannels": 64,
+        "outChannels": 64,
+        "inH": 32,
+        "inW": 32,
+        "kernelH": 3,
+        "kernelW": 3,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 64, 32, 32], "dtype": "float16", "dist": "normal", "seed": 211, "scale": 0.2 },
+        "w": { "shape": [64, 64, 3, 3], "dtype": "float16", "dist": "normal", "seed": 212, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 63, 63], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-b1c64m64-32x32-k3s2",
+      "preset": "smoke",
+      "vars": {
+        "batch": 1,
+        "inChannels": 64,
+        "outChannels": 64,
+        "inH": 32,
+        "inW": 32,
+        "kernelH": 3,
+        "kernelW": 3,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 64, 32, 32], "dtype": "float32", "dist": "normal", "seed": 211, "scale": 0.2 },
+        "w": { "shape": [64, 64, 3, 3], "dtype": "float32", "dist": "normal", "seed": 212, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 63, 63], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f16-large-upsampler-b1c128m128-64x64-k4s2",
+      "preset": "smoke",
+      "vars": {
+        "batch": 1,
+        "inChannels": 128,
+        "outChannels": 128,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 4,
+        "kernelW": 4,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 128, 64, 64], "dtype": "float16", "dist": "normal", "seed": 213, "scale": 0.2 },
+        "w": { "shape": [128, 128, 4, 4], "dtype": "float16", "dist": "normal", "seed": 214, "scale": 0.05 }
+      },
+      "outputs": { "y": { "shape": [1, 128, 128, 128], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-depthwise-g64-32x32-k4s2",
+      "preset": "smoke",
+      "vars": {
+        "batch": 1,
+        "inChannels": 64,
+        "outChannels": 64,
+        "inChannelsPerGroup": 1,
+        "inH": 32,
+        "inW": 32,
+        "kernelH": 4,
+        "kernelW": 4,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1,
+        "group": 64
+      },
+      "attrs": { "group": 64, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 64, 32, 32], "dtype": "float32", "dist": "normal", "seed": 215, "scale": 0.2 },
+        "w": { "shape": [64, 1, 4, 4], "dtype": "float32", "dist": "normal", "seed": 216, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 64, 64], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannelsPerGroup * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-depthwise-pointwise-g64-512x512-dispatch-fold",
+      "preset": "edge",
+      "provenance": {
+        "notes": "Realistic depthwise pointwise projection at 16,777,216 outputs. Exercises the two-dimensional dispatch fold while exposing scalar coordinate/division overhead in the generic grouped kernel."
+      },
+      "vars": { "count": 16777216 },
+      "attrs": { "group": 64 },
+      "inputs": {
+        "x": { "shape": [1, 64, 512, 512], "dtype": "float32", "dist": "normal", "seed": 221, "scale": 0.2 },
+        "w": { "shape": [64, 1, 1, 1], "dtype": "float32", "dist": "normal", "seed": 222, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 512, 512], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "(2 * args.count + 64) * 4" }] }
+    },
+    {
+      "name": "conv-transpose-f32-1d-b1c128m128-256-k4s2",
+      "preset": "smoke",
+      "vars": { "batch": 1, "inChannels": 128, "outChannels": 128, "inW": 256, "kernelW": 4, "strideW": 2, "padW": 1 },
+      "attrs": { "strides": [2], "pads": [1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 128, 256], "dtype": "float32", "dist": "normal", "seed": 217, "scale": 0.2 },
+        "w": { "shape": [128, 128, 4], "dtype": "float32", "dist": "normal", "seed": 218, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 128, 512], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * args.batch * args.outChannels * args.inW * args.inChannels * args.kernelW" }
+        ]
+      }
+    },
+    {
+      "name": "neural-audio-1d-f32-b1c256m256-w1024-k4s2-pathology",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for performance coverage",
+        "notes": "Realistic neural-audio upsampling stage for the direct NCW fallback. Every output sample serially scans 256 input channels and four kernel taps."
+      },
+      "vars": { "batch": 1, "inChannels": 256, "outChannels": 256, "inW": 1024, "kernelW": 4, "strideW": 2, "padW": 1 },
+      "attrs": { "strides": [2], "pads": [1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 256, 1024], "dtype": "float32", "dist": "normal", "seed": 219, "scale": 0.2 },
+        "w": { "shape": [256, 256, 4], "dtype": "float32", "dist": "normal", "seed": 220, "scale": 0.05 }
+      },
+      "outputs": { "y": { "shape": [1, 256, 2048], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * args.batch * args.outChannels * args.inW * args.inChannels * args.kernelW" }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-3d-naive-b1c16m16-16x16x16-k3",
+      "preset": "stress",
+      "vars": {
+        "batch": 1,
+        "inChannels": 16,
+        "outChannels": 16,
+        "inD": 16,
+        "inH": 16,
+        "inW": 16,
+        "kernelD": 3,
+        "kernelH": 3,
+        "kernelW": 3
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "shape": [1, 16, 16, 16, 16], "dtype": "float32", "dist": "normal", "seed": 331, "scale": 0.2 },
+        "w": { "shape": [16, 16, 3, 3, 3], "dtype": "float32", "dist": "normal", "seed": 332, "scale": 0.05 }
+      },
+      "outputs": { "y": { "shape": [1, 16, 18, 18, 18], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inD * args.inH * args.inW * args.inChannels * args.kernelD * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "medical-decoder-3d-f32-b1c32m32-32x32x32-k3-pathology",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for performance coverage",
+        "notes": "Realistic 3D decoder stage for the direct NCDHW fallback. Each output voxel serially scans all 32 input channels and the 3x3x3 kernel."
+      },
+      "vars": {
+        "batch": 1,
+        "inChannels": 32,
+        "outChannels": 32,
+        "inD": 32,
+        "inH": 32,
+        "inW": 32,
+        "kernelD": 3,
+        "kernelH": 3,
+        "kernelW": 3
+      },
+      "attrs": { "pads": [1, 1, 1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 32, 32, 32, 32], "dtype": "float32", "dist": "normal", "seed": 333, "scale": 0.2 },
+        "w": { "shape": [32, 32, 3, 3, 3], "dtype": "float32", "dist": "normal", "seed": 334, "scale": 0.05 }
+      },
+      "outputs": { "y": { "shape": [1, 32, 32, 32, 32], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inD * args.inH * args.inW * args.inChannels * args.kernelD * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-grouped-naive-g4-c32m64-48x48-k4s2",
+      "preset": "stress",
+      "vars": {
+        "batch": 1,
+        "inChannels": 32,
+        "outChannels": 64,
+        "inChannelsPerGroup": 8,
+        "inH": 48,
+        "inW": 48,
+        "kernelH": 4,
+        "kernelW": 4,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1,
+        "group": 4
+      },
+      "attrs": { "group": 4, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 32, 48, 48], "dtype": "float32", "dist": "normal", "seed": 341, "scale": 0.2 },
+        "w": { "shape": [32, 16, 4, 4], "dtype": "float32", "dist": "normal", "seed": 342, "scale": 0.05 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 96, 96], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannelsPerGroup * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "grouped-bias-g4-c32m64-48x48-k4s2-generic-pathology",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for variant coverage",
+        "notes": "Adding channel bias excludes nchw2d_grouped_stride_phase and selects nchw2d_grouped_bias; the paired no-bias case isolates this optional-input route boundary."
+      },
+      "vars": {
+        "batch": 1,
+        "inChannels": 32,
+        "outChannels": 64,
+        "inChannelsPerGroup": 8,
+        "inH": 48,
+        "inW": 48,
+        "kernelH": 4,
+        "kernelW": 4,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 1,
+        "padW": 1,
+        "group": 4
+      },
+      "attrs": { "group": 4, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "shape": [1, 32, 48, 48], "dtype": "float32", "dist": "normal", "seed": 5104, "scale": 0.2 },
+        "w": { "shape": [32, 16, 4, 4], "dtype": "float32", "dist": "normal", "seed": 5105, "scale": 0.05 },
+        "bias": { "shape": [64], "dtype": "float32", "dist": "normal", "seed": 5106, "scale": 0.05 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 96, 96], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannelsPerGroup * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-dpt-reassemble-48ch-37x49-k4s4",
+      "preset": "all",
+      "provenance": {
+        "source": "Depth Anything V2 Small neck reassemble",
+        "notes": "Depth Anything V2 Small's factor-4 reassemble at 518x686. ConvTranspose contracts the projected neck width (48), not the backbone width."
+      },
+      "vars": {
+        "batch": 1,
+        "inChannels": 48,
+        "outChannels": 48,
+        "inH": 37,
+        "inW": 49,
+        "kernelH": 4,
+        "kernelW": 4,
+        "strideH": 4,
+        "strideW": 4,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [4, 4] },
+      "inputs": {
+        "x": { "shape": [1, 48, 37, 49], "dtype": "float32", "dist": "normal", "seed": 9421, "scale": 0.2 },
+        "w": { "shape": [48, 48, 4, 4], "dtype": "float32", "dist": "normal", "seed": 9422, "scale": 0.1 },
+        "bias": { "shape": [48], "dtype": "float32", "dist": "normal", "seed": 9423, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 48, 148, 196], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-dpt-reassemble-96ch-37x49-k2s2",
+      "preset": "all",
+      "provenance": {
+        "source": "Depth Anything V2 Small neck reassemble",
+        "notes": "The factor-2 reassemble of the same model, likewise contracting over the neck width (96). Four taps per output element, one of them real."
+      },
+      "vars": {
+        "batch": 1,
+        "inChannels": 96,
+        "outChannels": 96,
+        "inH": 37,
+        "inW": 49,
+        "kernelH": 2,
+        "kernelW": 2,
+        "strideH": 2,
+        "strideW": 2,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [2, 2] },
+      "inputs": {
+        "x": { "shape": [1, 96, 37, 49], "dtype": "float32", "dist": "normal", "seed": 9431, "scale": 0.2 },
+        "w": { "shape": [96, 96, 2, 2], "dtype": "float32", "dist": "normal", "seed": 9432, "scale": 0.1 },
+        "bias": { "shape": [96], "dtype": "float32", "dist": "normal", "seed": 9433, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 96, 74, 98], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * args.inH * args.inW * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,402 @@

+// 1x1 / im2col Conv routed as a batched GEMM:
+// Y[b] (M x N) = W (M x K) * X[b] (K x N), with M = outChannels,
+// K = inChannels * kh * kw after im2col, and N = H * W. Register-blocked
+// workgroup lanes compute TM x TN micro-tiles. The defaults use 16x16 lanes with
+// an 8x4 micro-tile over a 64x64 output tile. Weight A is shared across the batch;
+// input B and the output use the dispatch.z batch offset. Accumulation is f32;
+// the optional epilogue applies bias, residual Z, and activation in the
+// accumulator domain on store. Every dimension is bounds-checked.
+// f16 operands remain packed in workgroup memory and widen only at the FMA site;
+// the f32 accumulation order is retained while shared traffic is halved.
+//
+// Both shared tiles are stored as vec4 groups along the axis each lane consumes
+// contiguously: tileA groups four K values for one output row, tileB groups four
+// N columns for one K. One micro-tile step therefore reads TM + 4 vector words
+// instead of 4 * (TM + TN) scalars, so a shared word feeds four times as many
+// FMAs and the K loop runs four accumulation steps per iteration.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% set hasActivation = hasActivation is defined and hasActivation %}
+{% if hasActivation %}
+// Apply the fused activation in the f32 accumulator before the single output
+// cast, avoiding an intermediate convolution tensor.
+fn fused_act(v: f32) -> f32 {
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+}
+{% endif %}
+{% set tileT = "f16" if usesF16 else "f32" -%}
+{%- set GEMM_BK = gemmKTile if gemmKTile is defined else 16 %}
+{% set GEMM_TM = gemmThreadRows if gemmThreadRows is defined else 8 %}
+{% set GEMM_WG_X = gemmWorkgroupX if gemmWorkgroupX is defined else 16 %}
+{% set GEMM_BM = gemmMTile if gemmMTile is defined else 64 %}
+{% set GEMM_BN = gemmNTile if gemmNTile is defined else GEMM_WG_X * 4 %}
+{% set GEMM_TN = gemmThreadColumns if gemmThreadColumns is defined else (GEMM_BN / GEMM_WG_X)|int %}
+{% set GEMM_WG_Y = gemmWorkgroupY if gemmWorkgroupY is defined else (GEMM_BM / GEMM_TM)|int %}
+{% set components = ["x", "y", "z", "w"] %}
+const BK: u32 = {{ GEMM_BK }}u;
+const BM: u32 = {{ GEMM_BM }}u;
+const BN: u32 = {{ GEMM_BN }}u;
+const TM: u32 = {{ GEMM_TM }}u;
+const TN: u32 = {{ GEMM_TN }}u;
+const WG_X: u32 = {{ GEMM_WG_X }}u;
+const WG_SIZE: u32 = {{ GEMM_WG_X * GEMM_WG_Y }}u;
+// Vector words per tile row: A groups K, B groups N.
+const AK_VECS: u32 = BK / 4u;
+const BN_VECS: u32 = BN / 4u;
+// implicitIm2col leaves B unmaterialized. While staging a tile, the kernel
+// decodes each (k, n) column-matrix coordinate into a raw NCHW input address.
+// Overlapping windows may reread input values, trading address arithmetic and
+// cache traffic for the storage and bandwidth of the expanded column matrix.
+{% set implicitIm2col = implicitIm2col is defined and implicitIm2col %}
+{% set fusedNarrowProjection = fusedNarrowProjection is defined and fusedNarrowProjection %}
+{% if fusedNarrowProjection %}
+{% set projectionChannels = projectionOutChannels %}
+{% set projectionInputAct = inputActivation %}
+{% set projectionOutputAct = outputActivation %}
+{% set projectionScaled = hasOutputScale %}
+{% set projectionBiased = hasProjectionBias %}
+{% endif %}
+{% set narrowProjectionTile = "tileB" if GEMM_BK >= GEMM_BM else "projectionTile" %}
+{% set splitKValue = splitK if splitK is defined else 1 %}
+{% set splitKPartial = splitKValue > 1 %}
+{% set gemmKLimit = "kEnd" if splitKPartial else "K" %}
+{% set implicitGatherMode = tunables.IMPLICIT_TILED_GATHER_MODE if tunables.IMPLICIT_TILED_GATHER_MODE is defined else 0 %}
+{% set implicitGatherWgSize = GEMM_WG_X * GEMM_WG_Y %}
+{% set implicitGatherColumnVecs = (GEMM_BN / 4)|int %}
+{% set implicitGatherKChunks = (implicitGatherWgSize / implicitGatherColumnVecs)|int %}
+{% set implicitGatherDivides = implicitGatherWgSize % implicitGatherColumnVecs == 0 and implicitGatherKChunks > 0 and GEMM_BK % implicitGatherKChunks == 0 %}
+{% set useCarriedImplicitGather = implicitIm2col and implicitGatherMode == 1 and implicitGatherDivides %}
+{% if useCarriedImplicitGather %}
+{% set implicitGatherLoadWidth = (GEMM_BK / implicitGatherKChunks)|int %}
+{% endif %}
+{% if implicitIm2col %}
+const CONV_KERNEL_H: u32 = {{ convKernelH }}u;
+const CONV_KERNEL_W: u32 = {{ convKernelW }}u;
+const CONV_KSIZE: u32 = CONV_KERNEL_H * CONV_KERNEL_W;
+const CONV_STRIDE_H: u32 = {{ convStrideH }}u;
+const CONV_STRIDE_W: u32 = {{ convStrideW }}u;
+const CONV_DILATION_H: u32 = {{ convDilationH }}u;
+const CONV_DILATION_W: u32 = {{ convDilationW }}u;
+const CONV_PAD_TOP: i32 = {{ convPadTop }};
+const CONV_PAD_LEFT: i32 = {{ convPadLeft }};
+const CONV_IN_H: u32 = {{ convInH }}u;
+const CONV_IN_W: u32 = {{ convInW }}u;
+const CONV_OUT_W: u32 = {{ convOutW }}u;
+const CONV_IN_CHANNELS: u32 = {{ convInChannels }}u;
+{% endif %}
+{% if fusedNarrowProjection %}
+const PROJECTION_OUT_C: u32 = {{ projectionChannels }}u;
+{% else %}{% set emitConvStoreOut = not splitKPartial %}{% if emitConvStoreOut | default(true) %}fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
+  if (m >= params.M || n >= params.N) {
+    return;
+  }
+  var v = raw;
+{% if hasBias %}
+  v = v + f32(bias[m]);
+{% endif %}
+{% if hasActivation %}
+  v = fused_act(v);
+{% endif %}
+  y[yBase + m * params.N + n] = {{ T }}(v);
+}
+{% endif %}
+{% endif %}
+var<workgroup> tileA: array<array<vec4<{{ tileT }}>, AK_VECS>, BM>;
+var<workgroup> tileB: array<array<vec4<{{ tileT }}>, BN_VECS>, BK>;
+{% if fusedNarrowProjection and GEMM_BK < GEMM_BM %}
+// BK16 cannot reuse the 16-row input tile to publish a BM32 intermediate.
+// The extra 32x64 f32 tile keeps the total at 14 KiB, below WebGPU's
+// guaranteed 16 KiB workgroup-storage floor.
+var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
+{% endif %}
+{% macro publish_projection_input(localRow, localVec, component, globalRow, globalColumn, raw) %}
+  if ({{ globalRow }} < M && {{ globalColumn }} < N) {
+    var projectionInput = {{ raw }};
+{% if hasBias %}
+    projectionInput = projectionInput + f32(bias[{{ globalRow }}]);
+{% endif %}
+{% if projectionInputAct == "relu" %}
+    projectionInput = projection_relu(projectionInput);
+{% endif %}
+    {{ narrowProjectionTile }}[{{ localRow }}][{{ localVec }}].{{ component }} = projectionInput;
+  } else {
+    {{ narrowProjectionTile }}[{{ localRow }}][{{ localVec }}].{{ component }} = 0.0;
+  }
+{%- endmacro %}
+{% macro load_a_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
+    for (var linear = li; linear < BM * AK_VECS; linear += WG_SIZE) {
+      let ar = linear / AK_VECS;
+      let ac4 = linear % AK_VECS;
+      let am = {{ rowBase }} + ar;
+      let ak = {{ columnBase }} + ac4 * 4u;
+      var av = vec4<{{ tileT }}>({{ tileT }}(0.0));
+      if (am < {{ rowLimit }}) {
+        let aBase = am * K + ak;
+        if (ak + 3u < {{ columnLimit }}) {
+          av = vec4<{{ tileT }}>({{ tileT }}(w[aBase]), {{ tileT }}(w[aBase + 1u]), {{ tileT }}(w[aBase + 2u]), {{ tileT }}(w[aBase + 3u]));
+        } else {
+          for (var t = 0u; t < 4u; t = t + 1u) {
+            if (ak + t < {{ columnLimit }}) {
+              av[t] = {{ tileT }}(w[aBase + t]);
+            }
+          }
+        }
+      }
+      tileA[ar][ac4] = av;
+    }
+{%- endmacro %}
+{% macro load_b_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
+    for (var linear = li; linear < BK * BN_VECS; linear += WG_SIZE) {
+      let br = linear / BN_VECS;
+      let bc4 = linear % BN_VECS;
+      let bk = {{ rowBase }} + br;
+      let bn = {{ columnBase }} + bc4 * 4u;
+      var bvec = vec4<{{ tileT }}>({{ tileT }}(0.0));
+      if (bk < {{ rowLimit }}) {
+        let bBase = xBatchBase + bk * N + bn;
+        if (bn + 3u < {{ columnLimit }}) {
+          bvec = vec4<{{ tileT }}>({{ tileT }}(xm[bBase]), {{ tileT }}(xm[bBase + 1u]), {{ tileT }}(xm[bBase + 2u]), {{ tileT }}(xm[bBase + 3u]));
+        } else {
+          for (var t = 0u; t < 4u; t = t + 1u) {
+            if (bn + t < {{ columnLimit }}) {
+              bvec[t] = {{ tileT }}(xm[bBase + t]);
+            }
+          }
+        }
+      }
+      tileB[br][bc4] = bvec;
+    }
+{%- endmacro %}
+{% macro load_b_implicit_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
+    for (var linear = li; linear < BK * BN_VECS; linear += WG_SIZE) {
+      let br = linear / BN_VECS;
+      let bc4 = linear % BN_VECS;
+      let bk = {{ rowBase }} + br;
+      var bvec = vec4<{{ tileT }}>({{ tileT }}(0.0));
+      if (bk < {{ rowLimit }}) {
+        // k is the OIHW weight flattening (ic * KH + kh) * KW + kw; n is the output position.
+        let ic = bk / CONV_KSIZE;
+        let kq = bk % CONV_KSIZE;
+        let kh = kq / CONV_KERNEL_W;
+        let kw = kq % CONV_KERNEL_W;
+        let planeBase = (batch * CONV_IN_CHANNELS + ic) * CONV_IN_H;
+        for (var t = 0u; t < 4u; t = t + 1u) {
+          let bn = {{ columnBase }} + bc4 * 4u + t;
+          if (bn < {{ columnLimit }}) {
+            let oh = bn / CONV_OUT_W;
+            let ow = bn % CONV_OUT_W;
+            let ih = i32(oh * CONV_STRIDE_H + kh * CONV_DILATION_H) - CONV_PAD_TOP;
+            let iw = i32(ow * CONV_STRIDE_W + kw * CONV_DILATION_W) - CONV_PAD_LEFT;
+            if (ih >= 0 && ih < i32(CONV_IN_H) && iw >= 0 && iw < i32(CONV_IN_W)) {
+              bvec[t] = {{ tileT }}(xm[(planeBase + u32(ih)) * CONV_IN_W + u32(iw)]);
+            }
+          }
+        }
+      }
+      tileB[br][bc4] = bvec;
+    }
+{%- endmacro %}
+{% macro load_b_implicit_carried(rowBase, columnBase, rowLimit, columnLimit) %}
+    let bColVec = li / {{ implicitGatherKChunks }}u;
+    let bChunk = li % {{ implicitGatherKChunks }}u;
+    let bLocalK = bChunk * {{ implicitGatherLoadWidth }}u;
+    let bStartK = {{ rowBase }} + bLocalK;
+    let bCol0 = {{ columnBase }} + bColVec * 4u;
+    let bCols = vec4<u32>(bCol0, bCol0 + 1u, bCol0 + 2u, bCol0 + 3u);
+    let inColumn = bCols < vec4<u32>({{ columnLimit }});
+    let ih0 = vec4<i32>((bCols / vec4<u32>(CONV_OUT_W)) * CONV_STRIDE_H) - vec4<i32>(CONV_PAD_TOP);
+    let iw0 = vec4<i32>((bCols % vec4<u32>(CONV_OUT_W)) * CONV_STRIDE_W) - vec4<i32>(CONV_PAD_LEFT);
+    var carriedK = bStartK;
+    var carriedIc = carriedK / CONV_KSIZE;
+    let carriedKq = carriedK % CONV_KSIZE;
+    var carriedKh = carriedKq / CONV_KERNEL_W;
+    var carriedKw = carriedKq % CONV_KERNEL_W;
+    var carriedIh = ih0 + vec4<i32>(i32(carriedKh * CONV_DILATION_H));
+    var carriedIw = iw0 + vec4<i32>(i32(carriedKw * CONV_DILATION_W));
+    // Interior words can carry four addresses. Padding, the N tail, and the
+    // final K tile retain the coordinate-checked loop below and publish exact
+    // zeros.
+    let interior = all(inColumn)
+      && bStartK + {{ implicitGatherLoadWidth }}u <= {{ rowLimit }}
+      && all(ih0 >= vec4<i32>(0))
+      && all(ih0 + vec4<i32>(i32((CONV_KERNEL_H - 1u) * CONV_DILATION_H)) < vec4<i32>(i32(CONV_IN_H)))
+      && all(iw0 >= vec4<i32>(0))
+      && all(iw0 + vec4<i32>(i32((CONV_KERNEL_W - 1u) * CONV_DILATION_W)) < vec4<i32>(i32(CONV_IN_W)));
+    if (interior) {
+      let batchBase = batch * CONV_IN_CHANNELS * CONV_IN_H * CONV_IN_W;
+      let colStep = i32(CONV_DILATION_W);
+      let rowStep = i32(CONV_DILATION_H * CONV_IN_W);
+      let planeStep = i32(CONV_IN_H * CONV_IN_W);
+      let kwWrap = colStep * i32(CONV_KERNEL_W) - rowStep;
+      let khWrap = rowStep * i32(CONV_KERNEL_H) - planeStep;
+      var carriedAddress = vec4<i32>(i32(carriedIc) * planeStep) + carriedIh * vec4<i32>(i32(CONV_IN_W)) + carriedIw;
+      for (var i = 0u; i < {{ implicitGatherLoadWidth }}u; i = i + 1u) {
+        tileB[bLocalK + i][bColVec] = vec4<{{ tileT }}>(
+{% for component in components %}
+          {{ tileT }}(xm[batchBase + u32(carriedAddress.{{ component }})]){% if not loop.last %},{% endif %}
+{% endfor %}
+        );
+        carriedAddress = carriedAddress + vec4<i32>(colStep);
+        carriedKw = carriedKw + 1u;
+        if (carriedKw == CONV_KERNEL_W) {
+          carriedKw = 0u;
+          carriedAddress = carriedAddress - vec4<i32>(kwWrap);
+          carriedKh = carriedKh + 1u;
+          if (carriedKh == CONV_KERNEL_H) {
+            carriedKh = 0u;
+            carriedAddress = carriedAddress - vec4<i32>(khWrap);
+          }
+        }
+      }
+    } else {
+      for (var i = 0u; i < {{ implicitGatherLoadWidth }}u; i = i + 1u) {
+        var carriedValue = vec4<{{ tileT }}>({{ tileT }}(0.0));
+        let carriedRow = (batch * CONV_IN_CHANNELS + carriedIc) * CONV_IN_H;
+        let carriedOk = inColumn
+          & (carriedIh >= vec4<i32>(0)) & (carriedIh < vec4<i32>(i32(CONV_IN_H)))
+          & (carriedIw >= vec4<i32>(0)) & (carriedIw < vec4<i32>(i32(CONV_IN_W)));
+        if (carriedK < {{ rowLimit }}) {
+{% for component in components %}
+          if (carriedOk.{{ component }}) {
+            carriedValue.{{ component }} = {{ tileT }}(xm[(carriedRow + u32(carriedIh.{{ component }})) * CONV_IN_W + u32(carriedIw.{{ component }})]);
+          }
+{% endfor %}
+        }
+        tileB[bLocalK + i][bColVec] = carriedValue;
+        carriedK = carriedK + 1u;
+        carriedKw = carriedKw + 1u;
+        carriedIw = carriedIw + vec4<i32>(i32(CONV_DILATION_W));
+        if (carriedKw == CONV_KERNEL_W) {
+          carriedKw = 0u;
+          carriedIw = iw0;
+          carriedKh = carriedKh + 1u;
+          carriedIh = carriedIh + vec4<i32>(i32(CONV_DILATION_H));
+          if (carriedKh == CONV_KERNEL_H) {
+            carriedKh = 0u;
+            carriedIh = ih0;
+            carriedIc = carriedIc + 1u;
+          }
+        }
+      }
+    }
+{%- endmacro %}
+@compute @workgroup_size({{ GEMM_WG_X }}, {{ GEMM_WG_Y }}, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let M = params.M;
+  let K = params.K;
+  let N = params.N;
+  let mBase = wg.y * BM;
+  let nBase = wg.x * BN;
+  let batch = wg.z;
+{% if not implicitIm2col %}
+  let xBatchBase = batch * K * N;
+{% endif %}
+  let li = lid.y * WG_X + lid.x;
+  // Explicit vectors give the backend fixed register indices and expose TM
+  // independent FMA chains.
+{% for row in range(GEMM_TM) %}
+  var acc{{ row }} = vec{{ GEMM_TN }}<f32>(0.0);
+{% endfor %}
+  let numTiles = (K + BK - 1u) / BK;
+  for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
+    let kBase = kt * BK;
+{{ load_a_vec4("mBase", "kBase", "M", gemmKLimit) }}
+{% if useCarriedImplicitGather %}
+{{ load_b_implicit_carried("kBase", "nBase", gemmKLimit, "N") }}
+{% elif implicitIm2col %}
+{{ load_b_implicit_vec4("kBase", "nBase", gemmKLimit, "N") }}
+{% else %}
+{{ load_b_vec4("kBase", "nBase", gemmKLimit, "N") }}
+{% endif %}
+    workgroupBarrier();
+    // One iteration consumes four K steps: TM vector words of A and four of B
+    // cover TM x TN x 4 FMAs.
+    let aRow = lid.y * TM;
+    for (var kv: u32 = 0u; kv < AK_VECS; kv = kv + 1u) {
+{% for row in range(GEMM_TM) %}
+      let av{{ row }} = vec4<f32>(tileA[aRow + {{ row }}u][kv]);
+{% endfor %}
+{% for sub in range(4) %}
+      let bv{{ sub }} = vec{{ GEMM_TN }}<f32>(tileB[kv * 4u + {{ sub }}u][lid.x]);
+{% for row in range(GEMM_TM) %}
+      acc{{ row }} += vec{{ GEMM_TN }}<f32>(av{{ row }}.{{ components[sub] }}) * bv{{ sub }};
+{% endfor %}
+{% endfor %}
+    }
+    workgroupBarrier();
+  }
+{% if not fusedNarrowProjection %}
+  let yBatchBase = batch * M * N;
+{% endif %}
+  let m0 = mBase + lid.y * TM;
+  let n0 = nBase + lid.x * TN;
+{% for row in range(GEMM_TM) %}
+{% for column in range(GEMM_TN) %}
+{% if fusedNarrowProjection %}
+{{ publish_projection_input("lid.y * TM + " ~ row ~ "u", "lid.x * " ~ ((GEMM_TN / 4)|int) ~ "u + " ~ ((column / 4)|int) ~ "u", components[column % 4], "m0 + " ~ row ~ "u", "n0 + " ~ column ~ "u", "acc" ~ row ~ "." ~ components[column]) }}
+{% else %}
+  store_out(m0 + {{ row }}u, n0 + {{ column }}u, yBatchBase, acc{{ row }}.{{ components[column] }});
+{% endif %}
+{% endfor %}
+{% endfor %}{% if fusedNarrowProjection %}
+  // Conv's final K-tile barrier makes tileB dead before it becomes the
+  // intermediate tile. Every lane publishes its unique micro-tile, then one
+  // lane per spatial vector word consumes all logical M rows in increasing
+  // order, four adjacent columns at a time. Invalid M/N-tail cells are
+  // initialized above, so the barrier is uniform and no lane can observe a
+  // stale input-tile value.
+  workgroupBarrier();
+  if (li < BN_VECS) {
+    let projectionNBase = nBase + li * 4u;
+{% for oc in range(projectionChannels) %}
+    var projectionAcc{{ oc }} = vec4<f32>(0.0);
+{% endfor %}
+    for (var channel = 0u; channel < M; channel = channel + 1u) {
+      let projectionValues = vec4<f32>({{ narrowProjectionTile }}[channel][li]);
+{% for oc in range(projectionChannels) %}
+      projectionAcc{{ oc }} = projectionAcc{{ oc }} + projectionValues * f32(projectionW[{{ oc }}u * M + channel]);
+{% endfor %}
+    }
+    let projectionYBase = batch * PROJECTION_OUT_C * N;
+{% for oc in range(projectionChannels) %}
+    let projected{{ oc }} = projectionAcc{{ oc }}{% if projectionBiased %} + vec4<f32>(f32(projectionBias[{{ oc }}u])){% endif %};
+{% endfor %}
+{% for column in range(4) %}
+    let projectionN{{ column }} = projectionNBase + {{ column }}u;
+    if (projectionN{{ column }} < N) {
+{% for oc in range(projectionChannels) %}
+{% if projectionOutputAct == "relu" %}
+      let activated{{ oc }}_{{ column }} = projection_relu(projected{{ oc }}.{{ components[column] }});
+{% elif projectionOutputAct == "sigmoid" %}
+      let activated{{ oc }}_{{ column }} = sigmoid_safe(projected{{ oc }}.{{ components[column] }});
+{% else %}
+      let activated{{ oc }}_{{ column }} = projected{{ oc }}.{{ components[column] }};
+{% endif %}
+{% if projectionScaled %}
+      y[projectionYBase + {{ oc }}u * N + projectionN{{ column }}] = activated{{ oc }}_{{ column }} * params.outputScale;
+{% else %}
+      y[projectionYBase + {{ oc }}u * N + projectionN{{ column }}] = activated{{ oc }}_{{ column }};
+{% endif %}
+{% endfor %}
+    }
+{% endfor %}
+  }
+{% endif %}
+}

build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,121 @@

+// 1x1 Conv routed as a batched GEMM: Y[b] (M x N) = W (M x K) * X[b] (K x N)
+// with M = outChannels, K = inChannels, N = H*W. For a 1x1 kernel with
+// group == 1, stride 1, no padding, and dilation 1, the NCHW input is already
+// the [K, N] matrix and the OIHW weight is already the [M, K] matrix — no
+// data movement, only index arithmetic.
+//
+// Register-blocked 32x32 shared-memory tile (16x16 threads, each computes a 2x2
+// micro-tile, with a selectable K tile (16 by default), the same structure as
+// the general tiled matmul path — each shared element feeds 2 FMAs
+// and a 32-wide column tile is staged once instead of re-reading W per output
+// column. The weight matrix (A) is shared across the batch (no batch stride);
+// the input (B) and output carry the batch offset via dispatch.z. f32
+// accumulation, cast to the output element type on store; the optional bias is
+// per output channel (M row), added after the full K accumulation per ONNX Conv
+// semantics. f16 operands stay packed in workgroup memory and widen only when
+// consumed. Fully bounds-checked (any M/N/K).
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% set hasActivation = hasActivation is defined and hasActivation %}
+{% if hasActivation %}
+// Apply the fused activation in the f32 accumulator before the single output
+// cast, avoiding an intermediate convolution tensor.
+fn fused_act(v: f32) -> f32 {
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+}
+{% endif %}
+{% set tileT = "f16" if usesF16 else "f32" -%}
+const BK: u32 = 16u;
+const BM: u32 = 32u;
+const BN: u32 = 32u;
+// Store one output element with the optional bias/residual/activation epilogue in the
+// f32 accumulator domain: Y = activation(conv + bias + Z).
+fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
+  if (m >= params.M || n >= params.N) {
+    return;
+  }
+  var v = raw;
+{% if hasBias %}
+  v = v + f32(bias[m]);
+{% endif %}
+{% if hasActivation %}
+  v = fused_act(v);
+{% endif %}
+  y[yBase + m * params.N + n] = {{ T }}(v);
+}
+var<workgroup> tileA: array<array<{{ tileT }}, BK>, BM>;
+var<workgroup> tileB: array<array<{{ tileT }}, BN>, BK>;
+@compute @workgroup_size(16, 16, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let M = params.M;
+  let K = params.K;
+  let N = params.N;
+  let mBase = wg.y * BM;
+  let nBase = wg.x * BN;
+  let batch = wg.z;
+  let xBatchBase = batch * K * N;
+  let li = lid.y * 16u + lid.x;
+  var acc00: f32 = 0.0;
+  var acc01: f32 = 0.0;
+  var acc10: f32 = 0.0;
+  var acc11: f32 = 0.0;
+  let numTiles = (K + BK - 1u) / BK;
+  for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
+    let kBase = kt * BK;
+    // Cooperative load: 32x16 A(=W) tile + 16x32 B(=X[batch]) tile, 256 threads x 2 each.
+    for (var e: u32 = 0u; e < (BM * BK) / 256u; e = e + 1u) {
+      let idx = li + e * 256u;
+      let ar = idx / BK;
+      let ac = idx % BK;
+      let am = mBase + ar;
+      let ak = kBase + ac;
+      if (am < M && ak < K) {
+        tileA[ar][ac] = {{ tileT }}(w[am * K + ak]);
+      } else {
+        tileA[ar][ac] = {{ tileT }}(0.0);
+      }
+      let br = idx / BN;
+      let bc = idx % BN;
+      let bk = kBase + br;
+      let bn = nBase + bc;
+      if (bk < K && bn < N) {
+        tileB[br][bc] = {{ tileT }}(xm[xBatchBase + bk * N + bn]);
+      } else {
+        tileB[br][bc] = {{ tileT }}(0.0);
+      }
+    }
+    workgroupBarrier();
+    for (var kk: u32 = 0u; kk < BK; kk = kk + 1u) {
+      let a0 = f32(tileA[lid.y * 2u][kk]);
+      let a1 = f32(tileA[lid.y * 2u + 1u][kk]);
+      let b0 = f32(tileB[kk][lid.x * 2u]);
+      let b1 = f32(tileB[kk][lid.x * 2u + 1u]);
+      acc00 = acc00 + a0 * b0;
+      acc01 = acc01 + a0 * b1;
+      acc10 = acc10 + a1 * b0;
+      acc11 = acc11 + a1 * b1;
+    }
+    workgroupBarrier();
+  }
+  let m0 = mBase + lid.y * 2u;
+  let m1 = m0 + 1u;
+  let n0 = nBase + lid.x * 2u;
+  let n1 = n0 + 1u;
+  let yBatchBase = batch * M * N;
+  store_out(m0, n0, yBatchBase, acc00);
+  store_out(m0, n1, yBatchBase, acc01);
+  store_out(m1, n0, yBatchBase, acc10);
+  store_out(m1, n1, yBatchBase, acc11);
+}

build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,306 @@

+// Convolution as subgroup-matrix GEMM. OIHW weights form row-major [M,K]; `xm`
+// is either an NCHW 1x1 view [K,N] or a materialized im2col matrix. Bias is
+// applied per output channel after the full f32 accumulation.
+//
+// Full 8x8 matrix loads require complete K/N tiles. The aligned path enforces
+// that geometry; `padded` zero-fills tails, and `implicitIm2col` gathers and
+// zero-fills directly from NCHW without materializing the column matrix.
+//
+// `splitK` assigns disjoint K ranges to workgroup_id.z. Each slice writes raw
+// sums to tile-padded [splitK,batch,M_PAD,N_PAD] scratch; a later pass combines
+// the slices and applies bias or an epilogue. Padded tail cells remain zero and
+// are never copied to the logical output.
+{% set fusedNarrowProjection = fusedNarrowProjection if fusedNarrowProjection is defined else false %}
+{% set polyphase = polyphaseConvTranspose is defined and polyphaseConvTranspose %}
+{% set splitKValue = splitK if splitK is defined else 1 %}
+{% set splitKPartial = splitKValue > 1 %}
+{% set kLoopVar = "K_LOOP" if padded else "K" %}
+{% set nColsVar = "N_COLS" if padded else "N" %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+enable subgroups;
+{% if pinSubgroupSize32 %}
+enable subgroup_size_control;
+{% endif %}
+enable chromium_experimental_subgroup_matrix;
+diagnostic(off, chromium.subgroup_matrix_uniformity);
+{{ env.wgsl.resourceDeclarations }}
+{% set operandScalar = fScalar %}
+{% set accScalar = "f32" %}
+{% set useDirectMatrixStore = directMatrixStore is defined and directMatrixStore %}
+{% set tileRowsValue = tileRows if tileRows is defined else 32 %}
+{% set tileColsValue = tileCols if tileCols is defined else 64 %}
+{% set workgroupThreadsValue = workgroupThreads if workgroupThreads is defined else 128 %}
+{% set subgroupRowsValue = subgroupRows if subgroupRows is defined else 2 %}
+{% set subgroupColsValue = subgroupCols if subgroupCols is defined else 2 %}
+{% set subRowsValue = (tileRowsValue / subgroupRowsValue)|int %}
+{% set subColsValue = (tileColsValue / subgroupColsValue)|int %}
+{% set bLoadWidth = (tileColsValue * 32 / workgroupThreadsValue)|int %}
+{% set bKChunks = (32 / bLoadWidth)|int %}
+const M: u32 = {{ M }}u;
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const B_BATCH_STRIDE: u32 = {{ kLoopVar }} * {{ nColsVar }};
+{% if not splitKPartial and not polyphase and not fusedNarrowProjection %}
+const C_BATCH_STRIDE: u32 = M * N;
+{% endif %}
+const TILE_COLS: u32 = {{ tileColsValue }}u;
+const TILE_ROWS: u32 = {{ tileRowsValue }}u;
+const TILE_K: u32 = 32u;
+const SUB_COLS: u32 = {{ subColsValue }}u;
+const SUB_ROWS: u32 = {{ subRowsValue }}u;
+var<workgroup> tile_A: array<{{ operandScalar }}, {{ tileRowsValue }} * 32>;
+var<workgroup> tile_B: array<{{ operandScalar }}, {{ tileColsValue }} * 32>;
+{% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
+var<workgroup> scratch: array<array<array<{{ accScalar }}, 64>, 4>, {{ (workgroupThreadsValue / 32)|int }}>;
+{% endif %}
+fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
+  let a_global = tile_base + row;
+  let col = c_idx * 8u;
+  for (var col_offset = 0u; col_offset < 8u; col_offset = col_offset + 1u) {
+    let k = k_idx + col + col_offset;
+    if (a_global < M) {
+{% set A_PHASE = "phase_base + " if polyphase else "" %}
+{% if operandScalar == "f16" %}
+      tile_A[row * TILE_K + col + col_offset] = f16(w[{{ A_PHASE }}a_global * K + k]);
+{% else %}
+      tile_A[row * TILE_K + col + col_offset] = f32(w[{{ A_PHASE }}a_global * K + k]);
+{% endif %}
+    } else {
+{% if operandScalar == "f16" %}
+      tile_A[row * TILE_K + col + col_offset] = 0.0h;
+{% else %}
+      tile_A[row * TILE_K + col + col_offset] = 0.0;
+{% endif %}
+    }
+  }
+}
+fn loadSHMB(b_base: u32, tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
+  let b_col = tile_base + row;
+  let col = c_idx * {{ bLoadWidth }}u;
+  for (var i = 0u; i < {{ bLoadWidth }}u; i = i + 1u) {
+    let k = k_idx + col + i;
+{% if operandScalar == "f16" %}
+    tile_B[row * TILE_K + col + i] = f16(xm[b_base + k * {{ nColsVar }} + b_col]);
+{% else %}
+    tile_B[row * TILE_K + col + i] = f32(xm[b_base + k * {{ nColsVar }} + b_col]);
+{% endif %}
+  }
+}
+{% set hasActivation = hasActivation is defined and hasActivation %}
+{% if hasActivation %}
+// Apply the fused activation in the f32 accumulator before the single output
+// cast, avoiding an intermediate convolution tensor.
+fn fused_act(v: f32) -> f32 {
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+}
+{% endif %}
+{% set biasAdd = " + bv" if hasBias else "" %}
+{% if hasActivation or hasZ %}
+// Fused epilogue: Y = activation(conv + bias + Z), applied at the output store.
+// `raw` is the f32 convolution-plus-bias accumulator; Z shares Y's NCHW layout.
+fn epi(raw: f32) -> {{ T }} {
+  var r = raw;
+{% if hasActivation %}
+  r = fused_act(r);
+{% endif %}
+  return {{ T }}(r);
+}
+{% endif %}
+{% macro store_val(valExpr, idxExpr) %}
+{% if hasActivation or hasZ %}epi({{ valExpr }})
+{%- else %}{{ T }}({{ valExpr }})
+{%- endif %}
+{% endmacro %}
+{% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
+fn storeOutput(offset: u32, {% if hasBias or polyphase %}row_base: u32, {% endif %}row: u32, col: u32, src_slot: u32, row_limit: i32) {
+  if (row_limit > 0 && row < u32(row_limit)) {
+    let col2 = col + 1u;
+{% if hasBias %}
+    let bv = {{ accScalar }}(bias[row_base + row]);
+{% endif %}
+    y[offset + row * N + col]       = {{ store_val("scratch[src_slot][0][row * 8u + col]" ~ biasAdd, "offset + row * N + col") }};
+    y[offset + row * N + col + 8u]  = {{ store_val("scratch[src_slot][1][row * 8u + col]" ~ biasAdd, "offset + row * N + col + 8u") }};
+    y[offset + row * N + col + 16u] = {{ store_val("scratch[src_slot][2][row * 8u + col]" ~ biasAdd, "offset + row * N + col + 16u") }};
+    y[offset + row * N + col + 24u] = {{ store_val("scratch[src_slot][3][row * 8u + col]" ~ biasAdd, "offset + row * N + col + 24u") }};
+    y[offset + row * N + col2]       = {{ store_val("scratch[src_slot][0][row * 8u + col2]" ~ biasAdd, "offset + row * N + col2") }};
+    y[offset + row * N + col2 + 8u]  = {{ store_val("scratch[src_slot][1][row * 8u + col2]" ~ biasAdd, "offset + row * N + col2 + 8u") }};
+    y[offset + row * N + col2 + 16u] = {{ store_val("scratch[src_slot][2][row * 8u + col2]" ~ biasAdd, "offset + row * N + col2 + 16u") }};
+    y[offset + row * N + col2 + 24u] = {{ store_val("scratch[src_slot][3][row * 8u + col2]" ~ biasAdd, "offset + row * N + col2 + 24u") }};
+  }
+}
+{% endif %}
+@compute @workgroup_size({{ workgroupThreadsValue }}, 1, 1){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
+fn main(
+  @builtin(workgroup_id) workgroup_id: vec3<u32>,
+  @builtin(local_invocation_index) local_idx: u32,
+{% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
+  @builtin(subgroup_invocation_id) sg_id: u32,
+{% endif %}
+  @builtin(subgroup_size) sg_size: u32
+) {
+  let batch = workgroup_id.z;
+  let b_base = batch * B_BATCH_STRIDE;
+{% if not fusedNarrowProjection %}
+  let c_base = batch * C_BATCH_STRIDE;
+{% endif %}
+  let a_global_base = workgroup_id.y * TILE_ROWS;
+  let b_global_base = workgroup_id.x * TILE_COLS;
+  let subtile_id = local_idx / sg_size;
+  let subtile_idx = subtile_id / {{ subgroupRowsValue }}u;
+  let subtile_idy = subtile_id % {{ subgroupRowsValue }}u;
+  let base_A = subtile_idy * SUB_ROWS;
+  let base_B = subtile_idx * SUB_COLS;
+  var matC00: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC01: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC02: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC03: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC10: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC11: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC12: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  var matC13: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+  for (var kidx = 0u; kidx < {{ kLoopVar }}; kidx = kidx + TILE_K) {
+    loadSHMA(a_global_base, kidx, local_idx / 4u, local_idx % 4u);
+    loadSHMB(b_base, b_global_base, kidx, local_idx / {{ bKChunks }}u, local_idx % {{ bKChunks }}u);
+    workgroupBarrier();
+    for (var step = 0u; step < TILE_K; step = step + 8u) {
+      let matrix_a_offset = subtile_idy * SUB_ROWS * TILE_K + step;
+      var matA0: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>>(&tile_A, matrix_a_offset, false, TILE_K);
+      var matA1: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>>(&tile_A, matrix_a_offset + 8u * TILE_K, false, TILE_K);
+      let matrix_b_offset = subtile_idx * SUB_COLS * TILE_K + step;
+      var matB0: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset, true, TILE_K);
+      var matB1: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + 8u * TILE_K, true, TILE_K);
+      var matB2: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + 16u * TILE_K, true, TILE_K);
+      var matB3: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + 24u * TILE_K, true, TILE_K);
+      matC00 = subgroupMatrixMultiplyAccumulate(matA0, matB0, matC00);
+      matC01 = subgroupMatrixMultiplyAccumulate(matA0, matB1, matC01);
+      matC02 = subgroupMatrixMultiplyAccumulate(matA0, matB2, matC02);
+      matC03 = subgroupMatrixMultiplyAccumulate(matA0, matB3, matC03);
+      matC10 = subgroupMatrixMultiplyAccumulate(matA1, matB0, matC10);
+      matC11 = subgroupMatrixMultiplyAccumulate(matA1, matB1, matC11);
+      matC12 = subgroupMatrixMultiplyAccumulate(matA1, matB2, matC12);
+      matC13 = subgroupMatrixMultiplyAccumulate(matA1, matB3, matC13);
+    }
+    workgroupBarrier();
+  }
+{% if fusedNarrowProjection %}
+  // The Conv producer owns exactly one 32-row tile. Its input tile is dead after
+  // the final K iteration, so reuse those 32x64 f32 cells as the only
+  // cross-subgroup handoff into the narrow projection. All collective stores
+  // remain subgroup-uniform, and every fragment has a disjoint destination.
+  let fused_tile_offset = base_A * TILE_COLS + base_B;
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 0u * TILE_COLS +  0u, matC00, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 0u * TILE_COLS +  8u, matC01, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 0u * TILE_COLS + 16u, matC02, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 0u * TILE_COLS + 24u, matC03, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 8u * TILE_COLS +  0u, matC10, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 8u * TILE_COLS +  8u, matC11, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 8u * TILE_COLS + 16u, matC12, false, TILE_COLS);
+  subgroupMatrixStore(&tile_B, fused_tile_offset + 8u * TILE_COLS + 24u, matC13, false, TILE_COLS);
+  workgroupBarrier();
+  // One invocation owns one spatial column and walks the producer channels in
+  // increasing order to preserve the scalar projection's accumulation order.
+  // The inactive half of the workgroup has no remaining collective operation
+  // to reach.
+  let fused_global_col = b_global_base + local_idx;
+  if (local_idx < TILE_COLS && fused_global_col < N) {
+{% for oc in range(projectionOutChannels) %}
+    var projected_acc{{ oc }} = f32(0.0);
+{% endfor %}
+    for (var channel = 0u; channel < M; channel = channel + 1u) {
+      let producer_value = tile_B[channel * TILE_COLS + local_idx]{% if hasBias %} + f32(bias[channel]){% endif %};
+      let activated_value = producer_value;
+{% for oc in range(projectionOutChannels) %}
+      projected_acc{{ oc }} = projected_acc{{ oc }} + activated_value * f32(projectionW[{{ oc }}u * M + channel]);
+{% endfor %}
+    }
+    let fused_y_base = batch * {{ projectionOutChannels }}u * N + fused_global_col;
+{% for oc in range(projectionOutChannels) %}
+    let projected{{ oc }} = projected_acc{{ oc }};
+    let activated{{ oc }} = projected{{ oc }};
+    y[fused_y_base + {{ oc }}u * N] = activated{{ oc }};
+{% endfor %}
+  }
+{% elif useDirectMatrixStore %}
+  // Full f32 output tiles can be published by the subgroup-matrix collectives
+  // directly. With no epilogue this removes scratch traffic and three barriers.
+  // When an epilogue is present, publish the raw accumulators first, then apply
+  // bias, residual Z, and activation cooperatively after one storage barrier.
+  let matrix_c_offset = c_base + (a_global_base + base_A) * N + b_global_base + base_B;
+  subgroupMatrixStore(&y, matrix_c_offset + 0u * N +  0u, matC00, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 0u * N +  8u, matC01, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 0u * N + 16u, matC02, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 0u * N + 24u, matC03, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 8u * N +  0u, matC10, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 8u * N +  8u, matC11, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 8u * N + 16u, matC12, false, N);
+  subgroupMatrixStore(&y, matrix_c_offset + 8u * N + 24u, matC13, false, N);
+{% if hasBias or hasActivation or hasZ %}
+  storageBarrier();
+  for (var tile_idx = local_idx; tile_idx < TILE_ROWS * TILE_COLS; tile_idx += {{ workgroupThreadsValue }}u) {
+    let local_row = tile_idx / TILE_COLS;
+    let local_col = tile_idx % TILE_COLS;
+    let global_row = a_global_base + local_row;
+    let global_col = b_global_base + local_col;
+    let y_index = c_base + global_row * N + global_col;
+    var raw = f32(y[y_index]);
+{% if hasBias %}
+    raw += f32(bias[global_row]);
+{% endif %}
+{% if hasActivation or hasZ %}
+    y[y_index] = epi(raw);
+{% else %}
+    y[y_index] = {{ T }}(raw);
+{% endif %}
+  }
+{% endif %}
+{% endif %}
+{% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
+  // The four scratch banks are reused across the two row-groups, and each is written
+  // by a collective subgroupMatrixStore then read CROSS-LANE by storeOutput. Barriers
+  // give the reads visibility of the store AND stop the second row-group's store from
+  // clobbering the first's still-in-flight readback when a partial final M-tile
+  // diverges storeOutput's guard. Without both barriers the last valid row can be corrupted.
+  subgroupMatrixStore(&scratch[subtile_id][0], 0u, matC00, false, 8u);
+  subgroupMatrixStore(&scratch[subtile_id][1], 0u, matC01, false, 8u);
+  subgroupMatrixStore(&scratch[subtile_id][2], 0u, matC02, false, 8u);
+  subgroupMatrixStore(&scratch[subtile_id][3], 0u, matC03, false, 8u);
+  workgroupBarrier();
+  let row = sg_id / 4u;
+  let col = (sg_id % 4u) * 2u;
+  let row_base = a_global_base + base_A;
+  var matrix_c_offset = c_base + row_base * N + b_global_base + base_B;
+  var row_limit = i32(M) - i32(row_base);
+  storeOutput(matrix_c_offset, {% if hasBias or polyphase %}row_base, {% endif %}row, col, subtile_id, row_limit);
+  workgroupBarrier();
+  subgroupMatrixStore(&scratch[subtile_id][0], 0u, matC10, false, 8u);
+  subgroupMatrixStore(&scratch[subtile_id][1], 0u, matC11, false, 8u);
+  subgroupMatrixStore(&scratch[subtile_id][2], 0u, matC12, false, 8u);
+  subgroupMatrixStore(&scratch[subtile_id][3], 0u, matC13, false, 8u);
+  workgroupBarrier();
+  matrix_c_offset = matrix_c_offset + 8u * N;
+  row_limit = i32(M) - i32(row_base + 8u);
+  storeOutput(matrix_c_offset, {% if hasBias or polyphase %}row_base + 8u, {% endif %}row, col, subtile_id, row_limit);
+{% endif %}
+}

build/webgpu/conv-transpose-empty-input.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,57 @@

+{% 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 }}
+const COUNT: u32 = {{ source.count }}u;
+{% if hasBias %}
+const OUT_CHANNELS: u32 = {{ source.outChannels }}u;
+const OUTPUT_SPATIAL: u32 = {{ source.outputSpatial }}u;
+{% endif %}
+// An empty contraction contributes the additive identity at every output
+// coordinate. The optional ConvTranspose bias is applied after that reduction.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", "COUNT", note="device-axis") }}
+{% if hasBias %}
+  let output_channel = (index / OUTPUT_SPATIAL) % OUT_CHANNELS;
+  y[index] = bias[output_channel];
+{% else %}
+  y[index] = {{ T }}(0.0);
+{% endif %}
+}

build/webgpu/conv-transpose1d-ncl.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,102 @@

+{% 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 %}
+{% macro tap(raw, coord, out, kernel, axis, indent) %}
+{{ indent }}let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
+{{ indent }}if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
+{{ indent }}let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+{{ indent }}if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
+{%- endmacro -%}
+{% if usesF16 %}
+// Preserve f32 accumulation for the deep channel-by-kernel sum: widen each
+// f16 product and narrow only once at the output store.
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="device-axis") }}
+  let ow = index % params.outW;
+  var t = index / params.outW;
+  let oc = t % params.outChannels;
+  let batch = t / params.outChannels;
+{% if grouped %}
+  let groupIndex = oc / params.outChannelsPerGroup;
+  let localOc = oc - groupIndex * params.outChannelsPerGroup;
+  let inputStart = groupIndex * params.inChannelsPerGroup;
+{% endif %}
+  var acc = 0.0;
+{% if hasBias %}
+{% if usesF16 %}
+  acc = f32(bias[oc]);
+{% else %}
+  acc = bias[oc];
+{% endif %}
+{% endif %}
+{% if grouped %}
+  for (var localIc = 0u; localIc < params.inChannelsPerGroup; localIc = localIc + 1u) {
+    let ic = inputStart + localIc;
+    for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+{{ tap("rawW", "iw", "ow", "kw", "W", "      ") }}
+      let xIndex = (batch * params.inChannels + ic) * params.inW + u32(iw);
+      let wIndex = (ic * params.outChannelsPerGroup + localOc) * params.kernelW + kw;
+{% if usesF16 %}
+      acc = acc + f32(x[xIndex]) * f32(w[wIndex]);
+{% else %}
+      acc = acc + x[xIndex] * w[wIndex];
+{% endif %}
+    }
+  }
+{% else %}
+  for (var ic = 0u; ic < params.inChannels; ic = ic + 1u) {
+    for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+{{ tap("rawW", "iw", "ow", "kw", "W", "      ") }}
+      let xIndex = (batch * params.inChannels + ic) * params.inW + u32(iw);
+      let wIndex = (ic * params.outChannels + oc) * params.kernelW + kw;
+{% if usesF16 %}
+      acc = acc + f32(x[xIndex]) * f32(w[wIndex]);
+{% else %}
+      acc = acc + x[xIndex] * w[wIndex];
+{% endif %}
+    }
+  }
+{% endif %}
+{% if usesF16 %}
+  y[index] = f16(acc);
+{% else %}
+  y[index] = acc;
+{% endif %}
+}

build/webgpu/conv-transpose1d-phase-tiled.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,102 @@

+{{ env.wgsl.resourceDeclarations }}
+// Stride-2, k=4 transposed convolution decomposes into two dense phases per
+// input position:
+//   y[2j]   = x[j] W1 + x[j-1] W3
+//   y[2j+1] = x[j+1] W0 + x[j] W2
+// A workgroup computes both phases for a tile of positions and output channels.
+// Input samples and four-tap weight vectors are staged once per input-channel
+// tile, so every staged value is reused across the tile's dot products.
+const IN_CHANNELS: u32 = {{ source.inChannels }}u;
+const IN_WIDTH: u32 = {{ source.inWidth }}u;
+const OUT_CHANNELS: u32 = {{ source.outChannels }}u;
+const POSITION_TILE: u32 = {{ source.positionTile }}u;
+const OUTPUT_TILE: u32 = {{ source.outputTile }}u;
+const OUTPUTS_PER_THREAD: u32 = {{ source.outputsPerThread }}u;
+const OUTPUT_GROUPS: u32 = OUTPUT_TILE / OUTPUTS_PER_THREAD;
+const K_TILE: u32 = {{ source.kTile }}u;
+const WORKGROUP_SIZE: u32 = {{ source.workgroupSize }}u;
+const INPUT_TILE_STRIDE: u32 = POSITION_TILE + 2u;
+var<workgroup> tile_x: array<f32, {{ source.xTileElements }}>;
+var<workgroup> tile_w: array<vec4<f32>, {{ source.weightTileElements }}>;
+@compute @workgroup_size({{ source.workgroupSize }})
+fn main(
+  @builtin(local_invocation_index) lane: u32,
+  @builtin(workgroup_id) wid: vec3<u32>
+) {
+  let position_base = wid.x * POSITION_TILE;
+  let output_base = wid.y * OUTPUT_TILE;
+  let batch = wid.z;
+  let local_position = lane % POSITION_TILE;
+  let local_output_group = lane / POSITION_TILE;
+  let position = position_base + local_position;
+{% for slot in range(source.outputsPerThread) %}
+  let output_channel_{{ slot }} = output_base + local_output_group + {{ slot }}u * OUTPUT_GROUPS;
+  let output_valid_{{ slot }} = position < IN_WIDTH && output_channel_{{ slot }} < OUT_CHANNELS;
+  var even_{{ slot }} = 0.0;
+  var odd_{{ slot }} = 0.0;
+{% if hasBias %}
+  if (output_channel_{{ slot }} < OUT_CHANNELS) {
+    even_{{ slot }} = bias[output_channel_{{ slot }}];
+    odd_{{ slot }} = even_{{ slot }};
+  }
+{% endif %}
+{% endfor %}
+  for (var k_base = 0u; k_base < IN_CHANNELS; k_base += K_TILE) {
+    for (var load_index = lane; load_index < {{ source.xTileElements }}u; load_index += WORKGROUP_SIZE) {
+      let local_k = load_index / INPUT_TILE_STRIDE;
+      let halo_position = load_index - local_k * INPUT_TILE_STRIDE;
+      let input_channel = k_base + local_k;
+      let input_position = i32(position_base + halo_position) - 1;
+      var value = 0.0;
+      if (input_channel < IN_CHANNELS && input_position >= 0 && input_position < i32(IN_WIDTH)) {
+        let input_index = (batch * IN_CHANNELS + input_channel) * IN_WIDTH + u32(input_position);
+        value = x[input_index];
+      }
+      tile_x[load_index] = value;
+    }
+    for (var load_index = lane; load_index < {{ source.weightTileElements }}u; load_index += WORKGROUP_SIZE) {
+      let local_k = load_index / OUTPUT_TILE;
+      let local_oc = load_index - local_k * OUTPUT_TILE;
+      let input_channel = k_base + local_k;
+      let oc = output_base + local_oc;
+      var value = vec4<f32>(0.0);
+      if (input_channel < IN_CHANNELS && oc < OUT_CHANNELS) {
+        value = w[input_channel * OUT_CHANNELS + oc];
+      }
+      tile_w[load_index] = value;
+    }
+    workgroupBarrier();
+    if (position < IN_WIDTH) {
+      for (var local_k = 0u; local_k < K_TILE; local_k += 1u) {
+        let x_base = local_k * INPUT_TILE_STRIDE + local_position;
+        let previous = tile_x[x_base];
+        let center = tile_x[x_base + 1u];
+        let next = tile_x[x_base + 2u];
+{% for slot in range(source.outputsPerThread) %}
+        if (output_channel_{{ slot }} < OUT_CHANNELS) {
+          let weights_{{ slot }} = tile_w[local_k * OUTPUT_TILE + local_output_group + {{ slot }}u * OUTPUT_GROUPS];
+          even_{{ slot }} = fma(center, weights_{{ slot }}.y, even_{{ slot }});
+          even_{{ slot }} = fma(previous, weights_{{ slot }}.w, even_{{ slot }});
+          odd_{{ slot }} = fma(next, weights_{{ slot }}.x, odd_{{ slot }});
+          odd_{{ slot }} = fma(center, weights_{{ slot }}.z, odd_{{ slot }});
+        }
+{% endfor %}
+      }
+    }
+    workgroupBarrier();
+  }
+{% for slot in range(source.outputsPerThread) %}
+  if (output_valid_{{ slot }}) {
+    // The output width is exactly 2*IN_WIDTH, so one vec2 owns both phases.
+    let output_index_{{ slot }} = (batch * OUT_CHANNELS + output_channel_{{ slot }}) * IN_WIDTH + position;
+    y[output_index_{{ slot }}] = vec2<f32>(even_{{ slot }}, odd_{{ slot }});
+  }
+{% endfor %}
+}

build/webgpu/conv-transpose2d-col2im.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,96 @@

+{% 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 %}
+{% macro tap(raw, coord, out, kernel, axis, indent) %}
+{{ indent }}let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
+{{ indent }}if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
+{{ indent }}let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+{{ indent }}if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
+{%- endmacro -%}
+// ConvTranspose col2im scatter (output-stationary) for the im2col-GEMM path.
+//
+// The GEMM pass produced cols[m = (oc*KH+kh)*KW+kw, n = ih*IW+iw] = sum_ic
+// Wt[m, ic] * X[ic, n], i.e. the per-tap partial products with the IC contraction
+// already folded in. This pass reconstructs each output element by gathering the
+// KH*KW taps that land on it — exactly the (oh,ow) -> (ih,iw) mapping the direct
+// kernel uses, but reading the precomputed cols instead of re-running the IC sum.
+// One thread per output element, f32 accumulation, deterministic (no atomics).
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="") }}
+  let ow = index % params.outW;
+  var t = index / params.outW;
+  let oh = t % params.outH;
+  t = t / params.outH;
+  let oc = t % params.outChannels;
+  let batch = t / params.outChannels;
+  // cols is [M = OC*KH*KW, N = IH*IW] per batch image.
+  let cols_row_stride = params.inH * params.inW;
+  let cols_batch_stride = params.outChannels * params.kernelH * params.kernelW * cols_row_stride;
+  let b_base = batch * cols_batch_stride;
+  var acc = 0.0;
+{% if hasBias %}
+  acc = f32(bias[oc]);
+{% endif %}
+{% if exactTiling is defined and exactTiling %}
+  // When kernel equals stride with no padding or dilation, the taps tile the
+  // output and exactly one (kh, kw) contributes to each element. Derive that
+  // tap directly; the output-shape contract keeps ih and iw in range.
+  let kh = oh % params.kernelH;
+  let ih = oh / params.kernelH;
+  let kw = ow % params.kernelW;
+  let iw = ow / params.kernelW;
+  let m = (oc * params.kernelH + kh) * params.kernelW + kw;
+  let n = ih * params.inW + iw;
+  acc = acc + f32(cols[b_base + m * cols_row_stride + n]);
+{% else %}
+  for (var kh = 0u; kh < params.kernelH; kh = kh + 1u) {
+{{ tap("raw_h", "ih", "oh", "kh", "H", "    ") }}
+    for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+{{ tap("raw_w", "iw", "ow", "kw", "W", "      ") }}
+      let m = (oc * params.kernelH + kh) * params.kernelW + kw;
+      let n = u32(ih) * params.inW + u32(iw);
+      acc = acc + f32(cols[b_base + m * cols_row_stride + n]);
+    }
+  }
+{% endif %}
+  y[index] = {{ T }}(acc);
+}

build/webgpu/conv-transpose2d-grouped-stride-phase.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,70 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// With stride > 1 only one kernel phase contributes to a given output pixel.
+// Start each kernel loop at that phase and advance by stride so it visits only
+// contributing taps and needs no per-tap modulo or divisibility test.
+const COUNT: u32 = {{ source.count }}u;
+const OUT_C: u32 = {{ source.outC }}u;
+const OUT_CPG: u32 = {{ source.outCPerGroup }}u;
+const OUT_H: u32 = {{ source.outH }}u;
+const OUT_W: u32 = {{ source.outW }}u;
+const IN_C: u32 = {{ source.inC }}u;
+const IN_CPG: u32 = {{ source.inCPerGroup }}u;
+const IN_H: i32 = {{ source.inH }};
+const IN_W: i32 = {{ source.inW }};
+const IN_H_U: u32 = {{ source.inH }}u;
+const IN_W_U: u32 = {{ source.inW }}u;
+const KH: i32 = {{ source.kernelH }};
+const KW: i32 = {{ source.kernelW }};
+const STRIDE_H: i32 = {{ source.strideH }};
+const STRIDE_W: i32 = {{ source.strideW }};
+const PAD_H: i32 = {{ source.padH }};
+const PAD_W: i32 = {{ source.padW }};
+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 index = gid.x + gid.y * nwg.x * WG;
+  if (index >= COUNT) { return; }
+  let ow = index % OUT_W;
+  var t = index / OUT_W;
+  let oh = t % OUT_H;
+  t /= OUT_H;
+  let oc = t % OUT_C;
+  let batch = t / OUT_C;
+  let group = oc / OUT_CPG;
+  let localOc = oc - group * OUT_CPG;
+  let inputStart = group * IN_CPG;
+  let rawHBase = i32(oh) + PAD_H;
+  let rawWBase = i32(ow) + PAD_W;
+  let khPhase = rawHBase % STRIDE_H;
+  let kwPhase = rawWBase % STRIDE_W;
+{% if source.hasBias %}
+  var acc = f32(bias[oc]);
+{% else %}
+  var acc = 0.0;
+{% endif %}
+  for (var localIc = 0u; localIc < IN_CPG; localIc++) {
+    let ic = inputStart + localIc;
+    for (var kh = khPhase; kh < KH; kh += STRIDE_H) {
+      let rawH = rawHBase - kh;
+      if (rawH < 0) { continue; }
+      let ih = rawH / STRIDE_H;
+      if (ih >= IN_H) { continue; }
+      for (var kw = kwPhase; kw < KW; kw += STRIDE_W) {
+        let rawW = rawWBase - kw;
+        if (rawW < 0) { continue; }
+        let iw = rawW / STRIDE_W;
+        if (iw >= IN_W) { continue; }
+        let xIndex = ((batch * IN_C + ic) * IN_H_U + u32(ih)) * IN_W_U + u32(iw);
+        let wIndex = ((ic * OUT_CPG + localOc) * u32(KH) + u32(kh)) * u32(KW) + u32(kw);
+        acc += f32(x[xIndex]) * f32(w[wIndex]);
+      }
+    }
+  }
+  y[index] = {{ scalar }}(acc);
+}

build/webgpu/conv-transpose2d-nchw.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,124 @@

+{% 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 %}
+{% macro tap(raw, coord, out, kernel, axis, indent) %}
+{{ indent }}let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
+{{ indent }}if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
+{{ indent }}let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+{{ indent }}if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
+{%- endmacro -%}
+{% if usesF16 %}
+// Preserve f32 accumulation for the deep channel-by-kernel sum: widen each
+// f16 product and narrow only once at the output store.
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if depthwisePointwiseVec4 %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("vector_index", "", note="") }}
+  let base = vector_index * 4u;
+  if (base >= params.count) {
+    return;
+  }
+  // The contract guarantees a whole number of vec4s per NCHW channel plane,
+  // so these four adjacent values share one depthwise weight and never cross
+  // a channel boundary. This removes four copies of the generic coordinate,
+  // stride, padding, and kernel-loop arithmetic from the pointwise case.
+  let plane = params.outH * params.outW;
+  let oc = (base / plane) % params.outChannels;
+  let xv = vec4<f32>(x[base], x[base + 1u], x[base + 2u], x[base + 3u]);
+  let result = vec4<f32>(0.0) + xv * vec4<f32>(w[oc]);
+  y[base] = result.x;
+  y[base + 1u] = result.y;
+  y[base + 2u] = result.z;
+  y[base + 3u] = result.w;
+}
+{% else %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="device-axis") }}
+  let ow = index % params.outW;
+  var t = index / params.outW;
+  let oh = t % params.outH;
+  t = t / params.outH;
+  let oc = t % params.outChannels;
+  let batch = t / params.outChannels;
+{% if grouped %}
+  let group_index = oc / params.outChannelsPerGroup;
+  let local_oc = oc - group_index * params.outChannelsPerGroup;
+  let input_start = group_index * params.inChannelsPerGroup;
+{% endif %}
+  var acc = 0.0;
+{% if hasBias %}
+{% if usesF16 %}
+  acc = f32(bias[oc]);
+{% else %}
+  acc = bias[oc];
+{% endif %}
+{% endif %}
+{% if grouped %}
+  for (var local_ic = 0u; local_ic < params.inChannelsPerGroup; local_ic = local_ic + 1u) {
+    let ic = input_start + local_ic;
+{% else %}
+  for (var ic = 0u; ic < params.inChannels; ic = ic + 1u) {
+{% endif %}
+    for (var kh = 0u; kh < params.kernelH; kh = kh + 1u) {
+{{ tap("raw_h", "ih", "oh", "kh", "H", "      ") }}
+      for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+{{ tap("raw_w", "iw", "ow", "kw", "W", "        ") }}
+        let x_index = ((batch * params.inChannels + ic) * params.inH + u32(ih)) * params.inW + u32(iw);
+{% if grouped %}
+        let w_index = ((ic * params.outChannelsPerGroup + local_oc) * params.kernelH + kh) * params.kernelW + kw;
+{% else %}
+        let w_index = ((ic * params.outChannels + oc) * params.kernelH + kh) * params.kernelW + kw;
+{% endif %}
+{% if usesF16 %}
+        acc = acc + f32(x[x_index]) * f32(w[w_index]);
+{% else %}
+        acc = acc + x[x_index] * w[w_index];
+{% endif %}
+      }
+    }
+  }
+{% if usesF16 %}
+  y[index] = f16(acc);
+{% else %}
+  y[index] = acc;
+{% endif %}
+}
+{% endif %}

build/webgpu/conv-transpose2d-phase-gemm-tiled.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,155 @@

+// Fused k4s2 ConvTranspose for the common exact-2x upsampling geometry.
+//
+// A stride-two output has four independent parity phases. Within one phase,
+// each output pixel receives exactly four kernel taps per input channel, so the
+// operation is an implicit GEMM:
+//
+//   Y_phase[OC, IH*IW] = W_phase[OC, IC*4] * X_neighborhood[IC*4, IH*IW]
+//
+// Materializing the conventional [OC*16, IH*IW] columns tensor costs one large
+// scratch write, one large scratch read, and a second col2im dispatch. This
+// kernel gathers the four shifted input samples while cooperatively loading the
+// GEMM B tile, then stores each register tile directly to its interleaved output
+// phase. Boundary samples are represented by zero in shared memory.
+enable f16;
+{{ env.wgsl.resourceDeclarations }}
+const IN_C: u32 = {{ source.inChannels }}u;
+const IN_H: u32 = {{ source.inHeight }}u;
+const IN_W: u32 = {{ source.inWidth }}u;
+const OUT_C: u32 = {{ source.outChannels }}u;
+const OUT_H: u32 = {{ source.outHeight }}u;
+const OUT_W: u32 = {{ source.outWidth }}u;
+const BATCH: u32 = {{ source.batch }}u;
+const PAD_H: u32 = {{ source.padH }}u;
+const PAD_W: u32 = {{ source.padW }}u;
+const PHASES: u32 = 4u;
+const TAPS_PER_PHASE: u32 = 4u;
+const PHASE_EDGE: u32 = 2u;
+const N: u32 = IN_H * IN_W;
+const K: u32 = IN_C * TAPS_PER_PHASE;
+const M_TILE: u32 = {{ source.mTile }}u;
+const N_TILE: u32 = {{ source.nTile }}u;
+const BK: u32 = {{ source.kTile }}u;
+const THREAD_ROWS: u32 = {{ source.threadRows }}u;
+const THREAD_COLUMNS: u32 = {{ source.threadColumns }}u;
+const WG_EDGE: u32 = {{ source.workgroupEdge }}u;
+const WG_SIZE: u32 = WG_EDGE * WG_EDGE;
+var<workgroup> tileA: array<array<f16, BK>, M_TILE>;
+var<workgroup> tileB: array<array<f16, N_TILE>, BK>;
+fn store_output(batch: u32, phase: u32, oc: u32, n: u32, value: f32) {
+  if (oc >= OUT_C || n >= N) {
+    return;
+  }
+  let phaseH = phase / PHASE_EDGE;
+  let phaseW = phase % PHASE_EDGE;
+  let oh = (n / IN_W) * PHASE_EDGE + phaseH;
+  let ow = (n % IN_W) * PHASE_EDGE + phaseW;
+{% if source.hasBias %}
+  let result = value + f32(bias[oc]);
+{% else %}
+  let result = value;
+{% endif %}
+  y[((batch * OUT_C + oc) * OUT_H + oh) * OUT_W + ow] = f16(result);
+}
+@compute @workgroup_size({{ source.workgroupEdge }}, {{ source.workgroupEdge }}, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let phase = wg.z % PHASES;
+  let batch = wg.z / PHASES;
+  if (batch >= BATCH) {
+    return;
+  }
+  let phaseH = phase / PHASE_EDGE;
+  let phaseW = phase % PHASE_EDGE;
+  let kernelPhaseH = (phaseH + PAD_H) % PHASE_EDGE;
+  let kernelPhaseW = (phaseW + PAD_W) % PHASE_EDGE;
+  let inputShiftH = (phaseH + PAD_H - kernelPhaseH) / PHASE_EDGE;
+  let inputShiftW = (phaseW + PAD_W - kernelPhaseW) / PHASE_EDGE;
+  let mBase = wg.y * M_TILE;
+  let nBase = wg.x * N_TILE;
+  let lane = lid.y * WG_EDGE + lid.x;
+  let aRow = lid.y * THREAD_ROWS;
+  let bCol = lid.x * THREAD_COLUMNS;
+{% for r in range(source.threadRows) %}
+  var acc{{ r }} = vec4<f32>(0.0);
+{% endfor %}
+  for (var kBase = 0u; kBase < K; kBase += BK) {
+    // The square A/B tiles contain the same number of values, so one loop
+    // publishes a coalesced share of both operands per lane.
+    for (var linear = lane; linear < M_TILE * BK; linear += WG_SIZE) {
+      let ar = linear / BK;
+      let ac = linear % BK;
+      let oc = mBase + ar;
+      let k = kBase + ac;
+      if (oc < OUT_C && k < K) {
+        let ic = k / TAPS_PER_PHASE;
+        let tap = k % TAPS_PER_PHASE;
+        let tapH = tap / PHASE_EDGE;
+        let tapW = tap % PHASE_EDGE;
+        let kh = kernelPhaseH + tapH * PHASE_EDGE;
+        let kw = kernelPhaseW + tapW * PHASE_EDGE;
+        let wIndex = ((ic * OUT_C + oc) * 4u + kh) * 4u + kw;
+        tileA[ar][ac] = w[wIndex];
+      } else {
+        tileA[ar][ac] = f16(0.0);
+      }
+      let br = linear / N_TILE;
+      let bc = linear % N_TILE;
+      let bk = kBase + br;
+      let n = nBase + bc;
+      if (bk < K && n < N) {
+        let ic = bk / TAPS_PER_PHASE;
+        let tap = bk % TAPS_PER_PHASE;
+        let tapH = tap / PHASE_EDGE;
+        let tapW = tap % PHASE_EDGE;
+        let ih = i32(n / IN_W) + i32(inputShiftH) - i32(tapH);
+        let iw = i32(n % IN_W) + i32(inputShiftW) - i32(tapW);
+        if (ih >= 0 && ih < i32(IN_H) && iw >= 0 && iw < i32(IN_W)) {
+          let xIndex = ((batch * IN_C + ic) * IN_H + u32(ih)) * IN_W + u32(iw);
+          tileB[br][bc] = x[xIndex];
+        } else {
+          tileB[br][bc] = f16(0.0);
+        }
+      } else {
+        tileB[br][bc] = f16(0.0);
+      }
+    }
+    workgroupBarrier();
+    for (var kk = 0u; kk < BK; kk++) {
+      let bv = vec4<f32>(
+        f32(tileB[kk][bCol]),
+        f32(tileB[kk][bCol + 1u]),
+        f32(tileB[kk][bCol + 2u]),
+        f32(tileB[kk][bCol + 3u])
+      );
+{% for r in range(source.threadRows) %}
+      acc{{ r }} += vec4<f32>(f32(tileA[aRow + {{ r }}u][kk])) * bv;
+{% endfor %}
+    }
+    workgroupBarrier();
+  }
+  let oc0 = mBase + aRow;
+  let n0 = nBase + bCol;
+{% for r in range(source.threadRows) %}
+  store_output(batch, phase, oc0 + {{ r }}u, n0, acc{{ r }}.x);
+  store_output(batch, phase, oc0 + {{ r }}u, n0 + 1u, acc{{ r }}.y);
+  store_output(batch, phase, oc0 + {{ r }}u, n0 + 2u, acc{{ r }}.z);
+  store_output(batch, phase, oc0 + {{ r }}u, n0 + 3u, acc{{ r }}.w);
+{% endfor %}
+}

build/webgpu/conv-transpose2d-weight-reorder.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,57 @@

+{% 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 %}
+// ConvTranspose weight reorder for the im2col-GEMM path.
+//
+// The ONNX ConvTranspose weight is laid out [IC, OC, KH, KW]. The GEMM that
+// produces the cols matrix (cols[OC*KH*KW, IH*IW] = Wt @ X) needs its A operand
+// as a row-major [M = OC*KH*KW, K = IC] matrix, i.e. the IC dimension contracted
+// to the inner (column) axis. This pass gathers W into that transposed layout.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="") }}
+  // Wt[m, ic] with m = (oc * KH + kh) * KW + kw, contracted dim ic.
+  let ic = index % params.inChannels;
+  var m = index / params.inChannels;
+  let kw = m % params.kernelW;
+  m = m / params.kernelW;
+  let kh = m % params.kernelH;
+  let oc = m / params.kernelH;
+  let w_index = ((ic * params.outChannels + oc) * params.kernelH + kh) * params.kernelW + kw;
+  wt[index] = wsrc[w_index];
+}

build/webgpu/conv-transpose3d-col2im.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,88 @@

+{% 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 %}
+{% macro tap(raw, coord, out, kernel, axis, indent) %}
+{{ indent }}let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
+{{ indent }}if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
+{{ indent }}let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+{{ indent }}if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
+{%- endmacro -%}
+// Output-stationary rank-5 col2im gather for ConvTranspose-as-GEMM.
+//
+// GEMM has already contracted input channels into
+// cols[oc,kd,kh,kw,id,ih,iw]. This pass gathers only the spatial taps that land
+// on each output coordinate, retaining deterministic f32 accumulation and the
+// full stride/dilation/padding contract without atomics.
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="device-axis") }}
+  let ow = index % params.outW;
+  var t = index / params.outW;
+  let oh = t % params.outH;
+  t /= params.outH;
+  let od = t % params.outD;
+  t /= params.outD;
+  let oc = t % params.outChannels;
+  let batch = t / params.outChannels;
+  let inputPlane = params.inH * params.inW;
+  let inputVolume = params.inD * inputPlane;
+  let kernelPlane = params.kernelH * params.kernelW;
+  let kernelVolume = params.kernelD * kernelPlane;
+  let colsBatchStride = params.outChannels * kernelVolume * inputVolume;
+  let batchBase = batch * colsBatchStride;
+  var acc = 0.0;
+{% if hasBias %}
+  acc = f32(bias[oc]);
+{% endif %}
+  for (var kd = 0u; kd < params.kernelD; kd += 1u) {
+{{ tap("rawD", "id", "od", "kd", "D", "    ") }}
+    for (var kh = 0u; kh < params.kernelH; kh += 1u) {
+{{ tap("rawH", "ih", "oh", "kh", "H", "      ") }}
+      for (var kw = 0u; kw < params.kernelW; kw += 1u) {
+{{ tap("rawW", "iw", "ow", "kw", "W", "        ") }}
+        let matrixRow = ((oc * params.kernelD + kd) * params.kernelH + kh) *
+                        params.kernelW + kw;
+        let matrixColumn = (u32(id) * params.inH + u32(ih)) * params.inW +
+                           u32(iw);
+        acc += f32(cols[batchBase + matrixRow * inputVolume + matrixColumn]);
+      }
+    }
+  }
+  y[index] = {{ T }}(acc);
+}

build/webgpu/conv-transpose3d-ncdhw.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,73 @@

+{% 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 %}
+{% macro tap(raw, coord, out, kernel, axis, indent) %}
+{{ indent }}let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
+{{ indent }}if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
+{{ indent }}let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+{{ indent }}if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
+{%- endmacro -%}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="device-axis") }}
+  let ow = index % params.outW;
+  var t = index / params.outW;
+  let oh = t % params.outH;
+  t = t / params.outH;
+  let od = t % params.outD;
+  t = t / params.outD;
+  let oc = t % params.outChannels;
+  let batch = t / params.outChannels;
+  var acc = 0.0;
+{% if hasBias %}
+  acc = bias[oc];
+{% endif %}
+  for (var ic = 0u; ic < params.inChannels; ic = ic + 1u) {
+    for (var kd = 0u; kd < params.kernelD; kd = kd + 1u) {
+{{ tap("raw_d", "id", "od", "kd", "D", "      ") }}
+      for (var kh = 0u; kh < params.kernelH; kh = kh + 1u) {
+{{ tap("raw_h", "ih", "oh", "kh", "H", "        ") }}
+        for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+{{ tap("raw_w", "iw", "ow", "kw", "W", "          ") }}
+          let x_index = (((batch * params.inChannels + ic) * params.inD + u32(id)) * params.inH + u32(ih)) * params.inW + u32(iw);
+          let w_index = (((ic * params.outChannels + oc) * params.kernelD + kd) * params.kernelH + kh) * params.kernelW + kw;
+          acc = acc + x[x_index] * w[w_index];
+        }
+      }
+    }
+  }
+  y[index] = acc;
+}

build/webgpu/conv-transpose3d-weight-reorder.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,56 @@

+{% 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 %}
+// Reorder ONNX ConvTranspose weights from [IC, OC, KD, KH, KW] into the
+// row-major GEMM operand [OC*KD*KH*KW, IC]. All extents come from the rank-5
+// tensor contract through the uniform parameters.
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", note="device-axis") }}
+  let ic = index % params.inChannels;
+  var m = index / params.inChannels;
+  let kw = m % params.kernelW;
+  m /= params.kernelW;
+  let kh = m % params.kernelH;
+  m /= params.kernelH;
+  let kd = m % params.kernelD;
+  let oc = m / params.kernelD;
+  let sourceIndex = (((ic * params.outChannels + oc) * params.kernelD + kd) *
+                     params.kernelH + kh) * params.kernelW + kw;
+  wt[index] = wsrc[sourceIndex];
+}

build/webgpu/manifest.json ADDED Viewed

The diff for this file is too large to render. See raw diff

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,31 @@

+{
+  "name": "ai.onnx.ConvTranspose",
+  "id": "_ai_onnx_convtranspose_webgpu_a7764a8",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "K/3QCWt4Y0n8hE7RNLjJqIdqSDzkpSnjvOzP4cYpsgU=",
+      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "hwyCsCvuGx3JXDj/OC0mVEhBqvrqpR5KqTWdUUdInZY=",
+      "conv-1x1-gemm-tiled.wgsl.jinja": "VeRKrzHV+9CvWm0Lrsgpt5C98Gkz5MubGVEnYApQhdk=",
+      "conv-1x1-subgroup-matrix.wgsl.jinja": "unfOK0d/deJ/p0Mzv12Z49OZHoZ89yHE8tQvQGTpOhI=",
+      "conv-transpose-empty-input.wgsl.jinja": "242WjbMQKXP51v/XyR8R/FORSaUohUJy0qMTCgcyd7Y=",
+      "conv-transpose1d-ncl.wgsl.jinja": "QV+rvtffoKy5nFid4T0vKcLbiRGeCVef87otPiRZs3Q=",
+      "conv-transpose1d-phase-tiled.wgsl.jinja": "nN+V8lCFxP3Omu1vLhUqZMpCt7Gb/LLQAVhPGntzcxk=",
+      "conv-transpose2d-col2im.wgsl.jinja": "TlUk+O29OeiRgVi+sCzFtm/pvJjQ4Pxs727+pS7Fg4c=",
+      "conv-transpose2d-grouped-stride-phase.wgsl.jinja": "VW/769KFvrA4NywQuNw9uawTGodB/iKUJdAE9vszJx8=",
+      "conv-transpose2d-nchw.wgsl.jinja": "VB/VsdNoFh6eVJJ9x1rHjIJYG1INTVMGO145elCU7ow=",
+      "conv-transpose2d-phase-gemm-tiled.wgsl.jinja": "uF07bIUt+NQdzc5Jq6IDLUDn51h6GASE++9HYbPGZms=",
+      "conv-transpose2d-weight-reorder.wgsl.jinja": "Gt8Qw4oUFXATVT33pGW7hvEIZ+VaYPrN70sFkFPNfGs=",
+      "conv-transpose3d-col2im.wgsl.jinja": "hKL2pQTqi81DFVEVuUZd5KNoikhK5Fcq+e6xPSBpM60=",
+      "conv-transpose3d-ncdhw.wgsl.jinja": "F9SYKafHOobKRTCa9yE/dUD4SCU/Fzgo+EdQdiAssCg=",
+      "conv-transpose3d-weight-reorder.wgsl.jinja": "CghFqX1j85v2ZY6DjTGhO2B+Es/lERkUB+QR+KoTzNA=",
+      "manifest.json": "X8AIY7imVWPKlvjr6NpBMaBGsONv+0kH68l2z0fwYZM=",
+      "test.json": "xRSBVV55KeohZ8pTZbAS5n3B0wwxBev9MxJbaE6/wF4="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.ConvTranspose" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1856 @@

+{
+  "op": "ai.onnx.ConvTranspose",
+  "fixtureArrays": {
+    "grouped_two_groups_no_bias_input_x": [1, 2, 3, 4, -1, -2, 0.5, 1.5, 2, -3, 4, -5, 0.25, -0.5, 0.75, -1],
+    "grouped_two_groups_no_bias_input_w": [1, 0, 0.5, -1, -0.25, 0.75, 1.5, 0, 0.5, -0.5, 1, 0.25, 1.25, 0, -0.75, 0.5, -1, 0.5, 0.25, 1, 0.75, -1.25, 0, 0.5, 0.25, 1.5, -0.5, 0, -0.75, 0.5, 1, -1],
+    "ort_default_strides_dilations_multi_channel_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17],
+    "onnx_backend_convtranspose_input_w": [1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1],
+    "f16_1d_nobias_input_x": [0.1, 0.5, -0.25, 0.75, 1, -0.5, 0.25, -1, 0.5, -0.75, 1.25, -0.25, 0, 0.5, -1, 0.75],
+    "f16_1d_nobias_input_w": [0.5, -0.25, 1, 0.75, -0.5, 0.25, -1, 0.5, 0.125, 1, -0.5, 0.25, 0.5, 0.75, -0.25, -0.75, 0.5, 1]
+  },
+  "cases": [
+    {
+      "name": "f32_subnormal_times_large_scale_1d_gpu_gap",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest",
+        "notes": "A 1D transposed convolution with a pointwise large finite weight should preserve subnormal input signal as normal finite output."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [1e-39, -1e-39, 1e-38, -1e-38] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "constant", "value": 1e+38 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 4], "tolerance": 0.000001 } },
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes f32 subnormals to zero (FTZ); subnormal-magnitude values cannot survive the GPU arithmetic path."
+      }
+    },
+    {
+      "name": "f32_subnormal_times_large_scale_2d_gpu_gap",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest",
+        "notes": "A 2D transposed convolution with a pointwise large finite weight should preserve subnormal input signal as normal finite output."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [1e-39, -1e-39, 1e-38, -1e-38] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 1], "data": { "kind": "constant", "value": 1e+38 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 2], "tolerance": 0.000001 } },
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes f32 subnormals to zero (FTZ); subnormal-magnitude values cannot survive the GPU arithmetic path."
+      }
+    },
+    {
+      "name": "f32_subnormal_times_large_scale_3d_gpu_gap",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest",
+        "notes": "A 3D transposed convolution with a pointwise large finite weight should preserve subnormal input signal as normal finite output."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1e-39, -1e-39, 2e-39, -2e-39, 5e-39, -5e-39, 1e-38, -1e-38] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 1, 1], "data": { "kind": "constant", "value": 1e+38 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 2, 2], "tolerance": 0.000001 } },
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes f32 subnormals to zero (FTZ); subnormal-magnitude values cannot survive the GPU arithmetic path."
+      }
+    },
+    {
+      "name": "dispatch_cliff_ncl1d",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 16776961],
+          "data": { "kind": "cycle", "values": [0.5, -0.25, 1.0, -0.75] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 1], "data": { "kind": "constant", "value": 2.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 16776961], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "phase_tiled_1d_k4s2_batch_channel_position_tails",
+      "provenance": {
+        "notes": "Exercises the neural-audio stride-2 phase decomposition with non-tile-aligned input channels, output channels, positions, and batch > 1."
+      },
+      "attrs": { "strides": [2], "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 37, 35],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17, "scale": 0.2 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [37, 13, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.1 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 13, 70], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "phase_tiled_1d_k4s2_bias_tails",
+      "provenance": { "notes": "Bias companion for the phase-tiled neural-audio path with channel and position tails." },
+      "attrs": { "strides": [2], "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 17, 33],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23, "scale": 0.2 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [17, 19, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.1 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [19],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 19, 66], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "phase_tiled_1d_register_tile_channel_position_tails",
+      "provenance": {
+        "notes": "Crosses the occupancy-derived two-output register-tile threshold while retaining input-channel, output-channel, and position tails."
+      },
+      "attrs": { "strides": [2], "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 17, 1025],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21, "scale": 0.2 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [17, 65, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.27, "scale": 0.1 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 65, 2050], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "nchw_no_bias",
+      "attrs": { "strides": [2, 1], "pads": [1, 0, 1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5, 5], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "nchw_bias",
+      "attrs": { "strides": [2, 1], "pads": [1, 0, 1, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5, 5], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "non_vector_channels_no_bias",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [3, 5, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.23 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 5, 3, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "non_vector_channels_bias",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [3, 5, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.23 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 5, 3, 3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "non_vector_channels_bias_stride2_compact",
+      "provenance": {
+        "notes": "Compact companion for bench case conv-transpose-f32-non-vector-bias-b1c3m5-32x32-k3s2, preserving non-vec channel counts, bias, stride=2, and padded k3 output geometry."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.17, "scale": 0.2 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [3, 5, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.23, "scale": 0.12 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.29, "scale": 0.1 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 5, 15, 15], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "output_padding_extends_shape",
+      "attrs": { "strides": [2, 2], "output_padding": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, -1.0, 2.0, 0.5] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "dilated_kernel_spreads_taps",
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "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] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 3.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "grouped_two_groups_no_bias",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/grouped_two_groups_no_bias_input_x" } }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [4, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/grouped_two_groups_no_bias_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_grouped2d_stride2_phase",
+      "provenance": {
+        "notes": "Compact grouped stride-2 fixture that exercises the phase-decomposed nchw2d_grouped_stride_phase fallback on every feature tier."
+      },
+      "attrs": { "group": 2, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 3, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17, "scale": 0.25 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [4, 3, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 6, 5, 7], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "grouped_two_groups_bias",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/grouped_two_groups_no_bias_input_x" } }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [4, 2, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/grouped_two_groups_no_bias_input_w" } }
+        },
+        "bias": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, -1.0, 1.5, -0.25] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_batch_dim",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.DimWithZero"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1], "output_padding": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 1, 3, 3], "data": { "kind": "values", "values": [] } },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [-0.06230065, 0.37932432, -0.25388849, 0.33878803, 0.43709868, -0.22477469, 0.04118127, -0.44696793, 0.06373066]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 1, 6, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_outputpadding_strides2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_outputpadding_strides2"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1], "output_padding": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.16857791, -0.15161794, 0.08540368, 0.1820628, -0.21746576, 0.08245695, 0.1431433, -0.43156421, 0.30591947]
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [-0.06230065, 0.37932432, -0.25388849, 0.33878803, 0.43709868, -0.22477469, 0.04118127, -0.44696793, 0.06373066]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 6, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_outputpadding_strides2_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_outputpadding_strides2"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1], "output_padding": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.16857791, -0.15161794, 0.08540368, 0.1820628, -0.21746576, 0.08245695, 0.1431433, -0.43156421, 0.30591947]
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [-0.06230065, 0.37932432, -0.25388849, 0.33878803, 0.43709868, -0.22477469, 0.04118127, -0.44696793, 0.06373066]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 6, 6], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_2d_c2_single_output_channel",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_C2"
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.43, 0.42871707, 0.29552766, 0.17258859, 0.68087016, 0.7090254, 0.60937387, 0.58646585, 0.84525721, 0.47011843, 0.95854213, 0.3972888, 0.0585452, 0.1206734, 0.76727852, 0.46040912, 0.83495316, 0.02409773]
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.25616416, 0.10246604, 0.08771133, 0.30770606, 0.84369617, 0.3010619, 0.44524362, 0.6056068]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_dilation2_kernel2_simple",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_1"
+      },
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [11.0, 12.0, 21.0, 22.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_dilation3_kernel2_mixed_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_4"
+      },
+      "attrs": { "dilations": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 6, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_dilation3_asymmetric_pads_top_left",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_AsymmetricPads_1"
+      },
+      "attrs": { "pads": [2, 2, 1, 1], "dilations": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [42.0, 6.0, 4.0, 1.0, 27.0, 72.0, 7.0, 81.0, 45.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_dilation3_asymmetric_pads_bottom_right",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_AsymmetricPads_2"
+      },
+      "attrs": { "pads": [1, 1, 2, 2], "dilations": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [35.0, 49.0, 18.0, 14.0, 42.0, 6.0, 8.0, 1.0, 27.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_dilation3_asymmetric_pads_crop_top_left",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_AsymmetricPads_3"
+      },
+      "attrs": { "pads": [2, 2, 0, 0], "dilations": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 4],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [42.0, 6.0, 4.0, 12.0, 1.0, 27.0, 72.0, 9.0, 7.0, 81.0, 45.0, 63.0, 6.0, 27.0, 18.0, 54.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_dilation3_asymmetric_pads_crop_bottom_right",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_AsymmetricPads_4"
+      },
+      "attrs": { "pads": [0, 0, 2, 2], "dilations": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 4],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [21.0, 56.0, 7.0, 6.0, 63.0, 35.0, 49.0, 18.0, 21.0, 14.0, 42.0, 6.0, 3.0, 8.0, 1.0, 27.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_default_strides_dilations_multi_channel",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_DefaultStridesAndDilations"
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_default_strides_dilations_multi_channel_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 2],
+          "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, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_stride_width2_dilation_width3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_NonDefaultStridesAndDilations"
+      },
+      "attrs": { "strides": [1, 2], "dilations": [1, 3] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 1, 2], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1, 12], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_stride_height2_dilation_height3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_NonDefaultStridesAndDilations_T"
+      },
+      "attrs": { "strides": [2, 1], "dilations": [3, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 2, 1], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "w": { "dtype": "float32", "shape": [1, 1, 4, 1], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 12, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_bias_3x3_padded",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Bias_1"
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 5, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.22572887, -0.07105902, -0.40399021, -0.14461157, 0.05367219, -0.08353302, 0.41023391, 0.42745841, -0.3769345, -0.42057109, -0.1372498, 0.05485916, 0.34602994, -0.06402895, -0.06000063, 0.07891446, -0.09410021, 0.26251942, -0.11043271, 0.47966552, 0.34682763, -0.04511502, 0.22414422, 0.24618894, -0.21480265]
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [-0.0962126, 0.19827795, 0.03667754, 0.36756599, -0.01076147, -0.11781135, -0.11574665, -0.38404959, 0.44403327]
+          }
+        },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.04676145] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_2d_bias_2x2_no_pad",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Bias_2"
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 7, 7],
+          "data": {
+            "kind": "values",
+            "values": [0.01270282, 0.09657472, -0.36909008, -0.08085269, 0.0242992, 0.40873009, -0.46927932, 0.34412372, -0.39574206, 0.26234281, 0.27352369, -0.22265741, 0.43270493, -0.24710381, -0.03418651, -0.04413456, -0.16414353, 0.3158558, 0.1087395, -0.38577938, -0.38986659, -0.09614426, 0.17591673, 0.40140027, -0.0869683, -0.47193506, -0.05010766, 0.29325962, 0.22680271, -0.0793834, -0.36764491, 0.20451134, 0.46361887, -0.12190259, 0.03413916, 0.12307656, 0.28569579, -0.392129, 0.17179191, 0.27161086, -0.12766263, 0.1371125, 0.28137422, -0.39899838, 0.23824286, -0.19693244, 0.32956779, 0.46209556, -0.46913007]
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [-0.34922412, 0.1114341, -0.01778314, 0.46861196] }
+        },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.17402864] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 8, 8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_1d_mixed_channels",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D"
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [0.1, 1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 6.0, 5.0, 4.0, 3.0, 2.0, 1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_1d_asymmetric_padding",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_AsymmetricPads"
+      },
+      "attrs": { "pads": [1, 0] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [3.0, 5.0, 7.0, 4.0, 3.0, 5.0, 7.0, 4.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_output_shape_2_bias_single_row",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_OutputShape_2",
+        "notes": "ORT uses the exact output_shape attribute; this fixture has the same natural output size."
+      },
+      "attrs": { "output_shape": [1, 14] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 1, 10], "data": { "kind": "linspace", "start": 0.0, "end": 9.0 } },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 5],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 2.0, 1.0] }
+        },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 1, 14], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_output_shape_with_batch_size_bias_single_row",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_OutputShapeWithBatchSize",
+        "notes": "ORT uses the exact output_shape attribute; this fixture has the same natural output size."
+      },
+      "attrs": { "output_shape": [1, 14] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 1, 1, 10], "data": { "kind": "linspace", "start": 0.0, "end": 19.0 } },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 5],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 2.0, 1.0] }
+        },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 1, 1, 14], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_output_shape_1_equivalent_symmetric_crop",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_OutputShape_1",
+        "notes": "Covers the exact spatial output_shape [4, 4]; ConvTranspose derives the required symmetric crop."
+      },
+      "attrs": { "output_shape": [4, 4] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 3, 4, 4], "data": { "kind": "constant", "value": 1.0 } },
+        "w": { "dtype": "float32", "shape": [3, 3, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dilation3_kernel2_repeated_tiles",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_2"
+      },
+      "attrs": { "dilations": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [11.0, 12.0, 21.0, 22.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dilation2_kernel2_mixed_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Dilation_3"
+      },
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_convtranspose",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose" }
+    },
+    {
+      "name": "onnx_backend_convtranspose_dilations",
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [3.0, 8.0, 1.0, 9.0, 5.0, 7.0, 3.0, 2.0, 6.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [7.0, 2.0, 1.0, 9.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_dilations" }
+    },
+    {
+      "name": "onnx_backend_convtranspose_group_2",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_default_strides_dilations_multi_channel_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_group_2" }
+    },
+    {
+      "name": "onnx_backend_convtranspose_group_2_image_3",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2, 3, 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, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 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, 15.0, 16.0, 17.0]
+          }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2, 5, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_group_2_image_3" }
+    },
+    {
+      "name": "onnx_backend_convtranspose_pad",
+      "attrs": { "strides": [3, 2], "output_padding": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 10, 8] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_pad" }
+    },
+    {
+      "name": "onnx_backend_convtranspose_pads",
+      "attrs": { "strides": [3, 2], "pads": [1, 2, 1, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 7, 3] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_pads" }
+    },
+    {
+      "name": "onnx_backend_convtranspose_autopad_same_explicit",
+      "attrs": { "strides": [2, 2], "auto_pad": "SAME_UPPER" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 6, 6], "tolerance": 0 } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_autopad_same",
+        "notes": "SAME_UPPER is represented by equivalent explicit symmetric pads plus output padding for this fixture."
+      }
+    },
+    {
+      "name": "onnx_backend_convtranspose_output_shape_as_output_padding",
+      "attrs": { "strides": [3, 2], "output_shape": [10, 8] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 10, 8], "tolerance": 0 } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_output_shape",
+        "notes": "Uses the official spatial output_shape attribute from the ONNX backend fixture."
+      }
+    },
+    {
+      "name": "onnx_backend_convtranspose_kernel_shape_output_padding",
+      "attrs": { "strides": [3, 2], "output_padding": [1, 1], "output_shape": [10, 8], "kernel_shape": [3, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_convtranspose_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 10, 8], "tolerance": 0 } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_kernel_shape",
+        "notes": "Uses the official kernel_shape, output_shape, and output_padding attributes from the ONNX backend fixture."
+      }
+    },
+    {
+      "name": "onnx_backend_convtranspose_1d",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_1d" },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "linspace", "start": 0.0, "end": 2.0 } },
+        "w": { "dtype": "float32", "shape": [1, 2, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_convtranspose_1d_autopad_same_upper",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_AutoPad_SameUpper",
+        "notes": "Covers one-dimensional ConvTranspose with auto_pad=SAME_UPPER."
+      },
+      "attrs": { "auto_pad": "SAME_UPPER" },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1.0, 3.0, 5.0, 7.0, 1.0, 3.0, 5.0, 7.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_convtranspose_1d_autopad_same_lower",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_AutoPad_SameLower",
+        "notes": "Covers one-dimensional ConvTranspose with auto_pad=SAME_LOWER."
+      },
+      "attrs": { "auto_pad": "SAME_LOWER" },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [3.0, 5.0, 7.0, 4.0, 3.0, 5.0, 7.0, 4.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_convtranspose_2d_autopad_same_upper_stride2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_AutoPad_with_non_default_strides",
+        "notes": "Covers two-dimensional stride-2 ConvTranspose with auto_pad=SAME_UPPER."
+      },
+      "attrs": { "auto_pad": "SAME_UPPER", "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 2, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 6, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_convtranspose_1d_grouped",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_OutputShape_1_group_2_for_transpose_path",
+        "notes": "Covers grouped one-dimensional ConvTranspose."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float32", "shape": [2, 1, 1], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [2.0, 4.0, 9.0, 12.0] }
+        }
+      }
+    },
+    {
+      "name": "ncl1d_grouped_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_OutputShape_1_group_2_for_transpose_path",
+        "notes": "Grouped (group=2) 1D ConvTranspose WITH bias. Exercises the ncl1d_grouped_bias variant."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float32", "shape": [2, 1, 1], "data": { "kind": "values", "values": [2.0, 3.0] } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_convtranspose_1d_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_2D_Bias_2",
+        "notes": "Covers one-dimensional ConvTranspose with the optional bias input."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "w": { "dtype": "float32", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, -1.0, 0.5] } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [1.25, 4.25, 7.25, 6.25, -1.5, -2.0, -2.5, 1.0] }
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_convtranspose_3d",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convtranspose_3d" },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 3, 4, 5], "data": { "kind": "linspace", "start": 0.0, "end": 59.0 } },
+        "w": { "dtype": "float32", "shape": [1, 2, 3, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 2, 5, 6, 7], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_convtranspose_3d_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_3D",
+        "notes": "Covers three-dimensional 1x1x1 ConvTranspose with the optional bias input."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, -1.0, -2.0, -3.0, -4.0] }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [1.0] } },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 1, 2, 2, 2],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [1.25, 2.25, 3.25, 4.25, -0.75, -1.75, -2.75, -3.75] }
+        }
+      }
+    },
+    {
+      "name": "nchw2d_bias_f16",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1], "output_padding": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.16857791, -0.15161794, 0.08540368, 0.1820628, -0.21746576, 0.08245695, 0.1431433, -0.43156421, 0.30591947]
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [1, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [-0.06230065, 0.37932432, -0.25388849, 0.33878803, 0.43709868, -0.22477469, 0.04118127, -0.44696793, 0.06373066]
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 1, 6, 6], "tolerance": 0.001 } }
+    },
+    {
+      "name": "f16_1d_nobias",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_1d_nobias_input_x" } }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_1d_nobias_input_w" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 10], "tolerance": 0.02 } }
+    },
+    {
+      "name": "f16_grouped2d_nobias",
+      "attrs": { "group": 2, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 4, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [4, 3, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 6, 15, 15], "tolerance": 0.03 } }
+    },
+    {
+      "name": "f16_autopad_same_upper_stride2",
+      "attrs": { "auto_pad": "SAME_UPPER", "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.23 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 8, 8], "tolerance": 0.03 } }
+    },
+    {
+      "name": "nchw2d_f16_nobias",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 6, 6], "tolerance": 0.02 } }
+    },
+    {
+      "name": "f16_phase_gemm_k4s2_channel_and_spatial_tails",
+      "provenance": {
+        "notes": "Compact route lock for the fused four-phase k4s2 GEMM path. Odd channel and spatial sizes exercise K, M, and N tile tails plus all padded boundaries."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 5, 7, 9],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17, "scale": 0.2 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [5, 11, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13, "scale": 0.1 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 11, 14, 18], "tolerance": 0.03 } }
+    },
+    {
+      "name": "f16_phase_gemm_k4s2_bias_tails",
+      "provenance": {
+        "notes": "Bias companion for the fused four-phase k4s2 path with non-tile-aligned channels and phase-grid dimensions."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 3, 5, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23, "scale": 0.2 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 7, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.1 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.31, "scale": 0.05 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 7, 10, 12], "tolerance": 0.03 } }
+    },
+    {
+      "name": "f16_gemm_2d_no_bias",
+      "provenance": {
+        "notes": "f16 2D ConvTranspose, group 1, IC%32==0 and IH*IW%64==0 -> exercises the im2col-GEMM path nchw2d_gemm (subgroup-matrix) AND the direct nchw2d_f16 kernel."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 16, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.11, "scale": 0.4 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 16, 15, 15], "tolerance": 0.03 } }
+    },
+    {
+      "name": "f16_gemm_2d_bias",
+      "provenance": {
+        "notes": "f16 2D ConvTranspose, group 1, IC%32==0 and IH*IW%64==0 -> exercises the im2col-GEMM path nchw2d_gemm_bias (subgroup-matrix) AND the direct nchw2d_f16 kernel."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 16, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.11, "scale": 0.4 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 0.3 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 16, 15, 15], "tolerance": 0.03 } }
+    },
+    {
+      "name": "f16_gemm_2d_batch2",
+      "provenance": {
+        "notes": "f16 2D ConvTranspose GEMM path with batch=2 — exercises the GEMM workgroup_id.z batching and col2im batch decode."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.061, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 16, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.11, "scale": 0.4 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 16, 15, 15], "tolerance": 0.03 } }
+    },
+    {
+      "name": "nchw2d_gemm_f32_reg_m512_n1024",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 },
+          "shape": [1, 8, 32, 32]
+        },
+        "w": {
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 },
+          "shape": [8, 128, 2, 2]
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 128, 33, 33], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "nchw2d_gemm_f32_bias_reg_m512_n1024",
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 },
+          "shape": [1, 8, 32, 32]
+        },
+        "w": {
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 },
+          "shape": [8, 128, 2, 2]
+        },
+        "bias": {
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.23 },
+          "shape": [128]
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 128, 33, 33], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "empty_input_spatial_dim_3d_nonempty_output",
+      "provenance": {
+        "notes": "A zero-depth input leaves a nonempty rank-5 output because the kernel is wider than the empty spatial extent. The contraction is empty, so every output is the additive identity."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 0, 3, 3] },
+        "w": { "dtype": "float32", "shape": [1, 1, 3, 3, 3], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 2, 5, 5], "data": { "kind": "constant", "value": 0.0 } } }
+    },
+    {
+      "name": "empty_input_spatial_dim_3d_nonempty_output_bias",
+      "provenance": {
+        "notes": "Bias is applied after an empty rank-5 contraction, so each nonempty output channel is exactly its bias value without reading X or W."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 0, 1, 1] },
+        "w": { "dtype": "float32", "shape": [1, 2, 2, 1, 1], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.25, -2.5] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 1, 1],
+          "data": { "kind": "values", "values": [1.25, -2.5] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "empty_input_spatial_dim_2d_grouped_f16_nonempty_output",
+      "provenance": {
+        "notes": "The empty-reduction fill route also covers grouped rank-4 float16 without allocating or reading an im2col scratch buffer."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 2, 0, 1] },
+        "w": { "dtype": "float16", "shape": [2, 1, 2, 1], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [1, 2, 1, 1], "data": { "kind": "constant", "value": 0.0 }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "empty_input_spatial_dim_1d_f16_nonempty_output_bias",
+      "provenance": {
+        "notes": "The float16 bias-broadcast source is rendered explicitly: a zero-width rank-3 reduction returns one output position per channel containing only bias."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 1, 0] },
+        "w": { "dtype": "float16", "shape": [1, 2, 2], "data": { "kind": "constant", "value": 1.0 } },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.5, -0.75] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float16",
+          "shape": [1, 2, 1],
+          "data": { "kind": "values", "values": [0.5, -0.75] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "zero_batch_dim_3d_empty_output",
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 1, 3, 3, 3] },
+        "w": { "dtype": "float32", "shape": [1, 2, 3, 3, 3] }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 2, 5, 5, 5] } }
+    },
+    {
+      "name": "zero_batch_dim_grouped_empty_output",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 4, 4, 4] },
+        "w": { "dtype": "float32", "shape": [4, 2, 2, 2] }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [0, 4, 5, 5] } }
+    },
+    {
+      "name": "grouped_dispatch_fold_depthwise_1x1_over_16m",
+      "attrs": { "group": 64 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 512, 512] },
+        "w": { "dtype": "float32", "shape": [64, 1, 1, 1], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 512, 512] } }
+    },
+    {
+      "name": "grouped_2d_f16_bias_routing_gap",
+      "attrs": { "group": 2, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 4, 8, 8] },
+        "w": { "dtype": "float16", "shape": [4, 3, 3, 3] },
+        "bias": { "dtype": "float16", "shape": [6] }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 6, 15, 15], "tolerance": 0.03 } }
+    },
+    {
+      "name": "convtranspose3d_depth_output_padding1_compact",
+      "provenance": {
+        "source": "ONNX ConvTranspose-22 volumetric output_padding semantics",
+        "notes": "Covers a non-default depth entry in the exact output_padding list attribute."
+      },
+      "attrs": { "strides": [2, 1, 1], "output_padding": [1, 0, 0] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "w": { "dtype": "float32", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [3.0, 4.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 1, 1], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "dispatch_cliff_batch_over_65535_bias_forces_nchw_scalar",
+      "provenance": {
+        "notes": "Batch 65536 exceeds maxComputeWorkgroupsPerDimension, so gemmDispatchFits is false and every col2im GEMM variant drops out; the scalar NCHW kernel with its bias epilogue is the only remaining f32 dense path."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [65536, 1, 1, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0009, "cosStep": 0.0007 }
+        },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [65536, 1, 1, 1], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "dispatch_cliff_batch_over_65535_bias_f16_forces_nchw_scalar",
+      "provenance": {
+        "notes": "f16 companion of the batch dispatch cliff: with gemmDispatchFits false the tiled and subgroup-matrix col2im GEMMs are ineligible, and the stride-2 phase GEMM geometry does not apply, leaving the scalar f16 NCHW bias kernel."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [65536, 1, 1, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0009, "cosStep": 0.0007 }
+        },
+        "w": { "dtype": "float16", "shape": [1, 1, 1, 1], "data": { "kind": "constant", "value": 1.25 } },
+        "bias": { "dtype": "float16", "shape": [1], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [65536, 1, 1, 1], "tolerance": 0.002, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "f16_grouped2d_stride1_no_bias",
+      "provenance": {
+        "notes": "Grouped f16 ConvTranspose at stride 1: the grouped stride-phase kernel requires strideH/strideW >= 2, so unit stride is what routes a grouped f16 request to the scalar grouped NCHW kernel."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.17 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [2, 1, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 3, 3], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_grouped2d_stride1_bias",
+      "provenance": {
+        "notes": "Bias companion of the unit-stride grouped f16 case: the grouped stride-phase bias kernel is gated on stride >= 2, so the bias epilogue of the scalar grouped f16 NCHW kernel is only reachable at stride 1."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [2, 1, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.11 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.29 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 3, 3], "tolerance": 0.01 } }
+    },
+    {
+      "name": "convtranspose3d_subgroup_matrix_col2im_overlap_f32",
+      "provenance": {
+        "notes": "General rank-5 matrix-path coverage with K=IC=32, N=ID*IH*IW=64, and M=OC*KD*KH*KW=32. Constant operands make the complete 2x2x2 overlap multiplicity explicit while exercising weight reorder, subgroup-matrix contraction, and 3D col2im."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 4, 4, 4], "data": { "kind": "constant", "value": 1.0 } },
+        "w": { "dtype": "float32", "shape": [32, 4, 2, 2, 2], "data": { "kind": "constant", "value": 0.125 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 4, 5, 5, 5],
+          "tolerance": 0.0002,
+          "relTolerance": 0.0002,
+          "data": {
+            "kind": "cycle",
+            "values": [4.0, 8.0, 8.0, 8.0, 4.0, 8.0, 16.0, 16.0, 16.0, 8.0, 8.0, 16.0, 16.0, 16.0, 8.0, 8.0, 16.0, 16.0, 16.0, 8.0, 4.0, 8.0, 8.0, 8.0, 4.0, 8.0, 16.0, 16.0, 16.0, 8.0, 16.0, 32.0, 32.0, 32.0, 16.0, 16.0, 32.0, 32.0, 32.0, 16.0, 16.0, 32.0, 32.0, 32.0, 16.0, 8.0, 16.0, 16.0, 16.0, 8.0, 8.0, 16.0, 16.0, 16.0, 8.0, 16.0, 32.0, 32.0, 32.0, 16.0, 16.0, 32.0, 32.0, 32.0, 16.0, 16.0, 32.0, 32.0, 32.0, 16.0, 8.0, 16.0, 16.0, 16.0, 8.0, 8.0, 16.0, 16.0, 16.0, 8.0, 16.0, 32.0, 32.0, 32.0, 16.0, 16.0, 32.0, 32.0, 32.0, 16.0, 16.0, 32.0, 32.0, 32.0, 16.0, 8.0, 16.0, 16.0, 16.0, 8.0, 4.0, 8.0, 8.0, 8.0, 4.0, 8.0, 16.0, 16.0, 16.0, 8.0, 8.0, 16.0, 16.0, 16.0, 8.0, 8.0, 16.0, 16.0, 16.0, 8.0, 4.0, 8.0, 8.0, 8.0, 4.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "convtranspose3d_subgroup_matrix_bias_f32",
+      "provenance": {
+        "notes": "Bias companion for the aligned rank-5 matrix path. A pointwise volumetric kernel isolates the bias epilogue while preserving context-derived M/K/N eligibility."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 4, 4, 4], "data": { "kind": "constant", "value": 1.0 } },
+        "w": { "dtype": "float32", "shape": [32, 32, 1, 1, 1], "data": { "kind": "constant", "value": 0.125 } },
+        "bias": { "dtype": "float32", "shape": [32], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 32, 4, 4, 4],
+          "tolerance": 0.0002,
+          "relTolerance": 0.0002,
+          "data": { "kind": "constant", "value": 4.5 }
+        }
+      }
+    },
+    {
+      "name": "convtranspose3d_tiled_reg_large_spatial_f32",
+      "provenance": {
+        "notes": "Portable register-tiled rank-5 coverage with N=8*32*32 and M=32. Pointwise weights isolate the large matrix path and 3D tensor plumbing while tile eligibility remains manifest-derived."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 8, 32, 32], "data": { "kind": "constant", "value": 1.0 } },
+        "w": { "dtype": "float32", "shape": [32, 32, 1, 1, 1], "data": { "kind": "constant", "value": 0.03125 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 32, 8, 32, 32],
+          "tolerance": 0.0002,
+          "relTolerance": 0.0002,
+          "data": { "kind": "constant", "value": 1.0 }
+        }
+      }
+    },
+    {
+      "name": "convtranspose3d_tiled_bias_reg_large_spatial_f32",
+      "provenance": {
+        "notes": "Bias companion for the portable register-tiled rank-5 path, validating that bias stays in the output-stationary col2im epilogue rather than the intermediate GEMM."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 8, 32, 32], "data": { "kind": "constant", "value": 1.0 } },
+        "w": { "dtype": "float32", "shape": [32, 32, 1, 1, 1], "data": { "kind": "constant", "value": 0.03125 } },
+        "bias": { "dtype": "float32", "shape": [32], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1, 32, 8, 32, 32],
+          "tolerance": 0.0002,
+          "relTolerance": 0.0002,
+          "data": { "kind": "constant", "value": 1.5 }
+        }
+      }
+    },
+    {
+      "name": "exact_tiling_stride4_kernel4_gemm_bias",
+      "provenance": {
+        "notes": "Kernel == stride with no padding, no dilation and no output padding: the taps TILE the output rather than overlapping it, so col2im takes its collapsed single-tap path instead of the KH*KW gather. Shaped after Depth Anything V2's factor-4 reassemble (x [1,384,37,49], w [384,48,4,4]). Without a case in this quadrant the collapsed branch had zero coverage: poisoning it left all 163 ConvTranspose tests green."
+      },
+      "attrs": { "strides": [4, 4] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 16, 5, 7],
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [16, 8, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.11, "scale": 0.4 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 0.3 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 20, 28], "tolerance": 0.0005 } }
+    },
+    {
+      "name": "exact_tiling_nonsquare_stride2x3_no_bias",
+      "provenance": {
+        "notes": "Kernel == stride on BOTH axes but with different strides per axis (2 x 3), so a kh/kw or inW/outW swap in the collapsed col2im tap cannot cancel out the way it would on a square kernel."
+      },
+      "attrs": { "strides": [2, 3] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.061, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [8, 6, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.047, "cosStep": 0.089, "scale": 0.4 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 6, 8, 15], "tolerance": 0.0005 } }
+    },
+    {
+      "name": "exact_tiling_stride2_kernel2_f16_bias",
+      "provenance": {
+        "notes": "The f16 im2col-GEMM family shares the same col2im shader, so the collapsed tap needs a case on that side of the dtype split too. Shaped after Depth Anything V2's factor-2 reassemble."
+      },
+      "attrs": { "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.07, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 16, 2, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.11, "scale": 0.4 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 0.3 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 16, 16, 16], "tolerance": 0.03 } }
+    },
+    {
+      "name": "cols_gemm_2d_bias_m_and_n_tail",
+      "provenance": {
+        "notes": "An f32 rank-4 im2col-GEMM whose channel count is a multiple of 32 and whose BOTH gemm axes are ragged: gemmM = 12*3*3 = 108 tails a 32-row tile and gemmN = 9*11 = 99 tails a 64-column tile. Written to cover a subgroup-matrix 2D path that was then reverted, and kept because it is the shape that CAUGHT it: the shared sgmat kernel's materialized B load has no N guard, so this case read maxAbs 6.6 wrong while the same shape at N = 64 was exact."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 32, 9, 11],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.067, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 12, 3, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.053, "cosStep": 0.089, "scale": 0.4 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "sinStep": 0.131, "cosStep": 0.197, "scale": 0.3 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 12, 17, 21], "tolerance": 0.0005 } }
+    },
+    {
+      "name": "cols_gemm_2d_no_bias_m_aligned_n_below_tile",
+      "provenance": {
+        "notes": "The companion: gemmM = 16*4*4 = 256 is tile-aligned and gemmN = 5*6 = 30 is BELOW one column tile, the narrowest an N tail gets. Same provenance — it is half of the pair that showed a ragged N is not safe on the shared subgroup-matrix kernel."
+      },
+      "attrs": { "strides": [4, 4] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 64, 5, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.061, "scale": 0.4 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 16, 4, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.047, "cosStep": 0.083, "scale": 0.4 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 16, 20, 24], "tolerance": 0.0005 } }
+    },
+    {
+      "name": "f16_1d_bias",
+      "provenance": {
+        "notes": "Covers rank-3 f16 ConvTranspose with bias. Exact half-valued inputs keep bias-indexing errors observable without f16 rounding ambiguity."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 2, 8],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_1d_nobias_input_x" } }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [2, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/f16_1d_nobias_input_w" } }
+        },
+        "bias": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.25, -0.5, 1.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 3, 10], "tolerance": 0.02 } }
+    },
+    {
+      "name": "f16_1d_grouped_nobias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_OutputShape_1_group_2_for_transpose_path",
+        "notes": "Covers grouped rank-3 f16 ConvTranspose without bias. Distinct per-group weights 2 and 3 make group-indexing errors observable."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float16", "shape": [2, 1, 1], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 2], "tolerance": 0.01 } }
+    },
+    {
+      "name": "f16_1d_grouped_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_transpose_op_test.cc",
+        "test": "ConvTransposeTest.ConvTranspose_1D_OutputShape_1_group_2_for_transpose_path",
+        "notes": "Covers grouped rank-3 f16 ConvTranspose with per-output-channel bias. The two bias entries differ in sign and magnitude, making a wrong channel index observable."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 2, 2], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "w": { "dtype": "float16", "shape": [2, 1, 1], "data": { "kind": "values", "values": [2.0, 3.0] } },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 2, 2], "tolerance": 0.01 } }
+    },
+    {
+      "name": "zero_input_spatial_dim_3d_empty_output",
+      "provenance": {
+        "notes": "A zero input-depth dimension with a unit kernel produces a zero output-depth dimension. This exercises the direct rank-5 fallback when the GEMM column dimension is zero and the nonempty-output reduction does not apply."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 0, 2, 2], "data": { "kind": "values", "values": [] } },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 0, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "zero_input_spatial_dim_3d_empty_output_bias",
+      "provenance": {
+        "notes": "Bias does not make an output tensor with a zero spatial dimension nonempty. This is the bias-bearing sibling for the direct rank-5 zero-column fallback."
+      },
+      "attrs": {},
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 1, 0, 2, 2], "data": { "kind": "values", "values": [] } },
+        "w": { "dtype": "float32", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [2.0] } },
+        "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.75] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 0, 2, 2], "tolerance": 0 } }
+    }
+  ]
+}