Xenova HF Staff commited on 21 days ago

Commit

2def0ac

verified ·

1 Parent(s): 900921f

sync 6fdf6301e2bb

Browse files

Files changed (18) hide show

README.md +4 -2
build/webgpu/bench.json +40 -1
build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja +37 -163
build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja +28 -48
build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja +16 -161
build/webgpu/conv-transpose-empty-input.wgsl.jinja +5 -30
build/webgpu/conv-transpose1d-ncl.wgsl.jinja +6 -35
build/webgpu/conv-transpose2d-col2im.wgsl.jinja +6 -36
build/webgpu/conv-transpose2d-grouped-stride-phase.wgsl.jinja +7 -2
build/webgpu/conv-transpose2d-nchw.wgsl.jinja +8 -35
build/webgpu/conv-transpose2d-phase-gemm-tiled.wgsl.jinja +43 -10
build/webgpu/conv-transpose2d-weight-reorder.wgsl.jinja +5 -34
build/webgpu/conv-transpose3d-col2im.wgsl.jinja +6 -32
build/webgpu/conv-transpose3d-ncdhw.wgsl.jinja +6 -32
build/webgpu/conv-transpose3d-weight-reorder.wgsl.jinja +5 -30
build/webgpu/manifest.json +222 -442
build/webgpu/metadata.json +23 -29
build/webgpu/test.json +376 -11

README.md CHANGED Viewed

@@ -59,6 +59,8 @@ One implementation is selected per call from the device capabilities, the reques
 - `ncdhw3d_implicit_sgmat_bias` — Stride-1 transposed convolution as a forward convolution of the unchanged input with tap-reversed weights and begin pads of kernel-1-pad: the implicit-gather GEMM reads input windows directly and stores output elements directly, so no column matrix is materialized and no scatter pass runs, while the reduction spans the full in-channel-by-tap depth.
 - `nchw2d_grouped_stride_phase_blocked` — Grouped stride-phase ConvTranspose where one invocation owns one output pixel and a block of that group's output channels: every gathered input sample feeds the whole block and the group's weight slab is staged once per workgroup. It serves grouped stride-two-or-more shapes whose per-group weights fit workgroup memory.
 - `nchw2d_grouped_stride_phase_blocked_bias` — Grouped stride-phase ConvTranspose where one invocation owns one output pixel and a block of that group's output channels: every gathered input sample feeds the whole block and the group's weight slab is staged once per workgroup. It serves grouped stride-two-or-more shapes whose per-group weights fit workgroup memory.
 ## Device requirements
@@ -69,7 +71,7 @@ Some implementation variants require `subgroup-matrix`, `shader-f16`, and `subgr
 - [`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
 - [`conv-1x1-gemm-tiled-reg.wgsl.jinja`](build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja)
 - [`conv-1x1-gemm-tiled.wgsl.jinja`](build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja)
 - [`conv-1x1-subgroup-matrix.wgsl.jinja`](build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja)
@@ -89,7 +91,7 @@ Some implementation variants require `subgroup-matrix`, `shader-f16`, and `subgr
 ## 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.

 - `ncdhw3d_implicit_sgmat_bias` — Stride-1 transposed convolution as a forward convolution of the unchanged input with tap-reversed weights and begin pads of kernel-1-pad: the implicit-gather GEMM reads input windows directly and stores output elements directly, so no column matrix is materialized and no scatter pass runs, while the reduction spans the full in-channel-by-tap depth.
 - `nchw2d_grouped_stride_phase_blocked` — Grouped stride-phase ConvTranspose where one invocation owns one output pixel and a block of that group's output channels: every gathered input sample feeds the whole block and the group's weight slab is staged once per workgroup. It serves grouped stride-two-or-more shapes whose per-group weights fit workgroup memory.
 - `nchw2d_grouped_stride_phase_blocked_bias` — Grouped stride-phase ConvTranspose where one invocation owns one output pixel and a block of that group's output channels: every gathered input sample feeds the whole block and the group's weight slab is staged once per workgroup. It serves grouped stride-two-or-more shapes whose per-group weights fit workgroup memory.
+- `nchw2d_exact_store_gemm_f32` — Exact-tiling f32 transposed convolution retains the register GEMM contraction and stores directly into output, removing the column intermediate and gather pass.
+- `nchw2d_exact_store_gemm_f32_bias` — Exact-tiling f32 transposed convolution retains the register GEMM contraction and stores directly into output, removing the column intermediate and gather pass.
 ## Device requirements
 - [`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 cases
 - [`conv-1x1-gemm-tiled-reg.wgsl.jinja`](build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja)
 - [`conv-1x1-gemm-tiled.wgsl.jinja`](build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja)
 - [`conv-1x1-subgroup-matrix.wgsl.jinja`](build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

build/webgpu/bench.json CHANGED Viewed

@@ -342,7 +342,7 @@
       "preset": "stress",
       "provenance": {
         "source": "synthetic benchmark",
-        "notes": "Adding channel bias excludes nchw2d_grouped_stride_phase and selects nchw2d_grouped_bias; the paired no-bias case isolates this optional-input route boundary."
       },
       "vars": {
         "batch": 1,
@@ -446,6 +446,45 @@
           }
         ]
       }
     }
   ]
 }

       "preset": "stress",
       "provenance": {
         "source": "synthetic benchmark",
+        "notes": "Grouped stride-two transposed convolution with channel bias; the paired no-bias case isolates the effect of the optional input."
       },
       "vars": {
         "batch": 1,
           }
         ]
       }
+    },
+    {
+      "name": "conv-transpose-f32-exact-store-b2c65m80-29x53-k3x2-nobias",
+      "preset": "edge",
+      "provenance": { "notes": "Rectangular nonoverlapping taps with batch and channel tails." },
+      "attrs": { "strides": [3, 2] },
+      "inputs": {
+        "x": { "shape": [2, 65, 29, 53], "dtype": "float32", "dist": "normal", "seed": 9421, "scale": 0.2 },
+        "w": { "shape": [65, 80, 3, 2], "dtype": "float32", "dist": "normal", "seed": 9422, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [2, 80, 87, 106], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * dim(shapes.x, 0) * dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "conv-transpose-f32-exact-store-b1c33m48-41x43-k2x3-bias",
+      "preset": "edge",
+      "provenance": { "notes": "Rectangular nonoverlapping taps with batch and channel tails." },
+      "attrs": { "strides": [2, 3] },
+      "inputs": {
+        "x": { "shape": [1, 33, 41, 43], "dtype": "float32", "dist": "normal", "seed": 9421, "scale": 0.2 },
+        "w": { "shape": [33, 48, 2, 3], "dtype": "float32", "dist": "normal", "seed": 9422, "scale": 0.1 },
+        "bias": { "shape": [48], "dtype": "float32", "dist": "normal", "seed": 9423, "scale": 0.1 }
+      },
+      "outputs": { "y": { "shape": [1, 48, 82, 129], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * dim(shapes.x, 0) * dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)"
+          }
+        ]
+      }
     }
   ]
 }

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

@@ -14,45 +14,14 @@
 // 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 %}
-{% set tileT = "f16" if usesF16 else "f32" -%}
-{%- set GEMM_BK = gemmKTile if gemmKTile is defined else 16 %}
-{% set GEMM_TM = gemmThreadRows if gemmThreadRows is defined else 8 %}
-{% set GEMM_WG_X = gemmWorkgroupX if gemmWorkgroupX is defined else 16 %}
-{% set GEMM_BM = gemmMTile if gemmMTile is defined else 64 %}
-{% set GEMM_BN = gemmNTile if gemmNTile is defined else GEMM_WG_X * 4 %}
-{% set GEMM_TN = gemmThreadColumns if gemmThreadColumns is defined else (GEMM_BN / GEMM_WG_X)|int %}
-{% set GEMM_WG_Y = gemmWorkgroupY if gemmWorkgroupY is defined else (GEMM_BM / GEMM_TM)|int %}
 {% set components = ["x", "y", "z", "w"] %}
 const BK: u32 = {{ GEMM_BK }}u;
 const BM: u32 = {{ GEMM_BM }}u;
@@ -70,18 +39,11 @@ const BN_VECS: u32 = BN / 4u;
 // 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 gemmKLimit = "kEnd" if splitKPartial else "K" %}
 {% set implicitGatherMode = tunables.IMPLICIT_TILED_GATHER_MODE if tunables.IMPLICIT_TILED_GATHER_MODE is defined else 0 %}
-{% set implicitGatherWgSize = GEMM_WG_X * GEMM_WG_Y %}
-{% set implicitGatherColumnVecs = (GEMM_BN / 4)|int %}
-{% set implicitGatherKChunks = (implicitGatherWgSize / implicitGatherColumnVecs)|int %}
 {% set implicitGatherDivides = implicitGatherWgSize % implicitGatherColumnVecs == 0 and implicitGatherKChunks > 0 and GEMM_BK % implicitGatherKChunks == 0 %}
 {% set useCarriedImplicitGather = implicitIm2col and implicitGatherMode == 1 and implicitGatherDivides %}
 {% if useCarriedImplicitGather %}
@@ -102,72 +64,36 @@ const CONV_IN_W: u32 = {{ convInW }}u;
 const CONV_OUT_W: u32 = {{ convOutW }}u;
 const CONV_IN_CHANNELS: u32 = {{ convInChannels }}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;
   }
   var v = raw;
-{% if hasZ %}
-  v = v + f32(zResidual[yBase + m * params.N + n]);
-{% endif %}
-{% if hasActivation %}
-  v = fused_act(v);
 {% endif %}
   y[yBase + m * params.N + n] = {{ T }}(v);
-}
-{% endif %}
 {% endif %}
 var<workgroup> tileA: array<array<vec4<{{ tileT }}>, AK_VECS>, BM>;
 var<workgroup> tileB: array<array<vec4<{{ tileT }}>, BN_VECS>, BK>;
-{% if fusedNarrowProjection and GEMM_BK < GEMM_BM %}
-// BK16 cannot reuse the 16-row input tile to publish a BM32 intermediate.
-// The extra 32x64 f32 tile keeps the total at 14 KiB, below WebGPU's
-// guaranteed 16 KiB workgroup-storage floor.
-var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
-{% endif %}
-{% macro publish_projection_input(localRow, localVec, component, globalRow, globalColumn, raw) %}
-  if ({{ globalRow }} < M && {{ globalColumn }} < N) {
-    var projectionInput = {{ raw }};
-{% if projectionInputAct == "relu" %}
-    projectionInput = projection_relu(projectionInput);
-{% endif %}
-    {{ narrowProjectionTile }}[{{ localRow }}][{{ localVec }}].{{ component }} = projectionInput;
-  } else {
-    {{ narrowProjectionTile }}[{{ localRow }}][{{ localVec }}].{{ component }} = 0.0;
-  }
-{%- endmacro %}
 {% macro load_a_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BM * AK_VECS; linear += WG_SIZE) {
       let ar = linear / AK_VECS;
@@ -188,8 +114,7 @@ var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
         }
       }
       tileA[ar][ac4] = av;
-    }
-{%- endmacro %}
 {% macro load_b_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BK * BN_VECS; linear += WG_SIZE) {
       let br = linear / BN_VECS;
@@ -210,8 +135,7 @@ var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
         }
       }
       tileB[br][bc4] = bvec;
-    }
-{%- endmacro %}
 {% macro load_b_implicit_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BK * BN_VECS; linear += WG_SIZE) {
       let br = linear / BN_VECS;
@@ -239,8 +163,7 @@ var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
         }
       }
       tileB[br][bc4] = bvec;
-    }
-{%- endmacro %}
 {% macro load_b_implicit_carried(rowBase, columnBase, rowLimit, columnLimit) %}
     let bColVec = li / {{ implicitGatherKChunks }}u;
     let bChunk = li % {{ implicitGatherKChunks }}u;
@@ -324,8 +247,7 @@ var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
           }
         }
       }
-    }
-{%- endmacro %}
 @compute @workgroup_size({{ GEMM_WG_X }}, {{ GEMM_WG_Y }}, 1)
 fn main(
@@ -377,59 +299,11 @@ fn main(
     workgroupBarrier();
   }
-{% if not fusedNarrowProjection %}
   let yBatchBase = batch * M * N;
-{% endif %}
   let m0 = mBase + lid.y * TM;
   let n0 = nBase + lid.x * TN;
 {% for row in range(GEMM_TM) %}
 {% for column in range(GEMM_TN) %}
-{% if fusedNarrowProjection %}
-{{ publish_projection_input("lid.y * TM + " ~ row ~ "u", "lid.x * " ~ ((GEMM_TN / 4)|int) ~ "u + " ~ ((column / 4)|int) ~ "u", components[column % 4], "m0 + " ~ row ~ "u", "n0 + " ~ column ~ "u", "acc" ~ row ~ "." ~ components[column]) }}
-{% else %}
   store_out(m0 + {{ row }}u, n0 + {{ column }}u, yBatchBase, acc{{ row }}.{{ components[column] }});
-{% endif %}
-{% endfor %}
-{% endfor %}{% if fusedNarrowProjection %}
-  // Conv's final K-tile barrier makes tileB dead before it becomes the
-  // intermediate tile. Every lane publishes its unique micro-tile, then one
-  // lane per spatial vector word consumes all logical M rows in increasing
-  // order, four adjacent columns at a time. Invalid M/N-tail cells are
-  // initialized above, so the barrier is uniform and no lane can observe a
-  // stale input-tile value.
-  workgroupBarrier();
-  if (li < BN_VECS) {
-    let projectionNBase = nBase + li * 4u;
-{% for oc in range(projectionChannels) %}
-    var projectionAcc{{ oc }} = vec4<f32>(0.0);
-{% endfor %}
-    for (var channel = 0u; channel < M; channel = channel + 1u) {
-      let projectionValues = vec4<f32>({{ narrowProjectionTile }}[channel][li]);
-{% for oc in range(projectionChannels) %}
-      projectionAcc{{ oc }} = projectionAcc{{ oc }} + projectionValues * f32(projectionW[{{ oc }}u * M + channel]);
 {% endfor %}
-    }
-    let projectionYBase = batch * PROJECTION_OUT_C * N;
-{% for oc in range(projectionChannels) %}
-    let projected{{ oc }} = projectionAcc{{ oc }};
-{% endfor %}
-{% for column in range(4) %}
-    let projectionN{{ column }} = projectionNBase + {{ column }}u;
-    if (projectionN{{ column }} < N) {
-{% for oc in range(projectionChannels) %}
-{% if projectionOutputAct == "relu" %}
-      let activated{{ oc }}_{{ column }} = projection_relu(projected{{ oc }}.{{ components[column] }});
-{% elif projectionOutputAct == "sigmoid" %}
-      let activated{{ oc }}_{{ column }} = sigmoid_safe(projected{{ oc }}.{{ components[column] }});
-{% else %}
-      let activated{{ oc }}_{{ column }} = projected{{ oc }}.{{ components[column] }};
-{% endif %}
-      y[projectionYBase + {{ oc }}u * N + projectionN{{ column }}] = activated{{ oc }}_{{ column }};
-{% endfor %}
-    }
-{% endfor %}
-  }
-{% endif %}
-}

 // 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 tileT = "f16" if usesF16 else "f32" %}
+{% set GEMM_BK = gemmKTile if gemmKTile is defined else 16 %}
+{% set GEMM_TM = 8 %}
+{% set GEMM_WG_X = 16 %}
+{% set GEMM_BM = 64 %}
+{% set GEMM_BN = 64 %}
+{% set GEMM_TN = 4 %}
+{% set GEMM_WG_Y = 8 %}
 {% set components = ["x", "y", "z", "w"] %}
 const BK: u32 = {{ GEMM_BK }}u;
 const BM: u32 = {{ GEMM_BM }}u;
 // 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 gemmKLimit = "K" %}
 {% set implicitGatherMode = tunables.IMPLICIT_TILED_GATHER_MODE if tunables.IMPLICIT_TILED_GATHER_MODE is defined else 0 %}
+{% set implicitGatherWgSize = 128 %}
+{% set implicitGatherColumnVecs = 16 %}
+{% set implicitGatherKChunks = 8 %}
 {% set implicitGatherDivides = implicitGatherWgSize % implicitGatherColumnVecs == 0 and implicitGatherKChunks > 0 and GEMM_BK % implicitGatherKChunks == 0 %}
 {% set useCarriedImplicitGather = implicitIm2col and implicitGatherMode == 1 and implicitGatherDivides %}
 {% if useCarriedImplicitGather %}
 const CONV_OUT_W: u32 = {{ convOutW }}u;
 const CONV_IN_CHANNELS: u32 = {{ convInChannels }}u;
 {% endif %}
+fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
   if (m >= params.M || n >= params.N) {
     return;
   }
+{% if convTransposeExactStore | default(false) %}
+{% set convTransposeKernelH = convTransposeKernelH | default(1) %}
+{% set convTransposeKernelW = convTransposeKernelW | default(1) %}
+{% set convTransposeInputW = convTransposeInputW | default(1) %}
+{% set convTransposeOutputH = convTransposeOutputH | default(1) %}
+{% set convTransposeOutputW = convTransposeOutputW | default(1) %}
+  // Nonoverlapping transposed-convolution taps map each matrix element to
+  // exactly one output. Keep the old col2im bias-plus-dot rounding order.
+  let kernelArea = {{ convTransposeKernelH }}u * {{ convTransposeKernelW }}u;
+  let oc = m / kernelArea;
+  let tap = m % kernelArea;
+  let oh = (n / {{ convTransposeInputW }}u) * {{ convTransposeKernelH }}u + tap / {{ convTransposeKernelW }}u;
+  let ow = (n % {{ convTransposeInputW }}u) * {{ convTransposeKernelW }}u + tap % {{ convTransposeKernelW }}u;
+  let outputIndex = yBase + (oc * {{ convTransposeOutputH }}u + oh) * {{ convTransposeOutputW }}u + ow;
+  y[outputIndex] = {% if hasBias %}f32(bias[oc]){% else %}0.0{% endif %} + raw;
+{% else %}
   var v = raw;
+{% if hasBias %}
+  v = v + f32(bias[m]);
 {% endif %}
   y[yBase + m * params.N + n] = {{ T }}(v);
 {% endif %}
+}
 var<workgroup> tileA: array<array<vec4<{{ tileT }}>, AK_VECS>, BM>;
 var<workgroup> tileB: array<array<vec4<{{ tileT }}>, BN_VECS>, BK>;
 {% macro load_a_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BM * AK_VECS; linear += WG_SIZE) {
       let ar = linear / AK_VECS;
         }
       }
       tileA[ar][ac4] = av;
+    }{% endmacro %}
 {% macro load_b_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BK * BN_VECS; linear += WG_SIZE) {
       let br = linear / BN_VECS;
         }
       }
       tileB[br][bc4] = bvec;
+    }{% endmacro %}
 {% macro load_b_implicit_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BK * BN_VECS; linear += WG_SIZE) {
       let br = linear / BN_VECS;
         }
       }
       tileB[br][bc4] = bvec;
+    }{% endmacro %}
 {% macro load_b_implicit_carried(rowBase, columnBase, rowLimit, columnLimit) %}
     let bColVec = li / {{ implicitGatherKChunks }}u;
     let bChunk = li % {{ implicitGatherKChunks }}u;
           }
         }
       }
+    }{% endmacro %}
 @compute @workgroup_size({{ GEMM_WG_X }}, {{ GEMM_WG_Y }}, 1)
 fn main(
     workgroupBarrier();
   }
   let yBatchBase = batch * M * N;
   let m0 = mBase + lid.y * TM;
   let n0 = nBase + lid.x * TN;
 {% for row in range(GEMM_TM) %}
 {% for column in range(GEMM_TN) %}
   store_out(m0 + {{ row }}u, n0 + {{ column }}u, yBatchBase, acc{{ row }}.{{ components[column] }});
 {% endfor %}
+{% endfor %}}

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

@@ -6,64 +6,45 @@
 // 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 %}
-{% set tileT = "f16" if usesF16 else "f32" -%}
 const BK: u32 = 16u;
 const BM: u32 = 32u;
 const BN: u32 = 32u;
 // Store one output element with the 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;
   }
   var v = raw;
-{% if hasZ %}
-  v = v + f32(zResidual[yBase + m * params.N + n]);
-{% endif %}
-{% if hasActivation %}
-  v = fused_act(v);
-{% endif %}
   y[yBase + m * params.N + n] = {{ T }}(v);
 }
 var<workgroup> tileA: array<array<{{ tileT }}, BK>, BM>;
 var<workgroup> tileB: array<array<{{ tileT }}, BN>, BK>;
-@compute @workgroup_size(16, 16, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>
@@ -75,7 +56,7 @@ fn main(
   let nBase = wg.x * BN;
   let batch = wg.z;
   let xBatchBase = batch * K * N;
-  let li = lid.y * 16u + lid.x;
   var acc00: f32 = 0.0;
   var acc01: f32 = 0.0;
@@ -85,8 +66,8 @@ fn main(
   for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
     let kBase = kt * BK;
     // Cooperative load: 32x16 A(=W) tile + 16x32 B(=X[batch]) tile, 256 threads x 2 each.
-    for (var e: u32 = 0u; e < (BM * BK) / 256u; e = e + 1u) {
-      let idx = li + e * 256u;
       let ar = idx / BK;
       let ac = idx % BK;
       let am = mBase + ar;
@@ -112,10 +93,9 @@ fn main(
       let a1 = f32(tileA[lid.y * 2u + 1u][kk]);
       let b0 = f32(tileB[kk][lid.x * 2u]);
       let b1 = f32(tileB[kk][lid.x * 2u + 1u]);
-      acc00 = acc00 + a0 * b0;
-      acc01 = acc01 + a0 * b1;
-      acc10 = acc10 + a1 * b0;
-      acc11 = acc11 + a1 * b1;
     }
     workgroupBarrier();
   }

 // f16 operands remain packed in workgroup memory and widen when consumed. All
 // M, N, and K accesses are bounds-checked.
 {{ env.wgsl.resourceDeclarations }}
+{% set accumulatorPrefix = "acc" %}
+{% set tileT = "f32" %}
+{% set wgX = 16 %}
+{% set wgY = 16 %}
 const BK: u32 = 16u;
 const BM: u32 = 32u;
 const BN: u32 = 32u;
 // Store one output element with the configured bias/residual/activation
 // epilogue in the f32 accumulator domain.
 fn store_out(m: u32, n: u32, yBase: u32, raw: f32) {
   if (m >= params.M || n >= params.N) {
     return;
   }
+{% if convTransposeExactStore | default(false) %}
+{% set convTransposeKernelH = convTransposeKernelH | default(1) %}
+{% set convTransposeKernelW = convTransposeKernelW | default(1) %}
+{% set convTransposeInputW = convTransposeInputW | default(1) %}
+{% set convTransposeOutputH = convTransposeOutputH | default(1) %}
+{% set convTransposeOutputW = convTransposeOutputW | default(1) %}
+  // Nonoverlapping transposed-convolution taps map each matrix element to
+  // exactly one output. Keep the old col2im bias-plus-dot rounding order.
+  let kernelArea = {{ convTransposeKernelH }}u * {{ convTransposeKernelW }}u;
+  let oc = m / kernelArea;
+  let tap = m % kernelArea;
+  let oh = (n / {{ convTransposeInputW }}u) * {{ convTransposeKernelH }}u + tap / {{ convTransposeKernelW }}u;
+  let ow = (n % {{ convTransposeInputW }}u) * {{ convTransposeKernelW }}u + tap % {{ convTransposeKernelW }}u;
+  let outputIndex = yBase + (oc * {{ convTransposeOutputH }}u + oh) * {{ convTransposeOutputW }}u + ow;
+  y[outputIndex] = 0.0 + raw;
+{% else %}
   var v = raw;
   y[yBase + m * params.N + n] = {{ T }}(v);
+{% endif %}
 }
 var<workgroup> tileA: array<array<{{ tileT }}, BK>, BM>;
 var<workgroup> tileB: array<array<{{ tileT }}, BN>, BK>;
+@compute @workgroup_size({{ wgX }}, {{ wgY }}, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>
   let nBase = wg.x * BN;
   let batch = wg.z;
   let xBatchBase = batch * K * N;
+  let li = lid.y * {{ wgX }}u + lid.x;
   var acc00: f32 = 0.0;
   var acc01: f32 = 0.0;
   for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
     let kBase = kt * BK;
     // Cooperative load: 32x16 A(=W) tile + 16x32 B(=X[batch]) tile, 256 threads x 2 each.
+    for (var e: u32 = 0u; e < (BM * BK) / {{ wgX * wgY }}u; e = e + 1u) {
+      let idx = li + e * {{ wgX * wgY }}u;
       let ar = idx / BK;
       let ac = idx % BK;
       let am = mBase + ar;
       let a1 = f32(tileA[lid.y * 2u + 1u][kk]);
       let b0 = f32(tileB[kk][lid.x * 2u]);
       let b1 = f32(tileB[kk][lid.x * 2u + 1u]);
+{% for i in range(2) %}{% for j in range(2) %}
+      {{ accumulatorPrefix }}{{ i }}{{ j }} = {{ accumulatorPrefix }}{{ i }}{{ j }} + a{{ i }} * b{{ j }};
+{% endfor %}{% endfor %}
     }
     workgroupBarrier();
   }

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

@@ -11,19 +11,15 @@
 // 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 implicit3d = implicit and (convSpatialDims is defined and convSpatialDims == 3) %}
-{% set splitKValue = splitK if splitK is defined else 1 %}
-{% set splitKPartial = splitKValue > 1 %}
 {% set kLoopVar = "K_LOOP" if padded else "K" %}
 {% set nColsVar = "N_COLS" if padded else "N" %}
-{% 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) %}
@@ -42,7 +38,7 @@
 {% set convPadFront = convPadFront | default(0) %}
 {% set convInD = convInD | default(0) %}
 {% set convOutH = convOutH | default(0) %}
-{% set nPadded = nPadded | default(0) %}
 enable subgroups;
 {% if pinSubgroupSize32 %}
 enable subgroup_size_control;
@@ -50,36 +46,7 @@ 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" %}
 {% set useDirectMatrixStore = directMatrixStore is defined and directMatrixStore %}
@@ -136,31 +103,26 @@ const N_COLS: u32 = {{ nPadded }}u;
 {% endif %}
 const B_BATCH_STRIDE: u32 = {{ kLoopVar }} * {{ nColsVar }};
 {% endif %}
-{% if not splitKPartial and not polyphase and not fusedNarrowProjection %}
 const C_BATCH_STRIDE: u32 = M * N;
-{% endif %}
 const TILE_COLS: u32 = {{ tileColsValue }}u;
 const TILE_ROWS: u32 = {{ tileRowsValue }}u;
 const TILE_K: u32 = 32u;
 const SUB_COLS: u32 = {{ subColsValue }}u;
 const SUB_ROWS: u32 = {{ subRowsValue }}u;
-{% 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;
   for (var col_offset = 0u; col_offset < 8u; col_offset = col_offset + 1u) {
     let k = k_idx + col + col_offset;
     if (a_global < M{% if padded %} && k < K{% endif %}) {
-{% set A_PHASE = "phase_base + " if polyphase else "" %}
 {% if operandScalar == "f16" %}
       tile_A[row * TILE_K + col + col_offset] = f16(w[{{ A_PHASE }}a_global * K + k]);
 {% else %}
@@ -309,59 +271,12 @@ fn loadSHMB(b_base: u32, tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
   }
 {% 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 %}
 {% set biasAdd = " + bv" if hasBias else "" %}
-{% if hasActivation or hasZ %}
-// Fused epilogue: Y = activation(conv + bias + Z), applied at the output store.
-// `raw` is the f32 convolution-plus-bias accumulator; Z shares Y's NCHW layout.
-fn epi(raw: f32{% if hasZ %}, yIndex: u32{% endif %}) -> {{ T }} {
-  var r = raw;
-{% if hasZ %}
-  r = r + f32(zResidual[yIndex]);
-{% endif %}
-{% if hasActivation %}
-  r = fused_act(r);
-{% endif %}
-  return {{ T }}(r);
-}
-{% endif %}
 {% macro store_val(valExpr, idxExpr) %}
-{% if hasActivation or hasZ %}epi({{ valExpr }}{% if hasZ %}, {{ idxExpr }}{% endif %})
-{%- else %}{{ T }}({{ valExpr }})
-{%- endif %}
-{% endmacro %}
 {% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
 fn storeOutput(offset: u32, {% if hasBias or polyphase %}row_base: u32, {% endif %}row: u32, col: u32, src_slot: u32, row_limit: i32{% if padded %}, col_anchor: u32{% endif %}) {
@@ -408,9 +323,7 @@ fn main(
 ) {
   let batch = workgroup_id.z;
   let b_base = batch * B_BATCH_STRIDE;
-{% if not fusedNarrowProjection %}
   let c_base = batch * C_BATCH_STRIDE;
-{% endif %}
   let a_global_base = workgroup_id.y * TILE_ROWS;
   let b_global_base = workgroup_id.x * TILE_COLS;
   let subtile_id = local_idx / sg_size;
@@ -429,22 +342,20 @@ fn main(
   var matC13: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
   for (var kidx = 0u; kidx < {{ kLoopVar }}; kidx = kidx + TILE_K) {
-{% 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);
@@ -456,62 +367,10 @@ fn main(
       matC12 = subgroupMatrixMultiplyAccumulate(matA1, matB2, matC12);
       matC13 = subgroupMatrixMultiplyAccumulate(matA1, matB3, matC13);
     }
-{% if not directMatrixInputs %}
     workgroupBarrier();
-{% endif %}
   }
-{% if fusedNarrowProjection %}
-  // The Conv producer owns exactly one 32-row tile. Its input tile is dead after
-  // the final K iteration, so reuse those 32x64 f32 cells as the only
-  // cross-subgroup handoff into the narrow projection. All collective stores
-  // remain subgroup-uniform, and every fragment has a disjoint destination.
-  let fused_tile_offset = base_A * TILE_COLS + base_B;
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 0u * TILE_COLS +  0u, matC00, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 0u * TILE_COLS +  8u, matC01, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 0u * TILE_COLS + 16u, matC02, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 0u * TILE_COLS + 24u, matC03, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 8u * TILE_COLS +  0u, matC10, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 8u * TILE_COLS +  8u, matC11, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 8u * TILE_COLS + 16u, matC12, TILE_COLS);
-  subgroupMatrixStore<row_major>(&tile_B, fused_tile_offset + 8u * TILE_COLS + 24u, matC13, TILE_COLS);
-  workgroupBarrier();
-  // One invocation owns one spatial column and walks the producer channels in
-  // increasing order to preserve the scalar projection's accumulation order.
-  // The inactive half of the workgroup has no remaining collective operation
-  // to reach.
-  let fused_global_col = b_global_base + local_idx;
-  if (local_idx < TILE_COLS && fused_global_col < N) {
-{% for oc in range(projectionOutChannels) %}
-    var projected_acc{{ oc }} = f32(0.0);
-{% endfor %}
-    for (var channel = 0u; channel < M; channel = channel + 1u) {
-      let producer_value = tile_B[channel * TILE_COLS + local_idx]{% if hasBias %} + f32(bias[channel]){% endif %};
-{% 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 useDirectMatrixStore %}
 {% if padded %}
   // Padded N only needs guarded publication for the final partial-width tile.
   // Full tiles can use the same direct collective stores as the aligned path;
@@ -543,11 +402,7 @@ fn main(
 {% if hasBias %}
     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 %}
   }
 {% endif %}
 {% if padded %}

 // 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 = false %}
+{% set polyphase = false %}
 {% set implicit = implicitIm2col is defined and implicitIm2col %}
 {% set implicit3d = implicit and (convSpatialDims is defined and convSpatialDims == 3) %}
+{% set splitKPartial = false %}
 {% set kLoopVar = "K_LOOP" if padded else "K" %}
 {% set nColsVar = "N_COLS" if padded else "N" %}
+{% set hasZ = false %}
+{% set hasActivation = false %}
 {% set convKernelH = convKernelH | default(0) %}
 {% set convKernelW = convKernelW | default(0) %}
 {% set convStrideH = convStrideH | default(0) %}
 {% set convPadFront = convPadFront | default(0) %}
 {% set convInD = convInD | default(0) %}
 {% set convOutH = convOutH | default(0) %}
+{% set nPadded = 0 %}
 enable subgroups;
 {% if pinSubgroupSize32 %}
 enable subgroup_size_control;
 enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 {% set operandScalar = fScalar %}
 {% set accScalar = "f32" %}
 {% set useDirectMatrixStore = directMatrixStore is defined and directMatrixStore %}
 {% endif %}
 const B_BATCH_STRIDE: u32 = {{ kLoopVar }} * {{ nColsVar }};
 {% endif %}
 const C_BATCH_STRIDE: u32 = M * N;
 const TILE_COLS: u32 = {{ tileColsValue }}u;
 const TILE_ROWS: u32 = {{ tileRowsValue }}u;
 const TILE_K: u32 = 32u;
 const SUB_COLS: u32 = {{ subColsValue }}u;
 const SUB_ROWS: u32 = {{ subRowsValue }}u;
 var<workgroup> tile_A: array<{{ operandScalar }}, {{ tileRowsValue }} * 32>;
 var<workgroup> tile_B: array<{{ operandScalar }}, {{ tileColsValue }} * 32>;
 {% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
 var<workgroup> scratch: array<array<array<{{ accScalar }}, 64>, 4>, {{ (workgroupThreadsValue / 32)|int }}>;
 {% endif %}
 fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
   let a_global = tile_base + row;
   let col = c_idx * 8u;
   for (var col_offset = 0u; col_offset < 8u; col_offset = col_offset + 1u) {
     let k = k_idx + col + col_offset;
     if (a_global < M{% if padded %} && k < K{% endif %}) {
+{% set A_PHASE = "" %}
 {% if operandScalar == "f16" %}
       tile_A[row * TILE_K + col + col_offset] = f16(w[{{ A_PHASE }}a_global * K + k]);
 {% else %}
   }
 {% endif %}
 }
+{% set hasActivation = false %}
 {% set biasAdd = " + bv" if hasBias else "" %}
 {% macro store_val(valExpr, idxExpr) %}
+{% if hasActivation or hasZ %}
+epi({{ valExpr }}){% else %}
+{{ T }}({{ valExpr }}){% endif %}{% endmacro %}
 {% if (not useDirectMatrixStore or padded) and not splitKPartial and not fusedNarrowProjection %}
 fn storeOutput(offset: u32, {% if hasBias or polyphase %}row_base: u32, {% endif %}row: u32, col: u32, src_slot: u32, row_limit: i32{% if padded %}, col_anchor: u32{% endif %}) {
 ) {
   let batch = workgroup_id.z;
   let b_base = batch * B_BATCH_STRIDE;
   let c_base = batch * C_BATCH_STRIDE;
   let a_global_base = workgroup_id.y * TILE_ROWS;
   let b_global_base = workgroup_id.x * TILE_COLS;
   let subtile_id = local_idx / sg_size;
   var matC13: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
   for (var kidx = 0u; kidx < {{ kLoopVar }}; kidx = kidx + TILE_K) {
     loadSHMA(a_global_base, kidx, local_idx / 4u, local_idx % 4u);
     loadSHMB(b_base, b_global_base, kidx, local_idx / {{ bKChunks }}u, local_idx % {{ bKChunks }}u);
     workgroupBarrier();
     for (var step = 0u; step < TILE_K; step = step + 8u) {
+{% set directInputs = false %}
+      let matrix_a_offset = subtile_idy * SUB_ROWS * TILE_K + step;
 {% 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 = subtile_idx * SUB_COLS * TILE_K + step;
 {% for c in range(4) %}
+      var matB{{ c }}: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>, col_major>(&{{ "xm" if directInputs else "tile_B" }}, matrix_b_offset{% if c > 0 %} + {{ c * 8 }}u * TILE_K{% endif %}, {{ "N" if directInputs else "TILE_K" }});
 {% endfor %}
       matC00 = subgroupMatrixMultiplyAccumulate(matA0, matB0, matC00);
       matC12 = subgroupMatrixMultiplyAccumulate(matA1, matB2, matC12);
       matC13 = subgroupMatrixMultiplyAccumulate(matA1, matB3, matC13);
     }
     workgroupBarrier();
   }
+{% if useDirectMatrixStore %}
 {% if padded %}
   // Padded N only needs guarded publication for the final partial-width tile.
   // Full tiles can use the same direct collective stores as the aligned path;
 {% if hasBias %}
     raw += f32(bias[global_row]);
 {% endif %}
     y[y_index] = {{ T }}(raw);
   }
 {% endif %}
 {% if padded %}

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

@@ -1,36 +1,11 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 const COUNT: u32 = {{ count }}u;
@@ -43,7 +18,7 @@ const OUTPUT_SPATIAL: u32 = {{ outputSpatial }}u;
 // coordinate. The optional ConvTranspose bias is applied after that reduction.
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", "COUNT", note="device-axis") }}
 {% if hasBias %}
   let output_channel = (index / OUTPUT_SPATIAL) % OUT_CHANNELS;
   y[index] = bias[output_channel];

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 const COUNT: u32 = {{ count }}u;
 // coordinate. The optional ConvTranspose bias is applied after that reduction.
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index", "COUNT") }}
 {% if hasBias %}
   let output_channel = (index / OUTPUT_SPATIAL) % OUT_CHANNELS;
   y[index] = bias[output_channel];

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

@@ -1,51 +1,22 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
-if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
-{%- endmacro -%}
-{% if usesF16 %}
 // Preserve f32 accumulation for the deep channel-by-kernel sum: widen each
 // f16 product and narrow only once at the output store.
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="device-axis") }}
   let ow = index % params.outW;
   var t = index / params.outW;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }{% endmacro %}
 // Preserve f32 accumulation for the deep channel-by-kernel sum: widen each
 // f16 product and narrow only once at the output store.
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
   let ow = index % params.outW;
   var t = index / params.outW;

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

@@ -1,42 +1,16 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
-if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
-{%- endmacro -%}
 // ConvTranspose col2im scatter (output-stationary) for the im2col-GEMM path.
 //
 // The GEMM pass produced cols[m = (oc*KH+kh)*KW+kw, n = ih*IW+iw] = sum_ic
@@ -45,14 +19,10 @@ if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
 // KH*KW taps whose inverse stride-and-padding coordinates land on it, reading the
 // precomputed cols without repeating the IC contraction.
 // One thread per output element, f32 accumulation, deterministic (no atomics).
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="") }}
   let ow = index % params.outW;
   var t = index / params.outW;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }{% endmacro %}
 // ConvTranspose col2im scatter (output-stationary) for the im2col-GEMM path.
 //
 // The GEMM pass produced cols[m = (oc*KH+kh)*KW+kw, n = ih*IW+iw] = sum_ic
 // KH*KW taps whose inverse stride-and-padding coordinates land on it, reading the
 // precomputed cols without repeating the IC contraction.
 // One thread per output element, f32 accumulation, deterministic (no atomics).
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
   let ow = index % params.outW;
   var t = index / params.outW;

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

@@ -1,3 +1,9 @@
 {{ env.wgsl.resourceDeclarations }}
 // With stride > 1 only one kernel phase contributes to a given output pixel.
@@ -24,8 +30,7 @@ const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
-  if (index >= COUNT) { return; }
   let ow = index % OUT_W;
   var t = index / OUT_W;
   let oh = t % OUT_H;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) { return; }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 // With stride > 1 only one kernel phase contributes to a given output pixel.
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("WG", "index", "COUNT", guardInline=true) }}
   let ow = index % OUT_W;
   var t = index / OUT_W;
   let oh = t % OUT_H;

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

@@ -1,52 +1,25 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
 {% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
-if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
-{%- endmacro -%}
-{% if usesF16 %}
 // Preserve f32 accumulation for the deep channel-by-kernel sum: widen each
 // f16 product and narrow only once at the output store.
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% if depthwisePointwiseVec4 %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("vector_index", "", note="") }}
   let base = vector_index * 4u;
   if (base >= params.count) {
     return;
@@ -68,7 +41,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% else %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="device-axis") }}
   let ow = index % params.outW;
   var t = index / params.outW;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endif %}{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }{% endmacro %}
 // Preserve f32 accumulation for the deep channel-by-kernel sum: widen each
 // f16 product and narrow only once at the output store.
 {{ env.wgsl.resourceDeclarations }}
 {% if depthwisePointwiseVec4 %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "vector_index", "") }}
   let base = vector_index * 4u;
   if (base >= params.count) {
     return;
 {% else %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
   let ow = index % params.outW;
   var t = index / params.outW;

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

@@ -11,8 +11,6 @@
 // kernel gathers the four shifted input samples while cooperatively loading the
 // GEMM B tile, then stores each register tile directly to its interleaved output
 // phase. Boundary samples are represented by zero in shared memory.
-enable f16;
 {{ env.wgsl.resourceDeclarations }}
 const IN_C: u32 = {{ inChannels }}u;
@@ -42,18 +40,58 @@ const WG_SIZE: u32 = WG_EDGE * WG_EDGE;
 var<workgroup> tileA: array<array<f16, BK>, M_TILE>;
 var<workgroup> tileB: array<array<f16, N_TILE>, BK>;
 fn store_output(batch: u32, phase: u32, oc: u32, n: u32, value: f32) {
   if (oc >= OUT_C || n >= N) {
     return;
   }
   let phaseH = phase / PHASE_EDGE;
   let phaseW = phase % PHASE_EDGE;
   let oh = (n / IN_W) * PHASE_EDGE + phaseH;
   let ow = (n % IN_W) * PHASE_EDGE + phaseW;
 {% if hasBias %}
-  let result = value + f32(bias[oc]);
 {% else %}
-  let result = value;
 {% endif %}
   y[((batch * OUT_C + oc) * OUT_H + oh) * OUT_W + ow] = f16(result);
 }
@@ -69,12 +107,7 @@ fn main(
     return;
   }
-  let phaseH = phase / PHASE_EDGE;
-  let phaseW = phase % PHASE_EDGE;
-  let kernelPhaseH = (phaseH + PAD_H) % PHASE_EDGE;
-  let kernelPhaseW = (phaseW + PAD_W) % PHASE_EDGE;
-  let inputShiftH = (phaseH + PAD_H - kernelPhaseH) / PHASE_EDGE;
-  let inputShiftW = (phaseW + PAD_W - kernelPhaseW) / PHASE_EDGE;
   let mBase = wg.y * M_TILE;
   let nBase = wg.x * N_TILE;

 // kernel gathers the four shifted input samples while cooperatively loading the
 // GEMM B tile, then stores each register tile directly to its interleaved output
 // phase. Boundary samples are represented by zero in shared memory.
 {{ env.wgsl.resourceDeclarations }}
 const IN_C: u32 = {{ inChannels }}u;
 var<workgroup> tileA: array<array<f16, BK>, M_TILE>;
 var<workgroup> tileB: array<array<f16, N_TILE>, BK>;
+{% macro phase_geometry() %}
+  let phaseH = phase / PHASE_EDGE;
+  let phaseW = phase % PHASE_EDGE;
+  let kernelPhaseH = (phaseH + PAD_H) % PHASE_EDGE;
+  let kernelPhaseW = (phaseW + PAD_W) % PHASE_EDGE;
+  let inputShiftH = (phaseH + PAD_H - kernelPhaseH) / PHASE_EDGE;
+  let inputShiftW = (phaseW + PAD_W - kernelPhaseW) / PHASE_EDGE;{% endmacro %}
+// A synthetic zero input may multiply a NaN/Inf weight in the tiled dot.
+// Reconstruct an affected perimeter output from its reachable taps only.
+fn boundary_sum(batch: u32, phase: u32, oc: u32, n: u32) -> f32 {
+{{ phase_geometry() }}
+  let row = i32(n / IN_W) + i32(inputShiftH);
+  let col = i32(n % IN_W) + i32(inputShiftW);
+  var sum = 0.0;
+  for (var ic = 0u; ic < IN_C; ic += 1u) {
+    for (var tap = 0u; tap < TAPS_PER_PHASE; tap += 1u) {
+      let tapH = tap / PHASE_EDGE;
+      let tapW = tap % PHASE_EDGE;
+      let ih = row - i32(tapH);
+      let iw = col - i32(tapW);
+      if (ih >= 0 && ih < i32(IN_H) && iw >= 0 && iw < i32(IN_W)) {
+        let kh = kernelPhaseH + tapH * PHASE_EDGE;
+        let kw = kernelPhaseW + tapW * PHASE_EDGE;
+        let xIndex = ((batch * IN_C + ic) * IN_H + u32(ih)) * IN_W + u32(iw);
+        let wIndex = ((ic * OUT_C + oc) * (2u * PHASE_EDGE) + kh) * (2u * PHASE_EDGE) + kw;
+        sum = fma(f32(x[xIndex]), f32(w[wIndex]), sum);
+      }
+    }
+  }
+  return sum;
+}
 fn store_output(batch: u32, phase: u32, oc: u32, n: u32, value: f32) {
   if (oc >= OUT_C || n >= N) {
     return;
   }
   let phaseH = phase / PHASE_EDGE;
   let phaseW = phase % PHASE_EDGE;
+  var sum = value;
+  if ((bitcast<u32>(value) & 0x7fffffffu) > 0x7f800000u) {
+    let row = n / IN_W;
+    let col = n % IN_W;
+    if (row == 0u || row + 1u == IN_H || col == 0u || col + 1u == IN_W) {
+      sum = boundary_sum(batch, phase, oc, n);
+    }
+  }
   let oh = (n / IN_W) * PHASE_EDGE + phaseH;
   let ow = (n % IN_W) * PHASE_EDGE + phaseW;
 {% if hasBias %}
+  let result = sum + f32(bias[oc]);
 {% else %}
+  let result = sum;
 {% endif %}
   y[((batch * OUT_C + oc) * OUT_H + oh) * OUT_W + ow] = f16(result);
 }
     return;
   }
+{{ phase_geometry() }}
   let mBase = wg.y * M_TILE;
   let nBase = wg.x * N_TILE;

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

@@ -1,50 +1,21 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 // ConvTranspose weight reorder for the im2col-GEMM path.
 //
 // The ONNX ConvTranspose weight is laid out [IC, OC, KH, KW]. The GEMM that
 // produces the cols matrix (cols[OC*KH*KW, IH*IW] = Wt @ X) needs its A operand
 // as a row-major [M = OC*KH*KW, K = IC] matrix, i.e. the IC dimension contracted
 // to the inner (column) axis. This pass gathers W into that transposed layout.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="") }}
   // Wt[m, ic] with m = (oc * KH + kh) * KW + kw, contracted dim ic.
   let ic = index % params.inChannels;
   var m = index / params.inChannels;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 // ConvTranspose weight reorder for the im2col-GEMM path.
 //
 // The ONNX ConvTranspose weight is laid out [IC, OC, KH, KW]. The GEMM that
 // produces the cols matrix (cols[OC*KH*KW, IH*IW] = Wt @ X) needs its A operand
 // as a row-major [M = OC*KH*KW, K = IC] matrix, i.e. the IC dimension contracted
 // to the inner (column) axis. This pass gathers W into that transposed layout.
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
   // Wt[m, ic] with m = (oc * KH + kh) * KW + kw, contracted dim ic.
   let ic = index % params.inChannels;
   var m = index / params.inChannels;

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

@@ -1,42 +1,16 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
-if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
-{%- endmacro -%}
 // Output-stationary rank-5 col2im gather for ConvTranspose-as-GEMM.
 //
 // GEMM has already contracted input channels into
@@ -47,7 +21,7 @@ if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="device-axis") }}
   let ow = index % params.outW;
   var t = index / params.outW;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }{% endmacro %}
 // Output-stationary rank-5 col2im gather for ConvTranspose-as-GEMM.
 //
 // GEMM has already contracted input channels into
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
   let ow = index % params.outW;
   var t = index / params.outW;

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

@@ -1,46 +1,20 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
-if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }
-{%- endmacro -%}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="device-axis") }}
   let ow = index % params.outW;
   var t = index / params.outW;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 {% macro tap(raw, coord, out, kernel, axis) %}
 let {{ raw }} = i32({{ out }}) + params.pad{{ axis }} - i32({{ kernel }} * params.dilation{{ axis }});
 if ({{ raw }} < 0 || {{ raw }} % i32(params.stride{{ axis }}) != 0) { continue; }
 let {{ coord }} = {{ raw }} / i32(params.stride{{ axis }});
+if ({{ coord }} < 0 || {{ coord }} >= i32(params.in{{ axis }})) { continue; }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
   let ow = index % params.outW;
   var t = index / params.outW;

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

@@ -1,36 +1,11 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
-  }
-{%- endif %}
-{% endmacro %}
 {% set FORWARD_FLIP = forwardFlip is defined and forwardFlip %}
 {% if FORWARD_FLIP %}
 // Reorder ONNX ConvTranspose weights from [IC, OC, KD, KH, KW] into the
@@ -48,7 +23,7 @@
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-{{ flat_index_2d("index", note="device-axis") }}
 {% if FORWARD_FLIP %}
   var m = index;

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
   if ({{ name }} >= {{ bound }}) {
     return;
+  }{% endmacro %}
 {% set FORWARD_FLIP = forwardFlip is defined and forwardFlip %}
 {% if FORWARD_FLIP %}
 // Reorder ONNX ConvTranspose weights from [IC, OC, KD, KH, KW] into the
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index") }}
 {% if FORWARD_FLIP %}
   var m = index;

build/webgpu/manifest.json CHANGED Viewed

@@ -50,7 +50,6 @@
     "stridesLengthOk": "not has(attrs, \"strides\") or (attrs.strides | length) == spatialRank",
     "dilationsLengthOk": "not has(attrs, \"dilations\") or (attrs.dilations | length) == spatialRank",
     "padsLengthOk": "not has(attrs, \"pads\") or (attrs.pads | length) == 2 * spatialRank",
-    "outputPaddingLengthOk": "not has(attrs, \"output_padding\") or (attrs.output_padding | length) == spatialRank",
     "kernelD": "attrs.kernel_shape[0] if kernelShapeLengthOk and has(attrs, \"kernel_shape\") and spatialRank == 3 else 1",
     "kernelH": "attrs.kernel_shape[spatialRank - 2] if kernelShapeLengthOk and has(attrs, \"kernel_shape\") and spatialRank >= 2 else 1",
     "kernelW": "attrs.kernel_shape[spatialRank - 1] if kernelShapeLengthOk and has(attrs, \"kernel_shape\") and spatialRank >= 1 else 1",
@@ -66,6 +65,7 @@
     "padBack": "attrs.pads[spatialRank] if padsLengthOk and has(attrs, \"pads\") and spatialRank == 3 else 0",
     "padBottom": "attrs.pads[2 * spatialRank - 2] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 2 else 0",
     "padRight": "attrs.pads[2 * spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
     "outputPaddingD": "attrs.output_padding[0] if outputPaddingLengthOk and has(attrs, \"output_padding\") and spatialRank == 3 else 0",
     "outputPaddingH": "attrs.output_padding[spatialRank - 2] if outputPaddingLengthOk and has(attrs, \"output_padding\") and spatialRank >= 2 else 0",
     "outputPaddingW": "attrs.output_padding[spatialRank - 1] if outputPaddingLengthOk and has(attrs, \"output_padding\") and spatialRank >= 1 else 0",
@@ -104,6 +104,7 @@
     "rank5DenseContract": "spatialAttributesOk and ranks.x == 5 and ranks.w == 5 and ranks.y == 5 and attrs.group == 1 and batchChannelContract and outputDispatchFits and dim(shapes.y, 1) == dim(shapes.w, 1)",
     "noBiasContract": "not present.bias",
     "biasContract": "present.bias and ranks.bias == 1 and dim(shapes.bias, 0) == dim(shapes.y, 1)",
     "f32Type": "dtypes.T == \"f32\"",
     "f16Type": "device.features.has(\"shader-f16\") and dtypes.T == \"f16\"",
     "gemmM": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
@@ -150,18 +151,31 @@
     "groupedBlockOc": "16 if groupedOutCPerGroup % 16 == 0 else (8 if groupedOutCPerGroup % 8 == 0 else (4 if groupedOutCPerGroup % 4 == 0 else (2 if groupedOutCPerGroup % 2 == 0 else 1)))",
     "groupedWeightBytes": "(dim(shapes.x, 1) / attrs.group) * groupedOutCPerGroup * dim(shapes.w, 2) * dim(shapes.w, 3) * 4",
     "groupedPlanes": "dim(shapes.x, 0) * attrs.group * (groupedOutCPerGroup / groupedBlockOc)",
-    "groupedStridePhaseBlockedFits": "groupedWeightBytes <= device.limits.maxComputeWorkgroupStorageSize and groupedPlanes <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), tunables.WORKGROUP_SIZE) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"
   },
   "bindings": {
-    "y": { "buffer": "storage", "elementType": "$T" },
-    "bias": { "buffer": "read-only-storage", "elementType": "$T" },
-    "x": { "buffer": "read-only-storage", "elementType": "$T" },
-    "w": { "buffer": "read-only-storage", "elementType": "$phaseWeightElement" },
-    "y_2": { "name": "y", "buffer": "storage", "elementType": "$phaseOutputElement" },
-    "wsrc": { "arg": "w", "buffer": "read-only-storage", "elementType": "$T" },
     "wt": { "buffer": "storage", "elementType": "$T" },
     "params": {
-      "buffer": "uniform",
       "struct": [
         { "name": "count", "type": "u32", "value": "numel(shapes.w)" },
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
@@ -171,13 +185,12 @@
         { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 4)" }
       ]
     },
-    "w_2": { "scratch": "wt", "name": "w", "buffer": "read-only-storage", "elementType": "$T" },
-    "xm": { "arg": "x", "buffer": "read-only-storage", "elementType": "$T" },
-    "y_3": { "scratch": "cols", "name": "y", "buffer": "storage", "elementType": "$T" },
     "cols": { "buffer": "read-only-storage", "elementType": "$T" },
-    "params_2": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
         { "name": "kernelD", "type": "u32", "value": "dim(shapes.w, 2)" },
@@ -201,9 +214,8 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_4": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "count", "type": "u32", "value": "numel(shapes.w)" },
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
@@ -212,9 +224,8 @@
         { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 3)" }
       ]
     },
-    "params_5": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
         { "name": "kernelH", "type": "u32", "value": "dim(shapes.w, 2)" },
@@ -232,19 +243,17 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_6": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" },
         { "name": "K", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "N", "type": "u32", "value": "dim(shapes.x, 2) * dim(shapes.x, 3)" }
       ]
     },
-    "w_4": { "name": "w", "buffer": "read-only-storage", "elementType": "$T" },
-    "params_7": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inW", "type": "u32", "value": "dim(shapes.x, 2)" },
@@ -257,9 +266,8 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_8": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
@@ -274,9 +282,8 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_9": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inH", "type": "u32", "value": "dim(shapes.x, 2)" },
@@ -295,9 +302,8 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_10": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inD", "type": "u32", "value": "dim(shapes.x, 2)" },
@@ -322,9 +328,8 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_11": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
@@ -345,9 +350,8 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "params_12": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "outChannels", "type": "u32", "value": "dim(shapes.y, 1)" },
         { "name": "outH", "type": "u32", "value": "dim(shapes.y, 2)" },
@@ -361,15 +365,12 @@
       "id": "empty_input_reduction",
       "priority": 100,
       "when": ["noBiasContract", "emptyInputReductionContract"],
-      "derive": { "hasBias": false, "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.EmptyInput",
           "shader": "conv-transpose-empty-input.wgsl.jinja",
           "derive": {
-            "count": "numel(shapes.y)",
-            "outChannels": "dim(shapes.y, 1)",
             "outputSpatial": "dim(shapes.y, 2) if ranks.y == 3 else (dim(shapes.y, 2) * dim(shapes.y, 3) if ranks.y == 4 else dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4))"
           },
           "bindings": ["y"],
@@ -385,15 +386,12 @@
       "id": "empty_input_reduction_bias",
       "priority": 100,
       "when": ["biasContract", "emptyInputReductionContract"],
-      "derive": { "hasBias": true, "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.EmptyInputBias",
           "shader": "conv-transpose-empty-input.wgsl.jinja",
           "derive": {
-            "count": "numel(shapes.y)",
-            "outChannels": "dim(shapes.y, 1)",
             "outputSpatial": "dim(shapes.y, 2) if ranks.y == 3 else (dim(shapes.y, 2) * dim(shapes.y, 3) if ranks.y == 4 else dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4))"
           },
           "bindings": ["bias", "y"],
@@ -410,7 +408,6 @@
       "priority": 40,
       "when": ["noBiasContract", "phase1dCovered"],
       "derive": {
-        "hasBias": false,
         "phaseWeightElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "phaseOutputElement": "\"vec2<\" ~ dtypes.T ~ \">\""
       },
@@ -420,9 +417,7 @@
           "name": "ConvTranspose1d.PhaseTiled",
           "shader": "conv-transpose1d-phase-tiled.wgsl.jinja",
           "derive": {
-            "inChannels": "dim(shapes.x, 1)",
             "inWidth": "dim(shapes.x, 2)",
-            "outChannels": "dim(shapes.y, 1)",
             "positionTile": "tunables.PHASE1D_POSITION_TILE",
             "outputTile": "phase1dOutputTile",
             "outputsPerThread": "phase1dOutputsPerThread",
@@ -431,7 +426,7 @@
             "xTileElements": "tunables.PHASE1D_K_TILE * (tunables.PHASE1D_POSITION_TILE + 2)",
             "weightTileElements": "tunables.PHASE1D_K_TILE * phase1dOutputTile"
           },
-          "bindings": ["x", "w", "y_2"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2), tunables.PHASE1D_POSITION_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), phase1dOutputTile)",
@@ -445,7 +440,6 @@
       "priority": 45,
       "when": ["biasContract", "phase1dCovered"],
       "derive": {
-        "hasBias": true,
         "phaseWeightElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "phaseOutputElement": "\"vec2<\" ~ dtypes.T ~ \">\""
       },
@@ -455,9 +449,7 @@
           "name": "ConvTranspose1d.PhaseTiledBias",
           "shader": "conv-transpose1d-phase-tiled.wgsl.jinja",
           "derive": {
-            "inChannels": "dim(shapes.x, 1)",
             "inWidth": "dim(shapes.x, 2)",
-            "outChannels": "dim(shapes.y, 1)",
             "positionTile": "tunables.PHASE1D_POSITION_TILE",
             "outputTile": "phase1dOutputTile",
             "outputsPerThread": "phase1dOutputsPerThread",
@@ -466,7 +458,7 @@
             "xTileElements": "tunables.PHASE1D_K_TILE * (tunables.PHASE1D_POSITION_TILE + 2)",
             "weightTileElements": "tunables.PHASE1D_K_TILE * phase1dOutputTile"
           },
-          "bindings": ["x", "w", "bias", "y_2"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2), tunables.PHASE1D_POSITION_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), phase1dOutputTile)",
@@ -484,9 +476,7 @@
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
-        "usesF16": false,
         "fScalar": "\"f32\"",
-        "hasBias": false,
         "directMatrixStore": "ctImplicitDirectStore",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "subgroupRows": "tunables.SGMAT_SUBGROUP_ROWS",
@@ -536,7 +526,7 @@
           "id": "gemm",
           "name": "ConvTranspose3dImplicitSubgroupMatrix",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
-          "bindings": ["w_2", "xm", "y"],
           "dispatch": {
             "x": "ceilDiv(ctImplicitOutVolume, tunables.SGMAT_TILE_COLS)",
             "y": "ctImplicitYWorkgroups",
@@ -554,9 +544,7 @@
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
-        "usesF16": false,
         "fScalar": "\"f32\"",
-        "hasBias": true,
         "directMatrixStore": "ctImplicitDirectStore",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "subgroupRows": "tunables.SGMAT_SUBGROUP_ROWS",
@@ -606,7 +594,7 @@
           "id": "gemm",
           "name": "ConvTranspose3dImplicitSubgroupMatrixBias",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
-          "bindings": ["w_2", "xm", "bias", "y"],
           "dispatch": {
             "x": "ceilDiv(ctImplicitOutVolume, tunables.SGMAT_TILE_COLS)",
             "y": "ctImplicitYWorkgroups",
@@ -624,9 +612,7 @@
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
-        "usesF16": false,
         "fScalar": "\"f32\"",
-        "hasBias": false,
         "directMatrixStore": "gemm3dM % tunables.SGMAT_TILE_ROWS == 0",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "tileCols": "tunables.SGMAT_TILE_COLS",
@@ -658,7 +644,7 @@
           "name": "ConvTranspose3dColsGemm",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3"],
           "dispatch": {
             "x": "ceilDiv(gemm3dN, tunables.SGMAT_TILE_COLS)",
             "y": "ceilDiv(gemm3dM, tunables.SGMAT_TILE_ROWS)",
@@ -669,7 +655,7 @@
           "id": "col2im",
           "name": "ConvTranspose3dCol2Im",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -687,9 +673,7 @@
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
-        "usesF16": false,
         "fScalar": "\"f32\"",
-        "hasBias": true,
         "directMatrixStore": "gemm3dM % tunables.SGMAT_TILE_ROWS == 0",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "tileCols": "tunables.SGMAT_TILE_COLS",
@@ -721,7 +705,7 @@
           "name": "ConvTranspose3dColsGemm",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3"],
           "dispatch": {
             "x": "ceilDiv(gemm3dN, tunables.SGMAT_TILE_COLS)",
             "y": "ceilDiv(gemm3dM, tunables.SGMAT_TILE_ROWS)",
@@ -732,7 +716,7 @@
           "id": "col2im",
           "name": "ConvTranspose3dCol2ImBias",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -745,7 +729,7 @@
       "id": "ncdhw3d_gemm_tiled_f32",
       "priority": 30,
       "when": ["noBiasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits"],
-      "derive": { "usesF16": false, "hasBias": false, "M": "gemm3dM", "K": "dim(shapes.x, 1)", "N": "gemm3dN" },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
@@ -768,9 +752,9 @@
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
-            "w_2",
             "xm",
-            "y_3",
             {
               "name": "params",
               "struct": [
@@ -786,7 +770,7 @@
           "id": "col2im",
           "name": "ConvTranspose3dCol2Im",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -799,14 +783,7 @@
       "id": "ncdhw3d_gemm_tiled_f32_reg",
       "priority": 32,
       "when": ["noBiasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits", "ceilDiv(gemm3dM, 64) * ceilDiv(gemm3dN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
-      "derive": {
-        "usesF16": false,
-        "hasBias": false,
-        "gemmKTile": "32 if dim(shapes.x, 1) % 32 == 0 else 16",
-        "M": "gemm3dM",
-        "K": "dim(shapes.x, 1)",
-        "N": "gemm3dN"
-      },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
@@ -829,9 +806,9 @@
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
-            "w_2",
             "xm",
-            "y_3",
             {
               "name": "params",
               "struct": [
@@ -847,7 +824,7 @@
           "id": "col2im",
           "name": "ConvTranspose3dCol2Im",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -860,7 +837,7 @@
       "id": "ncdhw3d_gemm_tiled_f32_bias",
       "priority": 35,
       "when": ["biasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits"],
-      "derive": { "usesF16": false, "hasBias": true, "M": "gemm3dM", "K": "dim(shapes.x, 1)", "N": "gemm3dN" },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
@@ -883,9 +860,9 @@
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
-            "w_2",
             "xm",
-            "y_3",
             {
               "name": "params",
               "struct": [
@@ -901,7 +878,7 @@
           "id": "col2im",
           "name": "ConvTranspose3dCol2ImBias",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -914,14 +891,7 @@
       "id": "ncdhw3d_gemm_tiled_f32_bias_reg",
       "priority": 37,
       "when": ["biasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits", "ceilDiv(gemm3dM, 64) * ceilDiv(gemm3dN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
-      "derive": {
-        "usesF16": false,
-        "hasBias": true,
-        "gemmKTile": "32 if dim(shapes.x, 1) % 32 == 0 else 16",
-        "M": "gemm3dM",
-        "K": "dim(shapes.x, 1)",
-        "N": "gemm3dN"
-      },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
@@ -944,9 +914,9 @@
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
-            "w_2",
             "xm",
-            "y_3",
             {
               "name": "params",
               "struct": [
@@ -962,7 +932,7 @@
           "id": "col2im",
           "name": "ConvTranspose3dCol2ImBias",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -980,13 +950,10 @@
         "subgroupMatrixConfigs": [{ "componentType": "f16", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
-        "usesF16": true,
         "fScalar": "\"f16\"",
-        "hasBias": false,
         "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
         "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
       },
       "intermediates": [
         {
@@ -1005,7 +972,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1032,7 +999,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1050,13 +1017,10 @@
         "subgroupMatrixConfigs": [{ "componentType": "f16", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
-        "usesF16": true,
         "fScalar": "\"f16\"",
-        "hasBias": true,
         "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
         "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
       },
       "intermediates": [
         {
@@ -1075,7 +1039,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1102,7 +1066,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1115,14 +1079,7 @@
       "id": "nchw2d_gemm_tiled",
       "priority": 30,
       "when": ["noBiasContract", "f16Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f16GemmScratchFits", "rank4OutputShapeOk"],
-      "derive": {
-        "usesF16": true,
-        "hasBias": false,
-        "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
-        "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
-      },
       "intermediates": [
         {
           "id": "wt",
@@ -1140,7 +1097,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1152,7 +1109,7 @@
           "name": "ConvTransposeColsGemmTiledReg",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
@@ -1163,7 +1120,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1176,14 +1133,7 @@
       "id": "nchw2d_gemm_tiled_bias",
       "priority": 35,
       "when": ["biasContract", "f16Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f16GemmScratchFits", "rank4OutputShapeOk"],
-      "derive": {
-        "usesF16": true,
-        "hasBias": true,
-        "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
-        "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
-      },
       "intermediates": [
         {
           "id": "wt",
@@ -1201,7 +1151,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1213,7 +1163,7 @@
           "name": "ConvTransposeColsGemmTiledReg",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
@@ -1224,7 +1174,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1237,14 +1187,7 @@
       "id": "nchw2d_gemm_tiled_f32",
       "priority": 30,
       "when": ["noBiasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk"],
-      "derive": {
-        "usesF16": false,
-        "hasBias": false,
-        "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
-        "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
-      },
       "intermediates": [
         {
           "id": "wt",
@@ -1262,7 +1205,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1274,7 +1217,7 @@
           "name": "ConvTransposeColsGemmTiled",
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 32)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 32)",
@@ -1285,7 +1228,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1298,14 +1241,7 @@
       "id": "nchw2d_gemm_tiled_f32_reg",
       "priority": 32,
       "when": ["noBiasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk", "ceilDiv(gemmM, 64) * ceilDiv(gemmN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
-      "derive": {
-        "usesF16": false,
-        "hasBias": false,
-        "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
-        "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
-      },
       "intermediates": [
         {
           "id": "wt",
@@ -1323,7 +1259,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1335,7 +1271,7 @@
           "name": "ConvTransposeColsGemmRegTiled",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
@@ -1346,7 +1282,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1359,14 +1295,7 @@
       "id": "nchw2d_gemm_tiled_f32_bias",
       "priority": 35,
       "when": ["biasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk"],
-      "derive": {
-        "usesF16": false,
-        "hasBias": true,
-        "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
-        "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
-      },
       "intermediates": [
         {
           "id": "wt",
@@ -1384,7 +1313,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1396,7 +1325,7 @@
           "name": "ConvTransposeColsGemmTiled",
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 32)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 32)",
@@ -1407,7 +1336,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1420,14 +1349,7 @@
       "id": "nchw2d_gemm_tiled_f32_bias_reg",
       "priority": 37,
       "when": ["biasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk", "ceilDiv(gemmM, 64) * ceilDiv(gemmN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
-      "derive": {
-        "usesF16": false,
-        "hasBias": true,
-        "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
-        "K": "dim(shapes.x, 1)",
-        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)",
-        "exactTiling": "exactTiling2d"
-      },
       "intermediates": [
         {
           "id": "wt",
@@ -1445,7 +1367,7 @@
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
-          "bindings": ["wsrc", "wt", "params_4"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1457,7 +1379,7 @@
           "name": "ConvTransposeColsGemmRegTiled",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
-          "bindings": ["w_2", "xm", "y_3", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
@@ -1468,7 +1390,7 @@
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
-          "bindings": ["cols", "bias", "y", "params_5"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1489,23 +1411,18 @@
           "shader": "conv-transpose2d-phase-gemm-tiled.wgsl.jinja",
           "derive": {
             "batch": "dim(shapes.x, 0)",
-            "inChannels": "dim(shapes.x, 1)",
             "inHeight": "dim(shapes.x, 2)",
             "inWidth": "dim(shapes.x, 3)",
-            "outChannels": "dim(shapes.y, 1)",
             "outHeight": "dim(shapes.y, 2)",
             "outWidth": "dim(shapes.y, 3)",
-            "padH": "effectivePadTop",
-            "padW": "effectivePadLeft",
             "mTile": "tunables.PHASE2D_GEMM_TILE",
             "nTile": "tunables.PHASE2D_GEMM_TILE",
             "kTile": "tunables.PHASE2D_GEMM_K_TILE",
             "threadRows": "tunables.PHASE2D_GEMM_THREAD_TILE",
             "threadColumns": "tunables.PHASE2D_GEMM_THREAD_TILE",
-            "workgroupEdge": "tunables.PHASE2D_GEMM_WORKGROUP_EDGE",
-            "hasBias": false
           },
-          "bindings": ["x", "w", "y_2"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2) * dim(shapes.x, 3), tunables.PHASE2D_GEMM_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), tunables.PHASE2D_GEMM_TILE)",
@@ -1526,23 +1443,18 @@
           "shader": "conv-transpose2d-phase-gemm-tiled.wgsl.jinja",
           "derive": {
             "batch": "dim(shapes.x, 0)",
-            "inChannels": "dim(shapes.x, 1)",
             "inHeight": "dim(shapes.x, 2)",
             "inWidth": "dim(shapes.x, 3)",
-            "outChannels": "dim(shapes.y, 1)",
             "outHeight": "dim(shapes.y, 2)",
             "outWidth": "dim(shapes.y, 3)",
-            "padH": "effectivePadTop",
-            "padW": "effectivePadLeft",
             "mTile": "tunables.PHASE2D_GEMM_TILE",
             "nTile": "tunables.PHASE2D_GEMM_TILE",
             "kTile": "tunables.PHASE2D_GEMM_K_TILE",
             "threadRows": "tunables.PHASE2D_GEMM_THREAD_TILE",
             "threadColumns": "tunables.PHASE2D_GEMM_THREAD_TILE",
-            "workgroupEdge": "tunables.PHASE2D_GEMM_WORKGROUP_EDGE",
-            "hasBias": true
           },
-          "bindings": ["x", "w", "bias", "y_2"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2) * dim(shapes.x, 3), tunables.PHASE2D_GEMM_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), tunables.PHASE2D_GEMM_TILE)",
@@ -1555,37 +1467,22 @@
       "id": "nchw2d_grouped_stride_phase_blocked",
       "priority": 25,
       "when": ["noBiasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW", "groupedStridePhaseBlockedFits"],
-      "derive": {
-        "usesF16": "dtypes.T == \"f16\"",
-        "scalar": "dtypes.T",
-        "phaseWeightElement": "dtypes.T",
-        "phaseOutputElement": "dtypes.T"
-      },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.GroupedStridePhaseBlocked",
           "shader": "conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja",
           "derive": {
-            "outC": "dim(shapes.y, 1)",
             "outCPerGroup": "groupedOutCPerGroup",
-            "outH": "dim(shapes.y, 2)",
-            "outW": "dim(shapes.y, 3)",
-            "inC": "dim(shapes.x, 1)",
-            "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
-            "inH": "dim(shapes.x, 2)",
-            "inW": "dim(shapes.x, 3)",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
             "strideWSpec": "strideW",
-            "padH": "effectivePadTop",
-            "padW": "effectivePadLeft",
             "groups": "attrs.group",
-            "ocBlock": "groupedBlockOc",
-            "hasBias": false
           },
-          "bindings": ["x", "w", "y_2"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), tunables.WORKGROUP_SIZE)",
             "y": "groupedPlanes"
@@ -1594,85 +1491,54 @@
       ]
     },
     {
-      "id": "nchw2d_grouped_stride_phase_blocked_bias",
-      "priority": 25,
-      "when": ["biasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW", "groupedStridePhaseBlockedFits"],
-      "derive": {
-        "usesF16": "dtypes.T == \"f16\"",
-        "scalar": "dtypes.T",
-        "phaseWeightElement": "dtypes.T",
-        "phaseOutputElement": "dtypes.T"
-      },
       "passes": [
         {
           "id": "main",
-          "name": "ConvTranspose.GroupedStridePhaseBlockedBias",
-          "shader": "conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja",
           "derive": {
-            "outC": "dim(shapes.y, 1)",
-            "outCPerGroup": "groupedOutCPerGroup",
-            "outH": "dim(shapes.y, 2)",
-            "outW": "dim(shapes.y, 3)",
-            "inC": "dim(shapes.x, 1)",
-            "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
-            "inH": "dim(shapes.x, 2)",
-            "inW": "dim(shapes.x, 3)",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
-            "strideWSpec": "strideW",
-            "padH": "effectivePadTop",
-            "padW": "effectivePadLeft",
-            "groups": "attrs.group",
-            "ocBlock": "groupedBlockOc",
-            "hasBias": true
           },
-          "bindings": ["x", "w", "bias", "y_2"],
           "dispatch": {
-            "x": "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), tunables.WORKGROUP_SIZE)",
-            "y": "groupedPlanes"
           }
         }
       ]
     },
     {
-      "id": "nchw2d_grouped_stride_phase",
-      "priority": 20,
-      "when": ["noBiasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW"],
-      "derive": {
-        "usesF16": "dtypes.T == \"f16\"",
-        "scalar": "dtypes.T",
-        "phaseWeightElement": "dtypes.T",
-        "phaseOutputElement": "dtypes.T"
-      },
       "passes": [
         {
           "id": "main",
-          "name": "ConvTranspose.GroupedStridePhase",
-          "shader": "conv-transpose2d-grouped-stride-phase.wgsl.jinja",
           "derive": {
-            "count": "numel(shapes.y)",
-            "outC": "dim(shapes.y, 1)",
-            "outCPerGroup": "dim(shapes.y, 1) / attrs.group",
-            "outH": "dim(shapes.y, 2)",
-            "outW": "dim(shapes.y, 3)",
-            "inC": "dim(shapes.x, 1)",
-            "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
-            "inH": "dim(shapes.x, 2)",
-            "inW": "dim(shapes.x, 3)",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
             "strideWSpec": "strideW",
-            "padH": "effectivePadTop",
-            "padW": "effectivePadLeft",
-            "hasBias": false
           },
-          "bindings": ["x", "w", "y_2"],
           "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
           }
         }
       ]
@@ -1681,36 +1547,20 @@
       "id": "nchw2d_grouped_stride_phase_bias",
       "priority": 20,
       "when": ["biasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW"],
-      "derive": {
-        "usesF16": "dtypes.T == \"f16\"",
-        "scalar": "dtypes.T",
-        "phaseWeightElement": "dtypes.T",
-        "phaseOutputElement": "dtypes.T"
-      },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.GroupedStridePhaseBias",
           "shader": "conv-transpose2d-grouped-stride-phase.wgsl.jinja",
           "derive": {
-            "count": "numel(shapes.y)",
-            "outC": "dim(shapes.y, 1)",
             "outCPerGroup": "dim(shapes.y, 1) / attrs.group",
-            "outH": "dim(shapes.y, 2)",
-            "outW": "dim(shapes.y, 3)",
-            "inC": "dim(shapes.x, 1)",
-            "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
-            "inH": "dim(shapes.x, 2)",
-            "inW": "dim(shapes.x, 3)",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
-            "strideWSpec": "strideW",
-            "padH": "effectivePadTop",
-            "padW": "effectivePadLeft",
-            "hasBias": true
           },
-          "bindings": ["x", "w", "bias", "y_2"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1722,33 +1572,14 @@
     {
       "id": "ncl1d",
       "priority": 0,
-      "when": ["noBiasContract", "f32Type", "rank3DenseContract", "(rank3OutputShapeOk or ((attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") and dim(shapes.y, 2) == dim(shapes.x, 2) * strideW))"],
-      "derive": {},
-      "passes": [
-        {
-          "id": "main",
-          "name": "ConvTranspose1d",
-          "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_7"],
-          "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
-          }
-        }
-      ]
-    },
-    {
-      "id": "ncl1d_f16",
-      "priority": 5,
-      "when": ["noBiasContract", "f16Type", "rank3DenseContract", "rank3OutputShapeOk"],
-      "derive": { "usesF16": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1d",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_7"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1760,33 +1591,14 @@
     {
       "id": "ncl1d_bias",
       "priority": 10,
-      "when": ["biasContract", "f32Type", "rank3DenseContract", "rank3OutputShapeOk"],
-      "derive": { "hasBias": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1dBias",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_7"],
-          "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
-          }
-        }
-      ]
-    },
-    {
-      "id": "ncl1d_bias_f16",
-      "priority": 10,
-      "when": ["biasContract", "f16Type", "rank3DenseContract", "rank3OutputShapeOk"],
-      "derive": { "hasBias": true, "usesF16": true },
-      "passes": [
-        {
-          "id": "main",
-          "name": "ConvTranspose1dBias",
-          "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_7"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1798,33 +1610,14 @@
     {
       "id": "ncl1d_grouped",
       "priority": 10,
-      "when": ["noBiasContract", "f32Type", "rank3GroupedContract", "rank3OutputShapeOk"],
-      "derive": { "hasBias": false, "grouped": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1dGrouped",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_8"],
-          "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
-          }
-        }
-      ]
-    },
-    {
-      "id": "ncl1d_grouped_f16",
-      "priority": 5,
-      "when": ["noBiasContract", "f16Type", "rank3GroupedContract", "rank3OutputShapeOk"],
-      "derive": { "hasBias": false, "usesF16": true, "grouped": true },
-      "passes": [
-        {
-          "id": "main",
-          "name": "ConvTranspose1dGrouped",
-          "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_8"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1836,33 +1629,14 @@
     {
       "id": "ncl1d_grouped_bias",
       "priority": 15,
-      "when": ["biasContract", "f32Type", "rank3GroupedContract", "rank3OutputShapeOk"],
-      "derive": { "hasBias": true, "grouped": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1dGroupedBias",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_8"],
-          "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
-          }
-        }
-      ]
-    },
-    {
-      "id": "ncl1d_grouped_bias_f16",
-      "priority": 8,
-      "when": ["biasContract", "f16Type", "rank3GroupedContract", "rank3OutputShapeOk"],
-      "derive": { "hasBias": true, "usesF16": true, "grouped": true },
-      "passes": [
-        {
-          "id": "main",
-          "name": "ConvTranspose1dGroupedBias",
-          "shader": "conv-transpose1d-ncl.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_8"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1874,14 +1648,14 @@
     {
       "id": "nchw2d",
       "priority": 0,
-      "when": ["noBiasContract", "f32Type", "rank4DenseContract", "(rank4OutputShapeOk or ((attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") and dim(shapes.y, 2) == dim(shapes.x, 2) * strideH and dim(shapes.y, 3) == dim(shapes.x, 3) * strideW))"],
-      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_9"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1893,14 +1667,14 @@
     {
       "id": "nchw2d_bias",
       "priority": 10,
-      "when": ["biasContract", "f32Type", "rank4DenseContract", "rank4OutputShapeOk"],
-      "derive": { "hasBias": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_9"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1919,7 +1693,7 @@
           "id": "main",
           "name": "ConvTranspose3d",
           "shader": "conv-transpose3d-ncdhw.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_10"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1932,13 +1706,13 @@
       "id": "ncdhw3d_bias",
       "priority": 10,
       "when": ["biasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk"],
-      "derive": { "hasBias": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose3dBias",
           "shader": "conv-transpose3d-ncdhw.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_10"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1950,14 +1724,14 @@
     {
       "id": "nchw2d_grouped",
       "priority": 10,
-      "when": ["noBiasContract", "f32Type", "rank4GroupedContract", "rank4OutputShapeOk"],
-      "derive": { "hasBias": false, "grouped": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTransposeGrouped",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_11"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1967,16 +1741,16 @@
       ]
     },
     {
-      "id": "nchw2d_grouped_f16",
-      "priority": 5,
-      "when": ["noBiasContract", "f16Type", "rank4GroupedContract", "rank4OutputShapeOk"],
-      "derive": { "hasBias": false, "usesF16": true, "grouped": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvTransposeGrouped",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_11"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -1986,97 +1760,103 @@
       ]
     },
     {
-      "id": "nchw2d_grouped_bias_f16",
-      "priority": 8,
-      "when": ["biasContract", "f16Type", "rank4GroupedContract", "rank4OutputShapeOk"],
-      "derive": { "hasBias": true, "usesF16": true, "grouped": true },
       "passes": [
         {
           "id": "main",
-          "name": "ConvTransposeGrouped",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_11"],
           "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
         }
       ]
     },
     {
-      "id": "nchw2d_grouped_bias",
-      "priority": 15,
-      "when": ["biasContract", "f32Type", "rank4GroupedContract", "rank4OutputShapeOk"],
-      "derive": { "hasBias": true, "grouped": true },
-      "passes": [
         {
-          "id": "main",
-          "name": "ConvTransposeGrouped",
-          "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_11"],
-          "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
-          }
         }
-      ]
-    },
-    {
-      "id": "nchw2d_f16",
-      "priority": 5,
-      "when": ["noBiasContract", "f16Type", "rank4DenseContract", "(rank4OutputShapeOk or ((attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") and dim(shapes.y, 2) == dim(shapes.x, 2) * strideH and dim(shapes.y, 3) == dim(shapes.x, 3) * strideW))"],
-      "derive": { "usesF16": true },
       "passes": [
         {
-          "id": "main",
-          "name": "ConvTranspose2dF16",
-          "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_9"],
           "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
         }
       ]
     },
     {
-      "id": "nchw2d_bias_f16",
-      "priority": 10,
-      "when": ["biasContract", "f16Type", "rank4DenseContract", "rank4OutputShapeOk"],
-      "derive": { "hasBias": true, "usesF16": true },
-      "passes": [
         {
-          "id": "main",
-          "name": "ConvTranspose2dBiasF16",
-          "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "bias", "y", "params_9"],
-          "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "z": 1
-          }
         }
-      ]
-    },
-    {
-      "id": "nchw2d_depthwise_pointwise_vec4",
-      "priority": 30,
-      "when": ["noBiasContract", "f32Type", "rank4DepthwisePointwiseVec4Contract"],
-      "derive": { "hasBias": false, "grouped": true, "depthwisePointwiseVec4": true },
       "passes": [
         {
-          "id": "main",
-          "name": "ConvTransposeDepthwisePointwiseVec4",
-          "shader": "conv-transpose2d-nchw.wgsl.jinja",
-          "bindings": ["x", "w_4", "y", "params_12"],
           "dispatch": {
-            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
         }
       ]
     }

     "stridesLengthOk": "not has(attrs, \"strides\") or (attrs.strides | length) == spatialRank",
     "dilationsLengthOk": "not has(attrs, \"dilations\") or (attrs.dilations | length) == spatialRank",
     "padsLengthOk": "not has(attrs, \"pads\") or (attrs.pads | length) == 2 * spatialRank",
     "kernelD": "attrs.kernel_shape[0] if kernelShapeLengthOk and has(attrs, \"kernel_shape\") and spatialRank == 3 else 1",
     "kernelH": "attrs.kernel_shape[spatialRank - 2] if kernelShapeLengthOk and has(attrs, \"kernel_shape\") and spatialRank >= 2 else 1",
     "kernelW": "attrs.kernel_shape[spatialRank - 1] if kernelShapeLengthOk and has(attrs, \"kernel_shape\") and spatialRank >= 1 else 1",
     "padBack": "attrs.pads[spatialRank] if padsLengthOk and has(attrs, \"pads\") and spatialRank == 3 else 0",
     "padBottom": "attrs.pads[2 * spatialRank - 2] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 2 else 0",
     "padRight": "attrs.pads[2 * spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
+    "outputPaddingLengthOk": "not has(attrs, \"output_padding\") or (attrs.output_padding | length) == spatialRank",
     "outputPaddingD": "attrs.output_padding[0] if outputPaddingLengthOk and has(attrs, \"output_padding\") and spatialRank == 3 else 0",
     "outputPaddingH": "attrs.output_padding[spatialRank - 2] if outputPaddingLengthOk and has(attrs, \"output_padding\") and spatialRank >= 2 else 0",
     "outputPaddingW": "attrs.output_padding[spatialRank - 1] if outputPaddingLengthOk and has(attrs, \"output_padding\") and spatialRank >= 1 else 0",
     "rank5DenseContract": "spatialAttributesOk and ranks.x == 5 and ranks.w == 5 and ranks.y == 5 and attrs.group == 1 and batchChannelContract and outputDispatchFits and dim(shapes.y, 1) == dim(shapes.w, 1)",
     "noBiasContract": "not present.bias",
     "biasContract": "present.bias and ranks.bias == 1 and dim(shapes.bias, 0) == dim(shapes.y, 1)",
+    "hasBias": "present.bias",
     "f32Type": "dtypes.T == \"f32\"",
     "f16Type": "device.features.has(\"shader-f16\") and dtypes.T == \"f16\"",
     "gemmM": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
     "groupedBlockOc": "16 if groupedOutCPerGroup % 16 == 0 else (8 if groupedOutCPerGroup % 8 == 0 else (4 if groupedOutCPerGroup % 4 == 0 else (2 if groupedOutCPerGroup % 2 == 0 else 1)))",
     "groupedWeightBytes": "(dim(shapes.x, 1) / attrs.group) * groupedOutCPerGroup * dim(shapes.w, 2) * dim(shapes.w, 3) * 4",
     "groupedPlanes": "dim(shapes.x, 0) * attrs.group * (groupedOutCPerGroup / groupedBlockOc)",
+    "groupedStridePhaseBlockedFits": "groupedWeightBytes <= device.limits.maxComputeWorkgroupStorageSize and groupedPlanes <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), tunables.WORKGROUP_SIZE) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "count": "numel(shapes.y)",
+    "exactTiling": "exactTiling2d",
+    "inC": "dim(shapes.x, 1)",
+    "inCPerGroup": "dim(shapes.x, 1) / attrs.group",
+    "inChannels": "dim(shapes.x, 1)",
+    "inH": "dim(shapes.x, 2)",
+    "inW": "dim(shapes.x, 3)",
+    "outC": "dim(shapes.y, 1)",
+    "outChannels": "dim(shapes.y, 1)",
+    "outH": "dim(shapes.y, 2)",
+    "outW": "dim(shapes.y, 3)",
+    "padH": "effectivePadTop",
+    "padW": "effectivePadLeft",
+    "scalar": "dtypes.T"
   },
   "bindings": {
+    "y": { "elementType": "$T" },
+    "bias": { "elementType": "$T" },
+    "x": { "elementType": "$T" },
+    "w": { "elementType": "$phaseWeightElement" },
+    "y_main": { "name": "y", "elementType": "$phaseOutputElement" },
+    "wsrc": { "arg": "w", "elementType": "$T" },
     "wt": { "buffer": "storage", "elementType": "$T" },
     "params": {
       "struct": [
         { "name": "count", "type": "u32", "value": "numel(shapes.w)" },
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
         { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 4)" }
       ]
     },
+    "w_t": { "scratch": "wt", "name": "w", "buffer": "read-only-storage", "elementType": "$T" },
+    "xm": { "arg": "x", "elementType": "$T" },
+    "y_t": { "scratch": "cols", "name": "y", "elementType": "$T" },
     "cols": { "buffer": "read-only-storage", "elementType": "$T" },
+    "params_col2im": {
       "name": "params",
       "struct": [
         { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
         { "name": "kernelD", "type": "u32", "value": "dim(shapes.w, 2)" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_reorder": {
       "name": "params",
       "struct": [
         { "name": "count", "type": "u32", "value": "numel(shapes.w)" },
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
         { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 3)" }
       ]
     },
+    "params__uniform": {
       "name": "params",
       "struct": [
         { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
         { "name": "kernelH", "type": "u32", "value": "dim(shapes.w, 2)" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_gemm": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" },
         { "name": "K", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "N", "type": "u32", "value": "dim(shapes.x, 2) * dim(shapes.x, 3)" }
       ]
     },
+    "w_main": { "name": "w", "elementType": "$T" },
+    "params_main": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inW", "type": "u32", "value": "dim(shapes.x, 2)" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_ncl1d_grouped": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_nchw2d": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inH", "type": "u32", "value": "dim(shapes.x, 2)" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_ncdhw3d": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inD", "type": "u32", "value": "dim(shapes.x, 2)" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_nchw2d_grouped": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "params_conv_transpose2d_nchw": {
       "name": "params",
       "struct": [
         { "name": "outChannels", "type": "u32", "value": "dim(shapes.y, 1)" },
         { "name": "outH", "type": "u32", "value": "dim(shapes.y, 2)" },
       "id": "empty_input_reduction",
       "priority": 100,
       "when": ["noBiasContract", "emptyInputReductionContract"],
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.EmptyInput",
           "shader": "conv-transpose-empty-input.wgsl.jinja",
           "derive": {
             "outputSpatial": "dim(shapes.y, 2) if ranks.y == 3 else (dim(shapes.y, 2) * dim(shapes.y, 3) if ranks.y == 4 else dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4))"
           },
           "bindings": ["y"],
       "id": "empty_input_reduction_bias",
       "priority": 100,
       "when": ["biasContract", "emptyInputReductionContract"],
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.EmptyInputBias",
           "shader": "conv-transpose-empty-input.wgsl.jinja",
           "derive": {
             "outputSpatial": "dim(shapes.y, 2) if ranks.y == 3 else (dim(shapes.y, 2) * dim(shapes.y, 3) if ranks.y == 4 else dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4))"
           },
           "bindings": ["bias", "y"],
       "priority": 40,
       "when": ["noBiasContract", "phase1dCovered"],
       "derive": {
         "phaseWeightElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "phaseOutputElement": "\"vec2<\" ~ dtypes.T ~ \">\""
       },
           "name": "ConvTranspose1d.PhaseTiled",
           "shader": "conv-transpose1d-phase-tiled.wgsl.jinja",
           "derive": {
             "inWidth": "dim(shapes.x, 2)",
             "positionTile": "tunables.PHASE1D_POSITION_TILE",
             "outputTile": "phase1dOutputTile",
             "outputsPerThread": "phase1dOutputsPerThread",
             "xTileElements": "tunables.PHASE1D_K_TILE * (tunables.PHASE1D_POSITION_TILE + 2)",
             "weightTileElements": "tunables.PHASE1D_K_TILE * phase1dOutputTile"
           },
+          "bindings": ["x", "w", "y_main"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2), tunables.PHASE1D_POSITION_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), phase1dOutputTile)",
       "priority": 45,
       "when": ["biasContract", "phase1dCovered"],
       "derive": {
         "phaseWeightElement": "\"vec4<\" ~ dtypes.T ~ \">\"",
         "phaseOutputElement": "\"vec2<\" ~ dtypes.T ~ \">\""
       },
           "name": "ConvTranspose1d.PhaseTiledBias",
           "shader": "conv-transpose1d-phase-tiled.wgsl.jinja",
           "derive": {
             "inWidth": "dim(shapes.x, 2)",
             "positionTile": "tunables.PHASE1D_POSITION_TILE",
             "outputTile": "phase1dOutputTile",
             "outputsPerThread": "phase1dOutputsPerThread",
             "xTileElements": "tunables.PHASE1D_K_TILE * (tunables.PHASE1D_POSITION_TILE + 2)",
             "weightTileElements": "tunables.PHASE1D_K_TILE * phase1dOutputTile"
           },
+          "bindings": ["x", "w", "bias", "y_main"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2), tunables.PHASE1D_POSITION_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), phase1dOutputTile)",
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
         "fScalar": "\"f32\"",
         "directMatrixStore": "ctImplicitDirectStore",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "subgroupRows": "tunables.SGMAT_SUBGROUP_ROWS",
           "id": "gemm",
           "name": "ConvTranspose3dImplicitSubgroupMatrix",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
+          "bindings": ["w_t", "xm", "y"],
           "dispatch": {
             "x": "ceilDiv(ctImplicitOutVolume, tunables.SGMAT_TILE_COLS)",
             "y": "ctImplicitYWorkgroups",
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
         "fScalar": "\"f32\"",
         "directMatrixStore": "ctImplicitDirectStore",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "subgroupRows": "tunables.SGMAT_SUBGROUP_ROWS",
           "id": "gemm",
           "name": "ConvTranspose3dImplicitSubgroupMatrixBias",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
+          "bindings": ["w_t", "xm", "bias", "y"],
           "dispatch": {
             "x": "ceilDiv(ctImplicitOutVolume, tunables.SGMAT_TILE_COLS)",
             "y": "ctImplicitYWorkgroups",
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
         "fScalar": "\"f32\"",
         "directMatrixStore": "gemm3dM % tunables.SGMAT_TILE_ROWS == 0",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "tileCols": "tunables.SGMAT_TILE_COLS",
           "name": "ConvTranspose3dColsGemm",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t"],
           "dispatch": {
             "x": "ceilDiv(gemm3dN, tunables.SGMAT_TILE_COLS)",
             "y": "ceilDiv(gemm3dM, tunables.SGMAT_TILE_ROWS)",
           "id": "col2im",
           "name": "ConvTranspose3dCol2Im",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params_col2im"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
         "fScalar": "\"f32\"",
         "directMatrixStore": "gemm3dM % tunables.SGMAT_TILE_ROWS == 0",
         "tileRows": "tunables.SGMAT_TILE_ROWS",
         "tileCols": "tunables.SGMAT_TILE_COLS",
           "name": "ConvTranspose3dColsGemm",
           "shader": "conv-1x1-subgroup-matrix.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t"],
           "dispatch": {
             "x": "ceilDiv(gemm3dN, tunables.SGMAT_TILE_COLS)",
             "y": "ceilDiv(gemm3dM, tunables.SGMAT_TILE_ROWS)",
           "id": "col2im",
           "name": "ConvTranspose3dCol2ImBias",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params_col2im"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "ncdhw3d_gemm_tiled_f32",
       "priority": 30,
       "when": ["noBiasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits"],
+      "derive": { "usesF16": false },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
+            "w_t",
             "xm",
+            "y_t",
             {
               "name": "params",
               "struct": [
           "id": "col2im",
           "name": "ConvTranspose3dCol2Im",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params_col2im"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "ncdhw3d_gemm_tiled_f32_reg",
       "priority": 32,
       "when": ["noBiasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits", "ceilDiv(gemm3dM, 64) * ceilDiv(gemm3dN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
+      "derive": { "usesF16": false, "gemmKTile": "32 if dim(shapes.x, 1) % 32 == 0 else 16" },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
+            "w_t",
             "xm",
+            "y_t",
             {
               "name": "params",
               "struct": [
           "id": "col2im",
           "name": "ConvTranspose3dCol2Im",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params_col2im"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "ncdhw3d_gemm_tiled_f32_bias",
       "priority": 35,
       "when": ["biasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits"],
+      "derive": { "usesF16": false },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
+            "w_t",
             "xm",
+            "y_t",
             {
               "name": "params",
               "struct": [
           "id": "col2im",
           "name": "ConvTranspose3dCol2ImBias",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params_col2im"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "ncdhw3d_gemm_tiled_f32_bias_reg",
       "priority": 37,
       "when": ["biasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk", "gemm3dN > 0", "gemm3dDispatchFits", "f32Gemm3dScratchFits", "ceilDiv(gemm3dM, 64) * ceilDiv(gemm3dN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
+      "derive": { "usesF16": false, "gemmKTile": "32 if dim(shapes.x, 1) % 32 == 0 else 16" },
       "intermediates": [
         { "id": "wt", "dtype": "float32", "shape": "[gemmWeightElements]" },
         { "id": "cols", "dtype": "float32", "shape": "[gemm3dColumnElements]" }
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
           "bindings": [
+            "w_t",
             "xm",
+            "y_t",
             {
               "name": "params",
               "struct": [
           "id": "col2im",
           "name": "ConvTranspose3dCol2ImBias",
           "shader": "conv-transpose3d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params_col2im"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
         "subgroupMatrixConfigs": [{ "componentType": "f16", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
         "fScalar": "\"f16\"",
         "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
         "K": "dim(shapes.x, 1)",
+        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)"
       },
       "intermediates": [
         {
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
         "subgroupMatrixConfigs": [{ "componentType": "f16", "M": 8, "N": 8, "K": 8 }]
       },
       "derive": {
         "fScalar": "\"f16\"",
         "M": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
         "K": "dim(shapes.x, 1)",
+        "N": "dim(shapes.x, 2) * dim(shapes.x, 3)"
       },
       "intermediates": [
         {
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "nchw2d_gemm_tiled",
       "priority": 30,
       "when": ["noBiasContract", "f16Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f16GemmScratchFits", "rank4OutputShapeOk"],
+      "derive": { "usesF16": true },
       "intermediates": [
         {
           "id": "wt",
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "ConvTransposeColsGemmTiledReg",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t", "params_gemm"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "nchw2d_gemm_tiled_bias",
       "priority": 35,
       "when": ["biasContract", "f16Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f16GemmScratchFits", "rank4OutputShapeOk"],
+      "derive": { "usesF16": true },
       "intermediates": [
         {
           "id": "wt",
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "ConvTransposeColsGemmTiledReg",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t", "params_gemm"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "nchw2d_gemm_tiled_f32",
       "priority": 30,
       "when": ["noBiasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk"],
+      "derive": { "usesF16": false },
       "intermediates": [
         {
           "id": "wt",
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "ConvTransposeColsGemmTiled",
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t", "params_gemm"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 32)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 32)",
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "nchw2d_gemm_tiled_f32_reg",
       "priority": 32,
       "when": ["noBiasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk", "ceilDiv(gemmM, 64) * ceilDiv(gemmN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
+      "derive": { "usesF16": false },
       "intermediates": [
         {
           "id": "wt",
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "ConvTransposeColsGemmRegTiled",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t", "params_gemm"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
           "id": "col2im",
           "name": "ConvTransposeCol2Im",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "nchw2d_gemm_tiled_f32_bias",
       "priority": 35,
       "when": ["biasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk"],
+      "derive": { "usesF16": false },
       "intermediates": [
         {
           "id": "wt",
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "ConvTransposeColsGemmTiled",
           "shader": "conv-1x1-gemm-tiled.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t", "params_gemm"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 32)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 32)",
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "nchw2d_gemm_tiled_f32_bias_reg",
       "priority": 37,
       "when": ["biasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk", "ceilDiv(gemmM, 64) * ceilDiv(gemmN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT"],
+      "derive": { "usesF16": false },
       "intermediates": [
         {
           "id": "wt",
           "id": "reorder",
           "name": "ConvTransposeWeightReorder",
           "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
           "name": "ConvTransposeColsGemmRegTiled",
           "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
           "derive": { "hasBias": false },
+          "bindings": ["w_t", "xm", "y_t", "params_gemm"],
           "dispatch": {
             "x": "ceil(dim(shapes.x, 2) * dim(shapes.x, 3) / 64)",
             "y": "ceil(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) / 64)",
           "id": "col2im",
           "name": "ConvTransposeCol2ImBias",
           "shader": "conv-transpose2d-col2im.wgsl.jinja",
+          "bindings": ["cols", "bias", "y", "params__uniform"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
           "shader": "conv-transpose2d-phase-gemm-tiled.wgsl.jinja",
           "derive": {
             "batch": "dim(shapes.x, 0)",
             "inHeight": "dim(shapes.x, 2)",
             "inWidth": "dim(shapes.x, 3)",
             "outHeight": "dim(shapes.y, 2)",
             "outWidth": "dim(shapes.y, 3)",
             "mTile": "tunables.PHASE2D_GEMM_TILE",
             "nTile": "tunables.PHASE2D_GEMM_TILE",
             "kTile": "tunables.PHASE2D_GEMM_K_TILE",
             "threadRows": "tunables.PHASE2D_GEMM_THREAD_TILE",
             "threadColumns": "tunables.PHASE2D_GEMM_THREAD_TILE",
+            "workgroupEdge": "tunables.PHASE2D_GEMM_WORKGROUP_EDGE"
           },
+          "bindings": ["x", "w", "y_main"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2) * dim(shapes.x, 3), tunables.PHASE2D_GEMM_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), tunables.PHASE2D_GEMM_TILE)",
           "shader": "conv-transpose2d-phase-gemm-tiled.wgsl.jinja",
           "derive": {
             "batch": "dim(shapes.x, 0)",
             "inHeight": "dim(shapes.x, 2)",
             "inWidth": "dim(shapes.x, 3)",
             "outHeight": "dim(shapes.y, 2)",
             "outWidth": "dim(shapes.y, 3)",
             "mTile": "tunables.PHASE2D_GEMM_TILE",
             "nTile": "tunables.PHASE2D_GEMM_TILE",
             "kTile": "tunables.PHASE2D_GEMM_K_TILE",
             "threadRows": "tunables.PHASE2D_GEMM_THREAD_TILE",
             "threadColumns": "tunables.PHASE2D_GEMM_THREAD_TILE",
+            "workgroupEdge": "tunables.PHASE2D_GEMM_WORKGROUP_EDGE"
           },
+          "bindings": ["x", "w", "bias", "y_main"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.x, 2) * dim(shapes.x, 3), tunables.PHASE2D_GEMM_TILE)",
             "y": "ceilDiv(dim(shapes.y, 1), tunables.PHASE2D_GEMM_TILE)",
       "id": "nchw2d_grouped_stride_phase_blocked",
       "priority": 25,
       "when": ["noBiasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW", "groupedStridePhaseBlockedFits"],
+      "derive": { "phaseWeightElement": "dtypes.T", "phaseOutputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.GroupedStridePhaseBlocked",
           "shader": "conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja",
           "derive": {
             "outCPerGroup": "groupedOutCPerGroup",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
             "strideWSpec": "strideW",
             "groups": "attrs.group",
+            "ocBlock": "groupedBlockOc"
           },
+          "bindings": ["x", "w", "y_main"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), tunables.WORKGROUP_SIZE)",
             "y": "groupedPlanes"
       ]
     },
     {
+      "id": "nchw2d_grouped_stride_phase",
+      "priority": 20,
+      "when": ["noBiasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW"],
+      "derive": { "phaseWeightElement": "dtypes.T", "phaseOutputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
+          "name": "ConvTranspose.GroupedStridePhase",
+          "shader": "conv-transpose2d-grouped-stride-phase.wgsl.jinja",
           "derive": {
+            "outCPerGroup": "dim(shapes.y, 1) / attrs.group",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
+            "strideWSpec": "strideW"
           },
+          "bindings": ["x", "w", "y_main"],
           "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
     },
     {
+      "id": "nchw2d_grouped_stride_phase_blocked_bias",
+      "priority": 25,
+      "when": ["biasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW", "groupedStridePhaseBlockedFits"],
+      "derive": { "phaseWeightElement": "dtypes.T", "phaseOutputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
+          "name": "ConvTranspose.GroupedStridePhaseBlockedBias",
+          "shader": "conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja",
           "derive": {
+            "outCPerGroup": "groupedOutCPerGroup",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
             "strideWSpec": "strideW",
+            "groups": "attrs.group",
+            "ocBlock": "groupedBlockOc"
           },
+          "bindings": ["x", "w", "bias", "y_main"],
           "dispatch": {
+            "x": "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), tunables.WORKGROUP_SIZE)",
+            "y": "groupedPlanes"
           }
         }
       ]
       "id": "nchw2d_grouped_stride_phase_bias",
       "priority": 20,
       "when": ["biasContract", "rank4GroupedContract", "f16Ok(dtypes.T)", "strideH >= 2", "strideW >= 2", "dilationH == 1", "dilationW == 1", "dim(shapes.w, 2) >= 2", "dim(shapes.w, 2) <= 7", "dim(shapes.w, 3) >= 2", "dim(shapes.w, 3) <= 7", "dim(shapes.y, 2) == (dim(shapes.x, 2) - 1) * strideH - (padTop + padBottom) + dim(shapes.w, 2) + outputPaddingH", "dim(shapes.y, 3) == (dim(shapes.x, 3) - 1) * strideW - (padLeft + padRight) + dim(shapes.w, 3) + outputPaddingW"],
+      "derive": { "phaseWeightElement": "dtypes.T", "phaseOutputElement": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose.GroupedStridePhaseBias",
           "shader": "conv-transpose2d-grouped-stride-phase.wgsl.jinja",
           "derive": {
             "outCPerGroup": "dim(shapes.y, 1) / attrs.group",
             "kernelHSpec": "dim(shapes.w, 2)",
             "kernelWSpec": "dim(shapes.w, 3)",
             "strideHSpec": "strideH",
+            "strideWSpec": "strideW"
           },
+          "bindings": ["x", "w", "bias", "y_main"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncl1d",
       "priority": 0,
+      "when": ["noBiasContract", "f32Type or f16Type", "rank3DenseContract", "rank3OutputShapeOk or (f32Type and ((attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") and dim(shapes.y, 2) == dim(shapes.x, 2) * strideW))"],
+      "derive": { "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1d",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
+          "bindings": ["x", "w_main", "y", "params_main"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncl1d_bias",
       "priority": 10,
+      "when": ["biasContract", "f32Type or f16Type", "rank3DenseContract", "rank3OutputShapeOk"],
+      "derive": { "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1dBias",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
+          "bindings": ["x", "w_main", "bias", "y", "params_main"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncl1d_grouped",
       "priority": 10,
+      "when": ["noBiasContract", "f32Type or f16Type", "rank3GroupedContract", "rank3OutputShapeOk"],
+      "derive": { "grouped": true, "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1dGrouped",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
+          "bindings": ["x", "w_main", "y", "params_ncl1d_grouped"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncl1d_grouped_bias",
       "priority": 15,
+      "when": ["biasContract", "f32Type or f16Type", "rank3GroupedContract", "rank3OutputShapeOk"],
+      "derive": { "grouped": true, "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose1dGroupedBias",
           "shader": "conv-transpose1d-ncl.wgsl.jinja",
+          "bindings": ["x", "w_main", "bias", "y", "params_ncl1d_grouped"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw2d",
       "priority": 0,
+      "when": ["noBiasContract", "f32Type or f16Type", "rank4DenseContract", "(rank4OutputShapeOk or ((attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") and dim(shapes.y, 2) == dim(shapes.x, 2) * strideH and dim(shapes.y, 3) == dim(shapes.x, 3) * strideW))"],
+      "derive": { "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w_main", "y", "params_nchw2d"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw2d_bias",
       "priority": 10,
+      "when": ["biasContract", "f32Type or f16Type", "rank4DenseContract", "rank4OutputShapeOk"],
+      "derive": { "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w_main", "bias", "y", "params_nchw2d"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
           "id": "main",
           "name": "ConvTranspose3d",
           "shader": "conv-transpose3d-ncdhw.wgsl.jinja",
+          "bindings": ["x", "w_main", "y", "params_ncdhw3d"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "id": "ncdhw3d_bias",
       "priority": 10,
       "when": ["biasContract", "f32Type", "rank5DenseContract", "rank5OutputShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvTranspose3dBias",
           "shader": "conv-transpose3d-ncdhw.wgsl.jinja",
+          "bindings": ["x", "w_main", "bias", "y", "params_ncdhw3d"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw2d_grouped",
       "priority": 10,
+      "when": ["noBiasContract", "f32Type or f16Type", "rank4GroupedContract", "rank4OutputShapeOk"],
+      "derive": { "grouped": true, "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTransposeGrouped",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w_main", "y", "params_nchw2d_grouped"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       ]
     },
     {
+      "id": "nchw2d_grouped_bias",
+      "priority": 15,
+      "when": ["biasContract", "f32Type or f16Type", "rank4GroupedContract", "rank4OutputShapeOk"],
+      "derive": { "grouped": true, "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "ConvTransposeGrouped",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w_main", "bias", "y", "params_nchw2d_grouped"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       ]
     },
     {
+      "id": "nchw2d_depthwise_pointwise_vec4",
+      "priority": 30,
+      "when": ["noBiasContract", "f32Type", "rank4DepthwisePointwiseVec4Contract"],
+      "derive": { "grouped": true, "depthwisePointwiseVec4": true },
       "passes": [
         {
           "id": "main",
+          "name": "ConvTransposeDepthwisePointwiseVec4",
           "shader": "conv-transpose2d-nchw.wgsl.jinja",
+          "bindings": ["x", "w_main", "y", "params_conv_transpose2d_nchw"],
           "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
         }
       ]
     },
     {
+      "id": "nchw2d_exact_store_gemm_f32",
+      "priority": 38,
+      "when": ["noBiasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk", "exactTiling2d", "ceilDiv(gemmM, 64) * ceilDiv(gemmN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT", "device.limits.maxComputeInvocationsPerWorkgroup >= 128", "device.limits.maxComputeWorkgroupSizeX >= 16", "device.limits.maxComputeWorkgroupSizeY >= 8", "device.limits.maxComputeWorkgroupStorageSize >= 8192", "device.limits.maxStorageBuffersPerShaderStage >= 4"],
+      "derive": {
+        "usesF16": false,
+        "convTransposeExactStore": true,
+        "convTransposeKernelH": "dim(shapes.w, 2)",
+        "convTransposeKernelW": "dim(shapes.w, 3)",
+        "convTransposeInputW": "dim(shapes.x, 3)",
+        "convTransposeOutputH": "dim(shapes.y, 2)",
+        "convTransposeOutputW": "dim(shapes.y, 3)"
+      },
+      "intermediates": [
         {
+          "id": "wt",
+          "dtype": "float32",
+          "shape": "[dim(shapes.w, 0) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)]"
         }
+      ],
       "passes": [
         {
+          "id": "reorder",
+          "name": "ConvTransposeWeightReorder",
+          "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
+        },
+        {
+          "id": "gemm",
+          "name": "ConvTransposeExactOutputGemm",
+          "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
+          "bindings": ["w_t", "xm", "y", "params_gemm"],
+          "dispatch": { "x": "ceilDiv(gemmN, 64)", "y": "ceilDiv(gemmM, 64)", "z": "dim(shapes.x, 0)" }
         }
       ]
     },
     {
+      "id": "nchw2d_exact_store_gemm_f32_bias",
+      "priority": 38,
+      "when": ["biasContract", "f32Type", "rank4DenseContract", "gemmN > 0", "gemmDispatchFits", "f32GemmScratchFits", "rank4OutputShapeOk", "exactTiling2d", "ceilDiv(gemmM, 64) * ceilDiv(gemmN, 64) >= tunables.GEMM_REG_MIN_TILE_PRODUCT", "device.limits.maxComputeInvocationsPerWorkgroup >= 128", "device.limits.maxComputeWorkgroupSizeX >= 16", "device.limits.maxComputeWorkgroupSizeY >= 8", "device.limits.maxComputeWorkgroupStorageSize >= 8192", "device.limits.maxStorageBuffersPerShaderStage >= 4"],
+      "derive": {
+        "usesF16": false,
+        "convTransposeExactStore": true,
+        "convTransposeKernelH": "dim(shapes.w, 2)",
+        "convTransposeKernelW": "dim(shapes.w, 3)",
+        "convTransposeInputW": "dim(shapes.x, 3)",
+        "convTransposeOutputH": "dim(shapes.y, 2)",
+        "convTransposeOutputW": "dim(shapes.y, 3)"
+      },
+      "intermediates": [
         {
+          "id": "wt",
+          "dtype": "float32",
+          "shape": "[dim(shapes.w, 0) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)]"
         }
+      ],
       "passes": [
         {
+          "id": "reorder",
+          "name": "ConvTransposeWeightReorder",
+          "shader": "conv-transpose2d-weight-reorder.wgsl.jinja",
+          "bindings": ["wsrc", "wt", "params_reorder"],
           "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.w)), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
+        },
+        {
+          "id": "gemm",
+          "name": "ConvTransposeExactOutputGemm",
+          "shader": "conv-1x1-gemm-tiled-reg.wgsl.jinja",
+          "bindings": ["w_t", "xm", "y", "bias", "params_gemm"],
+          "dispatch": { "x": "ceilDiv(gemmN, 64)", "y": "ceilDiv(gemmM, 64)", "z": "dim(shapes.x, 0)" }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,35 +1,35 @@
 {
   "name": "ai.onnx.ConvTranspose",
-  "id": "_ai_onnx_convtranspose_webgpu_56773ae",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "dDVq94oHWtxnmS2Wl6meZDLrI+Y9O97/LkYR48NM3V0=",
-      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "XGiW/6HmanaqmNZH0rmwfYTirJyt8+/EwXnirJFyBHc=",
-      "conv-1x1-gemm-tiled.wgsl.jinja": "iFg95DLXfwwm52C2vQZPQ2f2NwKnQ/YkeKfXTcqcRTw=",
-      "conv-1x1-subgroup-matrix.wgsl.jinja": "6F2l/7FfatK1TB/0lmXl/lQmOw2Tgq384e/Z4A9YTIY=",
-      "conv-transpose-empty-input.wgsl.jinja": "Ui7q8I1UJMw4zfQHWYUErVRg80V/l6f9UMUdsSKThG8=",
-      "conv-transpose1d-ncl.wgsl.jinja": "TVGioAump1OMC6qd8iyd/wuVQoQnMwCDh4+nWWjTKEA=",
       "conv-transpose1d-phase-tiled.wgsl.jinja": "bMqHC0PYMPoDgBX8QjGyF4amRjVyyl2tUgWWXSmJQeY=",
-      "conv-transpose2d-col2im.wgsl.jinja": "hQb4bZVoepiuY7w+IioYT6CPYQmtbmWgrSHjno6lG/w=",
       "conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja": "p4uIVP2LZItD34v/s79/GIozvNDnXb/PqS3tRGFQhe8=",
-      "conv-transpose2d-grouped-stride-phase.wgsl.jinja": "d1fM5Z2hJ8W2zIPwMfSj0udfo1J8VrMzaqfzZScb1ok=",
-      "conv-transpose2d-nchw.wgsl.jinja": "fZJUH7iDzqcapU1Ddjz7o1JP5J2uBB9qCWLE7QNuSAw=",
-      "conv-transpose2d-phase-gemm-tiled.wgsl.jinja": "FW6lWrHoqI76BAw1GtSzOPgFAZPMtxkxITAy/6X7ckM=",
-      "conv-transpose2d-weight-reorder.wgsl.jinja": "tmJqqJBwS6L4d9nqC0maYJyFI593/O3JyaGm5qBlAI8=",
-      "conv-transpose3d-col2im.wgsl.jinja": "IaBjirYXO0HWlomg6qMzalECCUcXO4zQiSHR/zct8WE=",
-      "conv-transpose3d-ncdhw.wgsl.jinja": "kzq3eqznM7d7n2aUgEkWqjog0zGW6Pdsxynwq7sCTws=",
-      "conv-transpose3d-weight-reorder.wgsl.jinja": "XBM0oifIImJjNT8OoMyutVpoPvE9iGebry251Gcanls=",
-      "manifest.json": "wY9St4yXt/42sQaWNvjxkzpY0slhkzuGR6yLd8/VtG0=",
-      "test.json": "l57YIRQh+ROu5iMw+EoPUOGz8/DplIWgGtvYu8ALGUw="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "empty_input_reduction": ["conv-transpose-empty-input.wgsl.jinja"],
       "empty_input_reduction_bias": ["conv-transpose-empty-input.wgsl.jinja"],
@@ -54,28 +54,22 @@
       "nchw2d_phase_gemm_tiled": ["conv-transpose2d-phase-gemm-tiled.wgsl.jinja"],
       "nchw2d_phase_gemm_tiled_bias": ["conv-transpose2d-phase-gemm-tiled.wgsl.jinja"],
       "nchw2d_grouped_stride_phase_blocked": ["conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja"],
-      "nchw2d_grouped_stride_phase_blocked_bias": ["conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja"],
       "nchw2d_grouped_stride_phase": ["conv-transpose2d-grouped-stride-phase.wgsl.jinja"],
       "nchw2d_grouped_stride_phase_bias": ["conv-transpose2d-grouped-stride-phase.wgsl.jinja"],
       "ncl1d": ["conv-transpose1d-ncl.wgsl.jinja"],
-      "ncl1d_f16": ["conv-transpose1d-ncl.wgsl.jinja"],
       "ncl1d_bias": ["conv-transpose1d-ncl.wgsl.jinja"],
-      "ncl1d_bias_f16": ["conv-transpose1d-ncl.wgsl.jinja"],
       "ncl1d_grouped": ["conv-transpose1d-ncl.wgsl.jinja"],
-      "ncl1d_grouped_f16": ["conv-transpose1d-ncl.wgsl.jinja"],
       "ncl1d_grouped_bias": ["conv-transpose1d-ncl.wgsl.jinja"],
-      "ncl1d_grouped_bias_f16": ["conv-transpose1d-ncl.wgsl.jinja"],
       "nchw2d": ["conv-transpose2d-nchw.wgsl.jinja"],
       "nchw2d_bias": ["conv-transpose2d-nchw.wgsl.jinja"],
       "ncdhw3d": ["conv-transpose3d-ncdhw.wgsl.jinja"],
       "ncdhw3d_bias": ["conv-transpose3d-ncdhw.wgsl.jinja"],
       "nchw2d_grouped": ["conv-transpose2d-nchw.wgsl.jinja"],
-      "nchw2d_grouped_f16": ["conv-transpose2d-nchw.wgsl.jinja"],
-      "nchw2d_grouped_bias_f16": ["conv-transpose2d-nchw.wgsl.jinja"],
       "nchw2d_grouped_bias": ["conv-transpose2d-nchw.wgsl.jinja"],
-      "nchw2d_f16": ["conv-transpose2d-nchw.wgsl.jinja"],
-      "nchw2d_bias_f16": ["conv-transpose2d-nchw.wgsl.jinja"],
-      "nchw2d_depthwise_pointwise_vec4": ["conv-transpose2d-nchw.wgsl.jinja"]
     }
   }
 }

 {
   "name": "ai.onnx.ConvTranspose",
+  "id": "_ai_onnx_convtranspose_webgpu_0e6933f",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "vJsBuZpz7kxmXlf9dInpQAZItXPZxN+5xPRrDmUUfrY=",
+      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "oOxUhgEAuPbZrr0h5rD0E7KnTqroZOk/lVT4AYQjngA=",
+      "conv-1x1-gemm-tiled.wgsl.jinja": "9xyemcOew6wTxIZw0vCzuFUNrx1kouVK4qrmKefko80=",
+      "conv-1x1-subgroup-matrix.wgsl.jinja": "rRCXC/HSy3WXK5TzvmRTWkNTAZ80EpnekwWU83oTrkI=",
+      "conv-transpose-empty-input.wgsl.jinja": "I5t4SroHcl3JqmeiggiIDwjUEGTPLyyX27ShKVhw33c=",
+      "conv-transpose1d-ncl.wgsl.jinja": "rMEEIJ8o3ag4cADd0WmCzaxMlTzKRvKvuvtyCMyknso=",
       "conv-transpose1d-phase-tiled.wgsl.jinja": "bMqHC0PYMPoDgBX8QjGyF4amRjVyyl2tUgWWXSmJQeY=",
+      "conv-transpose2d-col2im.wgsl.jinja": "fyDOZ6oCcj0h4R05wvmNKBnCKf+7xXXZpL7fQF15zMI=",
       "conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja": "p4uIVP2LZItD34v/s79/GIozvNDnXb/PqS3tRGFQhe8=",
+      "conv-transpose2d-grouped-stride-phase.wgsl.jinja": "3M2QhoCoMwtaATTUwz9jphQMRex0DEurNOimwlAW8Tc=",
+      "conv-transpose2d-nchw.wgsl.jinja": "C1ftXjB5oTtPuHLtKKVwshB1AMU+YXMArwmyS9NuTKo=",
+      "conv-transpose2d-phase-gemm-tiled.wgsl.jinja": "ftnzjnquOMpNm0hesCDaCp4hNwRrZwZp+b5RtBdUn9w=",
+      "conv-transpose2d-weight-reorder.wgsl.jinja": "KySJAoLUvLxoaTVjyEndrtcrS5hrlD/kYR8mWe/FSbY=",
+      "conv-transpose3d-col2im.wgsl.jinja": "GWbXZCcp7MorH/SUgtpiMlJKQEZK4Djq6VqqqYy3cVM=",
+      "conv-transpose3d-ncdhw.wgsl.jinja": "E82Z/wkc6hQYhWbuzEEU7YBptljE8DZpy0HPhOAfguc=",
+      "conv-transpose3d-weight-reorder.wgsl.jinja": "RJ5NvBPIZM3y61n+pvcQ7Vb+YpMEg6yR6gcjc1qUcQ0=",
+      "manifest.json": "BEwtw+Cge4ZoOyuZle8VlOGJjCo/tF01SCY4n6/CZj4=",
+      "test.json": "83npgDmD6vc3C8HyVruLUbPLuMIjDznYNU8oq2G9nek="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
       "empty_input_reduction": ["conv-transpose-empty-input.wgsl.jinja"],
       "empty_input_reduction_bias": ["conv-transpose-empty-input.wgsl.jinja"],
       "nchw2d_phase_gemm_tiled": ["conv-transpose2d-phase-gemm-tiled.wgsl.jinja"],
       "nchw2d_phase_gemm_tiled_bias": ["conv-transpose2d-phase-gemm-tiled.wgsl.jinja"],
       "nchw2d_grouped_stride_phase_blocked": ["conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja"],
       "nchw2d_grouped_stride_phase": ["conv-transpose2d-grouped-stride-phase.wgsl.jinja"],
+      "nchw2d_grouped_stride_phase_blocked_bias": ["conv-transpose2d-grouped-stride-phase-blocked.wgsl.jinja"],
       "nchw2d_grouped_stride_phase_bias": ["conv-transpose2d-grouped-stride-phase.wgsl.jinja"],
       "ncl1d": ["conv-transpose1d-ncl.wgsl.jinja"],
       "ncl1d_bias": ["conv-transpose1d-ncl.wgsl.jinja"],
       "ncl1d_grouped": ["conv-transpose1d-ncl.wgsl.jinja"],
       "ncl1d_grouped_bias": ["conv-transpose1d-ncl.wgsl.jinja"],
       "nchw2d": ["conv-transpose2d-nchw.wgsl.jinja"],
       "nchw2d_bias": ["conv-transpose2d-nchw.wgsl.jinja"],
       "ncdhw3d": ["conv-transpose3d-ncdhw.wgsl.jinja"],
       "ncdhw3d_bias": ["conv-transpose3d-ncdhw.wgsl.jinja"],
       "nchw2d_grouped": ["conv-transpose2d-nchw.wgsl.jinja"],
       "nchw2d_grouped_bias": ["conv-transpose2d-nchw.wgsl.jinja"],
+      "nchw2d_depthwise_pointwise_vec4": ["conv-transpose2d-nchw.wgsl.jinja"],
+      "nchw2d_exact_store_gemm_f32": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-transpose2d-weight-reorder.wgsl.jinja"],
+      "nchw2d_exact_store_gemm_f32_bias": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-transpose2d-weight-reorder.wgsl.jinja"]
     }
   }
 }

build/webgpu/test.json CHANGED Viewed

@@ -5,7 +5,12 @@
     "ort_default_strides_dilations_multi_channel_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17],
     "onnx_backend_convtranspose_input_w": [1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1],
     "f16_1d_nobias_input_x": [0.1, 0.5, -0.25, 0.75, 1, -0.5, 0.25, -1, 0.5, -0.75, 1.25, -0.25, 0, 0.5, -1, 0.75],
-    "f16_1d_nobias_input_w": [0.5, -0.25, 1, 0.75, -0.5, 0.25, -1, 0.5, 0.125, 1, -0.5, 0.25, 0.5, 0.75, -0.25, -0.75, 0.5, 1]
   },
   "cases": [
     {
@@ -133,7 +138,7 @@
     {
       "name": "phase_tiled_1d_register_tile_channel_position_tails",
       "provenance": {
-        "notes": "Crosses the occupancy-derived two-output register-tile threshold while retaining input-channel, output-channel, and position tails."
       },
       "attrs": { "strides": [2], "pads": [1, 1] },
       "inputs": {
@@ -1323,7 +1328,7 @@
     {
       "name": "f16_phase_gemm_k4s2_channel_and_spatial_tails",
       "provenance": {
-        "notes": "Compact route lock for the fused four-phase k4s2 GEMM path. Odd channel and spatial sizes exercise K, M, and N tile tails plus all padded boundaries; offset operands keep each accumulation proportional to its tap count."
       },
       "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -1368,7 +1373,7 @@
     {
       "name": "f16_gemm_2d_no_bias",
       "provenance": {
-        "notes": "f16 2D ConvTranspose, group 1, IC%32==0 and IH*IW%64==0 -> exercises the im2col-GEMM path nchw2d_gemm (subgroup-matrix) AND the direct nchw2d_f16 kernel."
       },
       "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -1388,7 +1393,7 @@
     {
       "name": "f16_gemm_2d_bias",
       "provenance": {
-        "notes": "f16 2D ConvTranspose, group 1, IC%32==0 and IH*IW%64==0 -> exercises the im2col-GEMM path nchw2d_gemm_bias (subgroup-matrix) AND the direct nchw2d_f16 kernel."
       },
       "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
       "inputs": {
@@ -1586,7 +1591,7 @@
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 1, 1], "tolerance": 0.000001 } }
     },
     {
-      "name": "dispatch_cliff_batch_over_65535_bias_forces_nchw_scalar",
       "provenance": {
         "notes": "Batch 65536 exceeds maxComputeWorkgroupsPerDimension, so gemmDispatchFits is false and every col2im GEMM variant drops out; the scalar NCHW kernel with its bias epilogue is the only remaining f32 dense path."
       },
@@ -1603,9 +1608,9 @@
       "outputs": { "y": { "dtype": "float32", "shape": [65536, 1, 1, 1], "tolerance": 0.00001 } }
     },
     {
-      "name": "dispatch_cliff_batch_over_65535_bias_f16_forces_nchw_scalar",
       "provenance": {
-        "notes": "For a float16 batch above the dispatch limit, matrix paths are ineligible and the geometry does not admit stride-2 phase GEMM, selecting the scalar NCHW bias path."
       },
       "attrs": {},
       "inputs": {
@@ -1690,7 +1695,7 @@
     {
       "name": "convtranspose3d_subgroup_matrix_bias_f32",
       "provenance": {
-        "notes": "An aligned rank-5 pointwise volumetric matrix path isolates the bias epilogue while preserving M/K/N tile eligibility."
       },
       "attrs": {},
       "inputs": {
@@ -1711,7 +1716,7 @@
     {
       "name": "convtranspose3d_tiled_reg_large_spatial_f32",
       "provenance": {
-        "notes": "A portable register-tiled rank-5 case with N=8*32*32 and M=32 uses pointwise weights to isolate the large matrix path and three-dimensional tensor indexing."
       },
       "attrs": {},
       "inputs": {
@@ -1774,7 +1779,7 @@
     {
       "name": "convtranspose3d_implicit_sgmat_k3_nopad_bias_f32",
       "provenance": {
-        "notes": "Zero-pad stride-1 k3 growth shape with bias: the forward equivalence uses full begin pads of kernel-1 on every axis and the bias epilogue lands on directly stored output tiles."
       },
       "inputs": {
         "x": {
@@ -2004,6 +2009,366 @@
         "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.75] } }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 0, 2, 2], "tolerance": 0 } }
     }
   ]
 }

     "ort_default_strides_dilations_multi_channel_input_x": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17],
     "onnx_backend_convtranspose_input_w": [1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1],
     "f16_1d_nobias_input_x": [0.1, 0.5, -0.25, 0.75, 1, -0.5, 0.25, -1, 0.5, -0.75, 1.25, -0.25, 0, 0.5, -1, 0.75],
+    "f16_1d_nobias_input_w": [0.5, -0.25, 1, 0.75, -0.5, 0.25, -1, 0.5, 0.125, 1, -0.5, 0.25, 0.5, 0.75, -0.25, -0.75, 0.5, 1],
+    "f16_phase_perimeter_h1_w1_first_nan_input_w": [0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, "NaN", 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125],
+    "f16_phase_perimeter_h1_w5_last_nan_input_x": [0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125],
+    "f16_phase_perimeter_h1_w5_last_infinity_input_w": [0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, "Infinity"],
+    "f16_phase_perimeter_h4_w1_first_infinity_input_x": [0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375],
+    "f16_phase_perimeter_h3_w5_first_nan_input_x": [0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375, 0.40625, 0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875, 0.21875, 0.25, 0.28125, 0.3125, 0.34375, 0.375]
   },
   "cases": [
     {
     {
       "name": "phase_tiled_1d_register_tile_channel_position_tails",
       "provenance": {
+        "notes": "Input-channel, output-channel and position tails exercise transposed convolution at a large output shape."
       },
       "attrs": { "strides": [2], "pads": [1, 1] },
       "inputs": {
     {
       "name": "f16_phase_gemm_k4s2_channel_and_spatial_tails",
       "provenance": {
+        "notes": "A stride-two transposed convolution with a four-tap kernel, odd channel and spatial sizes checks padding and reduction tails. Offset operands make a missing tap observable."
       },
       "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
       "inputs": {
     {
       "name": "f16_gemm_2d_no_bias",
       "provenance": {
+        "notes": "f16 2D ConvTranspose, group 1, IC%32==0 and IH*IW%64==0 -> exercises the im2col-GEMM path nchw2d_gemm (subgroup-matrix) AND the direct nchw2d kernel."
       },
       "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
       "inputs": {
     {
       "name": "f16_gemm_2d_bias",
       "provenance": {
+        "notes": "f16 2D ConvTranspose, group 1, IC%32==0 and IH*IW%64==0 -> exercises the im2col-GEMM path nchw2d_gemm_bias (subgroup-matrix) AND the direct nchw2d_bias kernel."
       },
       "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
       "inputs": {
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 5, 1, 1], "tolerance": 0.000001 } }
     },
     {
+      "name": "dispatch_cliff_batch_over_65535_bias_nchw_scalar",
       "provenance": {
         "notes": "Batch 65536 exceeds maxComputeWorkgroupsPerDimension, so gemmDispatchFits is false and every col2im GEMM variant drops out; the scalar NCHW kernel with its bias epilogue is the only remaining f32 dense path."
       },
       "outputs": { "y": { "dtype": "float32", "shape": [65536, 1, 1, 1], "tolerance": 0.00001 } }
     },
     {
+      "name": "dispatch_cliff_batch_over_65535_bias_f16_nchw_scalar",
       "provenance": {
+        "notes": "A batch of 65,536 float16 rows, beyond a single dispatch grid, checks a 1x1 transposed convolution with a single-element bias at default stride 1."
       },
       "attrs": {},
       "inputs": {
     {
       "name": "convtranspose3d_subgroup_matrix_bias_f32",
       "provenance": {
+        "notes": "A 3-D pointwise (1x1x1) transposed convolution with 32 input and 32 output channels uses constant inputs (1.0), weights (0.125) and bias (0.5); every output equals exactly 4.5, directly checking the accumulation and bias addition."
       },
       "attrs": {},
       "inputs": {
     {
       "name": "convtranspose3d_tiled_reg_large_spatial_f32",
       "provenance": {
+        "notes": "A rank-five pointwise transposed convolution with N=8*32*32 and M=32 checks a large spatial extent and three-dimensional tensor indexing."
       },
       "attrs": {},
       "inputs": {
     {
       "name": "convtranspose3d_implicit_sgmat_k3_nopad_bias_f32",
       "provenance": {
+        "notes": "A zero-pad, stride-one 3D transposed convolution grows each spatial axis and applies channel bias to the output."
       },
       "inputs": {
         "x": {
         "bias": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.75] } }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1, 1, 0, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_phase_perimeter_h1_w1_first_nan",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 1, 1],
+          "data": { "kind": "values", "values": [0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875] }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h1_w1_first_nan_input_w" }
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [2, 2, 2, 2], "tolerance": 0.001, "allowNaN": true, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h1_w1_first_neg_infinity",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 1, 1],
+          "data": { "kind": "values", "values": [0.03125, 0.0625, 0.09375, 0.125, 0.15625, 0.1875] }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, "-Infinity", 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125]
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 2, 2], "tolerance": 0.001, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h1_w5_last_nan",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 1, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h1_w5_last_nan_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, "NaN"]
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [2, 2, 2, 10], "tolerance": 0.001, "allowNaN": true, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h1_w5_last_infinity",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 1, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h1_w5_last_nan_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h1_w5_last_infinity_input_w" }
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 2, 10], "tolerance": 0.001, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h4_w1_first_infinity",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 1],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h4_w1_first_infinity_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, "Infinity", 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125]
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 8, 2], "tolerance": 0.001, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h4_w1_last_neg_infinity",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 4, 1],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h4_w1_first_infinity_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, 0.078125, 0.09375, 0.109375, 0.015625, 0.03125, 0.046875, 0.0625, "-Infinity"]
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 8, 2], "tolerance": 0.001, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h3_w5_first_nan",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h3_w5_first_nan_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h1_w1_first_nan_input_w" }
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float16", "shape": [2, 2, 6, 10], "tolerance": 0.001, "allowNaN": true, "relTolerance": 0.002 }
+      },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "f16_phase_perimeter_h3_w5_last_infinity",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h3_w5_first_nan_input_x" }
+          }
+        },
+        "w": {
+          "dtype": "float16",
+          "shape": [3, 2, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/f16_phase_perimeter_h1_w5_last_infinity_input_w" }
+          }
+        },
+        "bias": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.25, -0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2, 6, 10], "tolerance": 0.001, "relTolerance": 0.002 } },
+      "provenance": {
+        "notes": "A non-finite weight at the first or last tap of the final input/output channel is unreachable for some perimeter outputs. The phase GEMM must omit those taps instead of multiplying synthetic padding zeros by NaN/Infinity. Batch two and one-row/one-column cases cover phase addressing and degenerate axes."
+      }
+    },
+    {
+      "name": "exact_store_tail_b2_c17_m5_k2x3_nobias",
+      "attrs": { "strides": [2, 3] },
+      "tunables": { "GEMM_REG_MIN_TILE_PRODUCT": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 17, 7, 11],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125, 0.75, -1.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [17, 5, 2, 3],
+          "data": { "kind": "cycle", "values": [-0.5, 0.25, 0.75, -0.125, 0.5, 0.0625] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 14, 33], "tolerance": 0, "relTolerance": 0 } },
+      "provenance": {
+        "source": "derived",
+        "notes": "Nonoverlapping transposed convolution with batch, contraction, matrix-row and spatial tails. Dyadic inputs keep the CPU reference exact on a compact tensor."
+      }
+    },
+    {
+      "name": "exact_store_tail_b2_c17_m5_k2x3_bias",
+      "attrs": { "strides": [2, 3] },
+      "tunables": { "GEMM_REG_MIN_TILE_PRODUCT": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 17, 7, 11],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125, 0.75, -1.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [17, 5, 2, 3],
+          "data": { "kind": "cycle", "values": [-0.5, 0.25, 0.75, -0.125, 0.5, 0.0625] }
+        },
+        "bias": { "dtype": "float32", "shape": [5], "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 14, 33], "tolerance": 0, "relTolerance": 0 } },
+      "provenance": {
+        "source": "derived",
+        "notes": "Nonoverlapping transposed convolution with batch, contraction, matrix-row and spatial tails. Dyadic inputs keep the CPU reference exact on a compact tensor."
+      }
+    },
+    {
+      "name": "exact_store_tail_b1_c19_m7_k3x2_nobias",
+      "attrs": { "strides": [3, 2] },
+      "tunables": { "GEMM_REG_MIN_TILE_PRODUCT": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 19, 5, 13],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125, 0.75, -1.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [19, 7, 3, 2],
+          "data": { "kind": "cycle", "values": [-0.5, 0.25, 0.75, -0.125, 0.5, 0.0625] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 7, 15, 26], "tolerance": 0, "relTolerance": 0 } },
+      "provenance": {
+        "source": "derived",
+        "notes": "Nonoverlapping transposed convolution with batch, contraction, matrix-row and spatial tails. Dyadic inputs keep the CPU reference exact on a compact tensor."
+      }
+    },
+    {
+      "name": "exact_store_tail_b1_c19_m7_k3x2_bias",
+      "attrs": { "strides": [3, 2] },
+      "tunables": { "GEMM_REG_MIN_TILE_PRODUCT": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 19, 5, 13],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125, 0.75, -1.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [19, 7, 3, 2],
+          "data": { "kind": "cycle", "values": [-0.5, 0.25, 0.75, -0.125, 0.5, 0.0625] }
+        },
+        "bias": { "dtype": "float32", "shape": [7], "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 7, 15, 26], "tolerance": 0, "relTolerance": 0 } },
+      "provenance": {
+        "source": "derived",
+        "notes": "Nonoverlapping transposed convolution with batch, contraction, matrix-row and spatial tails. Dyadic inputs keep the CPU reference exact on a compact tensor."
+      }
+    },
+    {
+      "name": "exact_store_tail_kw2_last_infinity",
+      "attrs": { "strides": [2, 2] },
+      "tunables": { "GEMM_REG_MIN_TILE_PRODUCT": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 17, 3, 5],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125, 0.75, -1.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [17, 5, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, "Infinity"]
+          }
+        },
+        "bias": { "dtype": "float32", "shape": [5], "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 5, 6, 10], "tolerance": 0, "relTolerance": 0 } },
+      "provenance": {
+        "source": "derived",
+        "notes": "A nonfinite final weight exercises the partial contraction tile and spatial tap mapping without contaminating unrelated output channels or taps."
+      }
+    },
+    {
+      "name": "exact_store_tail_kw3_last_nan",
+      "attrs": { "strides": [2, 3] },
+      "tunables": { "GEMM_REG_MIN_TILE_PRODUCT": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 17, 3, 5],
+          "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125, 0.75, -1.0] }
+        },
+        "w": {
+          "dtype": "float32",
+          "shape": [17, 5, 2, 3],
+          "data": {
+            "kind": "values",
+            "values": [0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, 0.125, "NaN"]
+          }
+        },
+        "bias": { "dtype": "float32", "shape": [5], "data": { "kind": "cycle", "values": [0.25, -0.5, 0.125] } }
+      },
+      "outputs": {
+        "y": { "dtype": "float32", "shape": [2, 5, 6, 15], "tolerance": 0, "relTolerance": 0, "allowNaN": true }
+      },
+      "provenance": {
+        "source": "derived",
+        "notes": "A nonfinite final weight exercises the partial contraction tile and spatial tap mapping without contaminating unrelated output channels or taps."
+      }
     }
   ]
 }