sync 2e7068faf55e

Browse files

Files changed (10) hide show

README.md +76 -0
build/webgpu/average-pool2d-nchw-horizontal-reuse.wgsl.jinja +88 -0
build/webgpu/average-pool2d-nchw-w3s1-reuse.wgsl.jinja +86 -0
build/webgpu/bench.json +173 -0
build/webgpu/manifest.json +491 -0
build/webgpu/metadata.json +22 -0
build/webgpu/pool-global-reduction.wgsl.jinja +93 -0
build/webgpu/pool-window-nd.wgsl.jinja +91 -0
build/webgpu/pool-window-unroll.wgsl.jinja +80 -0
build/webgpu/test.json +0 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.AveragePool
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 19
+## Description
+Applies average pooling over a sliding kernel window on input tensor `X`, computing the mean of values within each window position and writing results to `Y`. Output spatial dimensions are determined by `kernel_shape`, `strides`, `dilations`, `pads`, and `ceil_mode`; padded positions are excluded from the average by default unless `count_include_pad` is set.
+See the [ONNX `AveragePool` spec](https://onnx.ai/onnx/operators/onnx__AveragePool.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Input data tensor of shape `(N x C x D1 x D2 ... Dn)`, where `N` is the batch size and `C` is the number of channels. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | derived; see description | Output tensor with pooled values; spatial dimensions vary based on kernel, stride, dilation, and pad settings. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `auto_pad` | `"NOTSET"` | Deprecated auto-padding mode: `NOTSET` (use explicit pads), `SAME_UPPER`, `SAME_LOWER` (pad so output size is `ceil(input / stride)`), or `VALID` (no padding). |
+| `count_include_pad` | `0` | When non-zero, pad pixels are counted in the divisor when computing the average; defaults to 0 (exclude pad). |
+| `ceil_mode` | `0` | When non-zero, uses ceiling instead of floor when computing the output spatial shape; defaults to 0. |
+| `kernel_shape` | — | Required kernel shape, with one positive value per spatial axis. |
+| `strides` | — | Stride along each spatial axis. When omitted, every stride is 1. |
+| `pads` | — | Padding at the beginning and end of each spatial axis, ordered as `[begin_0, ..., begin_n, end_0, ..., end_n]`. When omitted, every pad is 0. |
+| `dilations` | — | Dilation along each spatial axis. When omitted, every dilation is 1. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`average-pool2d-nchw-horizontal-reuse.wgsl.jinja`](build/webgpu/average-pool2d-nchw-horizontal-reuse.wgsl.jinja)
+- [`average-pool2d-nchw-w3s1-reuse.wgsl.jinja`](build/webgpu/average-pool2d-nchw-w3s1-reuse.wgsl.jinja)
+- [`pool-global-reduction.wgsl.jinja`](build/webgpu/pool-global-reduction.wgsl.jinja)
+- [`pool-window-nd.wgsl.jinja`](build/webgpu/pool-window-nd.wgsl.jinja)
+- [`pool-window-unroll.wgsl.jinja`](build/webgpu/pool-window-unroll.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.AveragePool", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [1, 3, 32] } }, {
+  attrs: { kernel_shape: [2] },
+});
+```

build/webgpu/average-pool2d-nchw-horizontal-reuse.wgsl.jinja ADDED Viewed

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

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Four adjacent windows share source columns. For a K-wide, dilation-D window
+// at stride S, their union spans (K-1)*D+1+3*S columns per row. Each output
+// retains its own accumulation order.
+const COUNT4: u32 = {{ source.count4 }}u;
+const OUT_W4: u32 = max(1u, {{ source.outW4 }}u);
+const OUT_H: u32 = max(1u, {{ source.outH }}u);
+{% if source.scalarOutput is defined and source.scalarOutput %}
+const OUT_W: u32 = {{ source.outW }}u;
+{% endif %}
+const IN_H: i32 = {{ source.inH }};
+const IN_W: i32 = {{ source.inW }};
+const IN_W_U: u32 = {{ source.inW }}u;
+const IN_PLANE: u32 = {{ source.inPlane }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+fn load_value(validRow: bool, base: u32, col: i32) -> f32 {
+  if (validRow && col >= 0 && col < IN_W) { return f32(x[base + u32(col)]); }
+  return 0.0;
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  let q = gid.x + gid.y * nwg.x * WG;
+  if (q >= COUNT4) { return; }
+  let ow0 = (q % OUT_W4) * 4u;
+  var t = q / OUT_W4;
+  let oh = t % OUT_H;
+  let plane = t / OUT_H;
+  let planeBase = plane * IN_PLANE;
+  let startW = i32(ow0 * {{ source.strideW }}u) - {{ source.padLeft }};
+  var acc0 = 0.0; var acc1 = 0.0; var acc2 = 0.0; var acc3 = 0.0;
+  // Window validity is separable: every valid kernel row combines with every
+  // valid kernel column. Count each axis once and combine the counts without
+  // retaining a count beside every reused value.
+  var validRows = 0u;
+{% for lane in range(4) %}
+  var validCols{{ lane }} = 0u;
+{% for kw in range(source.kernelW) %}
+  validCols{{ lane }} += select(
+    0u,
+    1u,
+    startW + {{ lane * source.strideW + kw * source.dilationW }} >= 0 &&
+      startW + {{ lane * source.strideW + kw * source.dilationW }} < IN_W,
+  );
+{% endfor %}
+{% endfor %}
+{% for kh in range(source.kernelH) %}
+  {
+    let ih = i32(oh * {{ source.strideH }}u) + {{ kh * source.dilationH - source.padTop }};
+    let validRow = ih >= 0 && ih < IN_H;
+    validRows += select(0u, 1u, validRow);
+    let rowBase = planeBase + u32(max(ih, 0)) * IN_W_U;
+{% for j in range((source.kernelW - 1) * source.dilationW + 1 + 3 * source.strideW) %}
+    let v{{ kh }}_{{ j }} = load_value(validRow, rowBase, startW + {{ j }});
+{% endfor %}
+{% for lane in range(4) %}
+{% for kw in range(source.kernelW) %}
+    acc{{ lane }} += v{{ kh }}_{{ lane * source.strideW + kw * source.dilationW }};
+{% endfor %}
+{% endfor %}
+  }
+{% endfor %}
+{% for lane in range(4) %}
+  let cnt{{ lane }} = validRows * validCols{{ lane }};
+{% endfor %}
+{% if source.scalarOutput is defined and source.scalarOutput %}
+  let outBase = (plane * OUT_H + oh) * OUT_W + ow0;
+  y[outBase] = {{ scalar }}(select(0.0, acc0 / f32(cnt0), cnt0 > 0u));
+  if (ow0 + 1u < OUT_W) { y[outBase + 1u] = {{ scalar }}(select(0.0, acc1 / f32(cnt1), cnt1 > 0u)); }
+  if (ow0 + 2u < OUT_W) { y[outBase + 2u] = {{ scalar }}(select(0.0, acc2 / f32(cnt2), cnt2 > 0u)); }
+  if (ow0 + 3u < OUT_W) { y[outBase + 3u] = {{ scalar }}(select(0.0, acc3 / f32(cnt3), cnt3 > 0u)); }
+{% else %}
+  y[q] = {{ vectorScalar }}(
+    {{ scalar }}(select(0.0, acc0 / f32(cnt0), cnt0 > 0u)),
+    {{ scalar }}(select(0.0, acc1 / f32(cnt1), cnt1 > 0u)),
+    {{ scalar }}(select(0.0, acc2 / f32(cnt2), cnt2 > 0u)),
+    {{ scalar }}(select(0.0, acc3 / f32(cnt3), cnt3 > 0u))
+  );
+{% endif %}
+}

build/webgpu/average-pool2d-nchw-w3s1-reuse.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,86 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// 2D NCHW AveragePool specialised for a 3-wide / stride-1 / pad-1 / dilation-1
+// W window (countIncludePad == 0). Each invocation produces four adjacent output
+// columns. With stride 1 the windows of adjacent columns overlap by 2, so the
+// direct evaluation rereads each input column up to three times. Each invocation
+// instead loads six contiguous columns per kernel row and slides the window in
+// registers. Accumulation remains kh-major then kw-major, and excludes padding
+// from the denominator.
+const OUT_W: u32 = {{ source.outW }}u;
+const OUT_H: u32 = {{ source.outH }}u;
+const OUT_W_Q: u32 = {{ source.outWQuarter }}u; // OUT_W / 4 (columns-of-four)
+const IN_H: i32 = {{ source.inH }};
+const IN_W: i32 = {{ source.inW }};
+const IN_W_U: u32 = {{ source.inW }}u;
+const IN_PLANE: u32 = {{ source.inPlane }}u;
+const PAD_TOP: i32 = {{ source.padTop }};
+// Load input element at (rowBase, col), returning 0 outside [0, IN_W). The row
+// is already gated, so only the horizontal bound is checked here.
+fn load_col(rowValid: bool, rowBase: u32, col: i32) -> f32 {
+  if (rowValid && col >= 0 && col < IN_W) {
+    return f32(x[rowBase + u32(col)]);
+  }
+  return 0.0;
+}
+fn col_valid(rowValid: bool, col: i32) -> u32 {
+  if (rowValid && col >= 0 && col < IN_W) { return 1u; }
+  return 0u;
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  let tile = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (tile >= params.tileCount) {
+    return;
+  }
+  let colTile = tile % OUT_W_Q;     // which group-of-four columns in this row
+  var t = tile / OUT_W_Q;
+  let oh = t % OUT_H;
+  let plane = t / OUT_H;            // batch * channels + channel
+  let planeBase = plane * IN_PLANE;
+  let ow0 = colTile * 4u;          // first output column owned by this thread
+  let startW = i32(ow0) - 1;        // PAD_LEFT == 1; leftmost input column read
+  let startH = i32(oh) - PAD_TOP;   // strideH == 1
+  var acc0 = 0.0; var acc1 = 0.0; var acc2 = 0.0; var acc3 = 0.0;
+  var cnt0 = 0u; var cnt1 = 0u; var cnt2 = 0u; var cnt3 = 0u;
+{% for kh in range(source.kernelH) %}
+  {
+    let ih = startH + {{ kh * source.dilationH }};
+    let rowValid = ih >= 0 && ih < IN_H;
+    let rowBase = planeBase + u32(max(ih, 0)) * IN_W_U;
+    // Six contiguous input columns feeding the four 3-wide windows.
+    let v0 = load_col(rowValid, rowBase, startW + 0);
+    let v1 = load_col(rowValid, rowBase, startW + 1);
+    let v2 = load_col(rowValid, rowBase, startW + 2);
+    let v3 = load_col(rowValid, rowBase, startW + 3);
+    let v4 = load_col(rowValid, rowBase, startW + 4);
+    let v5 = load_col(rowValid, rowBase, startW + 5);
+    let c0 = col_valid(rowValid, startW + 0);
+    let c1 = col_valid(rowValid, startW + 1);
+    let c2 = col_valid(rowValid, startW + 2);
+    let c3 = col_valid(rowValid, startW + 3);
+    let c4 = col_valid(rowValid, startW + 4);
+    let c5 = col_valid(rowValid, startW + 5);
+    acc0 = acc0 + v0 + v1 + v2; cnt0 = cnt0 + c0 + c1 + c2;
+    acc1 = acc1 + v1 + v2 + v3; cnt1 = cnt1 + c1 + c2 + c3;
+    acc2 = acc2 + v2 + v3 + v4; cnt2 = cnt2 + c2 + c3 + c4;
+    acc3 = acc3 + v3 + v4 + v5; cnt3 = cnt3 + c3 + c4 + c5;
+  }
+{% endfor %}
+  let outBase = (plane * OUT_H + oh) * OUT_W + ow0;
+  y[outBase + 0u] = {{ scalar }}(select(0.0, acc0 / f32(cnt0), cnt0 > 0u));
+  y[outBase + 1u] = {{ scalar }}(select(0.0, acc1 / f32(cnt1), cnt1 > 0u));
+  y[outBase + 2u] = {{ scalar }}(select(0.0, acc2 / f32(cnt2), cnt2 > 0u));
+  y[outBase + 3u] = {{ scalar }}(select(0.0, acc3 / f32(cnt3), cnt3 > 0u));
+}

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,173 @@

+{
+  "op": "ai.onnx.AveragePool",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "average-pool-f32-b1c32-64x64-k2s2",
+      "preset": "smoke",
+      "vars": { "batch": 1, "channels": 32, "inH": 64, "inW": 64 },
+      "attrs": { "kernel_shape": [2, 2], "strides": [2, 2], "pads": [0, 0, 0, 0] },
+      "inputs": { "x": { "shape": [1, 32, 64, 64], "dtype": "float32", "dist": "normal", "seed": 116, "scale": 0.2 } },
+      "outputs": { "y": { "shape": [1, 32, 32, 32], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4" }]
+      }
+    },
+    {
+      "name": "nchw_8x64x112x112_k3s1p1_unrolled",
+      "preset": "smoke",
+      "vars": { "batch": 8, "channels": 64, "inH": 112, "inW": 112 },
+      "attrs": { "kernel_shape": [3, 3], "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 64, 112, 112], "dist": "normal", "seed": 205 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 112, 112], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4 * 2" }]
+      }
+    },
+    {
+      "name": "nchw_f16_8x64x112x112_k3s1p1_unrolled",
+      "preset": "smoke",
+      "vars": { "batch": 8, "channels": 64, "inH": 112, "inW": 112 },
+      "attrs": { "kernel_shape": [3, 3], "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": { "x": { "dtype": "float16", "shape": [8, 64, 112, 112], "dist": "normal", "seed": 206 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [8, 64, 112, 112], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 2 * 2" }]
+      }
+    },
+    {
+      "name": "average-pool-f32-b1c64-224x224-k7s2p3-generic",
+      "preset": "smoke",
+      "vars": { "batch": 1, "channels": 64, "inH": 224, "inW": 224 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [7, 7], "strides": [2, 2], "pads": [3, 3, 3, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [1, 64, 224, 224], "dist": "normal", "seed": 311, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 112, 112], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4" }] }
+    },
+    {
+      "name": "average-pool-f16-b1c64-224x224-k7s2p3-generic",
+      "preset": "smoke",
+      "vars": { "batch": 1, "channels": 64, "inH": 224, "inW": 224 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [7, 7], "strides": [2, 2], "pads": [3, 3, 3, 3] },
+      "inputs": { "x": { "dtype": "float16", "shape": [1, 64, 224, 224], "dist": "normal", "seed": 312, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 64, 112, 112], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 2" }] }
+    },
+    {
+      "name": "average-pool-f32-b8c64-112x114-k3s1p1-unaligned-fallback",
+      "preset": "smoke",
+      "vars": { "batch": 8, "channels": 64, "inH": 112, "inW": 114 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [3, 3], "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 64, 112, 114], "dist": "normal", "seed": 313, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 64, 112, 114], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4 * 2" }]
+      }
+    },
+    {
+      "name": "average-pool-f32-ncdhw-b1c32-32x64x64-k2s2-3d",
+      "preset": "smoke",
+      "vars": { "batch": 1, "channels": 32, "inD": 32, "inH": 64, "inW": 64 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [2, 2, 2], "strides": [2, 2, 2], "pads": [0, 0, 0, 0, 0, 0] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 32, 64, 64], "dist": "normal", "seed": 314, "scale": 0.2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 16, 32, 32], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inD * args.inH * args.inW * 4" }]
+      }
+    },
+    {
+      "name": "average-pool-f32-ncl-b8c256-1024-k3s2-1d",
+      "preset": "smoke",
+      "vars": { "batch": 8, "channels": 256, "inW": 1024 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [3], "strides": [2], "pads": [0, 0] },
+      "inputs": { "x": { "dtype": "float32", "shape": [8, 256, 1024], "dist": "normal", "seed": 315, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 256, 511], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inW * 4" }] }
+    },
+    {
+      "name": "global_avg_pool_bigkernel_tiny_output_launch_starved_stress",
+      "preset": "stress",
+      "vars": { "batch": 1, "channels": 512, "inH": 64, "inW": 64 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [64, 64], "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": { "x": { "dtype": "float32", "shape": [1, 512, 64, 64], "dist": "normal", "seed": 4101, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 512, 1, 1], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4" }] }
+    },
+    {
+      "name": "global-avg-pool-f32-b2c256-96x96-segmentation-head",
+      "preset": "stress",
+      "vars": { "batch": 2, "channels": 256, "inH": 96, "inW": 96 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [96, 96], "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 256, 96, 96], "dist": "normal", "seed": 4103, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 256, 1, 1], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4" }]
+      }
+    },
+    {
+      "name": "nchw2d_generic_k7s1_overlap_reread_large_map_stress",
+      "preset": "stress",
+      "vars": { "batch": 4, "channels": 64, "inH": 128, "inW": 128 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [7, 7], "strides": [1, 1], "pads": [3, 3, 3, 3] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4, 64, 128, 128], "dist": "normal", "seed": 4102, "scale": 0.2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 64, 128, 128], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4 * 2" }]
+      }
+    },
+    {
+      "name": "average-pool-k7s1p3-f32-b2c128-96x96-segmentation-context",
+      "preset": "stress",
+      "vars": { "batch": 2, "channels": 128, "inH": 96, "inW": 96 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [7, 7], "strides": [1, 1], "pads": [3, 3, 3, 3] },
+      "inputs": { "x": { "dtype": "float32", "shape": [2, 128, 96, 96], "dist": "normal", "seed": 4104, "scale": 0.2 } },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 128, 96, 96], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4 * 2" }]
+      }
+    },
+    {
+      "name": "average-pool-k7s1p3-f32-b1c64-112x113-odd-width-cliff",
+      "preset": "stress",
+      "vars": { "batch": 1, "channels": 64, "inH": 112, "inW": 113 },
+      "attrs": { "count_include_pad": 0, "kernel_shape": [7, 7], "strides": [1, 1], "pads": [3, 3, 3, 3] },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 64, 112, 113], "dist": "normal", "seed": 4105, "scale": 0.2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 64, 112, 113], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inH * args.inW * 4 * 2" }]
+      }
+    },
+    {
+      "name": "global-averagepool-rank5-serial-reduction-pathology",
+      "preset": "stress",
+      "provenance": {
+        "source": "authored for variant coverage",
+        "notes": "Rank-5 global pooling over 131072 spatial values per plane verifies global_window_parallel coverage and its feature-tier fallbacks."
+      },
+      "vars": { "batch": 1, "channels": 32, "inD": 32, "inH": 64, "inW": 64 },
+      "attrs": {
+        "count_include_pad": 0,
+        "kernel_shape": [32, 64, 64],
+        "strides": [1, 1, 1],
+        "pads": [0, 0, 0, 0, 0, 0]
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1, 32, 32, 64, 64], "dist": "normal", "seed": 5101, "scale": 0.2 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 32, 1, 1, 1], "dist": "empty" } },
+      "bench": {
+        "metrics": [{ "type": "bandwidth", "value": "args.batch * args.channels * args.inD * args.inH * args.inW * 4" }]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,491 @@

+{
+  "domain": "ai.onnx",
+  "name": "AveragePool",
+  "sinceVersion": 19,
+  "description": "Applies average pooling over a sliding kernel window on input tensor `X`, computing the mean of values within each window position and writing results to `Y`. Output spatial dimensions are determined by `kernel_shape`, `strides`, `dilations`, `pads`, and `ceil_mode`; padded positions are excluded from the average by default unless `count_include_pad` is set.",
+  "inputs": [
+    {
+      "role": "X",
+      "dtype": "T",
+      "description": "Input data tensor of shape `(N x C x D1 x D2 ... Dn)`, where `N` is the batch size and `C` is the number of channels."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "description": "Output tensor with pooled values; spatial dimensions vary based on kernel, stride, dilation, and pad settings.",
+      "rank": "ranks.X",
+      "shape": "poolOutputShape"
+    }
+  ],
+  "attributes": { "auto_pad": "NOTSET", "count_include_pad": 0, "ceil_mode": 0 },
+  "attributeDescriptions": {
+    "auto_pad": "Deprecated auto-padding mode: `NOTSET` (use explicit pads), `SAME_UPPER`, `SAME_LOWER` (pad so output size is `ceil(input / stride)`), or `VALID` (no padding).",
+    "kernel_shape": "Required kernel shape, with one positive value per spatial axis.",
+    "strides": "Stride along each spatial axis. When omitted, every stride is 1.",
+    "pads": "Padding at the beginning and end of each spatial axis, ordered as `[begin_0, ..., begin_n, end_0, ..., end_n]`. When omitted, every pad is 0.",
+    "dilations": "Dilation along each spatial axis. When omitted, every dilation is 1.",
+    "count_include_pad": "When non-zero, pad pixels are counted in the divisor when computing the average; defaults to 0 (exclude pad).",
+    "ceil_mode": "When non-zero, uses ceiling instead of floor when computing the output spatial shape; defaults to 0."
+  },
+  "attributeConstraints": {
+    "auto_pad": { "values": ["NOTSET", "SAME_UPPER", "SAME_LOWER", "VALID"] },
+    "ceil_mode": { "values": [0, 1] },
+    "count_include_pad": { "values": [0, 1] },
+    "kernel_shape": { "required": true }
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "derive": {
+    "spatialRank": "ranks.X - 2",
+    "narrowSubgroupRange": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < device.adapterInfo.subgroupMaxSize and device.adapterInfo.subgroupMaxSize <= 16",
+    "kernelShapeLengthOk": "not has(attrs, \"kernel_shape\") or (attrs.kernel_shape | length) == spatialRank",
+    "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",
+    "strideD": "attrs.strides[0] if stridesLengthOk and has(attrs, \"strides\") and spatialRank == 3 else 1",
+    "strideH": "attrs.strides[spatialRank - 2] if stridesLengthOk and has(attrs, \"strides\") and spatialRank >= 2 else 1",
+    "strideW": "attrs.strides[spatialRank - 1] if stridesLengthOk and has(attrs, \"strides\") and spatialRank >= 1 else 1",
+    "dilationD": "attrs.dilations[0] if dilationsLengthOk and has(attrs, \"dilations\") and spatialRank == 3 else 1",
+    "dilationH": "attrs.dilations[spatialRank - 2] if dilationsLengthOk and has(attrs, \"dilations\") and spatialRank >= 2 else 1",
+    "dilationW": "attrs.dilations[spatialRank - 1] if dilationsLengthOk and has(attrs, \"dilations\") and spatialRank >= 1 else 1",
+    "padFront": "attrs.pads[0] if padsLengthOk and has(attrs, \"pads\") and spatialRank == 3 else 0",
+    "padTop": "attrs.pads[spatialRank - 2] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 2 else 0",
+    "padLeft": "attrs.pads[spatialRank - 1] if padsLengthOk and has(attrs, \"pads\") and spatialRank >= 1 else 0",
+    "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",
+    "poolBaseOk": "f16Ok(dtypes.T) and ranks.X >= 3 and ranks.X <= 5 and ranks.Y == ranks.X and has(attrs, \"kernel_shape\") and kernelShapeLengthOk and stridesLengthOk and dilationsLengthOk and padsLengthOk and kernelD > 0 and kernelH > 0 and kernelW > 0 and strideD > 0 and strideH > 0 and strideW > 0 and dilationD > 0 and dilationH > 0 and dilationW > 0 and padFront >= 0 and padTop >= 0 and padLeft >= 0 and padBack >= 0 and padBottom >= 0 and padRight >= 0 and (attrs.auto_pad == \"NOTSET\" or not has(attrs, \"pads\")) and dim(shapes.Y, 0) == dim(shapes.X, 0) and dim(shapes.Y, 1) == dim(shapes.X, 1)",
+    "explicitOutW": "min(ceil((dim(shapes.X, ranks.X - 1) + padLeft + padRight - ((kernelW - 1) * dilationW + 1)) / strideW) + 1, ceil((dim(shapes.X, ranks.X - 1) + padLeft) / strideW)) if attrs.ceil_mode != 0 else floor((dim(shapes.X, ranks.X - 1) + padLeft + padRight - ((kernelW - 1) * dilationW + 1)) / strideW) + 1",
+    "explicitOutH": "min(ceil((dim(shapes.X, ranks.X - 2) + padTop + padBottom - ((kernelH - 1) * dilationH + 1)) / strideH) + 1, ceil((dim(shapes.X, ranks.X - 2) + padTop) / strideH)) if attrs.ceil_mode != 0 else floor((dim(shapes.X, ranks.X - 2) + padTop + padBottom - ((kernelH - 1) * dilationH + 1)) / strideH) + 1",
+    "explicitOutD": "(min(ceil((dim(shapes.X, 2) + padFront + padBack - ((kernelD - 1) * dilationD + 1)) / strideD) + 1, ceil((dim(shapes.X, 2) + padFront) / strideD)) if attrs.ceil_mode != 0 else floor((dim(shapes.X, 2) + padFront + padBack - ((kernelD - 1) * dilationD + 1)) / strideD) + 1) if ranks.X == 5 else 0",
+    "inferredOutW": "ceil(dim(shapes.X, ranks.X - 1) / strideW) if attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\" else explicitOutW",
+    "inferredOutH": "ceil(dim(shapes.X, ranks.X - 2) / strideH) if attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\" else explicitOutH",
+    "inferredOutD": "ceil(dim(shapes.X, 2) / strideD) if attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\" else explicitOutD",
+    "poolOutputShape": "[dim(shapes.X, 0), dim(shapes.X, 1), inferredOutW] if ranks.X == 3 else ([dim(shapes.X, 0), dim(shapes.X, 1), inferredOutH, inferredOutW] if ranks.X == 4 else [dim(shapes.X, 0), dim(shapes.X, 1), inferredOutD, inferredOutH, inferredOutW])",
+    "autoPadTotalH": "max(0, (dim(shapes.Y, 2) - 1) * strideH + ((kernelH - 1) * dilationH + 1) - dim(shapes.X, 2)) if ranks.X == 4 else 0",
+    "autoPadTotalW": "max(0, (dim(shapes.Y, 3) - 1) * strideW + ((kernelW - 1) * dilationW + 1) - dim(shapes.X, 3)) if ranks.X == 4 else 0",
+    "effectivePadTop": "(floor(autoPadTotalH / 2) if attrs.auto_pad == \"SAME_UPPER\" else autoPadTotalH - floor(autoPadTotalH / 2)) if ranks.X == 4 and (attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") else padTop",
+    "effectivePadLeft": "(floor(autoPadTotalW / 2) if attrs.auto_pad == \"SAME_UPPER\" else autoPadTotalW - floor(autoPadTotalW / 2)) if ranks.X == 4 and (attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") else padLeft",
+    "effectivePadBottom": "autoPadTotalH - effectivePadTop if ranks.X == 4 and (attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") else padBottom",
+    "effectivePadRight": "autoPadTotalW - effectivePadLeft if ranks.X == 4 and (attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") else padRight",
+    "narrowF32W3UnrolledPreferred": "narrowSubgroupRange and dtypes.T == \"f32\"",
+    "globalWindowParallelCovered": "poolBaseOk and numel(shapes.Y) == dim(shapes.X, 0) * dim(shapes.X, 1) and numel(shapes.Y) <= device.limits.maxComputeWorkgroupsPerDimension and attrs.auto_pad == \"NOTSET\" and strideD == 1 and strideH == 1 and strideW == 1 and dilationD == 1 and dilationH == 1 and dilationW == 1 and padFront == 0 and padTop == 0 and padLeft == 0 and padBack == 0 and padBottom == 0 and padRight == 0 and ((ranks.X == 3 and kernelW == dim(shapes.X, 2)) or (ranks.X == 4 and kernelH == dim(shapes.X, 2) and kernelW == dim(shapes.X, 3)) or (ranks.X == 5 and kernelD == dim(shapes.X, 2) and kernelH == dim(shapes.X, 3) and kernelW == dim(shapes.X, 4))) and inner(shapes.X, 1) >= 256"
+  },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "bindingSets": {
+    "globalReduction": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$xElem" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "spatial", "type": "u32", "value": "inner(shapes.X, 1)" }] }
+      }
+    ],
+    "horizontal": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputElement" }
+    ],
+    "w3Reuse": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [{ "name": "tileCount", "type": "u32", "value": "numel(shapes.Y) / 4" }]
+        }
+      }
+    ],
+    "flatDirect": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" }] }
+      }
+    ],
+    "rank1Direct": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "channels", "type": "u32", "value": "dim(shapes.X, 1)" },
+            { "name": "inW", "type": "u32", "value": "dim(shapes.X, 2)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 2)" },
+            { "name": "kernelW", "type": "u32", "value": "kernelW" },
+            { "name": "strideW", "type": "u32", "value": "strideW" },
+            { "name": "dilationW", "type": "u32", "value": "dilationW" },
+            { "name": "padLeft", "type": "i32", "value": "padLeft" },
+            { "name": "padRight", "type": "i32", "value": "padRight" },
+            { "name": "countIncludePad", "type": "u32", "value": "attrs.count_include_pad" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" }
+          ]
+        }
+      }
+    ],
+    "rank3Direct": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "channels", "type": "u32", "value": "dim(shapes.X, 1)" },
+            { "name": "inD", "type": "u32", "value": "dim(shapes.X, 2)" },
+            { "name": "inH", "type": "u32", "value": "dim(shapes.X, 3)" },
+            { "name": "inW", "type": "u32", "value": "dim(shapes.X, 4)" },
+            { "name": "outD", "type": "u32", "value": "dim(shapes.Y, 2)" },
+            { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 3)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 4)" },
+            { "name": "kernelD", "type": "u32", "value": "kernelD" },
+            { "name": "kernelH", "type": "u32", "value": "kernelH" },
+            { "name": "kernelW", "type": "u32", "value": "kernelW" },
+            { "name": "strideD", "type": "u32", "value": "strideD" },
+            { "name": "strideH", "type": "u32", "value": "strideH" },
+            { "name": "strideW", "type": "u32", "value": "strideW" },
+            { "name": "dilationD", "type": "u32", "value": "dilationD" },
+            { "name": "dilationH", "type": "u32", "value": "dilationH" },
+            { "name": "dilationW", "type": "u32", "value": "dilationW" },
+            { "name": "padFront", "type": "i32", "value": "padFront" },
+            { "name": "padTop", "type": "i32", "value": "padTop" },
+            { "name": "padLeft", "type": "i32", "value": "padLeft" },
+            { "name": "padBack", "type": "i32", "value": "padBack" },
+            { "name": "padBottom", "type": "i32", "value": "padBottom" },
+            { "name": "padRight", "type": "i32", "value": "padRight" },
+            { "name": "countIncludePad", "type": "u32", "value": "attrs.count_include_pad" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" }
+          ]
+        }
+      }
+    ],
+    "rank2Direct": [
+      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "channels", "type": "u32", "value": "dim(shapes.X, 1)" },
+            { "name": "inH", "type": "u32", "value": "dim(shapes.X, 2)" },
+            { "name": "inW", "type": "u32", "value": "dim(shapes.X, 3)" },
+            { "name": "outH", "type": "u32", "value": "dim(shapes.Y, 2)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.Y, 3)" },
+            { "name": "kernelH", "type": "u32", "value": "kernelH" },
+            { "name": "kernelW", "type": "u32", "value": "kernelW" },
+            { "name": "strideH", "type": "u32", "value": "strideH" },
+            { "name": "strideW", "type": "u32", "value": "strideW" },
+            { "name": "dilationH", "type": "u32", "value": "dilationH" },
+            { "name": "dilationW", "type": "u32", "value": "dilationW" },
+            { "name": "padTop", "type": "i32", "value": "effectivePadTop" },
+            { "name": "padLeft", "type": "i32", "value": "effectivePadLeft" },
+            { "name": "padBottom", "type": "i32", "value": "effectivePadBottom" },
+            { "name": "padRight", "type": "i32", "value": "effectivePadRight" },
+            { "name": "countIncludePad", "type": "u32", "value": "attrs.count_include_pad" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.Y)" }
+          ]
+        }
+      }
+    ]
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "variants": [
+    {
+      "id": "global_window_parallel",
+      "priority": 50,
+      "when": ["globalWindowParallelCovered"],
+      "constants": {
+        "xElem": "(\"vec4<f16>\" if dtypes.T == \"f16\" else \"vec4<f32>\") if inner(shapes.X, 1) % 4 == 0 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.GlobalWindowParallel",
+          "source": {
+            "shader": "pool-global-reduction.wgsl.jinja",
+            "inputs": {
+              "op": "\"average\"",
+              "vec4": "inner(shapes.X, 1) % 4 == 0",
+              "workgroupSize": "tunables.WORKGROUP_SIZE",
+              "useSubgroups": "device.features.has(\"subgroups\")"
+            }
+          },
+          "bindings": "globalReduction",
+          "dispatch": { "x": "numel(shapes.Y)" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_k7s1_reuse_tail",
+      "priority": 34,
+      "when": ["poolBaseOk", "ranks.X == 4", "kernelH == 7", "kernelW == 7", "strideH == 1", "strideW == 1", "dilationH == 1", "dilationW == 1", "padTop == 3", "padBottom == 3", "padLeft == 3", "padRight == 3", "attrs.count_include_pad == 0", "attrs.ceil_mode == 0", "dim(shapes.X, 3) % 4 != 0", "dim(shapes.Y, 2) == dim(shapes.X, 2)", "dim(shapes.Y, 3) == dim(shapes.X, 3)", "dim(shapes.Y, 2) > 0", "dim(shapes.Y, 3) > 0"],
+      "constants": { "outputElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Nchw2dK7s1ReuseTail",
+          "source": {
+            "shader": "average-pool2d-nchw-horizontal-reuse.wgsl.jinja",
+            "inputs": {
+              "kernelH": 7,
+              "kernelW": 7,
+              "strideH": 1,
+              "strideW": 1,
+              "dilationH": 1,
+              "dilationW": 1,
+              "padTop": 3,
+              "padLeft": 3,
+              "inH": "dim(shapes.X, 2)",
+              "inW": "dim(shapes.X, 3)",
+              "inPlane": "dim(shapes.X, 2) * dim(shapes.X, 3)",
+              "outH": "dim(shapes.Y, 2)",
+              "outW": "dim(shapes.Y, 3)",
+              "outW4": "ceilDiv(dim(shapes.Y, 3), 4)",
+              "count4": "dim(shapes.Y, 0) * dim(shapes.Y, 1) * dim(shapes.Y, 2) * ceilDiv(dim(shapes.Y, 3), 4)",
+              "scalarOutput": true
+            }
+          },
+          "bindings": "horizontal",
+          "dispatch": {
+            "threads": "dim(shapes.Y, 0) * dim(shapes.Y, 1) * dim(shapes.Y, 2) * ceilDiv(dim(shapes.Y, 3), 4)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_w3s1_reuse",
+      "priority": 30,
+      "demoteWhen": ["narrowF32W3UnrolledPreferred"],
+      "when": ["poolBaseOk", "ranks.X == 4", "kernelW == 3", "strideW == 1", "dilationW == 1", "padLeft == 1", "padRight == 1", "kernelH >= 1", "kernelH <= 3", "strideH == 1", "dilationH >= 1", "attrs.count_include_pad == 0", "attrs.ceil_mode == 0", "dim(shapes.X, 3) % 4 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) + padTop + padBottom - ((kernelH - 1) * dilationH + 1) + 1", "dim(shapes.Y, 3) == dim(shapes.X, 3)", "dim(shapes.Y, 2) > 0", "dim(shapes.Y, 3) > 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Nchw2dW3s1Reuse",
+          "source": {
+            "shader": "average-pool2d-nchw-w3s1-reuse.wgsl.jinja",
+            "inputs": {
+              "kernelH": "kernelH",
+              "dilationH": "dilationH",
+              "padTop": "padTop",
+              "inH": "dim(shapes.X, 2)",
+              "inW": "dim(shapes.X, 3)",
+              "inPlane": "dim(shapes.X, 2) * dim(shapes.X, 3)",
+              "outH": "dim(shapes.Y, 2)",
+              "outW": "dim(shapes.Y, 3)",
+              "outWQuarter": "floor(dim(shapes.Y, 3) / 4)"
+            }
+          },
+          "bindings": "w3Reuse",
+          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_w3s1_reuse_tail",
+      "priority": 29,
+      "when": ["poolBaseOk", "ranks.X == 4", "kernelW == 3", "strideW == 1", "dilationW == 1", "padLeft == 1", "padRight == 1", "kernelH >= 1", "kernelH <= 3", "strideH == 1", "dilationH >= 1", "attrs.count_include_pad == 0", "attrs.ceil_mode == 0", "dim(shapes.X, 3) % 4 != 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) + padTop + padBottom - ((kernelH - 1) * dilationH + 1) + 1", "dim(shapes.Y, 3) == dim(shapes.X, 3)", "dim(shapes.Y, 2) > 0", "dim(shapes.Y, 3) > 0"],
+      "constants": { "outputElement": "dtypes.T" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Nchw2dW3s1ReuseTail",
+          "source": {
+            "shader": "average-pool2d-nchw-horizontal-reuse.wgsl.jinja",
+            "inputs": {
+              "kernelH": "kernelH",
+              "kernelW": 3,
+              "strideH": 1,
+              "strideW": 1,
+              "dilationH": "dilationH",
+              "dilationW": 1,
+              "padTop": "padTop",
+              "padLeft": 1,
+              "inH": "dim(shapes.X, 2)",
+              "inW": "dim(shapes.X, 3)",
+              "inPlane": "dim(shapes.X, 2) * dim(shapes.X, 3)",
+              "outH": "dim(shapes.Y, 2)",
+              "outW": "dim(shapes.Y, 3)",
+              "outW4": "ceilDiv(dim(shapes.Y, 3), 4)",
+              "count4": "dim(shapes.Y, 0) * dim(shapes.Y, 1) * dim(shapes.Y, 2) * ceilDiv(dim(shapes.Y, 3), 4)",
+              "scalarOutput": true
+            }
+          },
+          "bindings": "horizontal",
+          "dispatch": {
+            "threads": "dim(shapes.Y, 0) * dim(shapes.Y, 1) * dim(shapes.Y, 2) * ceilDiv(dim(shapes.Y, 3), 4)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_unrolled",
+      "priority": 25,
+      "when": ["poolBaseOk", "ranks.X == 4", "dim(shapes.Y, 2) == explicitOutH", "dim(shapes.Y, 3) == explicitOutW", "kernelH >= 1", "kernelH <= 3", "kernelW >= 1", "kernelW <= 3", "strideH >= 1", "strideW >= 1", "dilationH >= 1", "dilationW >= 1", "dim(shapes.Y, 2) > 0", "dim(shapes.Y, 3) > 0"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Nchw2dUnrolled",
+          "source": {
+            "shader": "pool-window-unroll.wgsl.jinja",
+            "inputs": {
+              "op": "\"average\"",
+              "kernelH": "kernelH",
+              "kernelW": "kernelW",
+              "strideH": "strideH",
+              "strideW": "strideW",
+              "dilationH": "dilationH",
+              "dilationW": "dilationW",
+              "padTop": "padTop",
+              "padLeft": "padLeft",
+              "inH": "dim(shapes.X, 2)",
+              "inW": "dim(shapes.X, 3)",
+              "outH": "dim(shapes.Y, 2)",
+              "outW": "dim(shapes.Y, 3)",
+              "padBottom": "padBottom",
+              "padRight": "padRight",
+              "countIncludePad": "attrs.count_include_pad != 0"
+            }
+          },
+          "bindings": "flatDirect",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncl1d",
+      "priority": 10,
+      "when": ["poolBaseOk", "ranks.X == 3", "dim(shapes.Y, 2) == explicitOutW"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Ncl1d",
+          "source": { "shader": "pool-window-nd.wgsl.jinja", "inputs": { "op": "\"average\"", "spatialRank": 1 } },
+          "bindings": "rank1Direct",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d",
+      "when": ["poolBaseOk", "ranks.X == 4", "((dim(shapes.Y, 2) == explicitOutH and dim(shapes.Y, 3) == explicitOutW) or ((attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\") and dim(shapes.Y, 2) == ceil(dim(shapes.X, 2) / strideH) and dim(shapes.Y, 3) == ceil(dim(shapes.X, 3) / strideW)))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "nchw2d",
+          "source": { "shader": "pool-window-nd.wgsl.jinja", "inputs": { "op": "\"average\"", "spatialRank": 2 } },
+          "bindings": "rank2Direct",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d",
+      "priority": 10,
+      "supersededBy": ["global_window_parallel"],
+      "when": ["poolBaseOk", "ranks.X == 5", "dim(shapes.Y, 2) == explicitOutD", "dim(shapes.Y, 3) == explicitOutH", "dim(shapes.Y, 4) == explicitOutW"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Ncdhw3d",
+          "source": { "shader": "pool-window-nd.wgsl.jinja", "inputs": { "op": "\"average\"", "spatialRank": 3 } },
+          "bindings": "rank3Direct",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_k7s2_reuse",
+      "priority": 36,
+      "when": ["poolBaseOk", "ranks.X == 4", "kernelH == 7", "kernelW == 7", "strideH == 2", "strideW == 2", "dilationH == 1", "dilationW == 1", "padTop == 3", "padBottom == 3", "padLeft == 3", "padRight == 3", "attrs.count_include_pad == 0", "attrs.ceil_mode == 0", "dim(shapes.Y, 2) == explicitOutH", "dim(shapes.Y, 3) == explicitOutW", "dim(shapes.Y, 3) % 4 == 0", "dim(shapes.Y, 2) > 0", "dim(shapes.Y, 3) > 0"],
+      "constants": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Nchw2dK7s2Reuse",
+          "source": {
+            "shader": "average-pool2d-nchw-horizontal-reuse.wgsl.jinja",
+            "inputs": {
+              "kernelH": 7,
+              "kernelW": 7,
+              "strideH": 2,
+              "strideW": 2,
+              "dilationH": 1,
+              "dilationW": 1,
+              "padTop": 3,
+              "padLeft": 3,
+              "inH": "dim(shapes.X, 2)",
+              "inW": "dim(shapes.X, 3)",
+              "inPlane": "dim(shapes.X, 2) * dim(shapes.X, 3)",
+              "outH": "dim(shapes.Y, 2)",
+              "outW": "dim(shapes.Y, 3)",
+              "outW4": "dim(shapes.Y, 3) / 4",
+              "count4": "numel(shapes.Y) / 4"
+            }
+          },
+          "bindings": "horizontal",
+          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_k7s1_reuse",
+      "priority": 35,
+      "when": ["poolBaseOk", "ranks.X == 4", "kernelH == 7", "kernelW == 7", "strideH == 1", "strideW == 1", "dilationH == 1", "dilationW == 1", "padTop == 3", "padBottom == 3", "padLeft == 3", "padRight == 3", "attrs.count_include_pad == 0", "attrs.ceil_mode == 0", "dim(shapes.X, 3) % 4 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2)", "dim(shapes.Y, 3) == dim(shapes.X, 3)", "dim(shapes.Y, 2) > 0", "dim(shapes.Y, 3) > 0"],
+      "constants": { "outputElement": "\"vec4<\" ~ dtypes.T ~ \">\"", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "AveragePool.Nchw2dK7s1Reuse",
+          "source": {
+            "shader": "average-pool2d-nchw-horizontal-reuse.wgsl.jinja",
+            "inputs": {
+              "kernelH": 7,
+              "kernelW": 7,
+              "strideH": 1,
+              "strideW": 1,
+              "dilationH": 1,
+              "dilationW": 1,
+              "padTop": 3,
+              "padLeft": 3,
+              "inH": "dim(shapes.X, 2)",
+              "inW": "dim(shapes.X, 3)",
+              "inPlane": "dim(shapes.X, 2) * dim(shapes.X, 3)",
+              "outH": "dim(shapes.Y, 2)",
+              "outW": "dim(shapes.Y, 3)",
+              "outW4": "dim(shapes.Y, 3) / 4",
+              "count4": "numel(shapes.Y) / 4"
+            }
+          },
+          "bindings": "horizontal",
+          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,22 @@

+{
+  "name": "ai.onnx.AveragePool",
+  "id": "_ai_onnx_averagepool_webgpu_208ddea",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "average-pool2d-nchw-horizontal-reuse.wgsl.jinja": "MuniH46O1hxMO9h44Pke8TPe9LdxPJxSfkWzpzH7ORA=",
+      "average-pool2d-nchw-w3s1-reuse.wgsl.jinja": "zPPcxTvNNFow8AZiF6f6aA528J72OfqiLD4r8DTNVVs=",
+      "bench.json": "YSXPXL9DWp91TDTmMZrcq81UYOnXi2L5BKFjPk+8MEU=",
+      "manifest.json": "6q/856PdHb80O+6DlOVhpkyox/A+DIvIcgNaiZBCN6E=",
+      "pool-global-reduction.wgsl.jinja": "kFAmHLC4tjXVW1gTRXNZHVeLo9o4BdkmCVhTIyB9doA=",
+      "pool-window-nd.wgsl.jinja": "WlKw9CY8+OR7l8jGO2lyJ72P0twJi3CIvechFVjC7jA=",
+      "pool-window-unroll.wgsl.jinja": "65NfjRJtzKQBVh23bYGyuwr7KUKKvdWsHnsBnGHSw7k=",
+      "test.json": "voBL9bHxxb4Sgoc7TC6W8uyXVsC/wEElhUg3BklfbL8="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.AveragePool" }
+}

build/webgpu/pool-global-reduction.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,93 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{% if source.useSubgroups %}
+enable subgroups;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Workgroup-parallel global pooling: one workgroup reduces one contiguous
+// (n, c) spatial plane (NC[D]HW layout makes the plane contiguous).
+// Threads stride the plane (vec4 loads when the plane size is divisible by 4),
+// then combine the per-thread partials. With subgroups: a subgroup collective
+// + a shared-memory fold of the per-subgroup partials. Without subgroups, a
+// full workgroup tree keeps workgroup-per-plane parallelism
+// instead of falling to the one-thread-per-plane serial scalar kernel. The
+// epilogue finalizes average / max / Lp. Sum reassociates;
+// max is exact.
+const WG: u32 = {{ source.workgroupSize }}u;
+{% if source.useSubgroups %}
+// WebGPU subgroups are at least 4 wide, so a workgroup holds at most WG / 4 subgroups.
+const MAX_SUBGROUPS: u32 = WG / 4u;
+var<workgroup> wg_partials: array<f32, MAX_SUBGROUPS>;
+{% else %}
+var<workgroup> wg_tree: array<f32, WG>;
+{% endif %}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>
+{%- if source.useSubgroups %},
+        @builtin(subgroup_size) sg_size: u32
+{%- endif %}) {
+  let plane = wid.x;
+  let tid = lid.x;
+{% if source.useSubgroups %}
+  // Seed cross-subgroup slots with the reduction identity so unwritten slots
+  // (when the workgroup has fewer than MAX_SUBGROUPS subgroups) are harmless.
+  if (tid < MAX_SUBGROUPS) {
+    wg_partials[tid] = 0.0;
+  }
+  workgroupBarrier();
+{% endif %}
+  var partial = 0.0;
+{% if source.vec4 %}
+  let spatialVec = params.spatial / 4u;
+  let base = plane * spatialVec;
+  var accv = vec4<f32>(0.0);
+  for (var i = tid; i < spatialVec; i = i + WG) {
+    let v = vec4<f32>(x[base + i]);
+    accv = accv + v;
+  }
+  partial = accv.x + accv.y + accv.z + accv.w;
+{% else %}
+  let base = plane * params.spatial;
+  for (var i = tid; i < params.spatial; i = i + WG) {
+    let v = f32(x[base + i]);
+    partial = partial + v;
+  }
+{% endif %}
+{% if source.useSubgroups %}
+  let sg_sum = subgroupAdd(partial);
+  if (subgroupElect()) {
+    wg_partials[tid / sg_size] = sg_sum;
+  }
+  workgroupBarrier();
+  if (tid == 0u) {
+    let nsg = min((WG + sg_size - 1u) / sg_size, MAX_SUBGROUPS);
+    var total = 0.0;
+    for (var i = 0u; i < nsg; i = i + 1u) {
+      total = total + wg_partials[i];
+    }
+    y[plane] = {{ scalar }}(total / f32(params.spatial));
+  }
+{% else %}
+  // No-subgroups fallback: full workgroup-tree reduction of the per-thread partials.
+  wg_tree[tid] = partial;
+  workgroupBarrier();
+  for (var step = WG >> 1u; step > 0u; step = step >> 1u) {
+    if (tid < step) {
+      wg_tree[tid] = wg_tree[tid] + wg_tree[tid + step];
+    }
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    let total = wg_tree[0];
+    y[plane] = {{ scalar }}(total / f32(params.spatial));
+  }
+{% endif %}
+}

build/webgpu/pool-window-nd.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,91 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Direct {{ source.spatialRank }}D {{ source.op }}-pool traversal for a channels-first tensor.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  let index = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (index >= params.count) { return; }
+  let ow = index % params.outW;
+  var t = index / params.outW;
+{% if source.spatialRank >= 2 %}
+  let oh = t % params.outH;
+  t /= params.outH;
+{% endif %}
+{% if source.spatialRank == 3 %}
+  let od = t % params.outD;
+  t /= params.outD;
+{% endif %}
+  let channel = t % params.channels;
+  let batch = t / params.channels;
+{% if source.spatialRank == 3 %}
+  let startD = i32(od * params.strideD) - params.padFront;
+{% endif %}
+{% if source.spatialRank >= 2 %}
+  let startH = i32(oh * params.strideH) - params.padTop;
+{% endif %}
+  let startW = i32(ow * params.strideW) - params.padLeft;
+  var acc = 0.0;
+  var seen = 0u;
+{% if source.spatialRank == 3 %}
+  for (var kd = 0u; kd < params.kernelD; kd += 1u) {
+    let id = startD + i32(kd * params.dilationD);
+    if (id < 0 || id >= i32(params.inD)) { continue; }
+{% endif %}
+{% if source.spatialRank >= 2 %}
+    for (var kh = 0u; kh < params.kernelH; kh += 1u) {
+      let ih = startH + i32(kh * params.dilationH);
+      if (ih < 0 || ih >= i32(params.inH)) { continue; }
+{% endif %}
+      for (var kw = 0u; kw < params.kernelW; kw += 1u) {
+        let iw = startW + i32(kw * params.dilationW);
+        if (iw < 0 || iw >= i32(params.inW)) { continue; }
+{% if source.spatialRank == 1 %}
+        let xIndex = (batch * params.channels + channel) * params.inW + u32(iw);
+{% elif source.spatialRank == 2 %}
+        let xIndex = ((batch * params.channels + channel) * params.inH + u32(ih)) * params.inW + u32(iw);
+{% else %}
+        let xIndex = (((batch * params.channels + channel) * params.inD + u32(id)) * params.inH + u32(ih)) * params.inW + u32(iw);
+{% endif %}
+        acc += f32(x[xIndex]);
+        seen += 1u;
+      }
+{% if source.spatialRank >= 2 %}
+    }
+{% endif %}
+{% if source.spatialRank == 3 %}
+  }
+{% endif %}
+  var denom = seen;
+  if (params.countIncludePad != 0u) {
+    denom = 0u;
+{% if source.spatialRank == 3 %}
+    for (var kd = 0u; kd < params.kernelD; kd += 1u) {
+      let id = startD + i32(kd * params.dilationD);
+      if (id < -params.padFront || id >= i32(params.inD) + params.padBack) { continue; }
+{% endif %}
+{% if source.spatialRank >= 2 %}
+      for (var kh = 0u; kh < params.kernelH; kh += 1u) {
+        let ih = startH + i32(kh * params.dilationH);
+        if (ih < -params.padTop || ih >= i32(params.inH) + params.padBottom) { continue; }
+{% endif %}
+        for (var kw = 0u; kw < params.kernelW; kw += 1u) {
+          let iw = startW + i32(kw * params.dilationW);
+          if (iw < -params.padLeft || iw >= i32(params.inW) + params.padRight) { continue; }
+          denom += 1u;
+        }
+{% if source.spatialRank >= 2 %}
+      }
+{% endif %}
+{% if source.spatialRank == 3 %}
+    }
+{% endif %}
+  }
+  y[index] = {{ scalar }}(select(0.0, acc / f32(denom), denom > 0u));
+}

build/webgpu/pool-window-unroll.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,80 @@

+// Compile-time-unrolled 2D NCHW pooling for small windows. Accumulation remains
+// kh-major/kw-minor, matching the runtime-loop fallbacks.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const OUT_W: u32 = {{ source.outW }}u;
+const OUT_H: u32 = {{ source.outH }}u;
+const IN_H: i32 = {{ source.inH }};
+const IN_W: i32 = {{ source.inW }};
+const IN_W_U: u32 = {{ source.inW }}u;
+const IN_PLANE: u32 = {{ source.inH * source.inW }}u;
+const STRIDE_H: u32 = {{ source.strideH }}u;
+const STRIDE_W: u32 = {{ source.strideW }}u;
+const PAD_TOP: i32 = {{ source.padTop }};
+const PAD_LEFT: i32 = {{ source.padLeft }};
+{% if source.op == "average" and source.countIncludePad %}
+const PAD_BOTTOM: i32 = {{ source.padBottom }};
+const PAD_RIGHT: i32 = {{ source.padRight }};
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  // Fold very large flat outputs into the dispatch's second dimension.
+  let index = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (index >= params.count) {
+    return;
+  }
+  let ow = index % OUT_W;
+  var plane = index / OUT_W;
+  let oh = plane % OUT_H;
+  plane /= OUT_H;
+  let planeBase = plane * IN_PLANE;
+  let startH = i32(oh * STRIDE_H) - PAD_TOP;
+  let startW = i32(ow * STRIDE_W) - PAD_LEFT;
+{% for kh in range(source.kernelH) %}
+  let ih_{{ kh }} = startH + {{ kh * source.dilationH }};
+{% endfor %}
+{% for kw in range(source.kernelW) %}
+  let iw_{{ kw }} = startW + {{ kw * source.dilationW }};
+{% endfor %}
+  var result = 0.0;
+  var seen = 0u;
+{% for kh in range(source.kernelH) %}
+  if (ih_{{ kh }} >= 0 && ih_{{ kh }} < IN_H) {
+    let rowBase{{ kh }} = planeBase + u32(ih_{{ kh }}) * IN_W_U;
+{% for kw in range(source.kernelW) %}
+    if (iw_{{ kw }} >= 0 && iw_{{ kw }} < IN_W) {
+      let value_{{ kh }}_{{ kw }} = f32(x[rowBase{{ kh }} + u32(iw_{{ kw }})]);
+      result += value_{{ kh }}_{{ kw }};
+      seen += 1u;
+    }
+{% endfor %}
+  }
+{% endfor %}
+{% if source.op == "average" and source.countIncludePad %}
+  var denominator = 0u;
+{% for kh in range(source.kernelH) %}
+  if (ih_{{ kh }} >= -PAD_TOP && ih_{{ kh }} < IN_H + PAD_BOTTOM) {
+{% for kw in range(source.kernelW) %}
+    if (iw_{{ kw }} >= -PAD_LEFT && iw_{{ kw }} < IN_W + PAD_RIGHT) {
+      denominator += 1u;
+    }
+{% endfor %}
+  }
+{% endfor %}
+{% else %}
+  let denominator = seen;
+{% endif %}
+  let output = select(0.0, result / f32(denominator), denominator > 0u);
+  y[index] = {{ scalar }}(output);
+}

build/webgpu/test.json ADDED Viewed

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