sync 91d990483a17

Browse files

Files changed (16) hide show

README.md +47 -11
build/webgpu/bench.json +600 -7
build/webgpu/conv-1x1-channel-reduce.wgsl.jinja +10 -9
build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja +73 -30
build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja +36 -22
build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja +118 -21
build/webgpu/conv-direct-nd.wgsl.jinja +41 -15
build/webgpu/conv-direct-unrolled.wgsl.jinja +45 -22
build/webgpu/conv-im2col-ncdhw.wgsl.jinja +12 -15
build/webgpu/conv-im2col-nchw.wgsl.jinja +13 -14
build/webgpu/conv-splitk-reduce.wgsl.jinja +42 -11
build/webgpu/conv1d-tiled-reg.wgsl.jinja +34 -22
build/webgpu/conv2d-grouped-large-w4.wgsl.jinja +139 -66
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +78 -17
build/webgpu/test.json +1649 -35

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX `Conv` spec](https://onnx.ai/onnx/operators/onnx__Conv.html) for t
 ## 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 the batch size and `C` is the number of channels. | required |
-| `W` | `w` | `T` | — | — | Convolution filter weights of shape `(M x C/group x k1 x ... x kn)`, where `M` is the number of output feature maps. | required |
-| `B` | `bias` | `T` | `1` | — | Optional 1D 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 determined by the kernel size, strides, dilations, and padding. | required |
 ## Attributes
@@ -37,8 +37,8 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `auto_pad` | `"NOTSET"` | Automatic padding mode. `NOTSET` uses `pads`; `SAME_UPPER` and `SAME_LOWER` choose padding so each output spatial size is `ceil(input / stride)`; `VALID` uses no padding. |
-| `group` | `1` | Number of groups that input and output channels are split into; defaults to 1. |
 | `dilations` | — | Optional dilation factors, one positive integer per spatial axis. Omission means all ones. |
 | `kernel_shape` | — | Optional kernel shape, one positive integer per spatial axis. When present, it must match the spatial dimensions of the weight tensor; omission infers the shape from the weights. |
 | `pads` | — | Optional explicit padding in ONNX order `[begin_axis_0, ..., begin_axis_n, end_axis_0, ..., end_axis_n]`. Omission means all zeros; it cannot be combined with an automatic padding mode. |
 | `strides` | — | Optional stride factors, one positive integer per spatial axis. Omission means all ones. |
@@ -49,13 +49,45 @@ Attributes and default values (overridable per request):
 | --- | --- |
 | `T` | `float32`, `float16` |
 ## Device requirements
-Some implementation variants require `subgroup-matrix` 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
@@ -73,10 +105,14 @@ Some implementation variants require `subgroup-matrix` and `subgroups`. These ar
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Inputs
+| Name | Upstream name | 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 the batch size and `C` is the number of channels. | required |
+| `w` | `W` | `T` | — | — | Convolution filter weights of shape `(M x C/group x k1 x ... x kn)`, where `M` is the number of output feature maps. | required |
+| `bias` | `B` | `T` | `1` | — | Optional 1D bias of length M added to each output channel. | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | same as `x` | derived | Output tensor whose spatial dimensions are determined by the kernel size, strides, dilations, and padding. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `auto_pad` | `"NOTSET"` | Automatic padding mode. `NOTSET` uses `pads`; `SAME_UPPER` and `SAME_LOWER` choose padding so each output spatial size is `ceil(input / stride)`; `VALID` uses no padding. |
 | `dilations` | — | Optional dilation factors, one positive integer per spatial axis. Omission means all ones. |
+| `group` | `1` | Number of groups that input and output channels are split into; defaults to 1. |
 | `kernel_shape` | — | Optional kernel shape, one positive integer per spatial axis. When present, it must match the spatial dimensions of the weight tensor; omission infers the shape from the weights. |
 | `pads` | — | Optional explicit padding in ONNX order `[begin_axis_0, ..., begin_axis_n, end_axis_0, ..., end_axis_n]`. Omission means all zeros; it cannot be combined with an automatic padding mode. |
 | `strides` | — | Optional stride factors, one positive integer per spatial axis. Omission means all ones. |
 | --- | --- |
 | `T` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `pointwise_channel_reduce` — Handles pointwise unit-stride convolution with few output channels by assigning one invocation to each spatial position and accumulating every output channel in registers. It does not materialize im2col scratch.
+- `pointwise_channel_reduce_bias` — Handles pointwise unit-stride convolution with few output channels by assigning one invocation to each spatial position and accumulating every output channel in registers. It does not materialize im2col scratch.
+- `gemm_1x1_subgroup_matrix` — Multiplies 1x1 weights directly by NCHW input. Complete output-channel tiles use direct subgroup-matrix loads, eliminating input staging and reduction-loop barriers; partial tiles retain guarded staging. Tile geometry follows the device's workgroup limits.
+- `gemm_1x1_subgroup_matrix_bias` — Multiplies 1x1 weights directly by NCHW input. Complete output-channel tiles use direct subgroup-matrix loads, eliminating input staging and reduction-loop barriers; partial tiles retain guarded staging. Tile geometry follows the device's workgroup limits.
+- `im2col_gemm_subgroup_matrix` — Materializes convolution windows as a column matrix, then multiplies it by the weights with subgroup-matrix operations. It applies when reduction and spatial extents are tile-aligned and column scratch fits.
+- `im2col_gemm_subgroup_matrix_bias` — Materializes convolution windows as a column matrix, then multiplies it by the weights with subgroup-matrix operations. It applies when reduction and spatial extents are tile-aligned and column scratch fits.
+- `im2col_direct_f32_subgroup_matrix` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
+- `im2col_direct_f16_subgroup_matrix` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
+- `im2col_direct_f32_subgroup_matrix_bias` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
+- `im2col_direct_f16_subgroup_matrix_bias` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
+- `implicit_im2col_subgroup_matrix` — Gathers each logical im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth that materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
+- `implicit_im2col_subgroup_matrix_bias` — Gathers each logical im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth that materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
+- `implicit_im2col3d_subgroup_matrix` — Gathers each logical 3-D im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth a volumetric materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
+- `implicit_im2col3d_subgroup_matrix_bias` — Gathers each logical 3-D im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth a volumetric materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
+- `implicit_im2col_subgroup_matrix_splitk` — Partitions the implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums from scratch. The combine reassociates floating-point addition relative to the unsplit route.
+- `implicit_im2col_subgroup_matrix_bias_splitk` — Partitions the implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums from scratch. The combine reassociates floating-point addition relative to the unsplit route.
+- `im2col_gemm_subgroup_matrix_padded` — Materializes a zero-padded column matrix before a subgroup-matrix multiply, allowing logical reduction and spatial extents that are not tile-aligned. Padded values contribute zero and are excluded from the logical output.
+- `im2col_gemm_subgroup_matrix_padded_bias` — Materializes a zero-padded column matrix before a subgroup-matrix multiply, allowing logical reduction and spatial extents that are not tile-aligned. Padded values contribute zero and are excluded from the logical output.
+- `im2col_gemm_subgroup_matrix_padded_splitk` — Materializes a zero-padded column matrix, partitions its subgroup-matrix reduction across workgroups, and combines raw partial sums from scratch. It preserves the padded route for low-parallelism outputs; the combine reassociates floating-point addition.
+- `im2col_gemm_subgroup_matrix_padded_bias_splitk` — Materializes a zero-padded column matrix, partitions its subgroup-matrix reduction across workgroups, and combines raw partial sums from scratch. It preserves the padded route for low-parallelism outputs; the combine reassociates floating-point addition.
+- `implicit_im2col_tiled_bias_reg_m32` — Gathers logical im2col elements directly from X while filling register-blocked GEMM tiles. It avoids column scratch and provides the portable implicit route without subgroup-matrix support.
+- `implicit_im2col_tiled_reg` — Gathers logical im2col elements directly from X while filling register-blocked GEMM tiles. It avoids column scratch and provides the portable implicit route without subgroup-matrix support.
+- `implicit_im2col_tiled_bias_reg` — Gathers logical im2col elements directly from X while filling register-blocked GEMM tiles. It avoids column scratch and provides the portable implicit route without subgroup-matrix support.
+- `implicit_im2col_tiled_reg_splitk` — Partitions the register-blocked implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
+- `implicit_im2col_tiled_bias_reg_splitk` — Partitions the register-blocked implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
+- `gemm_1x1_tiled_reg_splitk` — Partitions the register-blocked 1x1 GEMM reduction across additional workgroups when the output grid is small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
+- `gemm_1x1_tiled_bias_reg_splitk` — Partitions the register-blocked 1x1 GEMM reduction across additional workgroups when the output grid is small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 ## 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, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Conv",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
@@ -139,8 +138,9 @@
     {
       "name": "audio-feature-1d-f32-b1c256m256-w4096-k7-pathology",
       "preset": "stress",
       "provenance": {
-        "source": "authored for performance coverage",
         "notes": "Realistic long-sequence audio feature block with K and output width aligned for the materialized subgroup-matrix Conv1d path; reduced-feature tiers retain the general implicit-GEMM and direct NCW implementations."
       },
       "vars": { "batch": 1, "inChannels": 256, "outChannels": 256, "inW": 4096, "kernelW": 7, "strideW": 1, "padW": 3 },
@@ -198,7 +198,7 @@
       "name": "medical-segmentation-3d-f32-b1c32m32-32x32x32-k3-pathology",
       "preset": "stress",
       "provenance": {
-        "source": "authored for performance coverage",
         "notes": "Realistic 3D medical-segmentation feature block for the direct NCDHW kernel. Each output voxel serially scans 32 input channels and a 3x3x3 kernel."
       },
       "vars": {
@@ -547,8 +547,8 @@
     {
       "name": "conv-f32-group4-k5-64c-48x48-mid-kernel-band",
       "provenance": {
-        "source": "authored for performance coverage",
-        "notes": "Grouped 5x5 in the middle of the kernel-size band grouped_large_kernel_w4 used to refuse. No grouped benchmark existed between k3 and k9, so the k>=7 floor had nothing measured on either side of it."
       },
       "preset": "model",
       "vars": {
@@ -976,8 +976,8 @@
       "name": "manifest-gap-conv-f32-group8-k9-256c-64x63-w4-boundary",
       "preset": "stress",
       "provenance": {
-        "source": "authored for boundary coverage",
-        "notes": "A realistic odd-width grouped large-kernel convolution that falls outside grouped_large_kernel_w4's output-width % 4 == 0 guard; the paired W64 case isolates the route boundary."
       },
       "vars": {
         "batch": 1,
@@ -2163,6 +2163,599 @@
           }
         ]
       }
     }
   ]
 }

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
     {
       "name": "audio-feature-1d-f32-b1c256m256-w4096-k7-pathology",
       "preset": "stress",
+      "tunableSpace": { "CONV1D_TILE_M": [4, 8], "CONV1D_TILE_N": [4, 8], "CONV1D_K_TILE": [16, 32] },
       "provenance": {
+        "source": "repository-authored benchmark",
         "notes": "Realistic long-sequence audio feature block with K and output width aligned for the materialized subgroup-matrix Conv1d path; reduced-feature tiers retain the general implicit-GEMM and direct NCW implementations."
       },
       "vars": { "batch": 1, "inChannels": 256, "outChannels": 256, "inW": 4096, "kernelW": 7, "strideW": 1, "padW": 3 },
       "name": "medical-segmentation-3d-f32-b1c32m32-32x32x32-k3-pathology",
       "preset": "stress",
       "provenance": {
+        "source": "repository-authored benchmark",
         "notes": "Realistic 3D medical-segmentation feature block for the direct NCDHW kernel. Each output voxel serially scans 32 input channels and a 3x3x3 kernel."
       },
       "vars": {
     {
       "name": "conv-f32-group4-k5-64c-48x48-mid-kernel-band",
       "provenance": {
+        "source": "repository-authored benchmark",
+        "notes": "A grouped 5x5 convolution exercises the wide-output grouped route at an intermediate kernel size between the 3x3 and 9x9 benchmark shapes."
       },
       "preset": "model",
       "vars": {
       "name": "manifest-gap-conv-f32-group8-k9-256c-64x63-w4-boundary",
       "preset": "stress",
       "provenance": {
+        "source": "repository-authored benchmark",
+        "notes": "A realistic grouped 9x9 convolution produces 63 columns, so the final four-lane output group contains three valid columns and uses masked tail stores."
       },
       "vars": {
         "batch": 1,
           }
         ]
       }
+    },
+    {
+      "name": "conv-f32-depthwise-96c-56x56-k3-bias",
+      "preset": "edge",
+      "attrs": { "group": 96, "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 96, 56, 56], "dist": "normal", "seed": 7, "scale": 1 },
+        "w": { "dtype": "float32", "shape": [96, 1, 3, 3], "dist": "normal", "seed": 11, "scale": 0.2 },
+        "bias": { "dtype": "float32", "shape": [96], "dist": "normal", "seed": 13, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 96, 56, 56] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 96 * 56 * 56 * 9" }] }
+    },
+    {
+      "name": "conv-f32-depthwise-96c-56x56-k3-no-bias-control",
+      "preset": "edge",
+      "attrs": { "group": 96, "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 96, 56, 56], "dist": "normal", "seed": 7, "scale": 1 },
+        "w": { "dtype": "float32", "shape": [96, 1, 3, 3], "dist": "normal", "seed": 11, "scale": 0.2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 96, 56, 56] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 96 * 56 * 56 * 9" }] }
+    },
+    {
+      "name": "conv-f32-depthwise-576c-14x14-k3-bias",
+      "preset": "edge",
+      "attrs": { "group": 576, "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [32, 576, 14, 14], "dist": "normal", "seed": 7, "scale": 1 },
+        "w": { "dtype": "float32", "shape": [576, 1, 3, 3], "dist": "normal", "seed": 11, "scale": 0.2 },
+        "bias": { "dtype": "float32", "shape": [576], "dist": "normal", "seed": 13, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [32, 576, 14, 14] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 32 * 576 * 14 * 14 * 9" }] }
+    },
+    {
+      "name": "conv-f32-depthwise-960c-7x7-k3-bias",
+      "preset": "edge",
+      "attrs": { "group": 960, "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [32, 960, 7, 7], "dist": "normal", "seed": 7, "scale": 1 },
+        "w": { "dtype": "float32", "shape": [960, 1, 3, 3], "dist": "normal", "seed": 11, "scale": 0.2 },
+        "bias": { "dtype": "float32", "shape": [960], "dist": "normal", "seed": 13, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [32, 960, 7, 7] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 32 * 960 * 7 * 7 * 9" }] }
+    },
+    {
+      "name": "conv-f32-depthwise-192c-28x28-k3-bias",
+      "preset": "edge",
+      "attrs": { "group": 192, "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [32, 192, 28, 28], "dist": "normal", "seed": 7, "scale": 1 },
+        "w": { "dtype": "float32", "shape": [192, 1, 3, 3], "dist": "normal", "seed": 11, "scale": 0.2 },
+        "bias": { "dtype": "float32", "shape": [192], "dist": "normal", "seed": 13, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [32, 192, 28, 28] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 32 * 192 * 28 * 28 * 9" }] }
+    },
+    {
+      "name": "conv-f16-64c-56x56-k3-bias-sgmat",
+      "preset": "edge",
+      "attrs": { "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 64, 56, 56], "dist": "normal", "seed": 5, "scale": 1 },
+        "w": { "dtype": "float16", "shape": [64, 64, 3, 3], "dist": "normal", "seed": 9, "scale": 0.05 },
+        "bias": { "dtype": "float16", "shape": [64], "dist": "normal", "seed": 3, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 64, 56, 56] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 64 * 64 * 56 * 56 * 9" }] }
+    },
+    {
+      "name": "conv-f32-64c-56x56-k3-bias-sgmat-control",
+      "preset": "edge",
+      "attrs": { "kernel_shape": [3, 3], "pads": [1, 1, 1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 56, 56], "dist": "normal", "seed": 5, "scale": 1 },
+        "w": { "dtype": "float32", "shape": [64, 64, 3, 3], "dist": "normal", "seed": 9, "scale": 0.05 },
+        "bias": { "dtype": "float32", "shape": [64], "dist": "normal", "seed": 3, "scale": 0.1 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 56, 56] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 64 * 64 * 56 * 56 * 9" }] }
+    },
+    {
+      "name": "conv-reuse-float32-m32-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [32, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m32-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [32, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m64-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [64, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m64-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [64, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m96-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [96, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 96, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m96-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [96, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 96, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m128-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [128, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 128, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m128-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [128, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 128, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m256-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [256, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 256, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float32-m256-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float32", "shape": [256, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 256, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m32-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [32, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 32, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m32-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [32, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 32, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m64-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [64, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 64, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m64-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [64, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 64, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m96-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [96, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 96, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m96-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [96, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 96, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m128-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [128, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 128, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m128-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [128, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 128, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m256-c64-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 64, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [256, 64, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 256, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "conv-reuse-float16-m256-c256-n16384",
+      "preset": "edge",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1, 256, 128, 128], "dist": "normal", "seed": 821, "scale": 0.2 },
+        "w": { "dtype": "float16", "shape": [256, 256, 3, 3], "dist": "normal", "seed": 938, "scale": 0.05 }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 256, 128, 128] } },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "direct-1x1-float32-m63-k512-n4096",
+      "preset": "model",
+      "vars": {
+        "batch": 1,
+        "inChannels": 512,
+        "outChannels": 63,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 1,
+        "kernelW": 1,
+        "strideH": 1,
+        "strideW": 1,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "shape": [1, 512, 64, 64], "dtype": "float32", "dist": "normal", "seed": 354, "scale": 0.2 },
+        "w": { "shape": [63, 512, 1, 1], "dtype": "float32", "dist": "normal", "seed": 355, "scale": 0.02 }
+      },
+      "outputs": { "y": { "shape": [1, 63, 64, 64], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Row-tile boundary sweep: full, partial, and device-dependent complete row bands share the same reduction and spatial extent."
+      }
+    },
+    {
+      "name": "direct-1x1-float32-m64-k512-n4096",
+      "preset": "model",
+      "vars": {
+        "batch": 1,
+        "inChannels": 512,
+        "outChannels": 64,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 1,
+        "kernelW": 1,
+        "strideH": 1,
+        "strideW": 1,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "shape": [1, 512, 64, 64], "dtype": "float32", "dist": "normal", "seed": 354, "scale": 0.2 },
+        "w": { "shape": [64, 512, 1, 1], "dtype": "float32", "dist": "normal", "seed": 355, "scale": 0.02 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 64, 64], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Row-tile boundary sweep: full, partial, and device-dependent complete row bands share the same reduction and spatial extent."
+      }
+    },
+    {
+      "name": "direct-1x1-float32-m96-k512-n4096",
+      "preset": "model",
+      "vars": {
+        "batch": 1,
+        "inChannels": 512,
+        "outChannels": 96,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 1,
+        "kernelW": 1,
+        "strideH": 1,
+        "strideW": 1,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "shape": [1, 512, 64, 64], "dtype": "float32", "dist": "normal", "seed": 354, "scale": 0.2 },
+        "w": { "shape": [96, 512, 1, 1], "dtype": "float32", "dist": "normal", "seed": 355, "scale": 0.02 }
+      },
+      "outputs": { "y": { "shape": [1, 96, 64, 64], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Row-tile boundary sweep: full, partial, and device-dependent complete row bands share the same reduction and spatial extent."
+      }
+    },
+    {
+      "name": "direct-1x1-float16-m63-k512-n4096",
+      "preset": "model",
+      "vars": {
+        "batch": 1,
+        "inChannels": 512,
+        "outChannels": 63,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 1,
+        "kernelW": 1,
+        "strideH": 1,
+        "strideW": 1,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "shape": [1, 512, 64, 64], "dtype": "float16", "dist": "normal", "seed": 354, "scale": 0.2 },
+        "w": { "shape": [63, 512, 1, 1], "dtype": "float16", "dist": "normal", "seed": 355, "scale": 0.02 }
+      },
+      "outputs": { "y": { "shape": [1, 63, 64, 64], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Row-tile boundary sweep: full, partial, and device-dependent complete row bands share the same reduction and spatial extent."
+      }
+    },
+    {
+      "name": "direct-1x1-float16-m64-k512-n4096",
+      "preset": "model",
+      "vars": {
+        "batch": 1,
+        "inChannels": 512,
+        "outChannels": 64,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 1,
+        "kernelW": 1,
+        "strideH": 1,
+        "strideW": 1,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "shape": [1, 512, 64, 64], "dtype": "float16", "dist": "normal", "seed": 354, "scale": 0.2 },
+        "w": { "shape": [64, 512, 1, 1], "dtype": "float16", "dist": "normal", "seed": 355, "scale": 0.02 }
+      },
+      "outputs": { "y": { "shape": [1, 64, 64, 64], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Row-tile boundary sweep: full, partial, and device-dependent complete row bands share the same reduction and spatial extent."
+      }
+    },
+    {
+      "name": "direct-1x1-float16-m96-k512-n4096",
+      "preset": "model",
+      "vars": {
+        "batch": 1,
+        "inChannels": 512,
+        "outChannels": 96,
+        "inH": 64,
+        "inW": 64,
+        "kernelH": 1,
+        "kernelW": 1,
+        "strideH": 1,
+        "strideW": 1,
+        "padH": 0,
+        "padW": 0
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "shape": [1, 512, 64, 64], "dtype": "float16", "dist": "normal", "seed": 354, "scale": 0.2 },
+        "w": { "shape": [96, 512, 1, 1], "dtype": "float16", "dist": "normal", "seed": 355, "scale": 0.02 }
+      },
+      "outputs": { "y": { "shape": [1, 96, 64, 64], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic",
+        "notes": "Row-tile boundary sweep: full, partial, and device-dependent complete row bands share the same reduction and spatial extent."
+      }
     }
   ]
 }

build/webgpu/conv-1x1-channel-reduce.wgsl.jinja CHANGED Viewed

@@ -1,13 +1,14 @@
 // Pointwise convolution with few output channels is evaluated as a per-position
 // channel reduction instead of a tiled GEMM. One invocation walks the input
-// channels and accumulates every output channel directly in registers.
 // Consecutive invocations read consecutive positions from each NCHW channel
-// plane, so each channel remains a coalesced stream without shared-memory
-// staging or barriers.
 //
-// The output-channel accumulators are named and unrolled because dynamically
-// indexed local arrays may spill on some backends. Selection therefore limits
-// this route to cases whose complete accumulator set remains in registers.
 {% if usesF16 %}
 enable f16;
 {% endif %}
@@ -20,10 +21,10 @@ const COUNT: u32 = {{ batchCount }}u * PLANE;
 const WORKGROUP_SIZE: u32 = {{ channelReduceWorkgroupSize }}u;
 @compute @workgroup_size({{ channelReduceWorkgroupSize }}, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past
-  // maxComputeWorkgroupsPerDimension and reduces to the 1D form at y=0.
-  let idx = gid.x + gid.y * nwg.x * WORKGROUP_SIZE;
   if (idx >= COUNT) {
     return;
   }

 // Pointwise convolution with few output channels is evaluated as a per-position
 // channel reduction instead of a tiled GEMM. One invocation walks the input
+// channels and accumulates every output channel in a statically named local.
 // Consecutive invocations read consecutive positions from each NCHW channel
+// plane without shared-memory staging or barriers.
 //
+// The output-channel accumulators are named explicitly. Selection caps the
+// output-channel count, bounding the invocation-local accumulator storage.
+{% set usesF16 = usesF16 is defined and usesF16 %}
+{% set batchCount = batchCount | default(0) %}
+{% set yScalar = yScalar | default("f32") %}
 {% if usesF16 %}
 enable f16;
 {% endif %}
 const WORKGROUP_SIZE: u32 = {{ channelReduceWorkgroupSize }}u;
 @compute @workgroup_size({{ channelReduceWorkgroupSize }}, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past
+  // the per-axis dispatch fold width and reduces to the 1D form at y=0.
+  let idx = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WORKGROUP_SIZE;
   if (idx >= COUNT) {
     return;
   }

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

@@ -1,11 +1,10 @@
-// 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.
 //
@@ -14,16 +13,35 @@
 // 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 %}
@@ -50,19 +68,19 @@ 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 %}
@@ -91,7 +109,7 @@ const CONV_IN_CHANNELS: u32 = {{ convInChannels }}u;
 {% if splitKPartial %}
 // Split-K publishes raw f32 partials. A second pass owns the bias, optional
 // epilogue, output cast, and logical-shape store.
-const SPLIT_K: u32 = {{ splitK }}u;
 const K_TILES: u32 = {{ kTiles }}u;
 const PARTIAL_ROWS: u32 = {{ mPadded }}u;
 const PARTIAL_COLS: u32 = {{ nPadded }}u;
@@ -100,9 +118,36 @@ const PARTIAL_SLICE_STRIDE: u32 = {{ batchCount }}u * PARTIAL_BATCH_STRIDE;
 {% 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;
   }
@@ -110,6 +155,9 @@ const PROJECTION_OUT_C: u32 = {{ projectionChannels }}u;
 {% if hasBias %}
   v = v + f32(bias[m]);
 {% endif %}
 {% if hasActivation %}
   v = fused_act(v);
 {% endif %}
@@ -191,7 +239,7 @@ var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
       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;
@@ -322,8 +370,7 @@ fn main(
 {% 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 %}
@@ -415,7 +462,7 @@ fn main(
     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;
@@ -428,11 +475,7 @@ fn main(
 {% 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 %}

+// Register-blocked batched matrix product:
+// C[b] (M x N) = A (M x K) * B[b] (K x N). Workgroup lanes compute TM x TN
+// micro-tiles. The defaults use a 16x8 lane grid with an 8x4 micro-tile over a
+// 64x64 output tile. A is shared across the batch; B and C use the dispatch.z
+// batch offset. Accumulation is f32; an optional epilogue applies configured
+// bias, residual, and activation terms in the accumulator domain on store.
+// Every matrix 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.
 //
 // 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.
 {{ env.wgsl.resourceDeclarations }}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}
 // 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
+// repeated input loads for writing and rereading an expanded column matrix.
 {% set implicitIm2col = implicitIm2col is defined and implicitIm2col %}
 {% set fusedNarrowProjection = fusedNarrowProjection is defined and fusedNarrowProjection %}
+{% set projectionChannels = projectionOutChannels | default(0) %}
+{% set projectionInputAct = inputActivation | default("none") %}
+{% set projectionOutputAct = outputActivation | default("none") %}
 {% set narrowProjectionTile = "tileB" if GEMM_BK >= GEMM_BM else "projectionTile" %}
 {% set splitKValue = splitK if splitK is defined else 1 %}
 {% set splitKPartial = splitKValue > 1 %}
+{% set kTiles = kTiles | default(0) %}
+{% set mPadded = mPadded | default(0) %}
+{% set nPadded = nPadded | default(0) %}
+{% set batchCount = batchCount | default(0) %}
 {% 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 %}
 {% if splitKPartial %}
 // Split-K publishes raw f32 partials. A second pass owns the bias, optional
 // epilogue, output cast, and logical-shape store.
+const SPLIT_K: u32 = {{ splitKValue }}u;
 const K_TILES: u32 = {{ kTiles }}u;
 const PARTIAL_ROWS: u32 = {{ mPadded }}u;
 const PARTIAL_COLS: u32 = {{ nPadded }}u;
 {% endif %}
 {% if fusedNarrowProjection %}
+{% set inputActivation = inputActivation | default("none") %}
+{% set outputActivation = outputActivation | default("none") %}
+{% if inputActivation == "relu" or outputActivation == "relu" %}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+fn projection_relu(value: f32) -> f32 {
+  var out = max(value, 0.0);
+  if (is_nan_f32(value)) {
+    out = value;
+  }
+  return out;
+}
+{% endif %}
+{% if outputActivation == "sigmoid" %}
+fn sigmoid_safe(x: f32) -> f32 {
+  if (x >= 0.0) {
+    let z = exp(-x);
+    return 1.0 / (1.0 + z);
+  }
+  let z = exp(x);
+  return z / (1.0 + z);
+}
+{% endif %}
 const PROJECTION_OUT_C: u32 = {{ projectionChannels }}u;
+{% else %}{% set emitConvStoreOut = not splitKPartial %}{% set hasZ = hasZ is defined and hasZ %}
+{% if emitConvStoreOut | default(true) %}fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
   if (m >= params.M || n >= params.N) {
     return;
   }
 {% if hasBias %}
   v = v + f32(bias[m]);
 {% endif %}
+{% if hasZ %}
+  v = v + f32(zResidual[yBase + m * params.N + n]);
+{% endif %}
 {% if hasActivation %}
   v = fused_act(v);
 {% endif %}
       let bk = {{ rowBase }} + br;
       var bvec = vec4<{{ tileT }}>({{ tileT }}(0.0));
       if (bk < {{ rowLimit }}) {
+        // k flattens the input-channel and spatial-tap axes; n is the output position.
         let ic = bk / CONV_KSIZE;
         let kq = bk % CONV_KSIZE;
         let kh = kq / CONV_KERNEL_W;
 {% endif %}
   let li = lid.y * WG_X + lid.x;
+  // The statically named accumulator vectors form TM independent FMA chains.
 {% for row in range(GEMM_TM) %}
   var acc{{ row }} = vec{{ GEMM_TN }}<f32>(0.0);
 {% endfor %}
     let projectionYBase = batch * PROJECTION_OUT_C * N;
 {% for oc in range(projectionChannels) %}
+    let projected{{ oc }} = projectionAcc{{ oc }};
 {% endfor %}
 {% for column in range(4) %}
     let projectionN{{ column }} = projectionNBase + {{ column }}u;
 {% else %}
       let activated{{ oc }}_{{ column }} = projected{{ oc }}.{{ components[column] }};
 {% endif %}
       y[projectionYBase + {{ oc }}u * N + projectionN{{ column }}] = activated{{ oc }}_{{ column }};
 {% endfor %}
     }
 {% endfor %}

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

@@ -1,29 +1,39 @@
-// 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 %}
@@ -32,8 +42,9 @@ 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;
@@ -42,6 +53,9 @@ fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
 {% if hasBias %}
   v = v + f32(bias[m]);
 {% endif %}
 {% if hasActivation %}
   v = fused_act(v);
 {% endif %}

+// Tiled batched matrix product C[b] (M x N) = A (M x K) * B[b] (K x N).
+// A 16x16 workgroup computes a 32x32 output tile with a 2x2 micro-tile per
+// thread and a 16-element K tile. A is shared across the batch; B and C use the
+// dispatch.z batch offset. Accumulation is f32 and is cast to the output type on
+// store. A configured epilogue may add bias, residual, and activation terms.
+// f16 operands remain packed in workgroup memory and widen when consumed. All
+// M, N, and K accesses are bounds-checked.
 {{ env.wgsl.resourceDeclarations }}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}
 const BM: u32 = 32u;
 const BN: u32 = 32u;
+// Store one output element with the configured bias/residual/activation
+// epilogue in the f32 accumulator domain.
+{% set hasZ = hasZ is defined and hasZ %}
 fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
   if (m >= params.M || n >= params.N) {
     return;
 {% if hasBias %}
   v = v + f32(bias[m]);
 {% endif %}
+{% if hasZ %}
+  v = v + f32(zResidual[yBase + m * params.N + n]);
+{% endif %}
 {% if hasActivation %}
   v = fused_act(v);
 {% endif %}

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

@@ -1,6 +1,7 @@
-// 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
@@ -10,6 +11,7 @@
 // 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 implicit = implicitIm2col is defined and implicitIm2col %}
@@ -18,9 +20,32 @@
 {% 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;
@@ -28,7 +53,35 @@ enable subgroup_size_control;
 enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 {% set operandScalar = fScalar %}
 {% set accScalar = "f32" %}
@@ -92,7 +145,7 @@ const B_BATCH_STRIDE: u32 = {{ kLoopVar }} * {{ nColsVar }};
 const C_BATCH_STRIDE: u32 = M * N;
 {% endif %}
 {% if splitKPartial %}
-const SPLIT_K: u32 = {{ splitK }}u;
 // Whole 32-wide K tiles per slice, rounded up, with the last slice clamped to
 // K_LOOP. This avoids padding K to a multiple of SPLIT_K times the tile width.
 // A slice entirely past K_LOOP runs zero iterations and stores zero
@@ -108,12 +161,15 @@ 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;
@@ -270,12 +326,35 @@ fn loadSHMB(b_base: u32, tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
 {% 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 %}
@@ -283,8 +362,11 @@ fn fused_act(v: f32) -> f32 {
 {% 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 %}
@@ -292,7 +374,7 @@ fn epi(raw: f32) -> {{ T }} {
 }
 {% endif %}
 {% macro store_val(valExpr, idxExpr) %}
-{% if hasActivation or hasZ %}epi({{ valExpr }})
 {%- else %}{{ T }}({{ valExpr }})
 {%- endif %}
 {% endmacro %}
@@ -377,20 +459,23 @@ fn main(
 {% else %}
   for (var kidx = 0u; kidx < {{ kLoopVar }}; kidx = kidx + TILE_K) {
 {% endif %}
     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>, row_major>(&tile_A, matrix_a_offset, TILE_K);
-      var matA1: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, row_major>(&tile_A, matrix_a_offset + 8u * TILE_K, 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>, col_major>(&tile_B, matrix_b_offset, TILE_K);
-      var matB1: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>, col_major>(&tile_B, matrix_b_offset + 8u * TILE_K, TILE_K);
-      var matB2: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>, col_major>(&tile_B, matrix_b_offset + 16u * TILE_K, TILE_K);
-      var matB3: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>, col_major>(&tile_B, matrix_b_offset + 24u * TILE_K, TILE_K);
       matC00 = subgroupMatrixMultiplyAccumulate(matA0, matB0, matC00);
       matC01 = subgroupMatrixMultiplyAccumulate(matA0, matB1, matC01);
@@ -401,7 +486,9 @@ fn main(
       matC12 = subgroupMatrixMultiplyAccumulate(matA1, matB2, matC12);
       matC13 = subgroupMatrixMultiplyAccumulate(matA1, matB3, matC13);
     }
     workgroupBarrier();
   }
 {% if fusedNarrowProjection %}
@@ -431,7 +518,11 @@ fn main(
 {% 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 %}
@@ -439,9 +530,15 @@ fn main(
     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 splitKPartial %}
@@ -490,7 +587,7 @@ fn main(
     raw += f32(bias[global_row]);
 {% endif %}
 {% if hasActivation or hasZ %}
-    y[y_index] = epi(raw);
 {% else %}
     y[y_index] = {{ T }}(raw);
 {% endif %}

+// Batched subgroup-matrix product C[b] (M x N) = A (M x K) * B[b] (K x N).
+// A is row-major; B may be a direct matrix view, a materialized column matrix,
+// or an on-demand gather. Accumulation is f32. A configured epilogue may apply
+// bias or other output transforms after the full K reduction.
 //
 // Full 8x8 matrix loads require complete K/N tiles. The aligned path enforces
 // that geometry; `padded` zero-fills tails, and `implicitIm2col` gathers and
 // 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 directMatrixInputs = directMatrixInputs is defined and directMatrixInputs %}
 {% set fusedNarrowProjection = fusedNarrowProjection if fusedNarrowProjection is defined else false %}
 {% set polyphase = polyphaseConvTranspose is defined and polyphaseConvTranspose %}
 {% set implicit = implicitIm2col is defined and implicitIm2col %}
 {% set splitKPartial = splitKValue > 1 %}
 {% set kLoopVar = "K_LOOP" if padded else "K" %}
 {% set nColsVar = "N_COLS" if padded else "N" %}
+{% set hasZ = hasZ is defined and hasZ %}
+{% set projectionOutChannels = projectionOutChannels | default(0) %}
+{% set hasProjectionBias = hasProjectionBias is defined and hasProjectionBias %}
+{% set hasOutputScale = hasOutputScale is defined and hasOutputScale %}
+{% set convKernelH = convKernelH | default(0) %}
+{% set convKernelW = convKernelW | default(0) %}
+{% set convStrideH = convStrideH | default(0) %}
+{% set convStrideW = convStrideW | default(0) %}
+{% set convDilationH = convDilationH | default(0) %}
+{% set convDilationW = convDilationW | default(0) %}
+{% set convPadTop = convPadTop | default(0) %}
+{% set convPadLeft = convPadLeft | default(0) %}
+{% set convInH = convInH | default(0) %}
+{% set convInW = convInW | default(0) %}
+{% set convOutW = convOutW | default(0) %}
+{% set convInChannels = convInChannels | default(0) %}
+{% set convKernelD = convKernelD | default(0) %}
+{% set convStrideD = convStrideD | default(0) %}
+{% set convDilationD = convDilationD | default(0) %}
+{% set convPadFront = convPadFront | default(0) %}
+{% set convInD = convInD | default(0) %}
+{% set convOutH = convOutH | default(0) %}
+{% set nPadded = nPadded | default(0) %}
+{% set mPadded = mPadded | default(0) %}
+{% set kChunk = kChunk | default(0) %}
+{% set batchCount = batchCount | default(0) %}
 enable subgroups;
 {% if pinSubgroupSize32 %}
 enable subgroup_size_control;
 enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
+{% if fusedNarrowProjection %}{% set inputActivation = inputActivation | default("none") %}
+{% set outputActivation = outputActivation | default("none") %}
+{% if inputActivation == "relu" or outputActivation == "relu" %}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+fn projection_relu(value: f32) -> f32 {
+  var out = max(value, 0.0);
+  if (is_nan_f32(value)) {
+    out = value;
+  }
+  return out;
+}
+{% endif %}
+{% if outputActivation == "sigmoid" %}
+fn sigmoid_safe(x: f32) -> f32 {
+  if (x >= 0.0) {
+    let z = exp(-x);
+    return 1.0 / (1.0 + z);
+  }
+  let z = exp(x);
+  return z / (1.0 + z);
+}
+{% endif %}
+{% endif %}
 {% set operandScalar = fScalar %}
 {% set accScalar = "f32" %}
 const C_BATCH_STRIDE: u32 = M * N;
 {% endif %}
 {% if splitKPartial %}
+const SPLIT_K: u32 = {{ splitKValue }}u;
 // Whole 32-wide K tiles per slice, rounded up, with the last slice clamped to
 // K_LOOP. This avoids padding K to a multiple of SPLIT_K times the tile width.
 // A slice entirely past K_LOOP runs zero iterations and stores zero
 const SUB_COLS: u32 = {{ subColsValue }}u;
 const SUB_ROWS: u32 = {{ subRowsValue }}u;
+{% if not directMatrixInputs %}
 var<workgroup> tile_A: array<{{ operandScalar }}, {{ tileRowsValue }} * 32>;
 var<workgroup> tile_B: array<{{ operandScalar }}, {{ tileColsValue }} * 32>;
+{% endif %}
 {% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
 var<workgroup> scratch: array<array<array<{{ accScalar }}, 64>, 4>, {{ (workgroupThreadsValue / 32)|int }}>;
 {% endif %}
+{% if not directMatrixInputs %}
 fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
   let a_global = tile_base + row;
   let col = c_idx * 8u;
 {% endif %}
 }
+{% endif %}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}
 {% 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{% if hasZ %}, yIndex: u32{% endif %}) -> {{ T }} {
   var r = raw;
+{% if hasZ %}
+  r = r + f32(zResidual[yIndex]);
+{% endif %}
 {% if hasActivation %}
   r = fused_act(r);
 {% endif %}
 }
 {% endif %}
 {% macro store_val(valExpr, idxExpr) %}
+{% if hasActivation or hasZ %}epi({{ valExpr }}{% if hasZ %}, {{ idxExpr }}{% endif %})
 {%- else %}{{ T }}({{ valExpr }})
 {%- endif %}
 {% endmacro %}
 {% else %}
   for (var kidx = 0u; kidx < {{ kLoopVar }}; kidx = kidx + TILE_K) {
 {% endif %}
+{% if not directMatrixInputs %}
     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();
+{% endif %}
     for (var step = 0u; step < TILE_K; step = step + 8u) {
+{% set directInputs = directMatrixInputs is defined and directMatrixInputs %}
+      let matrix_a_offset = {% if directInputs %}(a_global_base + subtile_idy * SUB_ROWS) * K + kidx + step{% else %}subtile_idy * SUB_ROWS * TILE_K + step{% endif %};
+{% for r in range(2) %}
+      var matA{{ r }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, row_major>(&{{ "w" if directInputs else "tile_A" }}, matrix_a_offset{% if r > 0 %} + 8u * {{ "K" if directInputs else "TILE_K" }}{% endif %}, {{ "K" if directInputs else "TILE_K" }});
+{% endfor %}
+      let matrix_b_offset = {% if directInputs %}b_base + (kidx + step) * N + b_global_base + subtile_idx * SUB_COLS{% else %}subtile_idx * SUB_COLS * TILE_K + step{% endif %};
+{% for c in range(4) %}
+      var matB{{ c }}: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>, {% if directInputs %}row_major{% else %}col_major{% endif %}>(&{{ "xm" if directInputs else "tile_B" }}, matrix_b_offset{% if c > 0 %} + {{ c * 8 }}u{% if not directInputs %} * TILE_K{% endif %}{% endif %}, {{ "N" if directInputs else "TILE_K" }});
+{% endfor %}
       matC00 = subgroupMatrixMultiplyAccumulate(matA0, matB0, matC00);
       matC01 = subgroupMatrixMultiplyAccumulate(matA0, matB1, matC01);
       matC12 = subgroupMatrixMultiplyAccumulate(matA1, matB2, matC12);
       matC13 = subgroupMatrixMultiplyAccumulate(matA1, matB3, matC13);
     }
+{% if not directMatrixInputs %}
     workgroupBarrier();
+{% endif %}
   }
 {% if fusedNarrowProjection %}
 {% 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 %};
+{% if inputActivation == "relu" %}
+      let activated_value = projection_relu(producer_value);
+{% else %}
       let activated_value = producer_value;
+{% endif %}
 {% 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 }}{% if hasProjectionBias %} + f32(projectionBias[{{ oc }}u]){% endif %};
+{% if outputActivation == "relu" %}
+    let activated{{ oc }} = projection_relu(projected{{ oc }});
+{% elif outputActivation == "sigmoid" %}
+    let activated{{ oc }} = sigmoid_safe(projected{{ oc }});
+{% else %}
     let activated{{ oc }} = projected{{ oc }};
+{% endif %}
+    y[fused_y_base + {{ oc }}u * N] = activated{{ oc }}{% if hasOutputScale %} * params.outputScale{% endif %};
 {% endfor %}
   }
 {% elif splitKPartial %}
     raw += f32(bias[global_row]);
 {% endif %}
 {% if hasActivation or hasZ %}
+    y[y_index] = epi(raw{% if hasZ %}, y_index{% endif %});
 {% else %}
     y[y_index] = {{ T }}(raw);
 {% endif %}

build/webgpu/conv-direct-nd.wgsl.jinja CHANGED Viewed

@@ -1,36 +1,58 @@
-// Direct {{ source.spatialRank }}D convolution for channels-first tensors. Inputs are
 // accumulated in f32 and narrowed once at the output.
 // A fused epilogue may apply a residual input and activation before the store.
-{% if source.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 %}
 const WG: u32 = {{ convWorkgroupSize }}u;
 @compute @workgroup_size(WG)
-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 >= params.count) {
     return;
   }
   let ow = index % params.outW;
   var q = index / params.outW;
-{% if source.spatialRank >= 2 %}
   let oh = q % params.outH;
   q /= params.outH;
 {% endif %}
-{% if source.spatialRank == 3 %}
   let od = q % params.outD;
   q /= params.outD;
 {% endif %}
@@ -39,19 +61,19 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
   let group = oc / params.outChannelsPerGroup;
   var acc = 0.0;
-{% if source.hasBias %}
   acc = f32(bias[oc]);
 {% endif %}
   for (var ic = 0u; ic < params.weightInChannels; ic += 1u) {
     let inputChannel = group * params.inChannelsPerGroup + ic;
-{% if source.spatialRank == 3 %}
     for (var kd = 0u; kd < params.kernelD; kd += 1u) {
       let id = i32(od * params.strideD + kd * params.dilationD) - params.padD;
       if (id < 0 || id >= i32(params.inD)) {
         continue;
       }
 {% endif %}
-{% if source.spatialRank >= 2 %}
     for (var kh = 0u; kh < params.kernelH; kh += 1u) {
       let ih = i32(oh * params.strideH + kh * params.dilationH) - params.padH;
       if (ih < 0 || ih >= i32(params.inH)) {
@@ -63,10 +85,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
       if (iw < 0 || iw >= i32(params.inW)) {
         continue;
       }
-{% if source.spatialRank == 1 %}
       let xIndex = (batch * params.inChannels + inputChannel) * params.inW + u32(iw);
       let wIndex = (oc * params.weightInChannels + ic) * params.kernelW + kw;
-{% elif source.spatialRank == 2 %}
       let xIndex = ((batch * params.inChannels + inputChannel) * params.inH
                     + u32(ih)) * params.inW + u32(iw);
       let wIndex = ((oc * params.weightInChannels + ic) * params.kernelH
@@ -79,12 +101,16 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>,
 {% endif %}
       acc += f32(x[xIndex]) * f32(w[wIndex]);
     }
-{% if source.spatialRank >= 2 %}
     }
 {% endif %}
-{% if source.spatialRank == 3 %}
     }
 {% endif %}
   }
   y[index] = {{ T }}({% if hasActivation %}fused_act(acc){% else %}acc{% endif %});
 }

+// Direct N-dimensional convolution for channels-first tensors. Inputs are
 // accumulated in f32 and narrowed once at the output.
 // A fused epilogue may apply a residual input and activation before the store.
+{% if usesF16Spec %}
 enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}
+{% set hasZ = hasZ is defined and hasZ %}
 const WG: u32 = {{ convWorkgroupSize }}u;
 @compute @workgroup_size(WG)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   if (index >= params.count) {
     return;
   }
   let ow = index % params.outW;
   var q = index / params.outW;
+{% if spatialRankSpec >= 2 %}
   let oh = q % params.outH;
   q /= params.outH;
 {% endif %}
+{% if spatialRankSpec == 3 %}
   let od = q % params.outD;
   q /= params.outD;
 {% endif %}
   let group = oc / params.outChannelsPerGroup;
   var acc = 0.0;
+{% if hasBias %}
   acc = f32(bias[oc]);
 {% endif %}
   for (var ic = 0u; ic < params.weightInChannels; ic += 1u) {
     let inputChannel = group * params.inChannelsPerGroup + ic;
+{% if spatialRankSpec == 3 %}
     for (var kd = 0u; kd < params.kernelD; kd += 1u) {
       let id = i32(od * params.strideD + kd * params.dilationD) - params.padD;
       if (id < 0 || id >= i32(params.inD)) {
         continue;
       }
 {% endif %}
+{% if spatialRankSpec >= 2 %}
     for (var kh = 0u; kh < params.kernelH; kh += 1u) {
       let ih = i32(oh * params.strideH + kh * params.dilationH) - params.padH;
       if (ih < 0 || ih >= i32(params.inH)) {
       if (iw < 0 || iw >= i32(params.inW)) {
         continue;
       }
+{% if spatialRankSpec == 1 %}
       let xIndex = (batch * params.inChannels + inputChannel) * params.inW + u32(iw);
       let wIndex = (oc * params.weightInChannels + ic) * params.kernelW + kw;
+{% elif spatialRankSpec == 2 %}
       let xIndex = ((batch * params.inChannels + inputChannel) * params.inH
                     + u32(ih)) * params.inW + u32(iw);
       let wIndex = ((oc * params.weightInChannels + ic) * params.kernelH
 {% endif %}
       acc += f32(x[xIndex]) * f32(w[wIndex]);
     }
+{% if spatialRankSpec >= 2 %}
     }
 {% endif %}
+{% if spatialRankSpec == 3 %}
     }
 {% endif %}
   }
+{% if hasZ %}
+  // Optional residual epilogue: Y = activation(conv + B + Z); Z shares Y's layout.
+  acc = acc + f32(zResidual[index]);
+{% endif %}
   y[index] = {{ T }}({% if hasActivation %}fused_act(acc){% else %}acc{% endif %});
 }

build/webgpu/conv-direct-unrolled.wgsl.jinja CHANGED Viewed

@@ -1,36 +1,55 @@
 // Direct NCHW 2D convolution with kernel extent, strides, dilations, and
 // top/left pads compiled into the fully unrolled window below.
-// One thread per output element, same flat index unpack and the same
-// accumulation order (ic outer, then kh, kw ascending) as the
-// scalar NCHW fallback, so f32 results are bit-identical to it.
 // The kh/kw window offsets, kh*KW+kw weight offsets, and the per-channel
 // strides fold to literals; the per-row bounds check is hoisted per kh.
-// Supports grouped convolution exactly like the scalar fallback. f16 inputs
-// are widened to an f32 accumulator and narrowed once at store.
-{% 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 %}
-const KERNEL_AREA: u32 = {{ source.kernelH * source.kernelW }}u;
-const STRIDE_H: u32 = {{ source.strideH }}u;
-const STRIDE_W: u32 = {{ source.strideW }}u;
-const PAD_TOP: i32 = {{ source.padTop }};
-const PAD_LEFT: i32 = {{ source.padLeft }};
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past the
-  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
-  let index = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (index >= params.count) {
     return;
   }
@@ -54,16 +73,16 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   acc = f32(bias[oc]);
 {% endif %}
   for (var ic = 0u; ic < params.weightInChannels; ic = ic + 1u) {
-{% for kh in range(source.kernelH) %}
     {
-      let ih = ihBase + {{ kh * source.dilationH }}i;
       if (ih >= 0 && ih < i32(params.inH)) {
         let xRow = xBase + u32(ih) * params.inW;
-{% for kw in range(source.kernelW) %}
         {
-          let iw = iwBase + {{ kw * source.dilationW }}i;
           if (iw >= 0 && iw < i32(params.inW)) {
-            acc = acc + f32(x[xRow + u32(iw)]) * f32(w[wBase + {{ kh * source.kernelW + kw }}u]);
           }
         }
 {% endfor %}
@@ -73,5 +92,9 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
     xBase = xBase + xChannelStride;
     wBase = wBase + KERNEL_AREA;
   }
   y[index] = {{ T }}({% if hasActivation %}fused_act(acc){% else %}acc{% endif %});
 }

 // Direct NCHW 2D convolution with kernel extent, strides, dilations, and
 // top/left pads compiled into the fully unrolled window below.
+// One thread owns each output element and accumulates with input channels outer,
+// then kernel rows and columns in ascending order.
 // The kh/kw window offsets, kh*KW+kw weight offsets, and the per-channel
 // strides fold to literals; the per-row bounds check is hoisted per kh.
+// Group offsets are applied to both input channels and weights. f16 inputs are
+// widened to an f32 accumulator and narrowed once at store.
 {{ env.wgsl.resourceDeclarations }}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}
+{% set hasZ = hasZ is defined and hasZ %}
+const KERNEL_AREA: u32 = {{ kernelHSpec * kernelWSpec }}u;
+const STRIDE_H: u32 = {{ strideHSpec }}u;
+const STRIDE_W: u32 = {{ strideWSpec }}u;
+const PAD_TOP: i32 = {{ padTopSpec }};
+const PAD_LEFT: i32 = {{ padLeftSpec }};
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past the
+  // per-axis dispatch fold width (outputs > 16.7M elements).
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (index >= params.count) {
     return;
   }
   acc = f32(bias[oc]);
 {% endif %}
   for (var ic = 0u; ic < params.weightInChannels; ic = ic + 1u) {
+{% for kh in range(kernelHSpec) %}
     {
+      let ih = ihBase + {{ kh * dilationHSpec }}i;
       if (ih >= 0 && ih < i32(params.inH)) {
         let xRow = xBase + u32(ih) * params.inW;
+{% for kw in range(kernelWSpec) %}
         {
+          let iw = iwBase + {{ kw * dilationWSpec }}i;
           if (iw >= 0 && iw < i32(params.inW)) {
+            acc = acc + f32(x[xRow + u32(iw)]) * f32(w[wBase + {{ kh * kernelWSpec + kw }}u]);
           }
         }
 {% endfor %}
     xBase = xBase + xChannelStride;
     wBase = wBase + KERNEL_AREA;
   }
+{% if hasZ %}
+  // Optional residual epilogue: Y = activation(conv + B + Z); Z shares Y's layout.
+  acc = acc + f32(zResidual[index]);
+{% endif %}
   y[index] = {{ T }}({% if hasActivation %}fused_act(acc){% else %}acc{% endif %});
 }

build/webgpu/conv-im2col-ncdhw.wgsl.jinja CHANGED Viewed

@@ -6,25 +6,22 @@
 // Geometry comes from tensor shapes and Conv attributes. Compiling it into
 // constants removes invariant division and address work from
 // the inner GEMM while retaining the complete rank-5 Conv contract.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-const KERNEL_D: u32 = {{ source.kernelD }}u;
-const KERNEL_H: u32 = {{ source.kernelH }}u;
-const KERNEL_W: u32 = {{ source.kernelW }}u;
 const KERNEL_PLANE: u32 = KERNEL_H * KERNEL_W;
 const KERNEL_VOLUME: u32 = KERNEL_D * KERNEL_PLANE;
-const STRIDE_D: u32 = {{ source.strideD }}u;
-const STRIDE_H: u32 = {{ source.strideH }}u;
-const STRIDE_W: u32 = {{ source.strideW }}u;
-const DILATION_D: u32 = {{ source.dilationD }}u;
-const DILATION_H: u32 = {{ source.dilationH }}u;
-const DILATION_W: u32 = {{ source.dilationW }}u;
-const PAD_FRONT: i32 = {{ source.padFront }};
-const PAD_TOP: i32 = {{ source.padTop }};
-const PAD_LEFT: i32 = {{ source.padLeft }};
 @compute @workgroup_size({{ convWorkgroupSize }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {

 // Geometry comes from tensor shapes and Conv attributes. Compiling it into
 // constants removes invariant division and address work from
 // the inner GEMM while retaining the complete rank-5 Conv contract.
 {{ env.wgsl.resourceDeclarations }}
+const KERNEL_D: u32 = {{ kernelDSpec }}u;
+const KERNEL_H: u32 = {{ kernelHSpec }}u;
+const KERNEL_W: u32 = {{ kernelWSpec }}u;
 const KERNEL_PLANE: u32 = KERNEL_H * KERNEL_W;
 const KERNEL_VOLUME: u32 = KERNEL_D * KERNEL_PLANE;
+const STRIDE_D: u32 = {{ strideDSpec }}u;
+const STRIDE_H: u32 = {{ strideHSpec }}u;
+const STRIDE_W: u32 = {{ strideWSpec }}u;
+const DILATION_D: u32 = {{ dilationDSpec }}u;
+const DILATION_H: u32 = {{ dilationHSpec }}u;
+const DILATION_W: u32 = {{ dilationWSpec }}u;
+const PAD_FRONT: i32 = {{ padFrontSpec }};
+const PAD_TOP: i32 = {{ padTopSpec }};
+const PAD_LEFT: i32 = {{ padLeftSpec }};
 @compute @workgroup_size({{ convWorkgroupSize }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {

build/webgpu/conv-im2col-nchw.wgsl.jinja CHANGED Viewed

@@ -6,20 +6,19 @@
 // gid.x grid-strides over positions for coalesced reads and writes; gid.y fixes
 // k, allowing its ic/kh/kw decomposition to remain loop-invariant. The f32
 // scratch round-trips f16 inputs exactly when the GEMM narrows its tile loads.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-const KERNEL_H: u32 = {{ source.kernelH }}u;
-const KERNEL_W: u32 = {{ source.kernelW }}u;
 const KSIZE: u32 = KERNEL_H * KERNEL_W;
-const STRIDE_H: u32 = {{ source.strideH }}u;
-const STRIDE_W: u32 = {{ source.strideW }}u;
-const DILATION_H: u32 = {{ source.dilationH }}u;
-const DILATION_W: u32 = {{ source.dilationW }}u;
-const PAD_TOP: i32 = {{ source.padTop }};
-const PAD_LEFT: i32 = {{ source.padLeft }};
 {% if padded %}
 // Padded materialization for the subgroup-matrix GEMM: the cols buffer is
 // K -> ceil(K/32)*32 rows by N -> ceil(N/64)*64 columns, with rows >= kRows and
@@ -64,10 +63,10 @@ fn main(
     }
 {% if padded %}
     }
-    cols[(batch * K_PADDED + k) * COLS_STRIDE + pos] = value;
-{% else %}
-    cols[(batch * params.kRows + k) * params.outCount + pos] = value;
 {% endif %}
     pos = pos + step;
   }
 }

 // gid.x grid-strides over positions for coalesced reads and writes; gid.y fixes
 // k, allowing its ic/kh/kw decomposition to remain loop-invariant. The f32
 // scratch round-trips f16 inputs exactly when the GEMM narrows its tile loads.
+// A plan for f16 inputs may retain f16 columns to avoid that widening.
+{% set colsScalar = colsScalar | default("f32") %}
 {{ env.wgsl.resourceDeclarations }}
+const KERNEL_H: u32 = {{ kernelHSpec }}u;
+const KERNEL_W: u32 = {{ kernelWSpec }}u;
 const KSIZE: u32 = KERNEL_H * KERNEL_W;
+const STRIDE_H: u32 = {{ strideHSpec }}u;
+const STRIDE_W: u32 = {{ strideWSpec }}u;
+const DILATION_H: u32 = {{ dilationHSpec }}u;
+const DILATION_W: u32 = {{ dilationWSpec }}u;
+const PAD_TOP: i32 = {{ padTopSpec }};
+const PAD_LEFT: i32 = {{ padLeftSpec }};
 {% if padded %}
 // Padded materialization for the subgroup-matrix GEMM: the cols buffer is
 // K -> ceil(K/32)*32 rows by N -> ceil(N/64)*64 columns, with rows >= kRows and
     }
 {% if padded %}
     }
 {% endif %}
+{% set columnRows = "K_PADDED" if padded else "params.kRows" %}
+{% set columnStride = "COLS_STRIDE" if padded else "params.outCount" %}
+    cols[(batch * {{ columnRows }} + k) * {{ columnStride }} + pos] = {% if colsScalar == "f16" %}f16(value){% else %}value{% endif %};
     pos = pos + step;
   }
 }

build/webgpu/conv-splitk-reduce.wgsl.jinja CHANGED Viewed

@@ -8,15 +8,43 @@
 // contents never matter. One output element per thread, walking n fastest so
 // consecutive threads touch consecutive addresses in every slice.
 //
-// Summing the slices in index order gives a DIFFERENT rounding than the
-// single-pass kernel's one long K loop — the same reassociation any tiled GEMM
-// already makes, not a precision change, but it is why split-K variants are
-// compared against the op's tolerance and not pinned bit-for-bit to the
-// un-split path.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const M: u32 = {{ M }}u;
 const N: u32 = {{ N }}u;
@@ -28,10 +56,10 @@ const COUNT: u32 = {{ batchCount }}u * M * N;
 const WORKGROUP_SIZE: u32 = {{ reduceWorkgroupSize }}u;
 @compute @workgroup_size({{ reduceWorkgroupSize }}, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past
-  // maxComputeWorkgroupsPerDimension and reduces to the 1D form at y=0.
-  let idx = gid.x + gid.y * nwg.x * WORKGROUP_SIZE;
   if (idx >= COUNT) {
     return;
   }
@@ -47,6 +75,9 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   }
 {% if hasBias %}
   acc = acc + f32(bias[row]);
 {% endif %}
   y[idx] = {{ T }}(acc);
 }

 // contents never matter. One output element per thread, walking n fastest so
 // consecutive threads touch consecutive addresses in every slice.
 //
+// Summing slices in index order changes the f32 association relative to one
+// uninterrupted K loop, so the two orders need not be bit-identical.
 {{ env.wgsl.resourceDeclarations }}
+{% set applyActivation = hasActivation is defined and hasActivation %}
+{% if applyActivation %}
+{% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
+{% if hasActivation %}
+// Apply the fused activation in the f32 accumulator before the single output
+// cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
+}
+{% endif %}
+{% endif %}
 const M: u32 = {{ M }}u;
 const N: u32 = {{ N }}u;
 const WORKGROUP_SIZE: u32 = {{ reduceWorkgroupSize }}u;
 @compute @workgroup_size({{ reduceWorkgroupSize }}, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past
+  // the per-axis dispatch fold width and reduces to the 1D form at y=0.
+  let idx = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WORKGROUP_SIZE;
   if (idx >= COUNT) {
     return;
   }
   }
 {% if hasBias %}
   acc = acc + f32(bias[row]);
+{% endif %}
+{% if applyActivation %}
+  acc = fused_act(acc);
 {% endif %}
   y[idx] = {{ T }}(acc);
 }

build/webgpu/conv1d-tiled-reg.wgsl.jinja CHANGED Viewed

@@ -1,31 +1,43 @@
-// Implicit-GEMM register-tiled 1-D convolution over NCW: Y[b] (M x N) =
-// W (M x K) @ im2col(X[b]) (K x N) with M = outChannels, K = inChannels *
-// kernelW, N = outW. The im2col operand is never materialized: the X staging
-// loop decodes each GEMM row into (ic, kw) and samples x with stride /
-// dilation / left-pad applied, so one program covers general stride, padding,
-// and dilation (group == 1 only). Workgroup geometry is tunable: a WG_X x WG_Y
-// workgroup computes a (WG_Y*TM) x (WG_X*TN) output tile from BK-deep staged
-// slices, each thread owning a TM x TN register micro-tile. Cooperative loads
-// and stores are fully bounds-checked, so partial tiles on any edge stay
-// correct for every tunable setting. f32 accumulation; bias and the optional
-// activation epilogue are applied in the accumulator domain at
-// store. f16 operands stay packed in workgroup memory and widen at the FMA
-// site. Both tiles are indexed by their own output axis and group four K values
-// per vector word, so the micro-tile accumulates through dot() and the implicit
-// gather decodes one (ic, kw) tap per word, walking the remaining three. The
-// micro-tile is unrolled into named registers: a dynamically indexed private
-// array of vector words is large enough at the wide-M geometry that it may fail
-// to scalarize and spill the accumulators.
-{% 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 %}

+// Register-tiled implicit GEMM for group-1 NCW convolution:
+// Y[b] (M x N) = W (M x K) @ im2col(X[b]) (K x N), where M is outChannels, K
+// is inChannels * kernelW, and N is outW. The X tile is gathered directly from
+// NCW with the configured stride, dilation, and left padding; no column matrix
+// is materialized.
+//
+// A WG_X x WG_Y workgroup computes a (WG_Y * TM) x (WG_X * TN) output tile
+// from BK-deep slices, with one TM x TN register micro-tile per thread. Weight
+// and input tiles group four K values per workgroup-memory word and accumulate
+// through dot(). Bounds checks cover every partial tile. Accumulation, bias,
+// and the optional activation epilogue use f32; f16 operands widen at the FMA.
 {{ env.wgsl.resourceDeclarations }}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}

build/webgpu/conv2d-grouped-large-w4.wgsl.jinja CHANGED Viewed

@@ -1,13 +1,32 @@
-{% 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 %}
@@ -15,86 +34,130 @@ fn fused_act(v: f32) -> f32 {
 // value for neighboring output columns and input window for neighboring output
 // channels. One invocation computes four adjacent columns for OC_TILE channels,
 // sharing each input load across those channels. Shape and window geometry are
-// static, so all indexing divisors and kernel offsets are strength-reduced and
-// unrolled by the shader compiler.
-const COUNT_TILES: u32 = {{ source.countTiles }}u;
-{% if source.tailOutput %}const OUT_W: u32 = max(1u, {{ source.outW }}u);
 {% endif %}
-const OUT_W4: u32 = max(1u, {{ source.outW4 }}u);
-const OUT_H: u32 = max(1u, {{ source.outH }}u);
-const OUT_C: u32 = max(1u, {{ source.outC }}u);
-const OC_TILE: u32 = {{ source.ocTile }}u;
 const OUT_C_TILES: u32 = OUT_C / OC_TILE;
-const OUT_CPG: u32 = {{ source.outCPerGroup }}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 IN_PLANE: u32 = IN_H_U * IN_W_U;
-{% if not source.tailOutput %}const OUT_PLANE4: u32 = OUT_H * OUT_W4;
 {% endif %}
-const KAREA: u32 = {{ source.kernelH * source.kernelW }}u;
-const STRIDE_H: u32 = {{ source.strideH }}u;
-const STRIDE_W: i32 = {{ source.strideW }};
-const PAD_TOP: i32 = {{ source.padTop }};
-const PAD_LEFT: i32 = {{ source.padLeft }};
-const WG: u32 = {{ source.workgroupSize }}u;
-@compute @workgroup_size({{ source.workgroupSize }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
-  let q = gid.x + gid.y * nwg.x * WG;
   if (q >= COUNT_TILES) { return; }
   let ow4 = q % OUT_W4;
   var t = q / OUT_W4;
   let oh = t % OUT_H;
   t /= OUT_H;
   let ocBase = (t % OUT_C_TILES) * OC_TILE;
   let batch = t / OUT_C_TILES;
   let group = ocBase / OUT_CPG;
   let ihBase = i32(oh * STRIDE_H) - PAD_TOP;
   let iwBase = i32(ow4 * 4u) * STRIDE_W - PAD_LEFT;
   var xChannelBase = (batch * IN_C + group * IN_CPG) * IN_PLANE;
-{% for oct in range(source.ocTile) %}
   var wBase{{ oct }} = (ocBase + {{ oct }}u) * IN_CPG * KAREA;
-  var acc{{ oct }} = vec4<f32>(0.0);
 {% endfor %}
   for (var ic = 0u; ic < IN_CPG; ic++) {
-{% for kh in range(source.kernelH) %}
     {
-      let ih = ihBase + {{ kh * source.dilationH }};
       if (ih >= 0 && ih < IN_H) {
         let xRow = xChannelBase + u32(ih) * IN_W_U;
-{% if source.span <= source.spanCap %}
         // The four output columns and the kernel taps overlap heavily: tap kw of
         // lane l reads input column kw*dilationW + l*strideW, so kernelW*4 reads
         // land on only SPAN distinct columns. Load the row window once and let
-        // every tap take its value from a register. Out-of-range columns hold
-        // zero, which contributes nothing to the accumulation exactly as the
-        // per-tap bounds test did.
         //
-        // SPAN grows as (kernelW-1)*dilationW, so a widely dilated kernel would
-        // want more registers than a thread has and spill. Past spanCap the
-        // per-tap form below stays, where the register cost is fixed.
-{% for s in range(source.span) %}
 {% set used = namespace(value=false) %}
-{% for kw in range(source.kernelW) %}
-{% for lane in range(4) %}
-{% if s == kw * source.dilationW + lane * source.strideW %}{% set used.value = true %}{% endif %}
 {% endfor %}
 {% endfor %}
 {% if used.value %}
         var xw{{ s }} = 0.0;
 {% endif %}
 {% endfor %}
-{% for s in range(source.span) %}
 {% set used = namespace(value=false) %}
-{% for kw in range(source.kernelW) %}
-{% for lane in range(4) %}
-{% if s == kw * source.dilationW + lane * source.strideW %}{% set used.value = true %}{% endif %}
 {% endfor %}
 {% endfor %}
 {% if used.value %}
@@ -104,31 +167,31 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
         }
 {% endif %}
 {% endfor %}
-{% for kw in range(source.kernelW) %}
         {
-{% for oct in range(source.ocTile) %}
-          let weight{{ oct }} = f32(w[wBase{{ oct }} + {{ kh * source.kernelW + kw }}u]);
 {% endfor %}
-{% for lane in range(4) %}
-{% for oct in range(source.ocTile) %}
-          acc{{ oct }}.{{ ["x", "y", "z", "w"][lane] }} += xw{{ kw * source.dilationW + lane * source.strideW }} * weight{{ oct }};
 {% endfor %}
 {% endfor %}
         }
 {% endfor %}
 {% else %}
-{% for kw in range(source.kernelW) %}
         {
-{% for oct in range(source.ocTile) %}
-          let weight{{ oct }} = f32(w[wBase{{ oct }} + {{ kh * source.kernelW + kw }}u]);
 {% endfor %}
-          let iwK = iwBase + {{ kw * source.dilationW }};
-{% for lane in range(4) %}
-          let iwLane{{ lane }} = iwK + {{ lane * source.strideW }};
           if (iwLane{{ lane }} >= 0 && iwLane{{ lane }} < IN_W) {
             let xValue = f32(x[xRow + u32(iwLane{{ lane }})]);
-{% for oct in range(source.ocTile) %}
-            acc{{ oct }}.{{ ["x", "y", "z", "w"][lane] }} += xValue * weight{{ oct }};
 {% endfor %}
           }
 {% endfor %}
@@ -139,13 +202,23 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
     }
 {% endfor %}
     xChannelBase += IN_PLANE;
-{% for oct in range(source.ocTile) %}
     wBase{{ oct }} += KAREA;
 {% endfor %}
   }
-{% if source.tailOutput %}
   let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W + ow4 * 4u;
-{% for oct in range(source.ocTile) %}
   let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
 {% for lane in range(4) %}
   if (ow4 * 4u + {{ lane }}u < OUT_W) {
@@ -155,7 +228,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 {% endfor %}
 {% else %}
   let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W4 + ow4;
-{% for oct in range(source.ocTile) %}
   y[outBase + {{ oct }}u * OUT_PLANE4] = {{ vectorScalar }}({% if hasActivation %}vec4<f32>(fused_act(acc{{ oct }}.x), fused_act(acc{{ oct }}.y), fused_act(acc{{ oct }}.z), fused_act(acc{{ oct }}.w)){% else %}acc{{ oct }}{% endif %});
 {% endfor %}
 {% endif %}

 {{ env.wgsl.resourceDeclarations }}
 {% set hasActivation = hasActivation is defined and hasActivation %}
+{% set activation = activation | default("") %}
+{% set actAlpha = actAlpha | default(0.0) %}
+{% set actBeta = actBeta | default(0.0) %}
 {% if hasActivation %}
 // Apply the fused activation in the f32 accumulator before the single output
 // cast, avoiding an intermediate convolution tensor.
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
   return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(activation, actAlpha, actBeta) -}}
 }
 {% endif %}
 // value for neighboring output columns and input window for neighboring output
 // channels. One invocation computes four adjacent columns for OC_TILE channels,
 // sharing each input load across those channels. Shape and window geometry are
+// static, so indexing divisors and kernel offsets are shader constants.
+{% set dilatedLanes = dilatedLanes | default(false) %}
+{% set lanes = lanes | default(4) %}
+{% set quads = quads | default(1) %}
+{% set laneStep = dilationWSpec if dilatedLanes else strideWSpec %}
+{% set comps = ["x", "y", "z", "w"] %}
+{% macro acc_ref(oct, lane) %}{% if dilatedLanes %}acc{{ oct }}_{{ lane }}{% else %}acc{{ oct }}.{{ comps[lane] }}{% endif %}{% endmacro %}
+{% set usedCount = namespace(value=0) %}
+{% for s in range(span) %}
+{% set used = namespace(value=false) %}
+{% for kw in range(kernelWSpec) %}
+{% for lane in range(lanes) %}
+{% if s == kw * dilationWSpec + lane * laneStep %}{% set used.value = true %}{% endif %}
+{% endfor %}
+{% endfor %}
+{% if used.value %}{% set usedCount.value = usedCount.value + 1 %}{% endif %}
+{% endfor %}
+{% set registerForm = (usedCount.value <= spanCap) if dilatedLanes else (span <= spanCap) %}
+const COUNT_TILES: u32 = {{ countTiles }}u;
+{% if tailOutput %}const OUT_W: u32 = max(1u, {{ outW }}u);
+{% endif %}
+{% if dilatedLanes %}
+const OUT_W: u32 = max(1u, {{ outW }}u);
+const LANES: u32 = {{ lanes }}u;
+const RESIDUES: u32 = {{ dilationWSpec }}u;
+const TILES_PER_ROW: u32 = RESIDUES * {{ quads }}u;
+{% else %}
+const OUT_W4: u32 = max(1u, {{ outW4 }}u);
 {% endif %}
+const OUT_H: u32 = max(1u, {{ outH }}u);
+const OUT_C: u32 = max(1u, {{ outC }}u);
+const OC_TILE: u32 = {{ ocTile }}u;
 const OUT_C_TILES: u32 = OUT_C / OC_TILE;
+const OUT_CPG: u32 = {{ outCPerGroup }}u;
+const IN_C: u32 = {{ inC }}u;
+const IN_CPG: u32 = {{ inCPerGroup }}u;
+const IN_H: i32 = {{ inH }};
+const IN_W: i32 = {{ inW }};
+const IN_H_U: u32 = {{ inH }}u;
+const IN_W_U: u32 = {{ inW }}u;
 const IN_PLANE: u32 = IN_H_U * IN_W_U;
+{% if not tailOutput and not dilatedLanes %}const OUT_PLANE4: u32 = OUT_H * OUT_W4;
 {% endif %}
+const KAREA: u32 = {{ kernelHSpec * kernelWSpec }}u;
+const STRIDE_H: u32 = {{ strideHSpec }}u;
+const STRIDE_W: i32 = {{ strideWSpec }};
+const PAD_TOP: i32 = {{ padTopSpec }};
+const PAD_LEFT: i32 = {{ padLeftSpec }};
+const WG: u32 = {{ workgroupSizeSpec }}u;
+@compute @workgroup_size({{ workgroupSizeSpec }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let q = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   if (q >= COUNT_TILES) { return; }
+{% if dilatedLanes %}
+  let tile = q % TILES_PER_ROW;
+  var t = q / TILES_PER_ROW;
+{% else %}
   let ow4 = q % OUT_W4;
   var t = q / OUT_W4;
+{% endif %}
   let oh = t % OUT_H;
   t /= OUT_H;
   let ocBase = (t % OUT_C_TILES) * OC_TILE;
   let batch = t / OUT_C_TILES;
   let group = ocBase / OUT_CPG;
   let ihBase = i32(oh * STRIDE_H) - PAD_TOP;
+{% if dilatedLanes %}
+  // Lane l owns output column ow0 + l * dilationW: the residue class first,
+  // then the LANES-wide run within it. Runs past the row store nothing.
+  let ow0 = tile % RESIDUES + (tile / RESIDUES) * (LANES * RESIDUES);
+  let iwBase = i32(ow0) * STRIDE_W - PAD_LEFT;
+{% else %}
   let iwBase = i32(ow4 * 4u) * STRIDE_W - PAD_LEFT;
+{% endif %}
   var xChannelBase = (batch * IN_C + group * IN_CPG) * IN_PLANE;
+{% for oct in range(ocTile) %}
   var wBase{{ oct }} = (ocBase + {{ oct }}u) * IN_CPG * KAREA;
+{% if hasBias %}
+  // Seeding the accumulator with the channel bias replaces the zero seed and
+  // incorporates the bias before the reduction.
+  let bias{{ oct }} = f32(bias[ocBase + {{ oct }}u]);
+{% endif %}
+{% if dilatedLanes %}
+{% for lane in range(lanes) %}
+  var acc{{ oct }}_{{ lane }} = {% if hasBias %}bias{{ oct }}{% else %}0.0{% endif %};
+{% endfor %}
+{% else %}
+  var acc{{ oct }} = vec4<f32>({% if hasBias %}bias{{ oct }}{% else %}0.0{% endif %});
+{% endif %}
 {% endfor %}
   for (var ic = 0u; ic < IN_CPG; ic++) {
+{% for kh in range(kernelHSpec) %}
     {
+      let ih = ihBase + {{ kh * dilationHSpec }};
       if (ih >= 0 && ih < IN_H) {
         let xRow = xChannelBase + u32(ih) * IN_W_U;
+{% if registerForm %}
         // The four output columns and the kernel taps overlap heavily: tap kw of
         // lane l reads input column kw*dilationW + l*strideW, so kernelW*4 reads
         // land on only SPAN distinct columns. Load the row window once and let
+        // every tap take its value from a statically named local. Out-of-range
+        // columns contain zero and contribute nothing to the accumulation.
         //
+        // SPAN grows as (kernelW-1)*dilationW. `spanCap` bounds the number of
+        // staged row-window values; wider spans use the fixed-size per-tap form below.
+{% for s in range(span) %}
 {% set used = namespace(value=false) %}
+{% for kw in range(kernelWSpec) %}
+{% for lane in range(lanes) %}
+{% if s == kw * dilationWSpec + lane * laneStep %}{% set used.value = true %}{% endif %}
 {% endfor %}
 {% endfor %}
 {% if used.value %}
         var xw{{ s }} = 0.0;
 {% endif %}
 {% endfor %}
+{% for s in range(span) %}
 {% set used = namespace(value=false) %}
+{% for kw in range(kernelWSpec) %}
+{% for lane in range(lanes) %}
+{% if s == kw * dilationWSpec + lane * laneStep %}{% set used.value = true %}{% endif %}
 {% endfor %}
 {% endfor %}
 {% if used.value %}
         }
 {% endif %}
 {% endfor %}
+{% for kw in range(kernelWSpec) %}
         {
+{% for oct in range(ocTile) %}
+          let weight{{ oct }} = f32(w[wBase{{ oct }} + {{ kh * kernelWSpec + kw }}u]);
 {% endfor %}
+{% for lane in range(lanes) %}
+{% for oct in range(ocTile) %}
+          {{ acc_ref(oct, lane) }} += xw{{ kw * dilationWSpec + lane * laneStep }} * weight{{ oct }};
 {% endfor %}
 {% endfor %}
         }
 {% endfor %}
 {% else %}
+{% for kw in range(kernelWSpec) %}
         {
+{% for oct in range(ocTile) %}
+          let weight{{ oct }} = f32(w[wBase{{ oct }} + {{ kh * kernelWSpec + kw }}u]);
 {% endfor %}
+          let iwK = iwBase + {{ kw * dilationWSpec }};
+{% for lane in range(lanes) %}
+          let iwLane{{ lane }} = iwK + {{ lane * laneStep }};
           if (iwLane{{ lane }} >= 0 && iwLane{{ lane }} < IN_W) {
             let xValue = f32(x[xRow + u32(iwLane{{ lane }})]);
+{% for oct in range(ocTile) %}
+            {{ acc_ref(oct, lane) }} += xValue * weight{{ oct }};
 {% endfor %}
           }
 {% endfor %}
     }
 {% endfor %}
     xChannelBase += IN_PLANE;
+{% for oct in range(ocTile) %}
     wBase{{ oct }} += KAREA;
 {% endfor %}
   }
+{% if dilatedLanes %}
+  let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W;
+{% for oct in range(ocTile) %}
+  let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
+{% for lane in range(lanes) %}
+  if (ow0 + {{ lane * dilationWSpec }}u < OUT_W) {
+    y[outChannelBase{{ oct }} + ow0 + {{ lane * dilationWSpec }}u] = {{ scalar }}({% if hasActivation %}fused_act(acc{{ oct }}_{{ lane }}){% else %}acc{{ oct }}_{{ lane }}{% endif %});
+  }
+{% endfor %}
+{% endfor %}
+{% elif tailOutput %}
   let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W + ow4 * 4u;
+{% for oct in range(ocTile) %}
   let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
 {% for lane in range(4) %}
   if (ow4 * 4u + {{ lane }}u < OUT_W) {
 {% endfor %}
 {% else %}
   let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W4 + ow4;
+{% for oct in range(ocTile) %}
   y[outBase + {{ oct }}u * OUT_PLANE4] = {{ vectorScalar }}({% if hasActivation %}vec4<f32>(fused_act(acc{{ oct }}.x), fused_act(acc{{ oct }}.y), fused_act(acc{{ oct }}.z), fused_act(acc{{ oct }}.w)){% else %}acc{{ oct }}{% endif %});
 {% endfor %}
 {% endif %}

build/webgpu/manifest.json CHANGED Viewed

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

build/webgpu/metadata.json CHANGED Viewed

@@ -1,28 +1,89 @@
 {
   "name": "ai.onnx.Conv",
-  "id": "_ai_onnx_conv_webgpu_4822108",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "LNwpS8vDbGYX6SF0eYrHwA2PL9KohpIj3TEv6TF5iOQ=",
-      "conv-1x1-channel-reduce.wgsl.jinja": "N0hmEYctrx1dFezCieFBMSZw1AJkB9k1jqWbqwY73yw=",
-      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "uhup4og6kzn9fGBicWDhp8u7tRMWOxV1FP4h4B+1QC4=",
-      "conv-1x1-gemm-tiled.wgsl.jinja": "VeRKrzHV+9CvWm0Lrsgpt5C98Gkz5MubGVEnYApQhdk=",
-      "conv-1x1-subgroup-matrix.wgsl.jinja": "CTPjuc+j4mBdgIKG/LMdUSNTO/qeyS0810Gvzyl9waU=",
-      "conv-direct-nd.wgsl.jinja": "ANLdecurm5wdc72OIpjU8oiYtrMhL5a5z8XhDjYJXGg=",
-      "conv-direct-unrolled.wgsl.jinja": "ILQ1vB/EtaZvB65S5LwNliPq91fcT7LCzVoPDRoDolY=",
-      "conv-im2col-ncdhw.wgsl.jinja": "AsmaNPOjoGmh+YHO91mT5EIc7RSlGqgzm7uKASMS3vw=",
-      "conv-im2col-nchw.wgsl.jinja": "7IUuXo33elZx0pnR9M2vUXzXDjFuw2bJODkZViV/WQg=",
-      "conv-splitk-reduce.wgsl.jinja": "qAkIOXGzAj6HWYSscyL1cfNgM6A1Ym/1L/CjQ8ogCEA=",
-      "conv1d-tiled-reg.wgsl.jinja": "8xjz6LBfrQDZDFGDSYo8aV1m+yeVDvHMngjSAbOAmk8=",
-      "conv2d-grouped-large-w4.wgsl.jinja": "jt2Z54ENODhKgmScaR19Mq3Y3gAMybt2fP49zoZfwEA=",
-      "manifest.json": "hVZagZKncZQkjXOpaqdBUrFx0HXtV5u0K7rEqSuus1Q=",
-      "test.json": "ron7OZMUp/fWx5XDPPNom2Zg0q0lYcWgRzLyfgE4N5I="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Conv" }
 }

 {
   "name": "ai.onnx.Conv",
+  "id": "_ai_onnx_conv_webgpu_010d4b0",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "H1K2v5/L5nbBfjqTwogslpt8SrxTH6se0gj31ONSJeU=",
+      "conv-1x1-channel-reduce.wgsl.jinja": "Zjbu1oi+w5guLw54Y9C6TVJo0SntZDjgtsiie3o5md0=",
+      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "9yFpwXS1HzmWmm4W72n+aOIryLqc9+e88awZ1ogTe+Q=",
+      "conv-1x1-gemm-tiled.wgsl.jinja": "uqG5mRVvcHiUo0LDCUcVgqElVTvRNknx19Gm4oIALiU=",
+      "conv-1x1-subgroup-matrix.wgsl.jinja": "iDJs4RJk4iJarOZCjY7/5MFrHgR7CcqI6x5uwv2dYaA=",
+      "conv-direct-nd.wgsl.jinja": "FqwfAKVCF1BYBAtU7Y3rDcGrmZ5AYkH4FuTzdRlT2Rg=",
+      "conv-direct-unrolled.wgsl.jinja": "vrGUPwYkbYE1slII7nuK5wbERRcG5u0X1cW1kMc2sQs=",
+      "conv-im2col-ncdhw.wgsl.jinja": "vQtnzeEKC2wm23+kwcaWoX3VSaJSqno18ACe3w3trhI=",
+      "conv-im2col-nchw.wgsl.jinja": "IRpmtpeVQsLE6ud+Sf1230XaZVS1dxDtTZuSDwU3P8s=",
+      "conv-splitk-reduce.wgsl.jinja": "LbVU5DRL0Q51fvLjThmxJQZH8aPbZkvNzij4nY/cHmo=",
+      "conv1d-tiled-reg.wgsl.jinja": "ERqOPl+1RbnJ+CFqZvLSYw45qYMKgB7Je0teew5gyOM=",
+      "conv2d-grouped-large-w4.wgsl.jinja": "v8aAMQ2saiOgbQwLZAiVr3IeDT3h5YLO/v5tGL1x8pk=",
+      "manifest.json": "3/3/k9ZVj/vLw3xZ66XDMGWb86wzdZCoqDS0tKg9ziE=",
+      "test.json": "1OUEVPx0U/BTLtlxoSCFOWVgMfrlOfQb/TUY3VivfGw="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "ncw1d": ["conv-direct-nd.wgsl.jinja"],
+      "ncw1d_bias": ["conv-direct-nd.wgsl.jinja"],
+      "ncdhw3d": ["conv-direct-nd.wgsl.jinja"],
+      "ncdhw3d_bias": ["conv-direct-nd.wgsl.jinja"],
+      "nchw2d": ["conv-direct-nd.wgsl.jinja"],
+      "nchw2d_bias": ["conv-direct-nd.wgsl.jinja"],
+      "grouped_large_kernel_w4": ["conv2d-grouped-large-w4.wgsl.jinja"],
+      "grouped_large_kernel_w4_bias": ["conv2d-grouped-large-w4.wgsl.jinja"],
+      "grouped_large_kernel_w4_tail": ["conv2d-grouped-large-w4.wgsl.jinja"],
+      "grouped_large_kernel_w4_tail_bias": ["conv2d-grouped-large-w4.wgsl.jinja"],
+      "grouped_large_kernel_w4_dilated_lanes": ["conv2d-grouped-large-w4.wgsl.jinja"],
+      "grouped_large_kernel_w4_dilated_lanes_bias": ["conv2d-grouped-large-w4.wgsl.jinja"],
+      "pointwise_channel_reduce": ["conv-1x1-channel-reduce.wgsl.jinja"],
+      "pointwise_channel_reduce_bias": ["conv-1x1-channel-reduce.wgsl.jinja"],
+      "gemm_1x1_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
+      "gemm_1x1_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
+      "im2col_gemm_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_direct_f32_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_direct_f16_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_direct_f32_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_direct_f16_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "implicit_im2col_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
+      "implicit_im2col_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
+      "implicit_im2col3d_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
+      "implicit_im2col3d_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
+      "implicit_im2col_subgroup_matrix_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "implicit_im2col_subgroup_matrix_bias_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "im2col3d_gemm_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
+      "im2col3d_gemm_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
+      "im2col3d_gemm_tiled_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
+      "im2col3d_gemm_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
+      "im2col3d_gemm_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
+      "im2col3d_gemm_tiled_bias": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
+      "im2col_gemm_subgroup_matrix_padded": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_subgroup_matrix_padded_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_subgroup_matrix_padded_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "im2col_gemm_subgroup_matrix_padded_bias_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "implicit_im2col_tiled_bias_reg_m32": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
+      "implicit_im2col_tiled_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
+      "implicit_im2col_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
+      "implicit_im2col_tiled_reg_splitk": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "implicit_im2col_tiled_bias_reg_splitk": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "im2col_gemm_tiled_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled_bias": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "gemm_1x1_tiled_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
+      "gemm_1x1_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
+      "gemm_1x1_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja"],
+      "gemm_1x1_tiled_bias": ["conv-1x1-gemm-tiled.wgsl.jinja"],
+      "gemm_1x1_tiled_reg_splitk": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "gemm_1x1_tiled_bias_reg_splitk": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
+      "direct_unrolled": ["conv-direct-unrolled.wgsl.jinja"],
+      "direct_unrolled_bias": ["conv-direct-unrolled.wgsl.jinja"],
+      "conv1d_tiled_reg": ["conv1d-tiled-reg.wgsl.jinja"],
+      "conv1d_tiled_bias_reg": ["conv1d-tiled-reg.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Conv",
   "fixtureArrays": {
     "nchw_dilated_kernel_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24],
     "backend_stride2_with_padding_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34],
@@ -81,7 +80,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_op_test.cc",
         "test": "ConvTest",
-        "notes": "Zero-bias companion for the 1x1 Conv-as-GEMM subgroup-matrix path; the bias epilogue must preserve finite subnormal outputs."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
@@ -118,7 +117,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_op_test.cc",
         "test": "ConvTest",
-        "notes": "Zero-bias companion for the aligned im2col+subgroup-matrix path; the bias epilogue must preserve finite subnormal convolution outputs."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
@@ -155,7 +154,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_op_test.cc",
         "test": "ConvTest",
-        "notes": "Zero-bias companion for the padded im2col+subgroup-matrix path; padded K/N lanes and the bias epilogue must preserve finite subnormal real outputs."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
@@ -668,6 +667,51 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 16, 1], "tolerance": 0.0001 } },
       "attrs": {}
     },
     {
       "name": "ort_conv2d_1x1_splitk_with_bias",
       "provenance": {
@@ -879,8 +923,8 @@
     {
       "name": "conv1d_f16_audio_feature",
       "provenance": {
-        "source": "Manifest/schema gap regression",
-        "notes": "Compact half-precision audio-feature convolution covering the unified, dtype-generic NCW fallback."
       },
       "attrs": { "strides": [1], "dilations": [1], "pads": [1, 1] },
       "inputs": {
@@ -1210,19 +1254,19 @@
     {
       "name": "conv_3x3_im2col_gemm_dense_c256_8x8",
       "provenance": {
-        "notes": "Compact companion for wide 3x3 Conv bench cases such as flux2vae-dec-mid-512c-64x64-k3 and flux2vae-dec-up1-512c-128x128-k3, preserving the large K/M im2col-subgroup regime."
       },
       "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
         "x": {
           "dtype": "float32",
           "shape": [1, 256, 8, 8],
-          "data": { "kind": "fillFloat32", "scale": 0.16, "sinStep": 0.013, "cosStep": 0.031 }
         },
         "w": {
           "dtype": "float32",
           "shape": [256, 256, 3, 3],
-          "data": { "kind": "fillFloat32", "scale": 0.015, "sinStep": 0.007, "cosStep": 0.019 }
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 256, 8, 8], "tolerance": 0.006, "relTolerance": 0.006 } }
@@ -1534,7 +1578,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_fp16_test.cc",
         "test": "ConvFp16Test.Conv2D_Bias_2",
-        "notes": "Covers f16 bias with asymmetric explicit pads and non-unit strides from ORT's handwritten Conv fp16 suite."
       },
       "attrs": { "strides": [2, 3], "pads": [1, 2, 3, 1] },
       "inputs": {
@@ -1828,7 +1872,7 @@
       "name": "grouped_dilated_large_kernel_k11_oc4_nchw_fallback",
       "provenance": {
         "source": "Realistic grouped/dilated Conv stress geometry",
-        "notes": "Compact companion for conv-f32-group4-k11-dilated-128c-48x48; covers the K11 selection boundary with two OC4 tiles per group."
       },
       "attrs": { "group": 2, "strides": [1, 1], "dilations": [2, 2], "pads": [10, 10, 10, 10] },
       "inputs": {
@@ -1867,7 +1911,7 @@
       "name": "f16_grouped_large_kernel_k9_oc4_nchw_f16_fallback",
       "provenance": {
         "source": "Realistic grouped Conv stress geometry",
-        "notes": "Compact f16 OC4 companion for conv-f16-group8-k9-256c-64x64; two OC4 tiles per group exercise shared-input output-channel tiling."
       },
       "requires": { "features": ["shader-f16"] },
       "attrs": { "group": 2, "strides": [1, 1], "pads": [4, 4, 4, 4] },
@@ -1889,7 +1933,7 @@
       "name": "conv3d_depth_dilation2_compact",
       "provenance": {
         "source": "ONNX Conv-22 volumetric dilation semantics",
-        "notes": "Covers a non-default depth dilation. Existing 3D fixtures varied depth padding/stride but kept dilationD=1."
       },
       "attrs": { "strides": [1, 1, 1], "dilations": [2, 1, 1] },
       "inputs": {
@@ -1933,10 +1977,10 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 17], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
     {
-      "name": "grouped_k3_group4_w4_below_prior_min_kernel",
       "provenance": {
-        "source": "authored for route coverage",
-        "notes": "A 3x3 grouped convolution above the wide-workgroup spatial floor with a width divisible by four, so it takes grouped_large_kernel_w4. Kernel sizes below 7 could not reach this variant before, so nothing covered the w4 arm at k3."
       },
       "attrs": { "group": 4, "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -1946,10 +1990,10 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 20], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
     {
-      "name": "grouped_k3_depthwise_w4_below_prior_min_kernel",
       "provenance": {
-        "source": "authored for route coverage",
-        "notes": "Depthwise 3x3, the most common grouped shape in mobile vision backbones, on the same w4 arm."
       },
       "attrs": { "group": 8, "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -1959,15 +2003,20 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 20], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
     {
-      "name": "grouped_k3_group4_odd_width_tail_below_prior_min_kernel",
       "provenance": {
-        "source": "authored for route coverage",
-        "notes": "A 3x3 grouped convolution whose width is not divisible by four, covering the odd-width tail specialization at k3."
       },
       "attrs": { "group": 4, "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
         "x": { "dtype": "float32", "shape": [1, 8, 16, 17] },
-        "w": { "dtype": "float32", "shape": [8, 2, 3, 3] }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 17], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
@@ -1975,7 +2024,7 @@
       "name": "grouped_large_kernel_k9_dilation8_span_over_cap",
       "provenance": {
         "source": "ONNX Conv-22 dilations semantics (per-axis dilations [1, 8])",
-        "notes": "Grouped 7x9 conv whose retained input-column span (kernelW-1)*dilationW + 1 + 3*strideW = 65 exceeds GROUPED_MAX_REGISTER_SPAN (64), so the grouped wide-output route must render its per-tap bounds-tested column form instead of the row-window register form; every pre-existing grouped fixture has span <= 24. Asymmetric dilations are valid ONNX (dilations is per spatial axis). Geometry keeps all nine width taps live while leaving kw=0 out of range at ow<2 and kw=8 out of range at ow>=14, and the 7-tall window off both height edges, so the per-tap column and row guards are exercised in both directions."
       },
       "attrs": { "group": 2, "strides": [1, 1], "dilations": [1, 8], "pads": [3, 2, 3, 2] },
       "inputs": {
@@ -1996,7 +2045,7 @@
       "name": "grouped_large_kernel_k9_dilation8_stride2_span_over_cap_tail",
       "provenance": {
         "source": "ONNX Conv-22 strides/dilations semantics (strides [1, 2], dilations [1, 8])",
-        "notes": "Same over-cap span regime as grouped_large_kernel_k9_dilation8_span_over_cap (span = 8*8 + 1 + 3*2 = 71 > 64) but with strideW = 2 and an output width of 15. The per-tap form spaces its four lanes by lane*strideW, which folds to lane whenever strideW is 1, so strideW = 2 makes a dropped factor observable; the non-multiple-of-four width also exercises the masked tail epilogue. Output 18x15 = 270 keeps the wide-output route's spatial minimum."
       },
       "attrs": { "group": 2, "strides": [1, 2], "dilations": [1, 8], "pads": [3, 2, 3, 2] },
       "inputs": {
@@ -2016,7 +2065,7 @@
     {
       "name": "conv1d_aligned_sgmat_k3_s1_p1_c32m128_w1024",
       "provenance": {
-        "notes": "Whisper-stem-shaped NCW conv with K=96 and N=1024; covers aligned materialized subgroup-matrix Conv1d on the full tier and the register-tiled/scalar NCW fallbacks."
       },
       "attrs": { "strides": [1], "pads": [1, 1] },
       "inputs": {
@@ -2106,7 +2155,7 @@
     {
       "name": "conv1d_tiled_reg_f16_bias_k3_s1_p1_c16m128_w1024",
       "provenance": {
-        "notes": "f16 companion covering materialized subgroup-matrix Conv1d and the tiled NCW fallback: operands accumulate in f32 and narrow once at store."
       },
       "attrs": { "strides": [1], "pads": [1, 1] },
       "inputs": {
@@ -2151,7 +2200,7 @@
     {
       "name": "grouped_kernel12_wide_bias_forces_nchw_scalar",
       "provenance": {
-        "notes": "Bias companion of the oversized-kernel grouped conv: the bias epilogue must be exercised on the scalar NCHW direct kernel, which is the only path left once kernel width 12 and group > 1 reject the unrolled and GEMM variants."
       },
       "attrs": { "group": 2, "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
@@ -2196,7 +2245,7 @@
     {
       "name": "conv_8x8_im2col_gemm_tiled_bias_reg_m1_n8192",
       "provenance": {
-        "notes": "Bias companion of the N-axis register-tiled im2col GEMM: the 8x8 kernel keeps the 1x1 register-tiled shortcut ineligible, so the bias epilogue runs on the im2col register-tiled GEMM."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
@@ -2221,7 +2270,7 @@
     {
       "name": "conv3d_im2col_subgroup_matrix_aligned_f32",
       "provenance": {
-        "notes": "Rank-5 aligned matrix-path coverage. Shape-derived K=C*KD*KH*KW=32 and N=OD*OH*OW=64 exercise the NCDHW im2col producer and f32 subgroup-matrix consumer without relying on a model-specific size."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
@@ -2241,7 +2290,7 @@
     {
       "name": "conv3d_im2col_subgroup_matrix_aligned_bias_f16",
       "provenance": {
-        "notes": "F16+bias companion for the general aligned rank-5 matrix path. The geometry is derived from the tensor contract and also validates the scratch-store narrowing epilogue."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
@@ -2282,7 +2331,7 @@
     {
       "name": "conv3d_im2col_tiled_bias_reg_large_spatial_f32",
       "provenance": {
-        "notes": "Bias companion for the portable register-tiled rank-5 path, using the same shape-derived tile eligibility and a separately checkable epilogue."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
@@ -2348,7 +2397,7 @@
     {
       "name": "conv_3x3_implicit_im2col_m32_bk16_f16_bias_batch2",
       "provenance": {
-        "notes": "Portable M=32 carried-gather coverage for the remaining specialization edges: K=72 selects BK16, batch=2 exercises the z dispatch and batch base, and f16 X/W/Y plus bias exercise the half-storage path. The 128x256 plane is the conservative N=32768 specialization boundary established by the f32/f16 A/B controls."
       },
       "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -2459,7 +2508,7 @@
       "name": "conv_3x3_splitk_im2col_padded_k_n_tails",
       "tunables": { "TILED_SPLIT_K_MODE": 1 },
       "provenance": {
-        "notes": "Split-K over a materialized padded im2col on subgroup-matrix tiers and the portable implicit gather otherwise: 5x7 output is 35 positions, so the un-split GEMM would dispatch ONE workgroup and leave 39 of 40 cores idle. K = 432 and N = 35 exercise the reduction tails in both implementations."
       },
       "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -2651,7 +2700,7 @@
     {
       "name": "conv_1x1_implicit_im2col_splitk_aligned_starved",
       "provenance": {
-        "notes": "An ALIGNED 1x1 that is nonetheless starved: 32x32 is 16 column tiles over one 64-row tile, so the split outranks the dedicated aligned 1x1 kernel here even though that kernel is eligible and wins at any well-fed size."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
@@ -2790,6 +2839,1571 @@
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 91, 91], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     }
   ]
 }

 {
   "fixtureArrays": {
     "nchw_dilated_kernel_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24],
     "backend_stride2_with_padding_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34],
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_op_test.cc",
         "test": "ConvTest",
+        "notes": "A zero channel-bias tensor exercises the 1x1 matrix epilogue; finite subnormal convolution outputs must remain unchanged."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_op_test.cc",
         "test": "ConvTest",
+        "notes": "A zero channel-bias tensor exercises the aligned column-matrix epilogue; finite subnormal convolution outputs must remain unchanged."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_op_test.cc",
         "test": "ConvTest",
+        "notes": "A zero channel-bias tensor exercises the padded column-matrix epilogue; padded K/N lanes must not alter finite subnormal real outputs."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 16, 1], "tolerance": 0.0001 } },
       "attrs": {}
     },
+    {
+      "name": "gemm_1x1_tiled_reg_splitk_c128_16x16",
+      "provenance": {
+        "notes": "Route lock for the split-K 1x1 register tile: four 64x64 output tiles over K=128 take two K slices each, and the reduce sums the f32 partials."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 128, 16, 16],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.031 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 128, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.011, "cosStep": 0.023 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 16, 16], "tolerance": 0.0001 } },
+      "attrs": {}
+    },
+    {
+      "name": "gemm_1x1_tiled_reg_splitk_bias_c128_16x16",
+      "provenance": {
+        "notes": "The split-K 1x1 register tile produces two f32 partials per output tile; the reduce pass sums them before adding per-channel bias."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 128, 16, 16],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.019, "cosStep": 0.029 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 128, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.07, "cosStep": 0.05 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 16, 16], "tolerance": 0.0001 } },
+      "attrs": {}
+    },
     {
       "name": "ort_conv2d_1x1_splitk_with_bias",
       "provenance": {
     {
       "name": "conv1d_f16_audio_feature",
       "provenance": {
+        "source": "local fixture",
+        "notes": "A compact half-precision audio-feature convolution exercises the dtype-generic NCW implementation."
       },
       "attrs": { "strides": [1], "dilations": [1], "pads": [1, 1] },
       "inputs": {
     {
       "name": "conv_3x3_im2col_gemm_dense_c256_8x8",
       "provenance": {
+        "notes": "A compact 3x3 convolution retains a large K/M column-matrix regime while keeping the correctness tensor smaller than the model benchmarks. Offset operands keep each 2,304-tap accumulation proportional to its reduction length instead of cancelling toward zero, so a dropped tap or tile is observable."
       },
       "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
         "x": {
           "dtype": "float32",
           "shape": [1, 256, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.16, "sinStep": 0.013, "cosStep": 0.031, "offset": 0.2 }
         },
         "w": {
           "dtype": "float32",
           "shape": [256, 256, 3, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.015, "sinStep": 0.007, "cosStep": 0.019, "offset": 0.02 }
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 256, 8, 8], "tolerance": 0.006, "relTolerance": 0.006 } }
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/conv_fp16_test.cc",
         "test": "ConvFp16Test.Conv2D_Bias_2",
+        "notes": "This half-precision case combines channel bias, asymmetric explicit padding, and non-unit strides from the cited upstream Conv fixture."
       },
       "attrs": { "strides": [2, 3], "pads": [1, 2, 3, 1] },
       "inputs": {
       "name": "grouped_dilated_large_kernel_k11_oc4_nchw_fallback",
       "provenance": {
         "source": "Realistic grouped/dilated Conv stress geometry",
+        "notes": "A grouped, dilated 11x11 convolution produces two four-channel output tiles per group at the K11 selection boundary."
       },
       "attrs": { "group": 2, "strides": [1, 1], "dilations": [2, 2], "pads": [10, 10, 10, 10] },
       "inputs": {
       "name": "f16_grouped_large_kernel_k9_oc4_nchw_f16_fallback",
       "provenance": {
         "source": "Realistic grouped Conv stress geometry",
+        "notes": "A compact half-precision grouped 9x9 convolution produces two four-channel output tiles per group, exercising shared-input output-channel tiling."
       },
       "requires": { "features": ["shader-f16"] },
       "attrs": { "group": 2, "strides": [1, 1], "pads": [4, 4, 4, 4] },
       "name": "conv3d_depth_dilation2_compact",
       "provenance": {
         "source": "ONNX Conv-22 volumetric dilation semantics",
+        "notes": "Depth dilation 2 separates the two volumetric kernel planes while height and width dilation remain one."
       },
       "attrs": { "strides": [1, 1, 1], "dilations": [2, 1, 1] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 17], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
     {
+      "name": "grouped_k3_group4_w4",
       "provenance": {
+        "source": "repository-authored fixture",
+        "notes": "A grouped 3x3 convolution produces 20 columns and enough spatial work to select the four-column grouped specialization."
       },
       "attrs": { "group": 4, "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 20], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
     {
+      "name": "grouped_k3_depthwise_w4",
       "provenance": {
+        "source": "repository-authored fixture",
+        "notes": "A depthwise 3x3 convolution produces 20 columns and selects the four-column grouped specialization with one input channel per group."
       },
       "attrs": { "group": 8, "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 20], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
     {
+      "name": "grouped_k3_group4_odd_width_tail_bias",
       "provenance": {
+        "source": "repository-authored fixture",
+        "notes": "A grouped 3x3 convolution produces 17 columns, so the final four-lane output group contains one valid column. Mixed-sign channel bias exercises the biased masked-tail epilogue."
       },
       "attrs": { "group": 4, "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
         "x": { "dtype": "float32", "shape": [1, 8, 16, 17] },
+        "w": { "dtype": "float32", "shape": [8, 2, 3, 3] },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-0.5, 0.25, 0.75, -1.0, 1.25, -0.125, 0.375, -0.875] }
+        }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 17], "tolerance": 0.0001, "relTolerance": 0.0001 } }
     },
       "name": "grouped_large_kernel_k9_dilation8_span_over_cap",
       "provenance": {
         "source": "ONNX Conv-22 dilations semantics (per-axis dilations [1, 8])",
+        "notes": "For this grouped 7x9 convolution, (kernelW-1)*dilationW + 1 + 3*strideW = 65 exceeds GROUPED_MAX_REGISTER_SPAN (64), selecting per-tap bounds-tested column gathers. Asymmetric dilation keeps all nine width taps live while kw=0 is out of range at ow<2, kw=8 is out of range at ow>=14, and the seven-row window crosses both height edges."
       },
       "attrs": { "group": 2, "strides": [1, 1], "dilations": [1, 8], "pads": [3, 2, 3, 2] },
       "inputs": {
       "name": "grouped_large_kernel_k9_dilation8_stride2_span_over_cap_tail",
       "provenance": {
         "source": "ONNX Conv-22 strides/dilations semantics (strides [1, 2], dilations [1, 8])",
+        "notes": "The retained column span is 8*8 + 1 + 3*2 = 71, exceeding GROUPED_MAX_REGISTER_SPAN (64). strideW=2 makes each four-lane gather advance by twice the lane index, while output width 15 exercises the masked tail epilogue and 18x15 outputs satisfy the spatial work floor."
       },
       "attrs": { "group": 2, "strides": [1, 2], "dilations": [1, 8], "pads": [3, 2, 3, 2] },
       "inputs": {
     {
       "name": "conv1d_aligned_sgmat_k3_s1_p1_c32m128_w1024",
       "provenance": {
+        "notes": "A Whisper-stem-shaped NCW convolution with K=96 and N=1024 exercises aligned materialized matrix multiplication and the register-tiled and scalar NCW implementations."
       },
       "attrs": { "strides": [1], "pads": [1, 1] },
       "inputs": {
     {
       "name": "conv1d_tiled_reg_f16_bias_k3_s1_p1_c16m128_w1024",
       "provenance": {
+        "notes": "Half-precision operands exercise materialized matrix multiplication and tiled NCW paths; both accumulate in f32 and narrow once at the output store."
       },
       "attrs": { "strides": [1], "pads": [1, 1] },
       "inputs": {
     {
       "name": "grouped_kernel12_wide_bias_forces_nchw_scalar",
       "provenance": {
+        "notes": "Kernel width 12 with group > 1 selects the scalar NCHW implementation. Nonzero channel bias exercises its output epilogue."
       },
       "attrs": { "group": 2, "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
     {
       "name": "conv_8x8_im2col_gemm_tiled_bias_reg_m1_n8192",
       "provenance": {
+        "notes": "An 8x8 kernel excludes the 1x1 shortcut and selects the N-axis register-tiled column-matrix product. Nonzero channel bias exercises its output epilogue."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
     {
       "name": "conv3d_im2col_subgroup_matrix_aligned_f32",
       "provenance": {
+        "notes": "For this rank-5 tensor, K=C*KD*KH*KW=32 and N=OD*OH*OW=64. The aligned geometry exercises the NCDHW column producer and f32 matrix consumer without a model-specific size."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
     {
       "name": "conv3d_im2col_subgroup_matrix_aligned_bias_f16",
       "provenance": {
+        "notes": "Half-precision rank-5 input with channel bias exercises the aligned matrix path and its f32-to-f16 scratch-store epilogue."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
     {
       "name": "conv3d_im2col_tiled_bias_reg_large_spatial_f32",
       "provenance": {
+        "notes": "Rank-5 geometry selects the portable register-tiled path, and nonzero channel bias makes its epilogue independently observable."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
     {
       "name": "conv_3x3_implicit_im2col_m32_bk16_f16_bias_batch2",
       "provenance": {
+        "notes": "M=32 and K=72 select the BK16 carried-gather specialization. batch=2 exercises the z dispatch and batch offset; half-precision X, W, Y, and bias exercise packed storage at the N=128*256=32768 boundary."
       },
       "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
       "name": "conv_3x3_splitk_im2col_padded_k_n_tails",
       "tunables": { "TILED_SPLIT_K_MODE": 1 },
       "provenance": {
+        "notes": "A 5x7 output has 35 positions, so split-K adds independent workgroups to the single output tile. K=432 and N=35 exercise reduction tails in both the subgroup-matrix materialized-im2col route and the portable implicit-gather route."
       },
       "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
       "inputs": {
     {
       "name": "conv_1x1_implicit_im2col_splitk_aligned_starved",
       "provenance": {
+        "notes": "An aligned 1x1 convolution with a 32x32 output produces 16 column tiles over one 64-row tile, exercising split-K despite satisfying the dedicated aligned-kernel geometry."
       },
       "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
       "inputs": {
         }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 91, 91], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "grouped_dilated_lanes_k9_d4_g2_16x40",
+      "provenance": {
+        "notes": "Dilated-lane grouped route: four columns per invocation spaced by the dilation, so every tap reads a 12-register row window; width 40 is a whole number of lane runs in every residue class."
+      },
+      "attrs": { "group": 2, "strides": [1, 1], "dilations": [4, 4], "pads": [16, 16, 16, 16] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 16, 40],
+          "data": { "kind": "fillFloat32", "scale": 0.25, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [8, 4, 9, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 40], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "grouped_dilated_lanes_k9_d4_g2_odd_width_41_lane_tail",
+      "provenance": {
+        "notes": "Dilated-lane route with a width that is not a whole number of runs: residue classes hold 11, 10, 10 and 10 columns, so the last run of every class stores only its leading lanes."
+      },
+      "attrs": { "group": 2, "strides": [1, 1], "dilations": [4, 4], "pads": [16, 16, 16, 16] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 16, 41],
+          "data": { "kind": "fillFloat32", "scale": 0.25, "sinStep": 0.031, "cosStep": 0.019 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [8, 4, 9, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.023, "cosStep": 0.011 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 41], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "f16_grouped_dilated_lanes_k9_d5_g4_oc2_12x40",
+      "provenance": {
+        "notes": "f16 dilated-lane route at a two-channel output tile (four groups of two) and the smallest width the gate admits at dilation 5."
+      },
+      "attrs": { "group": 4, "strides": [1, 1], "dilations": [5, 5], "pads": [20, 20, 20, 20] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [1, 8, 12, 40],
+          "data": { "kind": "fillFloat32", "scale": 0.25, "sinStep": 0.041, "cosStep": 0.017 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [8, 2, 9, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.019, "cosStep": 0.031 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 8, 12, 40], "tolerance": 0.01, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "grouped_dilated_lanes_k9_d4_g2_16x40_bias",
+      "provenance": {
+        "notes": "The dilated-lanes grouped route carrying a bias: the accumulator is seeded per output channel before any tap is read, and the lane registers are the ones that must pick it up."
+      },
+      "attrs": { "group": 2, "strides": [1, 1], "dilations": [4, 4], "pads": [16, 16, 16, 16] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 8, 16, 40],
+          "data": { "kind": "fillFloat32", "scale": 0.25, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [8, 4, 9, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.5, -0.25, 1.0, -0.75, 0.125, 2.0, -1.5, 0.375] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8, 16, 40], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m32_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 32, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m32_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 32, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m32_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 32, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m32_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 32, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m64_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m64_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m64_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m64_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m96_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [96, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 96, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m96_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [96, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 96, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m96_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [96, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 96, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m96_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [96, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 96, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m128_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [128, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m128_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [128, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 128], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m128_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [128, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m128_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [128, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m31_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [31, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 31, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m33_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [33, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 33, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m65_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [65, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 65, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m64_c31_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 31, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 31, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 8, 8], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m64_c32_rank2_n7x9_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 7, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 7, 9], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float16_m32_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m32_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m32_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m32_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m96_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [96, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 96, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m96_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [96, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 96, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m96_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [96, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 96, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m96_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [96, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [96],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 96, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m128_c32_rank1_n128_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [128, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 128, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m128_c32_rank1_n128_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [128, 32, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 128, 128], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m128_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [128, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 128, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m128_c32_rank2_n8x8_bias1",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [128, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 128, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m31_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [31, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 31, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m33_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [33, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 33, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m65_c32_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [65, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 65, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c31_rank2_n8x8_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 31, 8, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 31, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 8, 8], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c32_rank2_n7x9_bias0",
+      "provenance": {
+        "notes": "Direct column matrix inputs cover both input scalar widths, full row bands, batched NCW/NCHW indexing and bias. A shallow reduction retains the unsplit accumulation order; incomplete rows/K/N exercise guarded fallbacks."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 7, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 7, 9], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float32_m128_c32_rank2_n32x32_bias1",
+      "provenance": {
+        "notes": "Full spatial tiles on each side of the output-channel reuse preference: retain implicit gathering at low reuse and select direct materialization at high reuse. Batched bias outputs retain the CPU oracle and default dtype tolerance."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 32, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [128, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 32, 32], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float32_m512_c32_rank2_n32x32_bias1",
+      "provenance": {
+        "notes": "Full spatial tiles on each side of the output-channel reuse preference: retain implicit gathering at low reuse and select direct materialization at high reuse. Batched bias outputs retain the CPU oracle and default dtype tolerance."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 32, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [512, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 512, 32, 32], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_columns_float16_m64_c32_rank2_n32x32_bias1",
+      "provenance": {
+        "notes": "Full spatial tiles on each side of the output-channel reuse preference: retain implicit gathering at low reuse and select direct materialization at high reuse. Batched bias outputs retain the CPU oracle and default dtype tolerance."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 32, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 32, 32], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_columns_float16_m256_c32_rank2_n32x32_bias1",
+      "provenance": {
+        "notes": "Full spatial tiles on each side of the output-channel reuse preference: retain implicit gathering at low reuse and select direct materialization at high reuse. Batched bias outputs retain the CPU oracle and default dtype tolerance."
+      },
+      "attrs": { "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 32, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.037, "cosStep": 0.013 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [256, 32, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.017, "cosStep": 0.029 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 256, 32, 32], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m32_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 32, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m32_k32_n64_mode1",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [32, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.23, "cosStep": 0.391 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 32, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m64_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m64_k32_n64_mode1",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        },
+        "bias": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.23, "cosStep": 0.391 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m33_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [33, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 33, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m65_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [65, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 65, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m96_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [96, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 96, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m64_k31_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 31, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 31, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m64_k33_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 33, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 33, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 1, 64], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m64_k32_n63_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 1, 63], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float32_m64_k32_n65_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 32, 1, 65],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 64, 1, 65], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m32_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m32_k32_n64_mode1",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [32, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.23, "cosStep": 0.391 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 32, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m64_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m64_k32_n64_mode1",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        },
+        "bias": {
+          "dtype": "float16",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.23, "cosStep": 0.391 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m33_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [33, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 33, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m65_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [65, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 65, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m96_k32_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [96, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 96, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m64_k31_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 31, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 31, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m64_k33_n64_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 33, 1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 33, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 1, 64], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m64_k32_n63_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 1, 63], "tolerance": 0.001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "direct_1x1_float16_m64_k32_n65_mode0",
+      "provenance": {
+        "notes": "Batched direct 1x1 matrix loads cover both scalar widths and fused epilogues. Neighboring incomplete M/K/N tiles retain guarded staging or the existing fallback; device limits determine the row tile."
+      },
+      "attrs": { "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 32, 1, 65],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.022099999999999998 }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [64, 32, 1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.0527 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64, 1, 65], "tolerance": 0.001, "relTolerance": 0.001 } }
     }
   ]
 }