sync 6fdf6301e2bb

Browse files

Files changed (14) hide show

README.md +22 -10
build/webgpu/bench.json +0 -0
build/webgpu/conv-1x1-channel-reduce.wgsl.jinja +9 -12
build/webgpu/conv-1x1-gemm-tiled-reg.wgsl.jinja +19 -158
build/webgpu/conv-1x1-gemm-tiled.wgsl.jinja +13 -50
build/webgpu/conv-1x1-subgroup-matrix.wgsl.jinja +12 -146
build/webgpu/conv-direct-nd.wgsl.jinja +10 -44
build/webgpu/conv-direct-unrolled.wgsl.jinja +10 -43
build/webgpu/conv-splitk-reduce.wgsl.jinja +9 -43
build/webgpu/conv1d-tiled-reg.wgsl.jinja +0 -34
build/webgpu/conv2d-grouped-large-w4.wgsl.jinja +17 -44
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +19 -17
build/webgpu/test.json +0 -0

README.md CHANGED Viewed

@@ -53,20 +53,24 @@ Attributes and default values (overridable per request):
 One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
 - `pointwise_channel_reduce` — Handles pointwise unit-stride convolution with few output channels by assigning one invocation to each spatial position and accumulating every output channel in registers. It does not materialize im2col scratch.
 - `pointwise_channel_reduce_bias` — Handles pointwise unit-stride convolution with few output channels by assigning one invocation to each spatial position and accumulating every output channel in registers. It does not materialize im2col scratch.
 - `gemm_1x1_subgroup_matrix` — Multiplies 1x1 weights directly by NCHW input. Complete output-channel tiles use direct subgroup-matrix loads, eliminating input staging and reduction-loop barriers; partial tiles retain guarded staging. Tile geometry follows the device's workgroup limits.
 - `gemm_1x1_subgroup_matrix_bias` — Multiplies 1x1 weights directly by NCHW input. Complete output-channel tiles use direct subgroup-matrix loads, eliminating input staging and reduction-loop barriers; partial tiles retain guarded staging. Tile geometry follows the device's workgroup limits.
 - `im2col_gemm_subgroup_matrix` — Materializes convolution windows as a column matrix, then multiplies it by the weights with subgroup-matrix operations. It applies when reduction and spatial extents are tile-aligned and column scratch fits.
 - `im2col_gemm_subgroup_matrix_bias` — Materializes convolution windows as a column matrix, then multiplies it by the weights with subgroup-matrix operations. It applies when reduction and spatial extents are tile-aligned and column scratch fits.
-- `im2col_direct_f32_subgroup_matrix` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
-- `im2col_direct_f16_subgroup_matrix` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
-- `im2col_direct_f32_subgroup_matrix_bias` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
-- `im2col_direct_f16_subgroup_matrix_bias` — Materialize tile-aligned convolution columns and load weights and columns directly into subgroup matrices, removing operand staging and K-loop barriers. Retain f32 accumulation, existing split-K reduction order, and the implicit gather preference below the measured channel-tile reuse threshold.
-- `implicit_im2col_subgroup_matrix` — Gathers each logical im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth that materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
-- `implicit_im2col_subgroup_matrix_bias` — Gathers each logical im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth that materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
-- `implicit_im2col3d_subgroup_matrix` — Gathers each logical 3-D im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth a volumetric materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
-- `implicit_im2col3d_subgroup_matrix_bias` — Gathers each logical 3-D im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth a volumetric materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
 - `implicit_im2col_subgroup_matrix_splitk` — Partitions the implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums from scratch. The combine reassociates floating-point addition relative to the unsplit route.
 - `implicit_im2col_subgroup_matrix_bias_splitk` — Partitions the implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums from scratch. The combine reassociates floating-point addition relative to the unsplit route.
 - `im2col_gemm_subgroup_matrix_padded` — Materializes a zero-padded column matrix before a subgroup-matrix multiply, allowing logical reduction and spatial extents that are not tile-aligned. Padded values contribute zero and are excluded from the logical output.
@@ -78,6 +82,14 @@ One implementation is selected per call from the device capabilities, the reques
 - `implicit_im2col_tiled_bias_reg` — Gathers logical im2col elements directly from X while filling register-blocked GEMM tiles. It avoids column scratch and provides the portable implicit route without subgroup-matrix support.
 - `implicit_im2col_tiled_reg_splitk` — Partitions the register-blocked implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 - `implicit_im2col_tiled_bias_reg_splitk` — Partitions the register-blocked implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 - `gemm_1x1_tiled_reg_splitk` — Partitions the register-blocked 1x1 GEMM reduction across additional workgroups when the output grid is small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 - `gemm_1x1_tiled_bias_reg_splitk` — Partitions the register-blocked 1x1 GEMM reduction across additional workgroups when the output grid is small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
@@ -90,7 +102,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-channel-reduce.wgsl.jinja`](build/webgpu/conv-1x1-channel-reduce.wgsl.jinja)
 - [`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)
@@ -106,7 +118,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.

 One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `grouped_large_kernel_w4` — Shares each input row window across output channels and columns with f32 accumulation. Multi-channel groups loop over kernel rows to limit generated instructions; depthwise groups retain unrolling. Dilation-aware lanes reuse overlapping taps. Unrolled large kernels are demoted when the adapter does not report a fixed 32-wide subgroup range.
+- `grouped_large_kernel_w4_bias` — Shares each input row window across output channels and columns with f32 accumulation. Multi-channel groups loop over kernel rows to limit generated instructions; depthwise groups retain unrolling. Dilation-aware lanes reuse overlapping taps. Unrolled large kernels are demoted when the adapter does not report a fixed 32-wide subgroup range.
+- `grouped_large_kernel_w4_tail` — Shares each input row window across output channels and columns with f32 accumulation. Multi-channel groups loop over kernel rows to limit generated instructions; depthwise groups retain unrolling. Dilation-aware lanes reuse overlapping taps. Unrolled large kernels are demoted when the adapter does not report a fixed 32-wide subgroup range.
+- `grouped_large_kernel_w4_tail_bias` — Shares each input row window across output channels and columns with f32 accumulation. Multi-channel groups loop over kernel rows to limit generated instructions; depthwise groups retain unrolling. Dilation-aware lanes reuse overlapping taps. Unrolled large kernels are demoted when the adapter does not report a fixed 32-wide subgroup range.
+- `grouped_large_kernel_w4_dilated_lanes` — Shares each input row window across output channels and columns with f32 accumulation. Multi-channel groups loop over kernel rows to limit generated instructions; depthwise groups retain unrolling. Dilation-aware lanes reuse overlapping taps. Unrolled large kernels are demoted when the adapter does not report a fixed 32-wide subgroup range.
+- `grouped_large_kernel_w4_dilated_lanes_bias` — Shares each input row window across output channels and columns with f32 accumulation. Multi-channel groups loop over kernel rows to limit generated instructions; depthwise groups retain unrolling. Dilation-aware lanes reuse overlapping taps. Unrolled large kernels are demoted when the adapter does not report a fixed 32-wide subgroup range.
 - `pointwise_channel_reduce` — Handles pointwise unit-stride convolution with few output channels by assigning one invocation to each spatial position and accumulating every output channel in registers. It does not materialize im2col scratch.
 - `pointwise_channel_reduce_bias` — Handles pointwise unit-stride convolution with few output channels by assigning one invocation to each spatial position and accumulating every output channel in registers. It does not materialize im2col scratch.
 - `gemm_1x1_subgroup_matrix` — Multiplies 1x1 weights directly by NCHW input. Complete output-channel tiles use direct subgroup-matrix loads, eliminating input staging and reduction-loop barriers; partial tiles retain guarded staging. Tile geometry follows the device's workgroup limits.
 - `gemm_1x1_subgroup_matrix_bias` — Multiplies 1x1 weights directly by NCHW input. Complete output-channel tiles use direct subgroup-matrix loads, eliminating input staging and reduction-loop barriers; partial tiles retain guarded staging. Tile geometry follows the device's workgroup limits.
 - `im2col_gemm_subgroup_matrix` — Materializes convolution windows as a column matrix, then multiplies it by the weights with subgroup-matrix operations. It applies when reduction and spatial extents are tile-aligned and column scratch fits.
 - `im2col_gemm_subgroup_matrix_bias` — Materializes convolution windows as a column matrix, then multiplies it by the weights with subgroup-matrix operations. It applies when reduction and spatial extents are tile-aligned and column scratch fits.
+- `im2col_direct_f32_subgroup_matrix` — Materializes tile-aligned convolution columns and loads weights and columns directly into subgroup matrices, with no operand staging or K-loop barriers. Accumulates in f32 in the split-K reduction order of the other column-matrix routes; below the channel-tile reuse threshold, implicit gathering is used instead.
+- `im2col_direct_f16_subgroup_matrix` — Materializes tile-aligned convolution columns and loads weights and columns directly into subgroup matrices, with no operand staging or K-loop barriers. Accumulates in f32 in the split-K reduction order of the other column-matrix routes; below the channel-tile reuse threshold, implicit gathering is used instead.
+- `im2col_direct_f32_subgroup_matrix_bias` — Materializes tile-aligned convolution columns and loads weights and columns directly into subgroup matrices, with no operand staging or K-loop barriers. Accumulates in f32 in the split-K reduction order of the other column-matrix routes; below the channel-tile reuse threshold, implicit gathering is used instead.
+- `im2col_direct_f16_subgroup_matrix_bias` — Materializes tile-aligned convolution columns and loads weights and columns directly into subgroup matrices, with no operand staging or K-loop barriers. Accumulates in f32 in the split-K reduction order of the other column-matrix routes; below the channel-tile reuse threshold, implicit gathering is used instead.
+- `implicit_im2col_subgroup_matrix` — Gathers each logical 2-D or 3-D im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth that materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
+- `implicit_im2col_subgroup_matrix_bias` — Gathers each logical 2-D or 3-D im2col tile directly from X instead of materializing a column matrix, trading extra address arithmetic for the scratch and bandwidth that materialization would cost. It applies when the output has few output-channel tiles, so the gather is not repeated across many tiles.
 - `implicit_im2col_subgroup_matrix_splitk` — Partitions the implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums from scratch. The combine reassociates floating-point addition relative to the unsplit route.
 - `implicit_im2col_subgroup_matrix_bias_splitk` — Partitions the implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums from scratch. The combine reassociates floating-point addition relative to the unsplit route.
 - `im2col_gemm_subgroup_matrix_padded` — Materializes a zero-padded column matrix before a subgroup-matrix multiply, allowing logical reduction and spatial extents that are not tile-aligned. Padded values contribute zero and are excluded from the logical output.
 - `implicit_im2col_tiled_bias_reg` — Gathers logical im2col elements directly from X while filling register-blocked GEMM tiles. It avoids column scratch and provides the portable implicit route without subgroup-matrix support.
 - `implicit_im2col_tiled_reg_splitk` — Partitions the register-blocked implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 - `implicit_im2col_tiled_bias_reg_splitk` — Partitions the register-blocked implicit-im2col reduction across workgroups when the unsplit output grid is too small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
+- `im2col_gemm_tiled_reg` — Materialized f32 columns feed the configured tiled GEMM with f32 accumulation.
+- `im2col_gemm_tiled` — Materialized f32 columns feed the configured tiled GEMM with f32 accumulation.
+- `im2col_gemm_tiled_bias_reg` — Materialized f32 columns feed the configured tiled GEMM with f32 accumulation. Bias is applied before the final output cast.
+- `im2col_gemm_tiled_bias` — Materialized f32 columns feed the configured tiled GEMM with f32 accumulation. Bias is applied before the final output cast.
+- `im2col_gemm_tiled_reg_f16_columns` — Exact f16 columns feed the register-tiled GEMM with f32 accumulation. Total batched output workgroups determine eligibility. The row tile rounds output channels to 32 and caps at 64, reducing padding for narrow outputs. Actual tile storage and invocation geometry must fit the WebGPU limits.
+- `im2col_gemm_tiled_f16_columns` — Exact f16 columns feed the smaller tiled GEMM with f32 accumulation. Short batched spatial workloads use the f32 column layout when it fits; this half-sized layout serves constrained buffers. The spatial-work crossover is independent of the gather workgroup size.
+- `im2col_gemm_tiled_bias_reg_f16_columns` — Exact f16 columns feed register-tiled GEMM with f32 accumulation. Batched output work determines eligibility; channel tiles round to 32 and cap at 64. Actual storage and invocation geometry must fit WebGPU limits. One bias load per microtile row feeds its columns before the final output cast.
+- `im2col_gemm_tiled_bias_f16_columns` — Exact f16 columns feed the smaller tiled GEMM with f32 accumulation. Short batched spatial workloads use the f32 column layout when it fits; this half-sized layout serves constrained buffers. The spatial-work crossover is independent of the gather workgroup size. Bias is applied before the final output cast.
 - `gemm_1x1_tiled_reg_splitk` — Partitions the register-blocked 1x1 GEMM reduction across additional workgroups when the output grid is small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 - `gemm_1x1_tiled_bias_reg_splitk` — Partitions the register-blocked 1x1 GEMM reduction across additional workgroups when the output grid is small, then combines raw partial sums. The combine reassociates floating-point addition relative to the unsplit route.
 - [`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-channel-reduce.wgsl.jinja`](build/webgpu/conv-1x1-channel-reduce.wgsl.jinja)
 - [`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)
 ## 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

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

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

@@ -1,3 +1,11 @@
 // Pointwise convolution with few output channels is evaluated as a per-position
 // channel reduction instead of a tiled GEMM. One invocation walks the input
 // channels and accumulates every output channel in a statically named local.
@@ -6,12 +14,6 @@
 //
 // The output-channel accumulators are named explicitly. Selection caps the
 // output-channel count, bounding the invocation-local accumulator storage.
-{% set usesF16 = usesF16 is defined and usesF16 %}
-{% set batchCount = batchCount | default(0) %}
-{% set yScalar = yScalar | default("f32") %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const IN_C: u32 = {{ inChannels }}u;
@@ -22,12 +24,7 @@ const WORKGROUP_SIZE: u32 = {{ channelReduceWorkgroupSize }}u;
 @compute @workgroup_size({{ channelReduceWorkgroupSize }}, 1, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past
-  // the per-axis dispatch fold width and reduces to the 1D form at y=0.
-  let idx = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WORKGROUP_SIZE;
-  if (idx >= COUNT) {
-    return;
-  }
   let batch = idx / PLANE;
   let pos = idx % PLANE;

+{% 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 %}
 // Pointwise convolution with few output channels is evaluated as a per-position
 // channel reduction instead of a tiled GEMM. One invocation walks the input
 // channels and accumulates every output channel in a statically named local.
 //
 // The output-channel accumulators are named explicitly. Selection caps the
 // output-channel count, bounding the invocation-local accumulator storage.
 {{ env.wgsl.resourceDeclarations }}
 const IN_C: u32 = {{ inChannels }}u;
 @compute @workgroup_size({{ channelReduceWorkgroupSize }}, 1, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("WORKGROUP_SIZE", "idx", "COUNT") }}
   let batch = idx / PLANE;
   let pos = idx % PLANE;

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

@@ -14,44 +14,13 @@
 // 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;
@@ -70,11 +39,6 @@ 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 kTiles = kTiles | default(0) %}
@@ -84,7 +48,7 @@ const BN_VECS: u32 = BN / 4u;
 {% 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 %}
@@ -116,78 +80,21 @@ const PARTIAL_COLS: u32 = {{ nPadded }}u;
 const PARTIAL_BATCH_STRIDE: u32 = PARTIAL_ROWS * PARTIAL_COLS;
 const PARTIAL_SLICE_STRIDE: u32 = {{ batchCount }}u * PARTIAL_BATCH_STRIDE;
 {% 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 hasBias %}
-  v = v + f32(bias[m]);
-{% endif %}
-{% if hasZ %}
-  v = v + f32(zResidual[yBase + m * params.N + n]);
-{% endif %}
-{% if hasActivation %}
-  v = fused_act(v);
 {% endif %}
   y[yBase + m * params.N + n] = {{ T }}(v);
 }
 {% endif %}
-{% endif %}
 var<workgroup> tileA: array<array<vec4<{{ tileT }}>, AK_VECS>, BM>;
 var<workgroup> tileB: array<array<vec4<{{ tileT }}>, BN_VECS>, BK>;
-{% if fusedNarrowProjection and GEMM_BK < GEMM_BM %}
-// BK16 cannot reuse the 16-row input tile to publish a BM32 intermediate.
-// The extra 32x64 f32 tile keeps the total at 14 KiB, below WebGPU's
-// guaranteed 16 KiB workgroup-storage floor.
-var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
-{% endif %}
-{% macro publish_projection_input(localRow, localVec, component, globalRow, globalColumn, raw) %}
-  if ({{ globalRow }} < M && {{ globalColumn }} < N) {
-    var projectionInput = {{ raw }};
-{% if hasBias %}
-    projectionInput = projectionInput + f32(bias[{{ globalRow }}]);
-{% endif %}
-{% if projectionInputAct == "relu" %}
-    projectionInput = projection_relu(projectionInput);
-{% endif %}
-    {{ narrowProjectionTile }}[{{ localRow }}][{{ localVec }}].{{ component }} = projectionInput;
-  } else {
-    {{ narrowProjectionTile }}[{{ localRow }}][{{ localVec }}].{{ component }} = 0.0;
-  }
-{%- endmacro %}
 {% macro load_a_vec4(rowBase, columnBase, rowLimit, columnLimit) %}
     for (var linear = li; linear < BM * AK_VECS; linear += WG_SIZE) {
       let ar = linear / AK_VECS;
@@ -208,8 +115,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;
@@ -230,8 +136,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;
@@ -259,8 +164,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;
@@ -344,8 +248,7 @@ var<workgroup> projectionTile: array<array<vec4<f32>, BN_VECS>, BM>;
           }
         }
       }
-    }
-{%- endmacro %}
 @compute @workgroup_size({{ GEMM_WG_X }}, {{ GEMM_WG_Y }}, 1)
 fn main(
@@ -423,62 +326,20 @@ fn main(
   let n0 = nBase + lid.x * TN;
   let partialBase = slice * PARTIAL_SLICE_STRIDE + batch * PARTIAL_BATCH_STRIDE;
 {% else %}
-{% if not fusedNarrowProjection %}
   let yBatchBase = batch * M * N;
-{% endif %}
   let m0 = mBase + lid.y * TM;
   let n0 = nBase + lid.x * TN;
 {% endif %}
 {% for row in range(GEMM_TM) %}
 {% for column in range(GEMM_TN) %}
 {% if splitKPartial %}
   y[partialBase + (m0 + {{ row }}u) * PARTIAL_COLS + n0 + {{ column }}u] = acc{{ row }}.{{ components[column] }};
-{% elif 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 = gemmThreadRows if gemmThreadRows is defined else 8 %}
+{% set GEMM_WG_X = 16 %}
 {% set GEMM_BM = gemmMTile if gemmMTile is defined else 64 %}
+{% set GEMM_BN = 64 %}
+{% set GEMM_TN = 4 %}
 {% 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;
 // 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 splitKValue = splitK if splitK is defined else 1 %}
 {% set splitKPartial = splitKValue > 1 %}
 {% set kTiles = kTiles | default(0) %}
 {% set gemmKLimit = "kEnd" if splitKPartial else "K" %}
 {% set implicitGatherMode = tunables.IMPLICIT_TILED_GATHER_MODE if tunables.IMPLICIT_TILED_GATHER_MODE is defined else 0 %}
 {% set implicitGatherWgSize = GEMM_WG_X * GEMM_WG_Y %}
+{% set implicitGatherColumnVecs = 16 %}
 {% 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 %}
 const PARTIAL_BATCH_STRIDE: u32 = PARTIAL_ROWS * PARTIAL_COLS;
 const PARTIAL_SLICE_STRIDE: u32 = {{ batchCount }}u * PARTIAL_BATCH_STRIDE;
 {% endif %}
+{% set emitConvStoreOut = not splitKPartial %}
+{% if emitConvStoreOut | default(true) %}fn store_out(m: u32, n: u32, yBase: u32, raw: f32{% if convPreloadBias | default(false) %}, rowBias: f32{% endif %}) {
   if (m >= params.M || n >= params.N) {
     return;
   }
   var v = raw;
 {% if hasBias %}
+  v = v + {% if convPreloadBias | default(false) %}rowBias{% else %}f32(bias[m]){% endif %};
 {% 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(
   let n0 = nBase + lid.x * TN;
   let partialBase = slice * PARTIAL_SLICE_STRIDE + batch * PARTIAL_BATCH_STRIDE;
 {% else %}
   let yBatchBase = batch * M * N;
   let m0 = mBase + lid.y * TM;
   let n0 = nBase + lid.x * TN;
 {% endif %}
 {% for row in range(GEMM_TM) %}
+{% if convPreloadBias | default(false) %}
+  var rowBias{{ row }} = 0.0;
+  if (m0 + {{ row }}u < M) { rowBias{{ row }} = f32(bias[m0 + {{ row }}u]); }
+{% endif %}
 {% for column in range(GEMM_TN) %}
 {% if splitKPartial %}
   y[partialBase + (m0 + {{ row }}u) * PARTIAL_COLS + n0 + {{ column }}u] = acc{{ row }}.{{ components[column] }};
 {% else %}
+  store_out(m0 + {{ row }}u, n0 + {{ column }}u, yBatchBase, acc{{ row }}.{{ components[column] }}{% if convPreloadBias | default(false) %}, rowBias{{ row }}{% endif %});
 {% endif %}
 {% endfor %}
+{% endfor %}}

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

@@ -6,67 +6,31 @@
 // 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 hasBias %}
-  v = v + f32(bias[m]);
-{% endif %}
-{% if hasZ %}
-  v = v + f32(zResidual[yBase + m * params.N + n]);
-{% endif %}
-{% if hasActivation %}
-  v = fused_act(v);
 {% endif %}
   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>
@@ -78,7 +42,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;
@@ -88,8 +52,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;
@@ -115,10 +79,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 = "f16" if usesF16 else "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 convPreloadBias | default(false) %}, rowBias: f32{% endif %}) {
   if (m >= params.M || n >= params.N) {
     return;
   }
   var v = raw;
 {% if hasBias %}
+  v = v + {% if convPreloadBias | default(false) %}rowBias{% else %}f32(bias[m]){% endif %};
 {% 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({{ 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

@@ -12,18 +12,16 @@
 // 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) %}
@@ -53,46 +51,16 @@ 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 %}
 {% set tileRowsValue = tileRows if tileRows is defined else 32 %}
-{% set tileColsValue = tileCols if tileCols is defined else 64 %}
 {% set workgroupThreadsValue = workgroupThreads if workgroupThreads is defined else 128 %}
 {% set subgroupRowsValue = subgroupRows if subgroupRows is defined else 2 %}
-{% set subgroupColsValue = subgroupCols if subgroupCols is defined else 2 %}
 {% set subRowsValue = (tileRowsValue / subgroupRowsValue)|int %}
-{% set subColsValue = (tileColsValue / subgroupColsValue)|int %}
 {% set bLoadWidth = (tileColsValue * 32 / workgroupThreadsValue)|int %}
 {% set bKChunks = (32 / bLoadWidth)|int %}
@@ -176,7 +144,7 @@ fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
   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 %}
@@ -325,59 +293,13 @@ 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 %}) {
@@ -431,9 +353,7 @@ fn main(
 {% else %}
   let batch = workgroup_id.z;
   let b_base = batch * B_BATCH_STRIDE;
-{% if not fusedNarrowProjection %}
   let c_base = batch * C_BATCH_STRIDE;
-{% endif %}
 {% endif %}
   let a_global_base = workgroup_id.y * TILE_ROWS;
   let b_global_base = workgroup_id.x * TILE_COLS;
@@ -491,57 +411,7 @@ fn main(
 {% 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 splitKPartial %}
   // Raw partials, no bias and no epilogue — the reduce pass owns both. The
   // scratch is whole tiles in both axes, so this needs none of the column or
   // row guards the real output store carries.
@@ -586,11 +456,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 %}

 // 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 = false %}
+{% set polyphase = false %}
 {% 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 = false %}
+{% set hasActivation = false %}
 {% set convKernelH = convKernelH | default(0) %}
 {% set convKernelW = convKernelW | default(0) %}
 {% set convStrideH = convStrideH | default(0) %}
 enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 {% set operandScalar = fScalar %}
 {% set accScalar = "f32" %}
 {% set useDirectMatrixStore = directMatrixStore is defined and directMatrixStore %}
 {% set tileRowsValue = tileRows if tileRows is defined else 32 %}
+{% set tileColsValue = 64 %}
 {% set workgroupThreadsValue = workgroupThreads if workgroupThreads is defined else 128 %}
 {% set subgroupRowsValue = subgroupRows if subgroupRows is defined else 2 %}
 {% set subRowsValue = (tileRowsValue / subgroupRowsValue)|int %}
+{% set subColsValue = 32 %}
 {% set bLoadWidth = (tileColsValue * 32 / workgroupThreadsValue)|int %}
 {% set bKChunks = (32 / bLoadWidth)|int %}
   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 %}
 }
 {% 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 %}) {
 {% else %}
   let batch = workgroup_id.z;
   let b_base = batch * B_BATCH_STRIDE;
   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;
 {% endif %}
   }
+{% if splitKPartial %}
   // Raw partials, no bias and no epilogue — the reduce pass owns both. The
   // scratch is whole tiles in both axes, so this needs none of the column or
   // row guards the real output store carries.
 {% if hasBias %}
     raw += f32(bias[global_row]);
 {% endif %}
     y[y_index] = {{ T }}(raw);
   }
 {% endif %}
 {% if padded %}

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

@@ -1,50 +1,20 @@
 // Direct N-dimensional convolution for channels-first tensors. Inputs are
 // accumulated in f32 and narrowed once at the output.
 // A fused epilogue may apply a residual input and activation before the store.
-{% if usesF16Spec %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% set hasActivation = hasActivation is defined and hasActivation %}
-{% set activation = activation | default("") %}
-{% set actAlpha = actAlpha | default(0.0) %}
-{% set actBeta = actBeta | default(0.0) %}
-{% if hasActivation %}
-// Apply the fused activation in the f32 accumulator before the single output
-// cast, avoiding an intermediate convolution tensor.
-{% macro fused_act_return(mode, alpha, beta) -%}
-{% if mode == "Relu" %}
-  return max(v, 0.0);
-{% elif mode == "Clip" %}
-  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
-{% elif mode == "LeakyRelu" %}
-  return select(v * f32({{ alpha }}), v, v >= 0.0);
-{% elif mode == "Sigmoid" %}
-  return 1.0 / (1.0 + exp(-v));
-{% elif mode == "Tanh" %}
-  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
-  // the builtin preserves that saturated result for larger accumulators.
-  return tanh(clamp(v, -10.0, 10.0));
-{% elif mode == "HardSigmoid" %}
-  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
-{% else %}
-  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
-{% endif %}
-{%- endmacro -%}
-fn fused_act(v: f32) -> f32 {
-{{ fused_act_return(activation, actAlpha, actBeta) -}}
-}
-{% endif %}
-{% set hasZ = hasZ is defined and hasZ %}
 const WG: u32 = {{ convWorkgroupSize }}u;
 @compute @workgroup_size(WG)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
-  if (index >= params.count) {
-    return;
-  }
   let ow = index % params.outW;
   var q = index / params.outW;
@@ -108,9 +78,5 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     }
 {% endif %}
   }
-{% if hasZ %}
-  // Optional residual epilogue: Y = activation(conv + B + Z); Z shares Y's layout.
-  acc = acc + f32(zResidual[index]);
-{% endif %}
-  y[index] = {{ T }}({% if hasActivation %}fused_act(acc){% else %}acc{% endif %});
 }

 // Direct N-dimensional convolution for channels-first tensors. Inputs are
 // accumulated in f32 and narrowed once at the output.
 // A fused epilogue may apply a residual input and activation before the store.
+{% 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 WG: u32 = {{ convWorkgroupSize }}u;
 @compute @workgroup_size(WG)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("WG", "index") }}
   let ow = index % params.outW;
   var q = index / params.outW;
     }
 {% endif %}
   }
+  y[index] = {{ T }}(acc);
 }

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

@@ -6,39 +6,15 @@
 // strides fold to literals; the per-row bounds check is hoisted per kh.
 // Group offsets are applied to both input channels and weights. f16 inputs are
 // widened to an f32 accumulator and narrowed once at store.
 {{ env.wgsl.resourceDeclarations }}
-{% set hasActivation = hasActivation is defined and hasActivation %}
-{% set activation = activation | default("") %}
-{% set actAlpha = actAlpha | default(0.0) %}
-{% set actBeta = actBeta | default(0.0) %}
-{% if hasActivation %}
-// Apply the fused activation in the f32 accumulator before the single output
-// cast, avoiding an intermediate convolution tensor.
-{% macro fused_act_return(mode, alpha, beta) -%}
-{% if mode == "Relu" %}
-  return max(v, 0.0);
-{% elif mode == "Clip" %}
-  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
-{% elif mode == "LeakyRelu" %}
-  return select(v * f32({{ alpha }}), v, v >= 0.0);
-{% elif mode == "Sigmoid" %}
-  return 1.0 / (1.0 + exp(-v));
-{% elif mode == "Tanh" %}
-  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
-  // the builtin preserves that saturated result for larger accumulators.
-  return tanh(clamp(v, -10.0, 10.0));
-{% elif mode == "HardSigmoid" %}
-  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
-{% else %}
-  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
-{% endif %}
-{%- endmacro -%}
-fn fused_act(v: f32) -> f32 {
-{{ fused_act_return(activation, actAlpha, actBeta) -}}
-}
-{% endif %}
-{% set hasZ = hasZ is defined and hasZ %}
 const KERNEL_AREA: u32 = {{ kernelHSpec * kernelWSpec }}u;
 const STRIDE_H: u32 = {{ strideHSpec }}u;
 const STRIDE_W: u32 = {{ strideWSpec }}u;
@@ -47,12 +23,7 @@ const PAD_LEFT: i32 = {{ padLeftSpec }};
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past the
-  // per-axis dispatch fold width (outputs > 16.7M elements).
-  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if (index >= params.count) {
-    return;
-  }
   let ow = index % params.outW;
   var t = index / params.outW;
@@ -92,9 +63,5 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     xBase = xBase + xChannelStride;
     wBase = wBase + KERNEL_AREA;
   }
-{% if hasZ %}
-  // Optional residual epilogue: Y = activation(conv + B + Z); Z shares Y's layout.
-  acc = acc + f32(zResidual[index]);
-{% endif %}
-  y[index] = {{ T }}({% if hasActivation %}fused_act(acc){% else %}acc{% endif %});
 }

 // strides fold to literals; the per-row bounds check is hoisted per kh.
 // Group offsets are applied to both input channels and weights. f16 inputs are
 // widened to an f32 accumulator and narrowed once at store.
+{% 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 KERNEL_AREA: u32 = {{ kernelHSpec * kernelWSpec }}u;
 const STRIDE_H: u32 = {{ strideHSpec }}u;
 const STRIDE_W: u32 = {{ strideWSpec }}u;
 @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;
     xBase = xBase + xChannelStride;
     wBase = wBase + KERNEL_AREA;
   }
+  y[index] = {{ T }}(acc);
 }

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

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

 //
 // Summing slices in index order changes the f32 association relative to one
 // uninterrupted K loop, so the two orders need not be bit-identical.
+{% 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 M: u32 = {{ M }}u;
 const N: u32 = {{ N }}u;
 @compute @workgroup_size({{ reduceWorkgroupSize }}, 1, 1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("WORKGROUP_SIZE", "idx", "COUNT") }}
   let col = idx % N;
   let row = (idx / N) % M;
   let batch = idx / (M * N);
   }
 {% if hasBias %}
   acc = acc + f32(bias[row]);
 {% endif %}
   y[idx] = {{ T }}(acc);
 }

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

@@ -10,37 +10,6 @@
 // through dot(). Bounds checks cover every partial tile. Accumulation, bias,
 // and the optional activation epilogue use f32; f16 operands widen at the FMA.
 {{ env.wgsl.resourceDeclarations }}
-{% set hasActivation = hasActivation is defined and hasActivation %}
-{% set activation = activation | default("") %}
-{% set actAlpha = actAlpha | default(0.0) %}
-{% set actBeta = actBeta | default(0.0) %}
-{% if hasActivation %}
-// Apply the fused activation in the f32 accumulator before the single output
-// cast, avoiding an intermediate convolution tensor.
-{% macro fused_act_return(mode, alpha, beta) -%}
-{% if mode == "Relu" %}
-  return max(v, 0.0);
-{% elif mode == "Clip" %}
-  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
-{% elif mode == "LeakyRelu" %}
-  return select(v * f32({{ alpha }}), v, v >= 0.0);
-{% elif mode == "Sigmoid" %}
-  return 1.0 / (1.0 + exp(-v));
-{% elif mode == "Tanh" %}
-  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
-  // the builtin preserves that saturated result for larger accumulators.
-  return tanh(clamp(v, -10.0, 10.0));
-{% elif mode == "HardSigmoid" %}
-  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
-{% else %}
-  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
-{% endif %}
-{%- endmacro -%}
-fn fused_act(v: f32) -> f32 {
-{{ fused_act_return(activation, actAlpha, actBeta) -}}
-}
-{% endif %}
 {% set tileT = "f16" if usesF16 else "f32" %}
 const WG_X: u32 = {{ convWgX }}u;
@@ -59,9 +28,6 @@ fn store_out(oc: u32, ow: u32, yBase: u32, raw: f32) {
   var v = raw;
 {% if hasBias %}
   v = v + f32(bias[oc]);
-{% endif %}
-{% if hasActivation %}
-  v = fused_act(v);
 {% endif %}
   y[yBase + oc * params.outW + ow] = {{ T }}(v);
 }

 // through dot(). Bounds checks cover every partial tile. Accumulation, bias,
 // and the optional activation epilogue use f32; f16 operands widen at the FMA.
 {{ env.wgsl.resourceDeclarations }}
 {% set tileT = "f16" if usesF16 else "f32" %}
 const WG_X: u32 = {{ convWgX }}u;
   var v = raw;
 {% if hasBias %}
   v = v + f32(bias[oc]);
 {% endif %}
   y[yBase + oc * params.outW + ow] = {{ T }}(v);
 }

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

@@ -1,43 +1,15 @@
 {{ 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 %}
 // Large grouped kernels are dominated by repeatedly fetching the same filter
 // value for neighboring output columns and input window for neighboring output
 // channels. One invocation computes four adjacent columns for OC_TILE channels,
 // sharing each input load across those channels. Shape and window geometry are
 // static, so indexing divisors and kernel offsets are shader constants.
-{% set dilatedLanes = dilatedLanes | default(false) %}
-{% set lanes = lanes | default(4) %}
-{% set quads = quads | default(1) %}
 {% set laneStep = dilationWSpec if dilatedLanes else strideWSpec %}
 {% set comps = ["x", "y", "z", "w"] %}
 {% macro acc_ref(oct, lane) %}{% if dilatedLanes %}acc{{ oct }}_{{ lane }}{% else %}acc{{ oct }}.{{ comps[lane] }}{% endif %}{% endmacro %}
@@ -86,8 +58,7 @@ const WG: u32 = {{ workgroupSizeSpec }}u;
 @compute @workgroup_size({{ workgroupSizeSpec }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  let q = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
-  if (q >= COUNT_TILES) { return; }
 {% if dilatedLanes %}
   let tile = q % TILES_PER_ROW;
@@ -128,9 +99,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% endfor %}
   for (var ic = 0u; ic < IN_CPG; ic++) {
-{% for kh in range(kernelHSpec) %}
     {
-      let ih = ihBase + {{ kh * dilationHSpec }};
       if (ih >= 0 && ih < IN_H) {
         let xRow = xChannelBase + u32(ih) * IN_W_U;
 {% if registerForm %}
@@ -170,7 +142,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% for kw in range(kernelWSpec) %}
         {
 {% for oct in range(ocTile) %}
-          let weight{{ oct }} = f32(w[wBase{{ oct }} + {{ kh * kernelWSpec + kw }}u]);
 {% endfor %}
 {% for lane in range(lanes) %}
 {% for oct in range(ocTile) %}
@@ -183,7 +155,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% for kw in range(kernelWSpec) %}
         {
 {% for oct in range(ocTile) %}
-          let weight{{ oct }} = f32(w[wBase{{ oct }} + {{ kh * kernelWSpec + kw }}u]);
 {% endfor %}
           let iwK = iwBase + {{ kw * dilationWSpec }};
 {% for lane in range(lanes) %}
@@ -201,7 +173,8 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
       }
     }
 {% endfor %}
-    xChannelBase += IN_PLANE;
 {% for oct in range(ocTile) %}
     wBase{{ oct }} += KAREA;
 {% endfor %}
@@ -212,7 +185,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
 {% for lane in range(lanes) %}
   if (ow0 + {{ lane * dilationWSpec }}u < OUT_W) {
-    y[outChannelBase{{ oct }} + ow0 + {{ lane * dilationWSpec }}u] = {{ scalar }}({% if hasActivation %}fused_act(acc{{ oct }}_{{ lane }}){% else %}acc{{ oct }}_{{ lane }}{% endif %});
   }
 {% endfor %}
 {% endfor %}
@@ -222,14 +195,14 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
 {% for lane in range(4) %}
   if (ow4 * 4u + {{ lane }}u < OUT_W) {
-    y[outChannelBase{{ oct }} + {{ lane }}u] = {{ scalar }}({% if hasActivation %}fused_act(acc{{ oct }}.{{ ["x", "y", "z", "w"][lane] }}){% else %}acc{{ oct }}.{{ ["x", "y", "z", "w"][lane] }}{% endif %});
   }
 {% endfor %}
 {% endfor %}
 {% else %}
   let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W4 + ow4;
 {% for oct in range(ocTile) %}
-  y[outBase + {{ oct }}u * OUT_PLANE4] = {{ vectorScalar }}({% if hasActivation %}vec4<f32>(fused_act(acc{{ oct }}.x), fused_act(acc{{ oct }}.y), fused_act(acc{{ oct }}.z), fused_act(acc{{ oct }}.w)){% else %}acc{{ oct }}{% endif %});
 {% endfor %}
 {% endif %}
 }

+{% 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 }}
 // Large grouped kernels are dominated by repeatedly fetching the same filter
 // value for neighboring output columns and input window for neighboring output
 // channels. One invocation computes four adjacent columns for OC_TILE channels,
 // sharing each input load across those channels. Shape and window geometry are
 // static, so indexing divisors and kernel offsets are shader constants.
 {% set laneStep = dilationWSpec if dilatedLanes else strideWSpec %}
 {% set comps = ["x", "y", "z", "w"] %}
 {% macro acc_ref(oct, lane) %}{% if dilatedLanes %}acc{{ oct }}_{{ lane }}{% else %}acc{{ oct }}.{{ comps[lane] }}{% endif %}{% endmacro %}
 @compute @workgroup_size({{ workgroupSizeSpec }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d("WG", "q", "COUNT_TILES", guardInline=true) }}
 {% if dilatedLanes %}
   let tile = q % TILES_PER_ROW;
 {% endfor %}
   for (var ic = 0u; ic < IN_CPG; ic++) {
+{% if rollKernelRows %}  for (var kernelRow = 0u; kernelRow < {{ kernelHSpec }}u; kernelRow++) {
+{% endif %}{% for kh in range(1 if rollKernelRows else kernelHSpec) %}
     {
+      let ih = ihBase + {% if rollKernelRows %}i32(kernelRow) * {{ dilationHSpec }}{% else %}{{ kh * dilationHSpec }}{% endif %};
       if (ih >= 0 && ih < IN_H) {
         let xRow = xChannelBase + u32(ih) * IN_W_U;
 {% if registerForm %}
 {% for kw in range(kernelWSpec) %}
         {
 {% for oct in range(ocTile) %}
+          let weight{{ oct }} = f32(w[wBase{{ oct }} + {% if rollKernelRows %}kernelRow * {{ kernelWSpec }}u + {{ kw }}u{% else %}{{ kh * kernelWSpec + kw }}u{% endif %}]);
 {% endfor %}
 {% for lane in range(lanes) %}
 {% for oct in range(ocTile) %}
 {% for kw in range(kernelWSpec) %}
         {
 {% for oct in range(ocTile) %}
+          let weight{{ oct }} = f32(w[wBase{{ oct }} + {% if rollKernelRows %}kernelRow * {{ kernelWSpec }}u + {{ kw }}u{% else %}{{ kh * kernelWSpec + kw }}u{% endif %}]);
 {% endfor %}
           let iwK = iwBase + {{ kw * dilationWSpec }};
 {% for lane in range(lanes) %}
       }
     }
 {% endfor %}
+{% if rollKernelRows %}  }
+{% endif %}    xChannelBase += IN_PLANE;
 {% for oct in range(ocTile) %}
     wBase{{ oct }} += KAREA;
 {% endfor %}
   let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
 {% for lane in range(lanes) %}
   if (ow0 + {{ lane * dilationWSpec }}u < OUT_W) {
+    y[outChannelBase{{ oct }} + ow0 + {{ lane * dilationWSpec }}u] = {{ scalar }}(acc{{ oct }}_{{ lane }});
   }
 {% endfor %}
 {% endfor %}
   let outChannelBase{{ oct }} = outBase + {{ oct }}u * OUT_H * OUT_W;
 {% for lane in range(4) %}
   if (ow4 * 4u + {{ lane }}u < OUT_W) {
+    y[outChannelBase{{ oct }} + {{ lane }}u] = {{ scalar }}(acc{{ oct }}.{{ ["x", "y", "z", "w"][lane] }});
   }
 {% endfor %}
 {% endfor %}
 {% else %}
   let outBase = ((batch * OUT_C + ocBase) * OUT_H + oh) * OUT_W4 + ow4;
 {% for oct in range(ocTile) %}
+  y[outBase + {{ oct }}u * OUT_PLANE4] = {{ vectorScalar }}(acc{{ oct }});
 {% endfor %}
 {% endif %}
 }

build/webgpu/manifest.json CHANGED Viewed

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

build/webgpu/metadata.json CHANGED Viewed

@@ -1,31 +1,31 @@
 {
   "name": "ai.onnx.Conv",
-  "id": "_ai_onnx_conv_webgpu_010d4b0",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "H1K2v5/L5nbBfjqTwogslpt8SrxTH6se0gj31ONSJeU=",
-      "conv-1x1-channel-reduce.wgsl.jinja": "Zjbu1oi+w5guLw54Y9C6TVJo0SntZDjgtsiie3o5md0=",
-      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "9yFpwXS1HzmWmm4W72n+aOIryLqc9+e88awZ1ogTe+Q=",
-      "conv-1x1-gemm-tiled.wgsl.jinja": "uqG5mRVvcHiUo0LDCUcVgqElVTvRNknx19Gm4oIALiU=",
-      "conv-1x1-subgroup-matrix.wgsl.jinja": "iDJs4RJk4iJarOZCjY7/5MFrHgR7CcqI6x5uwv2dYaA=",
-      "conv-direct-nd.wgsl.jinja": "FqwfAKVCF1BYBAtU7Y3rDcGrmZ5AYkH4FuTzdRlT2Rg=",
-      "conv-direct-unrolled.wgsl.jinja": "vrGUPwYkbYE1slII7nuK5wbERRcG5u0X1cW1kMc2sQs=",
       "conv-im2col-ncdhw.wgsl.jinja": "vQtnzeEKC2wm23+kwcaWoX3VSaJSqno18ACe3w3trhI=",
       "conv-im2col-nchw.wgsl.jinja": "IRpmtpeVQsLE6ud+Sf1230XaZVS1dxDtTZuSDwU3P8s=",
-      "conv-splitk-reduce.wgsl.jinja": "LbVU5DRL0Q51fvLjThmxJQZH8aPbZkvNzij4nY/cHmo=",
-      "conv1d-tiled-reg.wgsl.jinja": "ERqOPl+1RbnJ+CFqZvLSYw45qYMKgB7Je0teew5gyOM=",
-      "conv2d-grouped-large-w4.wgsl.jinja": "v8aAMQ2saiOgbQwLZAiVr3IeDT3h5YLO/v5tGL1x8pk=",
-      "manifest.json": "3/3/k9ZVj/vLw3xZ66XDMGWb86wzdZCoqDS0tKg9ziE=",
-      "test.json": "1OUEVPx0U/BTLtlxoSCFOWVgMfrlOfQb/TUY3VivfGw="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "ncw1d": ["conv-direct-nd.wgsl.jinja"],
       "ncw1d_bias": ["conv-direct-nd.wgsl.jinja"],
@@ -51,8 +51,6 @@
       "im2col_direct_f16_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
-      "implicit_im2col3d_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
-      "implicit_im2col3d_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix_bias_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
       "im2col3d_gemm_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
@@ -74,6 +72,10 @@
       "im2col_gemm_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "im2col_gemm_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "im2col_gemm_tiled_bias": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "gemm_1x1_tiled_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
       "gemm_1x1_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
       "gemm_1x1_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja"],

 {
   "name": "ai.onnx.Conv",
+  "id": "_ai_onnx_conv_webgpu_425dc7b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "qQlQzLmos636YZMqlseiLZNw3VMjVExXhvQhIkhNmso=",
+      "conv-1x1-channel-reduce.wgsl.jinja": "Ug2iTybBtLS8dnr7KOGbxUYKidE0/sb7LGLeAP3xka8=",
+      "conv-1x1-gemm-tiled-reg.wgsl.jinja": "fqy1AlbukRB6+EsP4NT8IsBbiAYWtxE7o8z87QN7nHI=",
+      "conv-1x1-gemm-tiled.wgsl.jinja": "mQWnnhk9Xlj20YcMc+6sQBuy67ZhdkgPUSdYgtauDaY=",
+      "conv-1x1-subgroup-matrix.wgsl.jinja": "/M/DUaLDk2rBCojLxieYQxQhpMei79mv4iM0shwJbl8=",
+      "conv-direct-nd.wgsl.jinja": "qooloadvj03r5vD7sSuCe1xGrWeuCxNx7B4lUjilOGQ=",
+      "conv-direct-unrolled.wgsl.jinja": "GsmGTy9UvercDDC0vNc3OOve9xcRtg2Hj9rSpx9YOSc=",
       "conv-im2col-ncdhw.wgsl.jinja": "vQtnzeEKC2wm23+kwcaWoX3VSaJSqno18ACe3w3trhI=",
       "conv-im2col-nchw.wgsl.jinja": "IRpmtpeVQsLE6ud+Sf1230XaZVS1dxDtTZuSDwU3P8s=",
+      "conv-splitk-reduce.wgsl.jinja": "jPCiZlPWzoQEoa9+ZwfZdSOHH+/hSYknvRFHKMvgkKM=",
+      "conv1d-tiled-reg.wgsl.jinja": "HvWuECHIzxqPrIyvfdlljhdOXpwzZo0jvAYuYDY+3kM=",
+      "conv2d-grouped-large-w4.wgsl.jinja": "mt2+nqnbFfJMac9s3sAsjTZlySD9VJFE+gHnkToVwt8=",
+      "manifest.json": "RtFRWBQdJ7bpJaP2G26O3a2GIKy1QaRpfZQqTFFmoYs=",
+      "test.json": "Y5+ejG2R11fFhIYo5ExoKSNeCtn278UgO3/JQgG1bY0="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
       "ncw1d": ["conv-direct-nd.wgsl.jinja"],
       "ncw1d_bias": ["conv-direct-nd.wgsl.jinja"],
       "im2col_direct_f16_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix_bias": ["conv-1x1-subgroup-matrix.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
       "implicit_im2col_subgroup_matrix_bias_splitk": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-splitk-reduce.wgsl.jinja"],
       "im2col3d_gemm_subgroup_matrix": ["conv-1x1-subgroup-matrix.wgsl.jinja", "conv-im2col-ncdhw.wgsl.jinja"],
       "im2col_gemm_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "im2col_gemm_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "im2col_gemm_tiled_bias": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled_reg_f16_columns": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled_f16_columns": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled_bias_reg_f16_columns": ["conv-1x1-gemm-tiled-reg.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
+      "im2col_gemm_tiled_bias_f16_columns": ["conv-1x1-gemm-tiled.wgsl.jinja", "conv-im2col-nchw.wgsl.jinja"],
       "gemm_1x1_tiled_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
       "gemm_1x1_tiled_bias_reg": ["conv-1x1-gemm-tiled-reg.wgsl.jinja"],
       "gemm_1x1_tiled": ["conv-1x1-gemm-tiled.wgsl.jinja"],

build/webgpu/test.json CHANGED Viewed

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