sync 2e7068faf55e

Browse files

Files changed (12) hide show

README.md +95 -0
build/webgpu/bench.json +293 -0
build/webgpu/conv-int-accumulate-spatial.wgsl.jinja +123 -0
build/webgpu/conv-int-im2col-spatial.wgsl.jinja +83 -0
build/webgpu/manifest.json +1845 -0
build/webgpu/metadata.json +24 -0
build/webgpu/qlinear-conv-dp4a-direct-3x3.wgsl.jinja +198 -0
build/webgpu/qlinear-conv-nchw-x4.wgsl.jinja +162 -0
build/webgpu/qlinear-conv-requantize.wgsl.jinja +76 -0
build/webgpu/quant-dp4a-matmul.wgsl.jinja +344 -0
build/webgpu/quant-pad-k-matrices.wgsl.jinja +50 -0
build/webgpu/test.json +1602 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.QLinearConv
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 10
+## Description
+Performs convolution on a quantized input using quantized weights, producing a quantized output. Inputs, weights, and output each carry their own scale and zero-point tensors (per-tensor or per-output-channel); an optional int32 bias must be pre-quantized with `scale = x_scale * w_scale` and zero point 0.
+See the [ONNX `QLinearConv` spec](https://onnx.ai/onnx/operators/onnx__QLinearConv.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `x` | `x` | `TX` | — | — | Quantized input data tensor shaped `(N x C x D1 x ... x Dn)`. | required |
+| `x_scale` | `x_scale` | `TF` | — | — | Per-tensor scale for input `x`. | required |
+| `x_zero_point` | `x_zero_point` | `TX` | — | — | Per-tensor zero point for input `x`. | required |
+| `w` | `w` | `TW` | — | — | Quantized weight tensor shaped `(M x C/group x k1 x ... x kn)`. | required |
+| `w_scale` | `w_scale` | `TF` | — | — | Scale for weights `w`; scalar for per-tensor or 1-D of length M for per-output-channel quantization. | required |
+| `w_zero_point` | `w_zero_point` | `TW` | — | — | Zero point for weights `w`; scalar or 1-D of length M matching `w_scale`. | required |
+| `y_scale` | `y_scale` | `TF` | — | — | Per-tensor scale for output `y`. | required |
+| `y_zero_point` | `y_zero_point` | `TY` | — | — | Per-tensor zero point for output `y`. | required |
+| `B` | `B` | `int32` | `1` | — | Optional 1-D bias of length `M`, pre-quantized with scale `x_scale * w_scale` and zero point 0. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `y` | `y` | `TY` | same as `x` | derived; see description | Quantized output tensor; shape is determined by kernel size, strides, and padding. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `auto_pad` | `"NOTSET"` | Automatic padding mode. `NOTSET` uses `pads`; `SAME_UPPER` and `SAME_LOWER` choose padding so each output spatial size is `ceil(input / stride)`; `VALID` uses no padding. |
+| `group` | `1` | Number of groups that input and output channels are split into; defaults to 1. |
+| `dilations` | — | Optional dilation factors, one positive integer per spatial axis. Omission means all ones. |
+| `kernel_shape` | — | Optional kernel shape, one positive integer per spatial axis. When present, it must match the spatial dimensions of the weight tensor; omission infers the shape from the weights. |
+| `pads` | — | Optional explicit padding in ONNX order `[begin_axis_0, ..., begin_axis_n, end_axis_0, ..., end_axis_n]`. Omission means all zeros; it cannot be combined with an automatic padding mode. |
+| `strides` | — | Optional stride factors, one positive integer per spatial axis. Omission means all ones. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `TX` | `uint8`, `int8` |
+| `TW` | `uint8`, `int8` |
+| `TY` | `uint8`, `int8` |
+| `TF` | `float32` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`conv-int-accumulate-spatial.wgsl.jinja`](build/webgpu/conv-int-accumulate-spatial.wgsl.jinja)
+- [`conv-int-im2col-spatial.wgsl.jinja`](build/webgpu/conv-int-im2col-spatial.wgsl.jinja)
+- [`qlinear-conv-dp4a-direct-3x3.wgsl.jinja`](build/webgpu/qlinear-conv-dp4a-direct-3x3.wgsl.jinja)
+- [`qlinear-conv-nchw-x4.wgsl.jinja`](build/webgpu/qlinear-conv-nchw-x4.wgsl.jinja)
+- [`qlinear-conv-requantize.wgsl.jinja`](build/webgpu/qlinear-conv-requantize.wgsl.jinja)
+- [`quant-dp4a-matmul.wgsl.jinja`](build/webgpu/quant-dp4a-matmul.wgsl.jinja)
+- [`quant-pad-k-matrices.wgsl.jinja`](build/webgpu/quant-pad-k-matrices.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.QLinearConv", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 1, 1] },
+  x_scale: { data: x_scaleData, shape: [1] },
+  x_zero_point: { data: x_zero_pointData, shape: [1] },
+  w: { data: wData, shape: [1, 1, 1] },
+  w_scale: { data: w_scaleData, shape: [1] },
+  w_zero_point: { data: w_zero_pointData, shape: [1] },
+  y_scale: { data: y_scaleData, shape: [1] },
+  y_zero_point: { data: y_zero_pointData, shape: [1] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,293 @@

+{
+  "op": "ai.onnx.QLinearConv",
+  "tunableSpace": {
+    "DP4A_DIRECT_SPATIAL_TILE": [8, 16],
+    "DP4A_DIRECT_OC_TILE": [8, 16],
+    "DP4A_DIRECT_IC_TILE": [16, 32, 64]
+  },
+  "cases": [
+    {
+      "name": "u8_nchw_1x16x32x32",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 16, 32, 32], "data": { "kind": "constant", "value": 127 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [16, 16, 3, 3], "data": { "kind": "constant", "value": 129 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.03] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.04] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 16, 30, 30] } },
+      "attrs": {}
+    },
+    {
+      "name": "u8s8_pointwise_1x64x56x56_oc128",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 56, 56], "data": { "kind": "constant", "value": 127 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 64, 1, 1], "data": { "kind": "constant", "value": -3 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.03] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 56, 56] } },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_cin64_aligned_1x56x56_oc128_u8s8",
+      "preset": "smoke",
+      "vars": { "batch": 1, "inChannels": 64, "outChannels": 128, "inH": 56, "inW": 56, "kernelH": 1, "kernelW": 1 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 56, 56], "dist": "randint", "seed": 7301, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 64, 1, 1], "dist": "randint", "seed": 7302, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 56, 56] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels"
+          }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_cin66_unaligned_k4_cliff_1x56x56_oc128_u8s8",
+      "preset": "smoke",
+      "vars": { "batch": 1, "inChannels": 66, "outChannels": 128, "inH": 56, "inW": 56, "kernelH": 1, "kernelW": 1 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 66, 56, 56], "dist": "randint", "seed": 7303, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 66, 1, 1], "dist": "randint", "seed": 7304, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 56, 56] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels"
+          }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv3x3_naive_resnet_1x128x56x56_oc128_u8s8",
+      "preset": "smoke",
+      "vars": { "batch": 1, "inChannels": 128, "outChannels": 128, "inH": 56, "inW": 56, "kernelH": 3, "kernelW": 3 },
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 128, 56, 56], "dist": "randint", "seed": 7305, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 128, 3, 3], "dist": "randint", "seed": 7306, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 56, 56] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "depthwise3x3_naive_mobilenet_1x256x28x28_g256_u8s8",
+      "preset": "smoke",
+      "vars": { "batch": 1, "inChannels": 256, "outChannels": 256, "inH": 28, "inW": 28, "kernelH": 3, "kernelW": 3 },
+      "attrs": { "group": 256, "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 256, 28, 28], "dist": "randint", "seed": 7307, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [256, 1, 3, 3], "dist": "randint", "seed": 7308, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 256, 28, 28] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bias3x3_forces_naive_accumulate_1x128x64x64_oc128_u8s8",
+      "preset": "stress",
+      "vars": { "batch": 1, "inChannels": 128, "outChannels": 128, "inH": 64, "inW": 64, "kernelH": 3, "kernelW": 3 },
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 128, 64, 64], "dist": "randint", "seed": 7401, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 128, 3, 3], "dist": "randint", "seed": 7402, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "B": { "dtype": "int32", "shape": [128], "dist": "randint", "seed": 7403, "min": -4096, "max": 4096 }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 64, 64] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "unaligned_k_padded_im2col_1x63x48x48_oc128_3x3_u8s8",
+      "preset": "stress",
+      "vars": { "batch": 1, "inChannels": 63, "outChannels": 128, "inH": 48, "inW": 48, "kernelH": 3, "kernelW": 3 },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 63, 48, 48], "dist": "randint", "seed": 7404, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 63, 3, 3], "dist": "randint", "seed": 7405, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 46, 46] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "bias_unaligned_k_baseline_fallback_1x63x48x48_oc128_3x3_u8s8",
+      "preset": "stress",
+      "vars": { "batch": 1, "inChannels": 63, "outChannels": 128, "inH": 48, "inW": 48, "kernelH": 3, "kernelW": 3 },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 63, 48, 48], "dist": "randint", "seed": 7410, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 63, 3, 3], "dist": "randint", "seed": 7411, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "B": { "dtype": "int32", "shape": [128], "dist": "randint", "seed": 7412, "min": -4096, "max": 4096 }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 46, 46] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * args.batch * args.outChannels * dim(shapes.y, 2) * dim(shapes.y, 3) * args.inChannels * args.kernelH * args.kernelW"
+          }
+        ]
+      }
+    },
+    {
+      "name": "u8s8_conv3d_dilated_depth_b1c16m32_16x32x32_k3",
+      "preset": "stress",
+      "attrs": { "strides": [1, 1, 1], "dilations": [2, 1, 1], "pads": [2, 1, 1, 2, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 16, 16, 32, 32], "dist": "randint", "seed": 7701, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [32, 16, 3, 3, 3], "dist": "randint", "seed": 7702, "min": -127, "max": 127 },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 32, 16, 32, 32], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4)"
+          }
+        ]
+      }
+    },
+    {
+      "name": "resnet_perchannel_3x3_b1c64_56x56_oc128_u8s8",
+      "preset": "stress",
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 56, 56], "dist": "randint", "seed": 7801, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 64, 3, 3], "dist": "randint", "seed": 7802, "min": -127, "max": 127 },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [0.0125, 0.015, 0.0175, 0.02] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [128], "data": { "kind": "cycle", "values": [-3, 0, 2, 5] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 56, 56] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "group2_feature_independent_oc4_cliff_b1c64_56x56_oc128_u8s8",
+      "preset": "stress",
+      "vars": { "batch": 1, "inChannels": 64, "outChannels": 128, "inH": 56, "inW": 56, "kernelH": 3, "kernelW": 3 },
+      "attrs": { "group": 2, "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 56, 56], "dist": "randint", "seed": 7811, "min": 0, "max": 255 },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [128, 32, 3, 3], "dist": "randint", "seed": 7812, "min": -127, "max": 127 },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [0.0125, 0.015, 0.0175, 0.02] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [128], "data": { "kind": "cycle", "values": [-3, 0, 2, 5] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 128, 56, 56] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/conv-int-accumulate-spatial.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,123 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% set depthIndent = "  " if source.spatialRank == 3 else "" %}
+{{ env.wgsl.resourceDeclarations }}
+fn read_x(index: u32) -> i32 {
+  {% if xUnsigned %}
+  return i32(x[index]);
+  {% else %}
+  return x[index];
+  {% endif %}
+}
+fn read_w(index: u32) -> i32 {
+  {% if wUnsigned %}
+  return i32(w[index]);
+  {% else %}
+  return w[index];
+  {% endif %}
+}
+fn read_x_zero() -> i32 {
+  {% if xUnsigned %}
+  return i32(x_zero_point[0]);
+  {% else %}
+  return x_zero_point[0];
+  {% endif %}
+}
+// A per-output-channel weight zero point applies at every spatial rank.
+fn read_w_zero({% if wZeroPerChannel %}oc: u32{% endif %}) -> i32 {
+  {% if wZeroPerChannel %}
+  {% if wUnsigned %}
+  return i32(w_zero_point[oc]);
+  {% else %}
+  return w_zero_point[oc];
+  {% endif %}
+  {% else %}
+  {% if wUnsigned %}
+  return i32(w_zero_point[0]);
+  {% else %}
+  return w_zero_point[0];
+  {% endif %}
+  {% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", guardInline=true) }}
+  let ow = index % params.outW;
+  var t = index / params.outW;
+  let oh = t % params.outH;
+  t = t / params.outH;
+  {% if source.spatialRank == 3 %}
+  let od = t % params.outD;
+  t = t / params.outD;
+  {% endif %}
+  let oc = t % params.outChannels;
+  let batch = t / params.outChannels;
+  let group = oc / params.outChannelsPerGroup;
+  let xzp = read_x_zero();
+  let wzp = read_w_zero({% if wZeroPerChannel %}oc{% endif %});
+  var acc = 0i;
+  for (var ic = 0u; ic < params.weightInChannels; ic = ic + 1u) {
+    let input_channel = group * params.inChannelsPerGroup + ic;
+    {% if source.spatialRank == 3 %}
+    for (var kd = 0u; kd < params.kernelD; kd = kd + 1u) {
+      let id = i32(od * params.strideD + kd * params.dilationD) - params.padD;
+      if (id < 0 || id >= i32(params.inD)) { continue; }
+    {% endif %}
+{{ depthIndent }}    for (var kh = 0u; kh < params.kernelH; kh = kh + 1u) {
+{{ depthIndent }}      let ih = i32(oh * params.strideH + kh * params.dilationH) - params.padH;
+{{ depthIndent }}      if (ih < 0 || ih >= i32(params.inH)) { continue; }
+{{ depthIndent }}      for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+{{ depthIndent }}        let iw = i32(ow * params.strideW + kw * params.dilationW) - params.padW;
+{{ depthIndent }}        if (iw < 0 || iw >= i32(params.inW)) { continue; }
+        {% if source.spatialRank == 3 %}
+{{ depthIndent }}        let x_index = (((batch * params.inChannels + input_channel) * params.inD + u32(id)) * params.inH + u32(ih)) * params.inW + u32(iw);
+{{ depthIndent }}        let w_index = (((oc * params.weightInChannels + ic) * params.kernelD + kd) * params.kernelH + kh) * params.kernelW + kw;
+        {% else %}
+{{ depthIndent }}        let x_index = ((batch * params.inChannels + input_channel) * params.inH + u32(ih)) * params.inW + u32(iw);
+{{ depthIndent }}        let w_index = ((oc * params.weightInChannels + ic) * params.kernelH + kh) * params.kernelW + kw;
+        {% endif %}
+{{ depthIndent }}        acc = acc + (read_x(x_index) - xzp) * (read_w(w_index) - wzp);
+{{ depthIndent }}      }
+{{ depthIndent }}    }
+    {% if source.spatialRank == 3 %}
+    }
+    {% endif %}
+  }
+  y[index] = acc;
+}

build/webgpu/conv-int-im2col-spatial.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,83 @@

+/* Widened-int8 im2col for convolution DP4A paths. The materialized K order is
+ * OIHW for 2D input and OIDHW for 3D input. Spatial padding is filled with the
+ * input zero point, so padded taps contribute exactly zero after centering.
+ * Adjacent x invocations cover adjacent output positions for coalesced writes.
+ * Optional extra K rows contain the raw input zero point to align the DP4A
+ * reduction. */
+{{ env.wgsl.resourceDeclarations }}
+{% if source.spatialRank == 3 %}
+const KERNEL_D: u32 = {{ source.kernelD }}u;
+{% endif %}
+const KERNEL_H: u32 = {{ source.kernelH }}u;
+const KERNEL_W: u32 = {{ source.kernelW }}u;
+const KERNEL_HW: u32 = KERNEL_H * KERNEL_W;
+{% if source.spatialRank == 3 %}
+const KSIZE: u32 = KERNEL_D * KERNEL_HW;
+const STRIDE_D: u32 = {{ source.strideD }}u;
+const DILATION_D: u32 = {{ source.dilationD }}u;
+const PAD_FRONT: i32 = {{ source.padFront }};
+{% else %}
+const KSIZE: u32 = KERNEL_HW;
+{% endif %}
+const STRIDE_H: u32 = {{ source.strideH }}u;
+const STRIDE_W: u32 = {{ source.strideW }}u;
+const DILATION_H: u32 = {{ source.dilationH }}u;
+const DILATION_W: u32 = {{ source.dilationW }}u;
+const PAD_TOP: i32 = {{ source.padTop }};
+const PAD_LEFT: i32 = {{ source.padLeft }};
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let pos = gid.x;
+  if (pos >= params.outCount) {
+    return;
+  }
+  let k = gid.y;
+  let batch = gid.z;
+{% if source.padK %}
+  if (k >= params.sourceKRows) {
+    cols[(batch * params.kRows + k) * params.outCount + pos] = x_zero_point[0];
+    return;
+  }
+{% endif %}
+  let ic = k / KSIZE;
+  let kq = k - ic * KSIZE;
+{% if source.spatialRank == 3 %}
+  let kd = kq / KERNEL_HW;
+  let khw = kq - kd * KERNEL_HW;
+  let kh = khw / KERNEL_W;
+  let kw = khw - kh * KERNEL_W;
+  let outHW = params.outH * params.outW;
+  let od = pos / outHW;
+  let ohw = pos - od * outHW;
+  let oh = ohw / params.outW;
+  let ow = ohw - oh * params.outW;
+  let id = i32(od * STRIDE_D + kd * DILATION_D) - PAD_FRONT;
+{% else %}
+  let kh = kq / KERNEL_W;
+  let kw = kq - kh * KERNEL_W;
+  let oh = pos / params.outW;
+  let ow = pos - oh * params.outW;
+{% endif %}
+  let ih = i32(oh * STRIDE_H + kh * DILATION_H) - PAD_TOP;
+  let iw = i32(ow * STRIDE_W + kw * DILATION_W) - PAD_LEFT;
+  var value: {{ bScalar }} = x_zero_point[0];
+{% if source.spatialRank == 3 %}
+  if (
+    id >= 0 && id < i32(params.inD)
+    && ih >= 0 && ih < i32(params.inH)
+    && iw >= 0 && iw < i32(params.inW)
+  ) {
+    value = x[
+      (((batch * params.inChannels + ic) * params.inD + u32(id)) * params.inH + u32(ih))
+      * params.inW + u32(iw)
+    ];
+  }
+{% else %}
+  if (ih >= 0 && ih < i32(params.inH) && iw >= 0 && iw < i32(params.inW)) {
+    value = x[((batch * params.inChannels + ic) * params.inH + u32(ih)) * params.inW + u32(iw)];
+  }
+{% endif %}
+  cols[(batch * params.kRows + k) * params.outCount + pos] = value;
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1845 @@

+{
+  "domain": "ai.onnx",
+  "name": "QLinearConv",
+  "sinceVersion": 10,
+  "description": "Performs convolution on a quantized input using quantized weights, producing a quantized output. Inputs, weights, and output each carry their own scale and zero-point tensors (per-tensor or per-output-channel); an optional int32 bias must be pre-quantized with `scale = x_scale * w_scale` and zero point 0.",
+  "inputs": [
+    { "role": "x", "dtype": "TX", "description": "Quantized input data tensor shaped `(N x C x D1 x ... x Dn)`." },
+    { "role": "x_scale", "dtype": "TF", "description": "Per-tensor scale for input `x`." },
+    { "role": "x_zero_point", "dtype": "TX", "description": "Per-tensor zero point for input `x`." },
+    { "role": "w", "dtype": "TW", "description": "Quantized weight tensor shaped `(M x C/group x k1 x ... x kn)`." },
+    {
+      "role": "w_scale",
+      "dtype": "TF",
+      "description": "Scale for weights `w`; scalar for per-tensor or 1-D of length M for per-output-channel quantization."
+    },
+    {
+      "role": "w_zero_point",
+      "dtype": "TW",
+      "description": "Zero point for weights `w`; scalar or 1-D of length M matching `w_scale`."
+    },
+    { "role": "y_scale", "dtype": "TF", "description": "Per-tensor scale for output `y`." },
+    { "role": "y_zero_point", "dtype": "TY", "description": "Per-tensor zero point for output `y`." },
+    {
+      "role": "B",
+      "dtype": "int32",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional 1-D bias of length `M`, pre-quantized with scale `x_scale * w_scale` and zero point 0."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "y",
+      "dtype": "TY",
+      "description": "Quantized output tensor; shape is determined by kernel size, strides, and padding.",
+      "rank": "ranks.x",
+      "shape": "[dim(shapes.x, 0), dim(shapes.w, 0), expectedOutputWidth] if ranks.x == 3 else ([dim(shapes.x, 0), dim(shapes.w, 0), expectedOutputHeight, expectedOutputWidth] if ranks.x == 4 else [dim(shapes.x, 0), dim(shapes.w, 0), expectedOutputDepth, expectedOutputHeight, expectedOutputWidth])"
+    }
+  ],
+  "attributes": { "auto_pad": "NOTSET", "group": 1 },
+  "attributeDescriptions": {
+    "auto_pad": "Automatic padding mode. `NOTSET` uses `pads`; `SAME_UPPER` and `SAME_LOWER` choose padding so each output spatial size is `ceil(input / stride)`; `VALID` uses no padding.",
+    "dilations": "Optional dilation factors, one positive integer per spatial axis. Omission means all ones.",
+    "group": "Number of groups that input and output channels are split into; defaults to 1.",
+    "kernel_shape": "Optional kernel shape, one positive integer per spatial axis. When present, it must match the spatial dimensions of the weight tensor; omission infers the shape from the weights.",
+    "pads": "Optional explicit padding in ONNX order `[begin_axis_0, ..., begin_axis_n, end_axis_0, ..., end_axis_n]`. Omission means all zeros; it cannot be combined with an automatic padding mode.",
+    "strides": "Optional stride factors, one positive integer per spatial axis. Omission means all ones."
+  },
+  "attributeConstraints": { "auto_pad": { "values": ["NOTSET", "SAME_UPPER", "SAME_LOWER", "VALID"] } },
+  "typeConstraints": { "TX": ["uint8", "int8"], "TW": ["uint8", "int8"], "TY": ["uint8", "int8"], "TF": ["float32"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "x", "role": "input" },
+    "x_scale": { "kind": "tensor", "semantic": "x_scale", "role": "input" },
+    "x_zero_point": { "kind": "tensor", "semantic": "x_zero_point", "role": "input" },
+    "w": { "kind": "tensor", "semantic": "w", "role": "input" },
+    "w_scale": { "kind": "tensor", "semantic": "w_scale", "role": "input" },
+    "w_zero_point": { "kind": "tensor", "semantic": "w_zero_point", "role": "input" },
+    "y_scale": { "kind": "tensor", "semantic": "y_scale", "role": "input" },
+    "y_zero_point": { "kind": "tensor", "semantic": "y_zero_point", "role": "input" },
+    "B": { "kind": "tensor", "semantic": "B", "role": "input", "required": false },
+    "y": { "kind": "tensor", "semantic": "y", "role": "output" }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 256,
+    "DP4A_DIRECT_SPATIAL_TILE": 8,
+    "DP4A_DIRECT_OC_TILE": 8,
+    "DP4A_DIRECT_IC_TILE": 16
+  },
+  "derive": {
+    "inputDepth": "dim(shapes.x, 2) if ranks.x == 5 else 1",
+    "inputHeight": "dim(shapes.x, 3) if ranks.x == 5 else (dim(shapes.x, 2) if ranks.x == 4 else 1)",
+    "inputWidth": "dim(shapes.x, ranks.x - 1) if ranks.x >= 3 else 1",
+    "outputDepth": "dim(shapes.y, 2) if ranks.y == 5 else 1",
+    "outputHeight": "dim(shapes.y, 3) if ranks.y == 5 else (dim(shapes.y, 2) if ranks.y == 4 else 1)",
+    "outputWidth": "dim(shapes.y, ranks.y - 1) if ranks.y >= 3 else 1",
+    "kernelDepth": "dim(shapes.w, 2) if ranks.w == 5 else 1",
+    "kernelHeight": "dim(shapes.w, 3) if ranks.w == 5 else (dim(shapes.w, 2) if ranks.w == 4 else 1)",
+    "kernelWidth": "dim(shapes.w, ranks.w - 1) if ranks.w >= 3 else 1",
+    "spatialRank": "ranks.w - 2",
+    "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",
+    "narrowSubgroupRange": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < device.adapterInfo.subgroupMaxSize and device.adapterInfo.subgroupMaxSize <= 16",
+    "autoPadSame": "attrs.auto_pad == \"SAME_UPPER\" or attrs.auto_pad == \"SAME_LOWER\"",
+    "autoPadValid": "attrs.auto_pad == \"VALID\"",
+    "samePadDepth": "max(0, (outputDepth - 1) * strideD + (kernelDepth - 1) * dilationD + 1 - inputDepth)",
+    "samePadHeight": "max(0, (outputHeight - 1) * strideH + (kernelHeight - 1) * dilationH + 1 - inputHeight)",
+    "samePadWidth": "max(0, (outputWidth - 1) * strideW + (kernelWidth - 1) * dilationW + 1 - inputWidth)",
+    "samePadFront": "floor(samePadDepth / 2) if attrs.auto_pad == \"SAME_UPPER\" else samePadDepth - floor(samePadDepth / 2)",
+    "samePadTop": "floor(samePadHeight / 2) if attrs.auto_pad == \"SAME_UPPER\" else samePadHeight - floor(samePadHeight / 2)",
+    "samePadLeft": "floor(samePadWidth / 2) if attrs.auto_pad == \"SAME_UPPER\" else samePadWidth - floor(samePadWidth / 2)",
+    "effectivePadFront": "samePadFront if autoPadSame else (0 if autoPadValid else padFront)",
+    "effectivePadTop": "samePadTop if autoPadSame else (0 if autoPadValid else padTop)",
+    "effectivePadLeft": "samePadLeft if autoPadSame else (0 if autoPadValid else padLeft)",
+    "expectedOutputDepth": "ceil(inputDepth / strideD) if autoPadSame else floor((inputDepth + (0 if autoPadValid else padFront + padBack) - ((kernelDepth - 1) * dilationD + 1)) / strideD) + 1",
+    "expectedOutputHeight": "ceil(inputHeight / strideH) if autoPadSame else floor((inputHeight + (0 if autoPadValid else padTop + padBottom) - ((kernelHeight - 1) * dilationH + 1)) / strideH) + 1",
+    "expectedOutputWidth": "ceil(inputWidth / strideW) if autoPadSame else floor((inputWidth + (0 if autoPadValid else padLeft + padRight) - ((kernelWidth - 1) * dilationW + 1)) / strideW) + 1",
+    "spatialAttributeLengthsOk": "kernelShapeLengthOk and stridesLengthOk and dilationsLengthOk and padsLengthOk",
+    "kernelShapeMatchesWeights": "not has(attrs, \"kernel_shape\") or (kernelW == kernelWidth and (spatialRank < 2 or kernelH == kernelHeight) and (spatialRank < 3 or kernelD == kernelDepth))",
+    "kernelExtentsOk": "kernelWidth >= 1 and (spatialRank < 2 or kernelHeight >= 1) and (spatialRank < 3 or kernelDepth >= 1)",
+    "stridesValuesOk": "not has(attrs, \"strides\") or (strideD >= 1 and floor(strideD) == strideD and strideH >= 1 and floor(strideH) == strideH and strideW >= 1 and floor(strideW) == strideW)",
+    "dilationsValuesOk": "not has(attrs, \"dilations\") or (dilationD >= 1 and floor(dilationD) == dilationD and dilationH >= 1 and floor(dilationH) == dilationH and dilationW >= 1 and floor(dilationW) == dilationW)",
+    "padsValuesOk": "padFront >= 0 and floor(padFront) == padFront and padTop >= 0 and floor(padTop) == padTop and padLeft >= 0 and floor(padLeft) == padLeft and padBack >= 0 and floor(padBack) == padBack and padBottom >= 0 and floor(padBottom) == padBottom and padRight >= 0 and floor(padRight) == padRight",
+    "explicitPadsOk": "attrs.auto_pad == \"NOTSET\" or not has(attrs, \"pads\")",
+    "spatialAttributesOk": "spatialRank >= 1 and spatialRank <= 3 and spatialAttributeLengthsOk and kernelShapeMatchesWeights and kernelExtentsOk and stridesValuesOk and dilationsValuesOk and padsValuesOk and explicitPadsOk",
+    "packedFeature": "device.wgslLanguageFeatures.has(\"packed_4x8_integer_dot_product\")",
+    "xyQuantContract": "(ranks.x_scale == 0 or (ranks.x_scale == 1 and dim(shapes.x_scale, 0) == 1)) and (ranks.y_scale == 0 or (ranks.y_scale == 1 and dim(shapes.y_scale, 0) == 1)) and (ranks.x_zero_point == 0 or (ranks.x_zero_point == 1 and dim(shapes.x_zero_point, 0) == 1)) and (ranks.y_zero_point == 0 or (ranks.y_zero_point == 1 and dim(shapes.y_zero_point, 0) == 1))",
+    "rank3TensorOk": "spatialAttributesOk and ranks.x == 3 and ranks.w == 3 and ranks.y == 3",
+    "rank4TensorOk": "spatialAttributesOk and (ranks.x == 4 and ranks.w == 4 and ranks.y == 4)",
+    "rank5TensorOk": "spatialAttributesOk and ranks.x == 5 and ranks.w == 5 and ranks.y == 5",
+    "outputBatchChannelsOk": "dim(shapes.y, 0) == dim(shapes.x, 0) and dim(shapes.y, 1) == dim(shapes.w, 0)",
+    "groupChannelsOk": "attrs.group >= 1 and dim(shapes.w, 1) * attrs.group == dim(shapes.x, 1) and dim(shapes.w, 0) % attrs.group == 0",
+    "output2dShapeOk": "dim(shapes.y, 2) == expectedOutputHeight and dim(shapes.y, 3) == expectedOutputWidth",
+    "scalarWeightQuantOk": "(ranks.w_scale == 0 or (ranks.w_scale == 1 and dim(shapes.w_scale, 0) == 1)) and (ranks.w_zero_point == 0 or (ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == 1))",
+    "flexibleWeightQuantOk": "(((ranks.w_scale == 0 or (ranks.w_scale == 1 and dim(shapes.w_scale, 0) == 1)) and (ranks.w_zero_point == 0 or (ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == 1))) or (ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) and ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)))",
+    "dp4a2dDispatchOk": "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3), 64) <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(dim(shapes.w, 0), 8) <= device.limits.maxComputeWorkgroupsPerDimension and dim(shapes.y, 0) <= device.limits.maxComputeWorkgroupsPerDimension",
+    "convM": "dim(shapes.w, 0)",
+    "convN": "inner(shapes.y, 1)",
+    "convK": "dim(shapes.w, 1) * inner(shapes.w, 1)",
+    "convBStride": "dim(shapes.w, 1) * inner(shapes.w, 1) * inner(shapes.y, 1)",
+    "convYStride": "dim(shapes.w, 0) * inner(shapes.y, 1)",
+    "paddedK": "ceilDiv(dim(shapes.w, 1) * inner(shapes.w, 1), 4) * 4",
+    "paddedBStride": "ceilDiv(dim(shapes.w, 1) * inner(shapes.w, 1), 4) * 4 * inner(shapes.y, 1)",
+    "paddedMatricesStorageOk": "dim(shapes.x, 0) * ceilDiv(dim(shapes.w, 1) * inner(shapes.w, 1), 4) * 4 * inner(shapes.y, 1) * 4 <= device.limits.maxStorageBufferBindingSize and dim(shapes.w, 0) * ceilDiv(dim(shapes.w, 1) * inner(shapes.w, 1), 4) * 4 * 4 <= device.limits.maxStorageBufferBindingSize",
+    "im2colStorageOk": "dim(shapes.x, 0) * dim(shapes.w, 1) * inner(shapes.w, 1) * inner(shapes.y, 1) * 4 <= device.limits.maxStorageBufferBindingSize",
+    "paddedPreparationDispatchOk": "ceilDiv(inner(shapes.y, 1), tunables.WORKGROUP_SIZE) <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(dim(shapes.w, 1) * inner(shapes.w, 1), 4) * 4 <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(dim(shapes.w, 0) * ceilDiv(dim(shapes.w, 1) * inner(shapes.w, 1), 4) * 4, tunables.WORKGROUP_SIZE) <= device.limits.maxComputeWorkgroupsPerDimension",
+    "direct3x3SharedBytes": "4 * ((tunables.DP4A_DIRECT_IC_TILE / 4) * (tunables.DP4A_DIRECT_SPATIAL_TILE + 2) * (tunables.DP4A_DIRECT_SPATIAL_TILE + 2) + tunables.DP4A_DIRECT_OC_TILE * (tunables.DP4A_DIRECT_IC_TILE / 4) * 9 + tunables.DP4A_DIRECT_OC_TILE)",
+    "direct3x3DispatchOk": "ceilDiv(dim(shapes.y, 3), tunables.DP4A_DIRECT_SPATIAL_TILE) <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(dim(shapes.y, 2), tunables.DP4A_DIRECT_SPATIAL_TILE) * ceilDiv(dim(shapes.w, 0), tunables.DP4A_DIRECT_OC_TILE) <= device.limits.maxComputeWorkgroupsPerDimension and dim(shapes.y, 0) <= device.limits.maxComputeWorkgroupsPerDimension",
+    "direct3x3DeviceOk": "tunables.DP4A_DIRECT_IC_TILE % 4 == 0 and tunables.DP4A_DIRECT_SPATIAL_TILE * tunables.DP4A_DIRECT_SPATIAL_TILE <= device.limits.maxComputeInvocationsPerWorkgroup and tunables.DP4A_DIRECT_SPATIAL_TILE <= device.limits.maxComputeWorkgroupSizeX and tunables.DP4A_DIRECT_SPATIAL_TILE <= device.limits.maxComputeWorkgroupSizeY and direct3x3SharedBytes <= device.limits.maxComputeWorkgroupStorageSize and direct3x3DispatchOk"
+  },
+  "bindingSets": {
+    "padPointwiseK": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      { "name": "b", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      { "name": "a_padded", "semantic": "w_padded", "buffer": { "type": "storage" }, "elementType": "i32" },
+      { "name": "b_padded", "semantic": "x_padded", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batch", "type": "u32", "value": "dim(shapes.x, 0)" },
+            { "name": "M", "type": "u32", "value": "dim(shapes.w, 0)" },
+            { "name": "N", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
+            { "name": "K", "type": "u32", "value": "dim(shapes.x, 1)" },
+            { "name": "KPadded", "type": "u32", "value": "ceilDiv(dim(shapes.x, 1), 4) * 4" }
+          ]
+        }
+      }
+    ],
+    "dp4aPaddedMain": [
+      { "name": "a", "semantic": "w_padded", "buffer": { "type": "read-only-storage" }, "elementType": "vec4<i32>" },
+      { "name": "b", "semantic": "x_padded", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "a_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "b_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "convM" },
+            { "name": "N", "type": "u32", "value": "convN" },
+            { "name": "K", "type": "u32", "value": "paddedK" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "paddedBStride" },
+            { "name": "yBatchStride", "type": "u32", "value": "convYStride" }
+          ]
+        }
+      }
+    ],
+    "dp4aPaddedIm2colMain": [
+      { "name": "a", "semantic": "w_padded", "buffer": { "type": "read-only-storage" }, "elementType": "vec4<i32>" },
+      { "name": "b", "semantic": "cols_padded", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      {
+        "name": "a_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "b_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "convM" },
+            { "name": "N", "type": "u32", "value": "convN" },
+            { "name": "K", "type": "u32", "value": "paddedK" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "paddedBStride" },
+            { "name": "yBatchStride", "type": "u32", "value": "convYStride" }
+          ]
+        }
+      }
+    ],
+    "im2colNchwPadded": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      { "name": "cols", "semantic": "cols_padded", "buffer": { "type": "storage" }, "elementType": "u32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "outCount", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
+            { "name": "inChannels", "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": "sourceKRows", "type": "u32", "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" },
+            {
+              "name": "kRows",
+              "type": "u32",
+              "value": "ceilDiv(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3), 4) * 4"
+            },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" }
+          ]
+        }
+      }
+    ],
+    "padWeightsK": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      { "name": "a_padded", "semantic": "w_padded", "buffer": { "type": "storage" }, "elementType": "i32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "dim(shapes.w, 0)" },
+            { "name": "K", "type": "u32", "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" },
+            {
+              "name": "KPadded",
+              "type": "u32",
+              "value": "ceilDiv(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3), 4) * 4"
+            }
+          ]
+        }
+      }
+    ],
+    "im2colNcdhw": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$bScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      { "name": "cols", "semantic": "cols3d", "buffer": { "type": "storage" }, "elementType": "$bScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "outCount", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4)" },
+            { "name": "inChannels", "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": "kRows",
+              "type": "u32",
+              "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4)"
+            },
+            { "name": "outH", "type": "u32", "value": "dim(shapes.y, 3)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.y, 4)" }
+          ]
+        }
+      }
+    ],
+    "im2colNchw": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$bScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      { "name": "cols", "semantic": "cols", "buffer": { "type": "storage" }, "elementType": "$bScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "outCount", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
+            { "name": "inChannels", "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": "kRows", "type": "u32", "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" }
+          ]
+        }
+      }
+    ],
+    "dp4aDirect3x3": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      {
+        "name": "w_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "u32",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "u32" }
+    ],
+    "dp4aMain": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$aVec4" },
+      {
+        "name": "b",
+        "arg": "cols",
+        "semantic": "cols",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar"
+      },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      {
+        "name": "a_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "b_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "convM" },
+            { "name": "N", "type": "u32", "value": "convN" },
+            { "name": "K", "type": "u32", "value": "convK" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "convBStride" },
+            { "name": "yBatchStride", "type": "u32", "value": "convYStride" }
+          ]
+        }
+      }
+    ],
+    "dp4aMain3d": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$aVec4" },
+      { "name": "b", "semantic": "cols3d", "buffer": { "type": "read-only-storage" }, "elementType": "$bScalar" },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      {
+        "name": "a_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "b_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "convM" },
+            { "name": "N", "type": "u32", "value": "convN" },
+            { "name": "K", "type": "u32", "value": "convK" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "convBStride" },
+            { "name": "yBatchStride", "type": "u32", "value": "convYStride" }
+          ]
+        }
+      }
+    ],
+    "dp4aPointwiseMain": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$aVec4" },
+      { "name": "b", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$bScalar" },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      {
+        "name": "a_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "b_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "convM" },
+            { "name": "N", "type": "u32", "value": "convN" },
+            { "name": "K", "type": "u32", "value": "convK" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "convBStride" },
+            { "name": "yBatchStride", "type": "u32", "value": "convYStride" }
+          ]
+        }
+      }
+    ],
+    "dp4aMainBias": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$aVec4" },
+      {
+        "name": "b",
+        "arg": "cols",
+        "semantic": "cols",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar"
+      },
+      {
+        "name": "a_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      {
+        "name": "a_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "b_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "bias", "arg": "B", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "M", "type": "u32", "value": "convM" },
+            { "name": "N", "type": "u32", "value": "convN" },
+            { "name": "K", "type": "u32", "value": "convK" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "convBStride" },
+            { "name": "yBatchStride", "type": "u32", "value": "convYStride" }
+          ]
+        }
+      }
+    ],
+    "accumulateNcw": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar"
+      },
+      { "name": "y", "semantic": "accum", "buffer": { "type": "storage" }, "elementType": "i32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
+            { "name": "inH", "type": "u32", "value": 1 },
+            { "name": "inW", "type": "u32", "value": "dim(shapes.x, 2)" },
+            { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+            { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
+            { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
+            { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.w, 0) / attrs.group" },
+            { "name": "kernelH", "type": "u32", "value": 1 },
+            { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 2)" },
+            { "name": "outH", "type": "u32", "value": 1 },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.y, 2)" },
+            { "name": "strideH", "type": "u32", "value": 1 },
+            { "name": "strideW", "type": "u32", "value": "strideW" },
+            { "name": "dilationH", "type": "u32", "value": 1 },
+            { "name": "dilationW", "type": "u32", "value": "dilationW" },
+            { "name": "padH", "type": "i32", "value": 0 },
+            { "name": "padW", "type": "i32", "value": "effectivePadLeft" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
+          ]
+        }
+      }
+    ],
+    "accumulateNcdhw": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar"
+      },
+      { "name": "y", "semantic": "accum", "buffer": { "type": "storage" }, "elementType": "i32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "inChannels", "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": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+            { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
+            { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
+            { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.w, 0) / attrs.group" },
+            { "name": "kernelD", "type": "u32", "value": "dim(shapes.w, 2)" },
+            { "name": "kernelH", "type": "u32", "value": "dim(shapes.w, 3)" },
+            { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 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": "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": "padD", "type": "i32", "value": "effectivePadFront" },
+            { "name": "padH", "type": "i32", "value": "effectivePadTop" },
+            { "name": "padW", "type": "i32", "value": "effectivePadLeft" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
+          ]
+        }
+      }
+    ],
+    "fusedNchwX4": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "w_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batch", "type": "u32", "value": "dim(shapes.x, 0)" },
+            { "name": "inChannels", "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": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+            { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
+            { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
+            { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.w, 0) / attrs.group" },
+            { "name": "kernelH", "type": "u32", "value": "dim(shapes.w, 2)" },
+            { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 3)" },
+            { "name": "outH", "type": "u32", "value": "dim(shapes.y, 2)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" },
+            { "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": "padH", "type": "i32", "value": "effectivePadTop" },
+            { "name": "padW", "type": "i32", "value": "effectivePadLeft" }
+          ]
+        }
+      }
+    ],
+    "accumulateNchw": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar"
+      },
+      { "name": "y", "semantic": "accum", "buffer": { "type": "storage" }, "elementType": "i32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "inChannels", "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": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+            { "name": "weightInChannels", "type": "u32", "value": "dim(shapes.w, 1)" },
+            { "name": "inChannelsPerGroup", "type": "u32", "value": "dim(shapes.x, 1) / attrs.group" },
+            { "name": "outChannelsPerGroup", "type": "u32", "value": "dim(shapes.w, 0) / attrs.group" },
+            { "name": "kernelH", "type": "u32", "value": "dim(shapes.w, 2)" },
+            { "name": "kernelW", "type": "u32", "value": "dim(shapes.w, 3)" },
+            { "name": "outH", "type": "u32", "value": "dim(shapes.y, 2)" },
+            { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" },
+            { "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": "padH", "type": "i32", "value": "effectivePadTop" },
+            { "name": "padW", "type": "i32", "value": "effectivePadLeft" },
+            { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
+          ]
+        }
+      }
+    ],
+    "requantizeScalar": [
+      { "name": "accum", "semantic": "accum", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "w_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+      }
+    ],
+    "requantizePerChannel": [
+      { "name": "accum", "semantic": "accum", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "w_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
+            { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+            {
+              "name": "outImageSize",
+              "type": "u32",
+              "value": "numel(shapes.y) / (dim(shapes.y, 0) * dim(shapes.y, 1))"
+            }
+          ]
+        }
+      }
+    ],
+    "requantizeBias": [
+      { "name": "accum", "semantic": "accum", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "w_scale",
+        "arg": "w_scale",
+        "semantic": "w_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "y_scale",
+        "arg": "y_scale",
+        "semantic": "y_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": 1
+      },
+      {
+        "name": "y_zero_point",
+        "arg": "y_zero_point",
+        "semantic": "y_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$yScalar",
+        "length": 1
+      },
+      { "name": "B", "arg": "B", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "i32" },
+      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
+            { "name": "outChannels", "type": "u32", "value": "dim(shapes.w, 0)" },
+            {
+              "name": "outImageSize",
+              "type": "u32",
+              "value": "numel(shapes.y) / (dim(shapes.y, 0) * dim(shapes.y, 1))"
+            }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "dp4a_pointwise_1x1_padded_k_baseline",
+      "priority": 21,
+      "when": ["packedFeature", "not present.B", "not device.features.has(\"shader-f16\")", "rank4TensorOk", "tensorDtypes.x == \"uint8\"", "tensorDtypes.w == \"int8\"", "tensorDtypes.y == \"uint8\"", "xyQuantContract", "flexibleWeightQuantOk", "attrs.group == 1", "dim(shapes.x, 1) % 4 != 0", "dim(shapes.x, 1) >= 4", "dim(shapes.w, 0) >= 8", "dim(shapes.y, 2) * dim(shapes.y, 3) >= 256", "dim(shapes.w, 2) == 1", "dim(shapes.w, 3) == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "dim(shapes.w, 1) == dim(shapes.x, 1)", "outputBatchChannelsOk", "dim(shapes.y, 2) == dim(shapes.x, 2)", "dim(shapes.y, 3) == dim(shapes.x, 3)", "paddedMatricesStorageOk", "dp4a2dDispatchOk"],
+      "constants": {
+        "aZeroPerRow": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "aScalar": "\"i32\"",
+        "bScalar": "\"u32\"",
+        "yScalar": "\"u32\"",
+        "aUnsigned": false,
+        "bUnsigned": true,
+        "yUnsigned": true,
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": true,
+        "bZeroPerColumn": false,
+        "requant": "\"per_row\" if ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) else \"per_tensor\"",
+        "scaleScalar": "\"f32\"",
+        "qMin": 0,
+        "qMax": 255,
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32,
+        "padB": true
+      },
+      "intermediates": [
+        { "id": "w_padded", "dtype": "int32", "shape": "[dim(shapes.w, 0) * ceilDiv(dim(shapes.x, 1), 4) * 4]" },
+        {
+          "id": "x_padded",
+          "dtype": "uint32",
+          "shape": "[dim(shapes.x, 0) * ceilDiv(dim(shapes.x, 1), 4) * 4 * dim(shapes.y, 2) * dim(shapes.y, 3)]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "pad_k",
+          "name": "QLinearConv.PadPointwiseK",
+          "shader": "quant-pad-k-matrices.wgsl.jinja",
+          "bindings": "padPointwiseK",
+          "dispatch": {
+            "threads": "max(dim(shapes.w, 0) * ceilDiv(dim(shapes.x, 1), 4) * 4, dim(shapes.x, 0) * ceilDiv(dim(shapes.x, 1), 4) * 4 * dim(shapes.y, 2) * dim(shapes.y, 3))",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        },
+        {
+          "id": "main",
+          "name": "QLinearConv.Dp4aPointwisePaddedK",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aPaddedMain",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
+            "y": "ceil(dim(shapes.w, 0) / 8)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_im2col_nchw_padded_k_baseline",
+      "priority": 21,
+      "when": ["packedFeature", "not present.B", "not device.features.has(\"shader-f16\")", "rank4TensorOk", "tensorDtypes.x == \"uint8\"", "tensorDtypes.w == \"int8\"", "tensorDtypes.y == \"uint8\"", "xyQuantContract", "flexibleWeightQuantOk", "attrs.group == 1", "strideH >= 1", "strideW >= 1", "dilationH >= 1", "dilationW >= 1", "(dim(shapes.w, 2) > 1 or dim(shapes.w, 3) > 1)", "dim(shapes.w, 1) == dim(shapes.x, 1)", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) % 4 != 0", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) >= 64", "dim(shapes.w, 0) >= 8", "dim(shapes.y, 2) * dim(shapes.y, 3) >= 256", "numel(shapes.y) >= 8192", "outputBatchChannelsOk", "output2dShapeOk", "paddedMatricesStorageOk", "paddedPreparationDispatchOk", "dp4a2dDispatchOk"],
+      "constants": {
+        "aZeroPerRow": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "aScalar": "\"i32\"",
+        "bScalar": "\"u32\"",
+        "yScalar": "\"u32\"",
+        "aUnsigned": false,
+        "bUnsigned": true,
+        "yUnsigned": true,
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": true,
+        "bZeroPerColumn": false,
+        "requant": "\"per_row\" if ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) else \"per_tensor\"",
+        "scaleScalar": "\"f32\"",
+        "qMin": 0,
+        "qMax": 255,
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32,
+        "padB": false
+      },
+      "intermediates": [
+        {
+          "id": "w_padded",
+          "dtype": "int32",
+          "shape": "[dim(shapes.w, 0) * ceilDiv(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3), 4) * 4]"
+        },
+        {
+          "id": "cols_padded",
+          "dtype": "uint32",
+          "shape": "[dim(shapes.x, 0) * ceilDiv(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3), 4) * 4 * dim(shapes.y, 2) * dim(shapes.y, 3)]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "im2col",
+          "name": "QLinearConv.Im2ColPaddedK",
+          "source": {
+            "shader": "conv-int-im2col-spatial.wgsl.jinja",
+            "inputs": {
+              "spatialRank": 2,
+              "padK": true,
+              "kernelH": "dim(shapes.w, 2)",
+              "kernelW": "dim(shapes.w, 3)",
+              "strideH": "strideH",
+              "strideW": "strideW",
+              "dilationH": "dilationH",
+              "dilationW": "dilationW",
+              "padTop": "effectivePadTop",
+              "padLeft": "effectivePadLeft"
+            }
+          },
+          "bindings": "im2colNchwPadded",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / tunables.WORKGROUP_SIZE)",
+            "y": "ceilDiv(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3), 4) * 4",
+            "z": "dim(shapes.x, 0)"
+          }
+        },
+        {
+          "id": "pad_w",
+          "name": "QLinearConv.PadIm2ColWeightsK",
+          "shader": "quant-pad-k-matrices.wgsl.jinja",
+          "bindings": "padWeightsK",
+          "dispatch": {
+            "threads": "dim(shapes.w, 0) * ceilDiv(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3), 4) * 4",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        },
+        {
+          "id": "main",
+          "name": "QLinearConv.Im2ColPaddedKDp4a",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aPaddedIm2colMain",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
+            "y": "ceil(dim(shapes.w, 0) / 8)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_direct_3x3_nchw_u8s8",
+      "priority": 30,
+      "demoteWhen": ["narrowSubgroupRange"],
+      "when": ["packedFeature", "not present.B", "rank4TensorOk", "tensorDtypes.x == \"uint8\"", "tensorDtypes.w == \"int8\"", "tensorDtypes.y == \"uint8\"", "xyQuantContract", "flexibleWeightQuantOk", "attrs.group == 1", "dim(shapes.w, 1) == dim(shapes.x, 1)", "dim(shapes.w, 2) == 3", "dim(shapes.w, 3) == 3", "strideH == 1", "strideW == 1", "dilationH == 1", "dilationW == 1", "padTop == 1", "padLeft == 1", "dim(shapes.y, 2) == dim(shapes.x, 2)", "dim(shapes.y, 3) == dim(shapes.x, 3)", "outputBatchChannelsOk", "output2dShapeOk", "dim(shapes.x, 1) % tunables.DP4A_DIRECT_IC_TILE == 0", "dim(shapes.w, 0) >= tunables.DP4A_DIRECT_OC_TILE", "direct3x3DeviceOk"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "QLinearConv.Direct3x3Dp4a",
+          "source": {
+            "shader": "qlinear-conv-dp4a-direct-3x3.wgsl.jinja",
+            "inputs": {
+              "batch": "dim(shapes.x, 0)",
+              "inChannels": "dim(shapes.x, 1)",
+              "inHeight": "dim(shapes.x, 2)",
+              "inWidth": "dim(shapes.x, 3)",
+              "outChannels": "dim(shapes.w, 0)",
+              "outHeight": "dim(shapes.y, 2)",
+              "outWidth": "dim(shapes.y, 3)",
+              "spatialTile": "tunables.DP4A_DIRECT_SPATIAL_TILE",
+              "ocTile": "tunables.DP4A_DIRECT_OC_TILE",
+              "icTile": "tunables.DP4A_DIRECT_IC_TILE",
+              "wScalePerChannel": "ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0)",
+              "wZeroPerChannel": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)"
+            }
+          },
+          "bindings": "dp4aDirect3x3",
+          "dispatch": {
+            "x": "ceilDiv(dim(shapes.y, 3), tunables.DP4A_DIRECT_SPATIAL_TILE)",
+            "y": "ceilDiv(dim(shapes.y, 2), tunables.DP4A_DIRECT_SPATIAL_TILE) * ceilDiv(dim(shapes.w, 0), tunables.DP4A_DIRECT_OC_TILE)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1",
+      "priority": 20,
+      "when": ["packedFeature", "not present.B", "rank4TensorOk", "flexibleWeightQuantOk", "xyQuantContract", "attrs.group == 1", "(device.features.has(\"shader-f16\") or dim(shapes.x, 1) % 4 == 0)", "dim(shapes.w, 2) == 1", "dim(shapes.w, 3) == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "dim(shapes.w, 1) == dim(shapes.x, 1)", "outputBatchChannelsOk", "dim(shapes.y, 2) == dim(shapes.x, 2)", "dim(shapes.y, 3) == dim(shapes.x, 3)", "dp4a2dDispatchOk"],
+      "constants": {
+        "aScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "bScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "aVec4": "(\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\") if dim(shapes.x, 1) % 4 == 0 else (\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\")",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": "dim(shapes.x, 1) % 4 == 0",
+        "aZeroPerRow": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "bZeroPerColumn": false,
+        "requant": "\"per_row\" if ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) else \"per_tensor\"",
+        "scaleScalar": "\"f32\"",
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127",
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "QLinearConv.Dp4aPointwiseFused",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aPointwiseMain",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
+            "y": "ceil(dim(shapes.w, 0) / 8)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_feature_independent_x4",
+      "priority": 10,
+      "when": ["not present.B", "rank4TensorOk", "flexibleWeightQuantOk", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk", "output2dShapeOk", "dim(shapes.y, 3) >= 4", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) >= 4", "numel(shapes.y) >= 32"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "wScalePerChannel": "ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0)",
+        "wZeroPerChannel": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "ocTile": "4 if dim(shapes.w, 0) >= attrs.group * 4 and dim(shapes.w, 0) % (attrs.group * 4) == 0 else 1",
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "QLinearConv.FeatureIndependentX4",
+          "shader": "qlinear-conv-nchw-x4.wgsl.jinja",
+          "bindings": "fusedNchwX4",
+          "dispatch": {
+            "threads": "dim(shapes.y, 0) * ceil(dim(shapes.y, 1) / (4 if dim(shapes.w, 0) >= attrs.group * 4 and dim(shapes.w, 0) % (attrs.group * 4) == 0 else 1)) * dim(shapes.y, 2) * ceil(dim(shapes.y, 3) / 4)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d_per_channel_quantization",
+      "when": ["not present.B", "rank5TensorOk and dim(shapes.y, 2) == expectedOutputDepth and dim(shapes.y, 3) == expectedOutputHeight and dim(shapes.y, 4) == expectedOutputWidth", "not scalarWeightQuantOk and ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) and ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": false,
+        "wScalePerChannel": true,
+        "wZeroPerChannel": true,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.AccumulatePerChannel3d",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 3 } },
+          "bindings": "accumulateNcdhw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.Requantize",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizePerChannel",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncw1d_scalar_quantization",
+      "when": ["not present.B", "rank3TensorOk and dim(shapes.y, 2) == expectedOutputWidth", "scalarWeightQuantOk", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": false,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.Accumulate1d",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 2 } },
+          "bindings": "accumulateNcw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.Requantize",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizeScalar",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d_scalar_quantization",
+      "when": ["not present.B", "rank5TensorOk and dim(shapes.y, 2) == expectedOutputDepth and dim(shapes.y, 3) == expectedOutputHeight and dim(shapes.y, 4) == expectedOutputWidth", "scalarWeightQuantOk", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": false,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.Accumulate3d",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 3 } },
+          "bindings": "accumulateNcdhw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.Requantize",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizeScalar",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_scalar_quantization",
+      "when": ["not present.B", "rank4TensorOk and output2dShapeOk", "scalarWeightQuantOk", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": false,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.Accumulate",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 2 } },
+          "bindings": "accumulateNchw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.Requantize",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizeScalar",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_im2col_ncdhw3d_scalar_quantization",
+      "priority": 20,
+      "when": ["packedFeature", "not present.B", "rank5TensorOk", "attrs.group == 1", "dim(shapes.w, 1) == dim(shapes.x, 1)", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4) % 4 == 0", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4) >= 32", "dim(shapes.w, 0) >= 8", "dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4) >= 256", "flexibleWeightQuantOk", "xyQuantContract", "outputBatchChannelsOk", "dim(shapes.y, 2) == expectedOutputDepth", "dim(shapes.y, 3) == expectedOutputHeight", "dim(shapes.y, 4) == expectedOutputWidth", "im2colStorageOk", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4) <= device.limits.maxComputeWorkgroupsPerDimension", "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4), tunables.WORKGROUP_SIZE) <= device.limits.maxComputeWorkgroupsPerDimension", "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4), 64) <= device.limits.maxComputeWorkgroupsPerDimension", "ceilDiv(dim(shapes.w, 0), 8) <= device.limits.maxComputeWorkgroupsPerDimension", "dim(shapes.x, 0) <= device.limits.maxComputeWorkgroupsPerDimension"],
+      "constants": {
+        "aScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "bScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": true,
+        "aZeroPerRow": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "bZeroPerColumn": false,
+        "requant": "\"per_row\" if ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) else \"per_tensor\"",
+        "scaleScalar": "\"f32\"",
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127",
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32,
+        "hasBias": false
+      },
+      "intermediates": [
+        {
+          "id": "cols3d",
+          "dtype": "int32",
+          "shape": "[dim(shapes.x, 0) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4) * dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4)]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "im2col",
+          "name": "QLinearConv.Im2Col3d",
+          "source": {
+            "shader": "conv-int-im2col-spatial.wgsl.jinja",
+            "inputs": {
+              "spatialRank": 3,
+              "kernelD": "dim(shapes.w, 2)",
+              "kernelH": "dim(shapes.w, 3)",
+              "kernelW": "dim(shapes.w, 4)",
+              "strideD": "strideD",
+              "strideH": "strideH",
+              "strideW": "strideW",
+              "dilationD": "dilationD",
+              "dilationH": "dilationH",
+              "dilationW": "dilationW",
+              "padFront": "effectivePadFront",
+              "padTop": "effectivePadTop",
+              "padLeft": "effectivePadLeft"
+            }
+          },
+          "bindings": "im2colNcdhw",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4) / tunables.WORKGROUP_SIZE)",
+            "y": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4)",
+            "z": "dim(shapes.x, 0)"
+          }
+        },
+        {
+          "id": "main",
+          "name": "QLinearConv.Im2Col3dDp4a",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aMain3d",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4) / 64)",
+            "y": "ceil(dim(shapes.w, 0) / 8)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_im2col_nchw_bias_quantization",
+      "priority": 25,
+      "when": ["packedFeature", "present.B", "ranks.B == 1", "dim(shapes.B, 0) == dim(shapes.w, 0)", "rank4TensorOk", "attrs.group == 1", "dim(shapes.w, 1) == dim(shapes.x, 1)", "outputBatchChannelsOk", "flexibleWeightQuantOk", "xyQuantContract"],
+      "constants": {
+        "aScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "bScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "aVec4": "(\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\") if dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) % 4 == 0 else (\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\")",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) % 4 == 0",
+        "aZeroPerRow": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "bZeroPerColumn": false,
+        "requant": "\"per_row\" if ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) else \"per_tensor\"",
+        "scaleScalar": "\"f32\"",
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127",
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32,
+        "hasBias": true
+      },
+      "intermediates": [
+        {
+          "id": "cols",
+          "dtype": "int32",
+          "shape": "[dim(shapes.x, 0) * (dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)) * (dim(shapes.y, 2) * dim(shapes.y, 3))]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "im2col",
+          "name": "QLinearConv.Im2Col",
+          "source": {
+            "shader": "conv-int-im2col-spatial.wgsl.jinja",
+            "inputs": {
+              "spatialRank": 2,
+              "kernelH": "dim(shapes.w, 2)",
+              "kernelW": "dim(shapes.w, 3)",
+              "strideH": "strideH",
+              "strideW": "strideW",
+              "dilationH": "dilationH",
+              "dilationW": "dilationW",
+              "padTop": "effectivePadTop",
+              "padLeft": "effectivePadLeft"
+            }
+          },
+          "bindings": "im2colNchw",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / tunables.WORKGROUP_SIZE)",
+            "y": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
+            "z": "dim(shapes.x, 0)"
+          }
+        },
+        {
+          "id": "main",
+          "name": "QLinearConv.Im2ColDp4aBias",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aMainBias",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
+            "y": "ceil(dim(shapes.w, 0) / 8)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_im2col_nchw_scalar_quantization",
+      "priority": 20,
+      "when": ["packedFeature", "not present.B", "flexibleWeightQuantOk", "xyQuantContract", "rank4TensorOk", "attrs.group == 1", "(device.features.has(\"shader-f16\") or dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) % 4 == 0)", "(dim(shapes.w, 2) > 1 or dim(shapes.w, 3) > 1)", "dim(shapes.w, 1) == dim(shapes.x, 1)", "outputBatchChannelsOk", "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) >= 64", "numel(shapes.y) >= 8192", "dp4a2dDispatchOk"],
+      "constants": {
+        "aScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "bScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "aVec4": "(\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\") if dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) % 4 == 0 else (\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\")",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) % 4 == 0",
+        "aZeroPerRow": "ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)",
+        "bZeroPerColumn": false,
+        "requant": "\"per_row\" if ranks.w_scale == 1 and dim(shapes.w_scale, 0) == dim(shapes.w, 0) else \"per_tensor\"",
+        "scaleScalar": "\"f32\"",
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127",
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32
+      },
+      "intermediates": [
+        {
+          "id": "cols",
+          "dtype": "int32",
+          "shape": "[dim(shapes.x, 0) * (dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)) * (dim(shapes.y, 2) * dim(shapes.y, 3))]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "im2col",
+          "name": "QLinearConv.Im2Col",
+          "source": {
+            "shader": "conv-int-im2col-spatial.wgsl.jinja",
+            "inputs": {
+              "spatialRank": 2,
+              "kernelH": "dim(shapes.w, 2)",
+              "kernelW": "dim(shapes.w, 3)",
+              "strideH": "strideH",
+              "strideW": "strideW",
+              "dilationH": "dilationH",
+              "dilationW": "dilationW",
+              "padTop": "effectivePadTop",
+              "padLeft": "effectivePadLeft"
+            }
+          },
+          "bindings": "im2colNchw",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / tunables.WORKGROUP_SIZE)",
+            "y": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)",
+            "z": "dim(shapes.x, 0)"
+          }
+        },
+        {
+          "id": "main",
+          "name": "QLinearConv.Im2ColDp4a",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aMain",
+          "dispatch": {
+            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
+            "y": "ceil(dim(shapes.w, 0) / 8)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_per_channel_quantization",
+      "priority": 5,
+      "when": ["not present.B", "rank4TensorOk", "ranks.w_scale == 1", "dim(shapes.w_scale, 0) == dim(shapes.w, 0)", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk", "output2dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": false,
+        "wScalePerChannel": true,
+        "wZeroPerChannel": true,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.AccumulatePerChannel",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 2 } },
+          "bindings": "accumulateNchw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.RequantizePerChannel",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizePerChannel",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_per_channel_quantization_bias",
+      "when": ["present.B", "ranks.B == 1", "dim(shapes.B, 0) == dim(shapes.w, 0)", "rank4TensorOk", "ranks.w_scale == 1", "dim(shapes.w_scale, 0) == dim(shapes.w, 0)", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk", "output2dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": true,
+        "wScalePerChannel": true,
+        "wZeroPerChannel": true,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.AccumulatePerChannel",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 2 } },
+          "bindings": "accumulateNchw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.RequantizePerChannelBias",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizeBias",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_scalar_quantization_bias",
+      "when": ["present.B", "ranks.B == 1", "dim(shapes.B, 0) == dim(shapes.w, 0)", "rank4TensorOk", "scalarWeightQuantOk", "xyQuantContract", "groupChannelsOk", "outputBatchChannelsOk", "output2dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "yScalar": "\"u32\" if tensorDtypes.y == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "yUnsigned": "tensorDtypes.y == \"uint8\"",
+        "hasBias": true,
+        "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+        "qMax": "255 if tensorDtypes.y == \"uint8\" else 127"
+      },
+      "intermediates": [{ "id": "accum", "dtype": "int32", "shape": "shapes.y" }],
+      "passes": [
+        {
+          "id": "accumulate",
+          "name": "QLinearConv.Accumulate",
+          "source": { "shader": "conv-int-accumulate-spatial.wgsl.jinja", "inputs": { "spatialRank": 2 } },
+          "bindings": "accumulateNchw",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "requantize",
+          "name": "QLinearConv.RequantizeBias",
+          "shader": "qlinear-conv-requantize.wgsl.jinja",
+          "bindings": "requantizeBias",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,24 @@

+{
+  "name": "ai.onnx.QLinearConv",
+  "id": "_ai_onnx_qlinearconv_webgpu_97f4388",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "7QcaEvhkU7vodLHPSkDnPEs+l3fyHE5M7iPreEs8EBs=",
+      "conv-int-accumulate-spatial.wgsl.jinja": "+RLzIvFAtuFaDFjixU0amSlvCZ4sOlphzPpkf+7Klis=",
+      "conv-int-im2col-spatial.wgsl.jinja": "fdcyh0E6YDzmlp/uWUnejeervSfidZul2GNQBrZUG+4=",
+      "manifest.json": "GD1vHqnoBAsn2BJePTNWNRgtBJUcvPp6xTHgX0xEsLM=",
+      "qlinear-conv-dp4a-direct-3x3.wgsl.jinja": "irlapLZkfGOdQZgLgddD34Djo8z+RSVCqUj0OrBH3Yw=",
+      "qlinear-conv-nchw-x4.wgsl.jinja": "UZfNJ4L5YVesOPAv/9ugYz8tJOhujdafH/BtcnPjrKQ=",
+      "qlinear-conv-requantize.wgsl.jinja": "w4e3cepxUD19iM8XgzpXUN7dyHCzzrxfrWjzrOaPps4=",
+      "quant-dp4a-matmul.wgsl.jinja": "yzJ3cN0cvihhQ8P3joswjg12M0p4ZrLQgfxXq9epORo=",
+      "quant-pad-k-matrices.wgsl.jinja": "UNPxeSYR99nYYGs6qmsHnO3/uHcDkdQB4jBZwZw2zmw=",
+      "test.json": "Q1nMc5F1AYDSVoAvj5w8rpbPrOw+6hZhLvwo5wV9yJ4="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.QLinearConv" }
+}

build/webgpu/qlinear-conv-dp4a-direct-3x3.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,198 @@

+// Direct tiled u8 x s8 3x3 convolution using packed 4-lane integer dots.
+//
+// An 8x8 output workgroup cooperatively stages a 10x10 halo for sixteen input
+// channels. Four channels are packed into each shared u32, so every staged
+// input word is reused by all output-channel accumulators and neighboring
+// output pixels. Weights are packed in the same channel groups, and the
+// convolution operates directly without a materialized im2col tensor.
+{{ env.wgsl.resourceDeclarations }}
+// Centralize the signed packed-int8 dot expression so unrolled callers cannot
+// collide in lowering-generated temporaries.
+fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
+  return dot4I8Packed(a_word, b_word);
+}
+const IN_C: u32 = {{ source.inChannels }}u;
+const IN_H: u32 = {{ source.inHeight }}u;
+const IN_W: u32 = {{ source.inWidth }}u;
+const OUT_C: u32 = {{ source.outChannels }}u;
+const OUT_H: u32 = {{ source.outHeight }}u;
+const OUT_W: u32 = {{ source.outWidth }}u;
+const BATCH: u32 = {{ source.batch }}u;
+const SPATIAL_TILE: u32 = {{ source.spatialTile }}u;
+const OC_TILE: u32 = {{ source.ocTile }}u;
+const IC_TILE: u32 = {{ source.icTile }}u;
+const IC_WORDS: u32 = IC_TILE / 4u;
+const PATCH_EDGE: u32 = SPATIAL_TILE + 2u;
+const PATCH_AREA: u32 = PATCH_EDGE * PATCH_EDGE;
+const X_TILE_WORDS: u32 = IC_WORDS * PATCH_AREA;
+const W_TILE_WORDS: u32 = OC_TILE * IC_WORDS * 9u;
+const WG_SIZE: u32 = SPATIAL_TILE * SPATIAL_TILE;
+const OUT_TILE_ROWS: u32 = (OUT_H + SPATIAL_TILE - 1u) / SPATIAL_TILE;
+const K: i32 = i32(IN_C * 9u);
+var<workgroup> xTile: array<u32, X_TILE_WORDS>;
+var<workgroup> wTile: array<u32, W_TILE_WORDS>;
+var<workgroup> wSums: array<atomic<i32>, OC_TILE>;
+fn pack_x(v: vec4<u32>) -> u32 {
+  return ((v.x ^ 0x80u) & 0xFFu)
+    | (((v.y ^ 0x80u) & 0xFFu) << 8u)
+    | (((v.z ^ 0x80u) & 0xFFu) << 16u)
+    | (((v.w ^ 0x80u) & 0xFFu) << 24u);
+}
+fn pack_w(v: vec4<i32>) -> u32 {
+  return (u32(v.x) & 0xFFu)
+    | ((u32(v.y) & 0xFFu) << 8u)
+    | ((u32(v.z) & 0xFFu) << 16u)
+    | ((u32(v.w) & 0xFFu) << 24u);
+}
+fn signed_byte(v: u32, shift: u32) -> i32 {
+  let b = i32((v >> shift) & 0xFFu);
+  return select(b, b - 256, b >= 128);
+}
+fn sum4_packed(v: u32) -> i32 {
+  return signed_byte(v, 0u) + signed_byte(v, 8u)
+    + signed_byte(v, 16u) + signed_byte(v, 24u);
+}
+fn read_w_zero({% if source.wZeroPerChannel %}oc: u32{% endif %}) -> i32 {
+{% if source.wZeroPerChannel %}
+  return w_zero_point[oc];
+{% else %}
+  return w_zero_point[0];
+{% endif %}
+}
+fn read_w_scale({% if source.wScalePerChannel %}oc: u32{% endif %}) -> f32 {
+{% if source.wScalePerChannel %}
+  return w_scale[oc];
+{% else %}
+  return w_scale[0];
+{% endif %}
+}
+fn requantize(acc: i32{% if source.wScalePerChannel %}, oc: u32{% endif %}) -> u32 {
+  let scaled = f32(acc) * x_scale[0] * read_w_scale({% if source.wScalePerChannel %}oc{% endif %}) / y_scale[0];
+  let q = clamp(
+    i32(clamp(round(scaled), -2.0e9, 2.0e9)) + i32(y_zero_point[0]),
+    0,
+    255,
+  );
+  return u32(q);
+}
+@compute @workgroup_size({{ source.spatialTile }}, {{ source.spatialTile }}, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let lane = lid.y * SPATIAL_TILE + lid.x;
+  let tileY = wg.y % OUT_TILE_ROWS;
+  let ocBlock = wg.y / OUT_TILE_ROWS;
+  let ocBase = ocBlock * OC_TILE;
+  let oh = tileY * SPATIAL_TILE + lid.y;
+  let ow = wg.x * SPATIAL_TILE + lid.x;
+  let outputValid = wg.z < BATCH && oh < OUT_H && ow < OUT_W;
+  let xzp = x_zero_point[0];
+  let cb = 128 - i32(xzp);
+  if (lane < OC_TILE) {
+    atomicStore(&wSums[lane], 0);
+  }
+  workgroupBarrier();
+{% for oc in range(source.ocTile) %}
+  var acc{{ oc }} = 0i;
+{% endfor %}
+  var sumX = 0i;
+  for (var icBase = 0u; icBase < IN_C; icBase += IC_TILE) {
+    // Pack four NCHW channels for each halo position. Padding uses the input
+    // zero point, so the zero-point correction below cancels it exactly.
+    for (var item = lane; item < X_TILE_WORDS; item += WG_SIZE) {
+      let channelWord = item / PATCH_AREA;
+      let patchIndex = item % PATCH_AREA;
+      let py = patchIndex / PATCH_EDGE;
+      let px = patchIndex % PATCH_EDGE;
+      let iy = i32(tileY * SPATIAL_TILE + py) - 1;
+      let ix = i32(wg.x * SPATIAL_TILE + px) - 1;
+      var xv = vec4<u32>(xzp);
+      if (iy >= 0 && iy < i32(IN_H) && ix >= 0 && ix < i32(IN_W)) {
+        let spatial = u32(iy) * IN_W + u32(ix);
+        let ic0 = icBase + channelWord * 4u;
+        let batchBase = wg.z * IN_C * IN_H * IN_W;
+        xv = vec4<u32>(
+          x[batchBase + (ic0 + 0u) * IN_H * IN_W + spatial],
+          x[batchBase + (ic0 + 1u) * IN_H * IN_W + spatial],
+          x[batchBase + (ic0 + 2u) * IN_H * IN_W + spatial],
+          x[batchBase + (ic0 + 3u) * IN_H * IN_W + spatial],
+        );
+      }
+      xTile[item] = pack_x(xv);
+    }
+    // Pack the same four-channel groups for every output channel and tap.
+    for (var item = lane; item < W_TILE_WORDS; item += WG_SIZE) {
+      let tap = item % 9u;
+      let channelWord = (item / 9u) % IC_WORDS;
+      let localOc = item / (9u * IC_WORDS);
+      let oc = ocBase + localOc;
+      var wv = vec4<i32>(0);
+      if (oc < OUT_C) {
+        let kh = tap / 3u;
+        let kw = tap % 3u;
+        let ic0 = icBase + channelWord * 4u;
+        wv = vec4<i32>(
+          w[((oc * IN_C + ic0 + 0u) * 3u + kh) * 3u + kw],
+          w[((oc * IN_C + ic0 + 1u) * 3u + kh) * 3u + kw],
+          w[((oc * IN_C + ic0 + 2u) * 3u + kh) * 3u + kw],
+          w[((oc * IN_C + ic0 + 3u) * 3u + kh) * 3u + kw],
+        );
+        atomicAdd(&wSums[localOc], wv.x + wv.y + wv.z + wv.w);
+      }
+      wTile[item] = pack_w(wv);
+    }
+    workgroupBarrier();
+    if (outputValid) {
+      for (var channelWord = 0u; channelWord < IC_WORDS; channelWord++) {
+        for (var kh = 0u; kh < 3u; kh++) {
+          for (var kw = 0u; kw < 3u; kw++) {
+            let tap = kh * 3u + kw;
+            let xp = xTile[channelWord * PATCH_AREA + (lid.y + kh) * PATCH_EDGE + lid.x + kw];
+            sumX += sum4_packed(xp);
+{% for oc in range(source.ocTile) %}
+            acc{{ oc }} += dot4_packed(wTile[({{ oc }}u * IC_WORDS + channelWord) * 9u + tap], xp);
+{% endfor %}
+          }
+        }
+      }
+    }
+    workgroupBarrier();
+  }
+  if (!outputValid) {
+    return;
+  }
+{% for oc in range(source.ocTile) %}
+  {
+    let outOc = ocBase + {{ oc }}u;
+    if (outOc < OUT_C) {
+      let ca = -read_w_zero({% if source.wZeroPerChannel %}outOc{% endif %});
+      let corrected = acc{{ oc }}
+        + cb * atomicLoad(&wSums[{{ oc }}u])
+        + ca * sumX
+        + K * ca * cb;
+      let outIndex = ((wg.z * OUT_C + outOc) * OUT_H + oh) * OUT_W + ow;
+      y[outIndex] = requantize(corrected{% if source.wScalePerChannel %}, outOc{% endif %});
+    }
+  }
+{% endfor %}
+}

build/webgpu/qlinear-conv-nchw-x4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,162 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+fn x_zero() -> i32 {
+{% if xUnsigned %}
+  return i32(x_zero_point[0]);
+{% else %}
+  return x_zero_point[0];
+{% endif %}
+}
+fn w_zero({% if wZeroPerChannel %}oc: u32{% endif %}) -> i32 {
+{% if wZeroPerChannel %}
+  let i = oc;
+{% else %}
+  let i = 0u;
+{% endif %}
+{% if wUnsigned %}
+  return i32(w_zero_point[i]);
+{% else %}
+  return w_zero_point[i];
+{% endif %}
+}
+fn centered_x(i: u32, zp: i32) -> i32 {
+{% if xUnsigned %}
+  return i32(x[i]) - zp;
+{% else %}
+  return x[i] - zp;
+{% endif %}
+}
+fn centered_w(i: u32, zp: i32) -> i32 {
+{% if wUnsigned %}
+  return i32(w[i]) - zp;
+{% else %}
+  return w[i] - zp;
+{% endif %}
+}
+fn y_zero() -> i32 {
+{% if yUnsigned %}
+  return i32(y_zero_point[0]);
+{% else %}
+  return y_zero_point[0];
+{% endif %}
+}
+fn requant(acc: i32{% if wScalePerChannel %}, oc: u32{% endif %}) -> i32 {
+{% if wScalePerChannel %}
+  let ws = w_scale[oc];
+{% else %}
+  let ws = w_scale[0];
+{% endif %}
+  let scaled = f32(acc) * x_scale[0] * ws / y_scale[0];
+  return clamp(i32(clamp(round(scaled), -2.0e9, 2.0e9)) + y_zero(), {{ qMin }}, {{ qMax }});
+}
+fn write_y(i: u32, value: i32) {
+{% if yUnsigned %}
+  y[i] = u32(value);
+{% else %}
+  y[i] = value;
+{% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("flat", "", note="") }}
+  let groups_per_row = (params.outW + 3u) / 4u;
+  let oc_blocks = (params.outChannels + {{ ocTile - 1 }}u) / {{ ocTile }}u;
+  let group_count = params.batch * oc_blocks * params.outH * groups_per_row;
+  if (flat >= group_count) { return; }
+  let ow0 = (flat % groups_per_row) * 4u;
+  let row = flat / groups_per_row;
+  let oh = row % params.outH;
+  let oc0 = ((row / params.outH) % oc_blocks) * {{ ocTile }}u;
+  let n = row / (params.outH * oc_blocks);
+  let group = oc0 / params.outChannelsPerGroup;
+  let in_start = group * params.inChannelsPerGroup;
+  let xzp = x_zero();
+{% for ot in range(ocTile) %}
+  var acc{{ ot }} = vec4<i32>(0);
+{% endfor %}
+  for (var icg = 0u; icg < params.inChannelsPerGroup; icg = icg + 1u) {
+    let ic = in_start + icg;
+    for (var kh = 0u; kh < params.kernelH; kh = kh + 1u) {
+      let ih = i32(oh * params.strideH + kh * params.dilationH) - params.padH;
+      for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+        var x0 = 0i;
+        var x1 = 0i;
+        var x2 = 0i;
+        var x3 = 0i;
+        if (ih >= 0 && ih < i32(params.inH)) {
+          let iw0 = i32(ow0 * params.strideW + kw * params.dilationW) - params.padW;
+          let iw1 = i32((ow0 + 1u) * params.strideW + kw * params.dilationW) - params.padW;
+          let iw2 = i32((ow0 + 2u) * params.strideW + kw * params.dilationW) - params.padW;
+          let iw3 = i32((ow0 + 3u) * params.strideW + kw * params.dilationW) - params.padW;
+          let base = (n * params.inChannels + ic) * params.inH * params.inW + u32(ih) * params.inW;
+          if (ow0 < params.outW && iw0 >= 0 && iw0 < i32(params.inW)) { x0 = centered_x(base + u32(iw0), xzp); }
+          if (ow0 + 1u < params.outW && iw1 >= 0 && iw1 < i32(params.inW)) { x1 = centered_x(base + u32(iw1), xzp); }
+          if (ow0 + 2u < params.outW && iw2 >= 0 && iw2 < i32(params.inW)) { x2 = centered_x(base + u32(iw2), xzp); }
+          if (ow0 + 3u < params.outW && iw3 >= 0 && iw3 < i32(params.inW)) { x3 = centered_x(base + u32(iw3), xzp); }
+        }
+        let xv = vec4<i32>(x0, x1, x2, x3);
+{% for ot in range(ocTile) %}
+        let oc{{ ot }} = oc0 + {{ ot }}u;
+        if (oc{{ ot }} < params.outChannels) {
+          let wi{{ ot }} = (((oc{{ ot }} * params.weightInChannels + icg) * params.kernelH + kh) * params.kernelW + kw);
+          let wc{{ ot }} = centered_w(wi{{ ot }}, w_zero({% if wZeroPerChannel %}oc{{ ot }}{% endif %}));
+          acc{{ ot }} = acc{{ ot }} + xv * vec4<i32>(wc{{ ot }});
+        }
+{% endfor %}
+      }
+    }
+  }
+{% for ot in range(ocTile) %}
+  let out_oc{{ ot }} = oc0 + {{ ot }}u;
+  if (out_oc{{ ot }} < params.outChannels) {
+    let out_base{{ ot }} = ((n * params.outChannels + out_oc{{ ot }}) * params.outH + oh) * params.outW + ow0;
+    write_y(out_base{{ ot }}, requant(acc{{ ot }}.x{% if wScalePerChannel %}, out_oc{{ ot }}{% endif %}));
+    if (ow0 + 1u < params.outW) { write_y(out_base{{ ot }} + 1u, requant(acc{{ ot }}.y{% if wScalePerChannel %}, out_oc{{ ot }}{% endif %})); }
+    if (ow0 + 2u < params.outW) { write_y(out_base{{ ot }} + 2u, requant(acc{{ ot }}.z{% if wScalePerChannel %}, out_oc{{ ot }}{% endif %})); }
+    if (ow0 + 3u < params.outW) { write_y(out_base{{ ot }} + 3u, requant(acc{{ ot }}.w{% if wScalePerChannel %}, out_oc{{ ot }}{% endif %})); }
+  }
+{% endfor %}
+}

build/webgpu/qlinear-conv-requantize.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,76 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+fn read_y_zero() -> i32 {
+  {% if yUnsigned %}
+  return i32(y_zero_point[0]);
+  {% else %}
+  return y_zero_point[0];
+  {% endif %}
+}
+fn write_y(index: u32, value: i32) {
+  {% if yUnsigned %}
+  y[index] = u32(value);
+  {% else %}
+  y[index] = value;
+  {% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d(guardInline=true) }}
+{% if hasBias or wScalePerChannel %}
+  let oc = (i / params.outImageSize) % params.outChannels;
+{% endif %}
+  {% if hasBias %}
+  // Optional int32 bias is per output channel, already in (x_scale * w_scale)
+  // units, so it is added to the int32 accumulator before requantization.
+  let raw = accum[i] + B[oc];
+  {% else %}
+  let raw = accum[i];
+  {% endif %}
+{% if wScalePerChannel %}
+  let w_scale_value = w_scale[oc];
+{% else %}
+  let w_scale_value = w_scale[0];
+{% endif %}
+  let scaled = f32(raw) * x_scale[0] * w_scale_value / y_scale[0];
+  // Clamp the rounded float into an i32-representable range before the cast so a
+  // huge-finite / Inf requantized value saturates (the final clamp pins it to
+  // [qMin,qMax]) instead of invoking undefined behaviour in i32(huge_float).
+  let q = clamp(i32(clamp(round(scaled), -2.0e9, 2.0e9)) + read_y_zero(), {{ qMin }}, {{ qMax }});
+  write_y(i, q);
+}

build/webgpu/quant-dp4a-matmul.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,344 @@

+{% if requant != "none" and scaleScalar == "f16" %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const TILE_M: u32 = {{ tileM }}u;
+const TILE_N: u32 = {{ tileN }}u;
+const KT_WORDS: u32 = {{ ktWords }}u;
+const A_TILE_WORDS: u32 = TILE_M * KT_WORDS;
+{% set GEMV_UNROLL = gemvKUnroll if gemvKUnroll is defined else 4 %}
+{% if tileM == 1 %}
+const GEMV_K_UNROLL: u32 = {{ GEMV_UNROLL }}u;
+{% endif %}
+{% set needsASum = hasBZero or bUnsigned %}
+{% set needsBSum = hasAZero or aUnsigned %}
+var<workgroup> a_tile: array<u32, A_TILE_WORDS>;
+{% if needsASum %}
+var<workgroup> a_row_sum: array<atomic<i32>, TILE_M>;
+{% endif %}
+// Pack 4 widened A values into one u32 (byte 0 = lowest k).
+fn pack_a(v: vec4<{{ aScalar }}>) -> u32 {
+{% if aUnsigned %}
+  // u8 -> biased i8: (u - 128) two's-complement low byte == u ^ 0x80.
+  return ((v.x ^ 0x80u) & 0xFFu) | (((v.y ^ 0x80u) & 0xFFu) << 8u) | (((v.z ^ 0x80u) & 0xFFu) << 16u) | (((v.w ^ 0x80u) & 0xFFu) << 24u);
+{% else %}
+  return (u32(v.x) & 0xFFu) | ((u32(v.y) & 0xFFu) << 8u) | ((u32(v.z) & 0xFFu) << 16u) | ((u32(v.w) & 0xFFu) << 24u);
+{% endif %}
+}
+{% if needsASum %}
+// Sum of the 4 bias-shifted A values of one packed word.
+fn sum4_a(v: vec4<{{ aScalar }}>) -> i32 {
+{% if aUnsigned %}
+  return i32(v.x + v.y + v.z + v.w) - 512;
+{% else %}
+  return v.x + v.y + v.z + v.w;
+{% endif %}
+}
+{% endif %}
+fn pack_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> u32 {
+{% if bUnsigned %}
+  return ((b0 ^ 0x80u) & 0xFFu) | (((b1 ^ 0x80u) & 0xFFu) << 8u) | (((b2 ^ 0x80u) & 0xFFu) << 16u) | (((b3 ^ 0x80u) & 0xFFu) << 24u);
+{% else %}
+  return (u32(b0) & 0xFFu) | ((u32(b1) & 0xFFu) << 8u) | ((u32(b2) & 0xFFu) << 16u) | ((u32(b3) & 0xFFu) << 24u);
+{% endif %}
+}
+// Centralize the signed packed-int8 dot expression so unrolled callers cannot
+// collide in lowering-generated temporaries.
+fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
+  return dot4I8Packed(a_word, b_word);
+}
+{% if needsBSum %}
+fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> i32 {
+{% if bUnsigned %}
+  return i32(b0 + b1 + b2 + b3) - 512;
+{% else %}
+  return b0 + b1 + b2 + b3;
+{% endif %}
+}
+{% endif %}
+{% if hasAZero %}
+fn read_a_zero({% if aZeroPerRow %}row: u32{% endif %}) -> i32 {
+{% if aZeroPerRow %}
+{% if aUnsigned %}
+  return i32(a_zero_point[row]);
+{% else %}
+  return a_zero_point[row];
+{% endif %}
+{% else %}
+{% if aUnsigned %}
+  return i32(a_zero_point[0]);
+{% else %}
+  return a_zero_point[0];
+{% endif %}
+{% endif %}
+}
+{% endif %}
+{% if hasBZero %}
+fn read_b_zero({% if bZeroPerColumn %}col: u32{% endif %}) -> i32 {
+{% if bZeroPerColumn %}
+{% if bUnsigned %}
+  return i32(b_zero_point[col]);
+{% else %}
+  return b_zero_point[col];
+{% endif %}
+{% else %}
+{% if bUnsigned %}
+  return i32(b_zero_point[0]);
+{% else %}
+  return b_zero_point[0];
+{% endif %}
+{% endif %}
+}
+{% endif %}
+{% if requant != "none" %}
+fn read_y_zero() -> i32 {
+{% if yUnsigned %}
+  return i32(y_zero_point[0]);
+{% else %}
+  return y_zero_point[0];
+{% endif %}
+}
+{% endif %}
+@compute @workgroup_size(TILE_N, 1, 1)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
+  let batch = wg.z;
+  let a_base4 = batch * params.aBatchStride4;
+  let b_base = batch * params.bBatchStride;
+  let y_base = batch * params.yBatchStride;
+  let m_block = wg.y * TILE_M;
+  let n_block = wg.x * TILE_N;
+  let tid = lid.x;
+  let col = n_block + tid;
+  let n_valid = col < params.N;
+{% if aPacked == false %}
+  // Tail variants keep the hot loop on complete words. Packing the one partial
+  // word afterward prevents its bounds checks from infecting every full-word load.
+  let k_words_total = params.K / 4u;
+  let tail_lanes = params.K - k_words_total * 4u;
+{% else %}
+  let k_words_total = params.K / 4u;
+{% endif %}
+{% if needsASum %}
+  if (tid < TILE_M) {
+    atomicStore(&a_row_sum[tid], 0);
+  }
+  workgroupBarrier();
+{% endif %}
+{% if tileM == 1 %}
+  var r0: i32 = 0;
+{% for u in range(1, GEMV_UNROLL) %}
+  var r0_{{ u }}: i32 = 0;
+{% endfor %}
+{% else %}
+{% for m in range(tileM) %}
+  var r{{ m }}: i32 = 0;
+{% endfor %}
+{% endif %}
+{% if needsBSum %}
+  var sum_b: i32 = 0;
+{% if tileM == 1 %}
+{% for u in range(1, GEMV_UNROLL) %}
+  var sum_b_{{ u }}: i32 = 0;
+{% endfor %}
+{% endif %}
+{% endif %}
+  for (var kw0 = 0u; kw0 < k_words_total; kw0 = kw0 + KT_WORDS) {
+    let words = min(KT_WORDS, k_words_total - kw0);
+    // Cooperative vec4 load + pack of the A tile (TILE_M rows x `words` u32).
+    for (var i = tid; i < TILE_M * words; i = i + TILE_N) {
+      let mi = i / words;
+      let ki = i % words;
+      let src_m = m_block + mi;
+      var packed = 0u;
+      if (src_m < params.M) {
+{% if aPacked == false %}
+        let k4 = (kw0 + ki) * 4u;
+        let a_row = a_base4 + src_m * params.K;
+        let v = vec4<{{ aScalar }}>(
+          a[a_row + k4],
+          a[a_row + k4 + 1u],
+          a[a_row + k4 + 2u],
+          a[a_row + k4 + 3u],
+        );
+{% else %}
+        let v = a[a_base4 + src_m * k_words_total + kw0 + ki];
+{% endif %}
+        packed = pack_a(v);
+{% if needsASum %}
+        atomicAdd(&a_row_sum[mi], sum4_a(v));
+{% endif %}
+      }
+      a_tile[mi * KT_WORDS + ki] = packed;
+    }
+    workgroupBarrier();
+    if (n_valid) {
+{% if tileM == 1 %}
+      // M=1 has only one output accumulator, so a scalar K loop forms a long
+      // dependent DP4A chain. Interleave several words while preserving the
+      // coalesced mapping of neighboring threads to neighboring B columns.
+      let unrolled_words = words - words % GEMV_K_UNROLL;
+      for (var kk = 0u; kk < unrolled_words; kk += GEMV_K_UNROLL) {
+{% for u in range(GEMV_UNROLL) %}
+        let k4_{{ u }} = (kw0 + kk + {{ u }}u) * 4u;
+        let b_row_{{ u }} = b_base + k4_{{ u }} * params.N + col;
+        let b0_{{ u }} = b[b_row_{{ u }}];
+        let b1_{{ u }} = b[b_row_{{ u }} + params.N];
+        let b2_{{ u }} = b[b_row_{{ u }} + 2u * params.N];
+        let b3_{{ u }} = b[b_row_{{ u }} + 3u * params.N];
+{% if needsBSum %}
+        sum_b{{ "" if u == 0 else "_" ~ u }} += sum4_b(b0_{{ u }}, b1_{{ u }}, b2_{{ u }}, b3_{{ u }});
+{% endif %}
+        let bp_{{ u }} = pack_b(b0_{{ u }}, b1_{{ u }}, b2_{{ u }}, b3_{{ u }});
+        r0{{ "" if u == 0 else "_" ~ u }} += dot4_packed(a_tile[kk + {{ u }}u], bp_{{ u }});
+{% endfor %}
+      }
+      for (var kk = unrolled_words; kk < words; kk++) {
+        let k4 = (kw0 + kk) * 4u;
+        let b_row = b_base + k4 * params.N + col;
+        let b0 = b[b_row];
+        let b1 = b[b_row + params.N];
+        let b2 = b[b_row + 2u * params.N];
+        let b3 = b[b_row + 3u * params.N];
+{% if needsBSum %}
+        sum_b += sum4_b(b0, b1, b2, b3);
+{% endif %}
+        let bp = pack_b(b0, b1, b2, b3);
+        r0 += dot4_packed(a_tile[kk], bp);
+      }
+{% else %}
+      for (var kk = 0u; kk < words; kk = kk + 1u) {
+        let k4 = (kw0 + kk) * 4u;
+        let b_row = b_base + k4 * params.N + col;
+        let b0 = b[b_row];
+        let b1 = b[b_row + params.N];
+        let b2 = b[b_row + 2u * params.N];
+        let b3 = b[b_row + 3u * params.N];
+{% if needsBSum %}
+        sum_b = sum_b + sum4_b(b0, b1, b2, b3);
+{% endif %}
+        let bp = pack_b(b0, b1, b2, b3);
+{% for m in range(tileM) %}
+        r{{ m }} = r{{ m }} + dot4_packed(a_tile[{{ m }}u * KT_WORDS + kk], bp);
+{% endfor %}
+      }
+{% endif %}
+    }
+    workgroupBarrier();
+  }
+{% if aPacked == false %}
+  // Exactly one uniform guarded word handles K % 4. Missing lanes use the raw
+  // bias-domain zero (128 for u8, 0 for i8), so their shifted values, sums, and
+  // dot-product contributions are all zero while params.K remains unpadded for
+  // the zero-point correction below.
+  if (tail_lanes != 0u) {
+    if (tid < TILE_M) {
+      let src_m = m_block + tid;
+      var packed = 0u;
+      if (src_m < params.M) {
+        let k4 = k_words_total * 4u;
+        let a_row = a_base4 + src_m * params.K;
+        let a0 = a[a_row + k4];
+        var a1: {{ aScalar }} = {{ "128u" if aUnsigned else "0" }};
+        var a2: {{ aScalar }} = {{ "128u" if aUnsigned else "0" }};
+        if (tail_lanes > 1u) { a1 = a[a_row + k4 + 1u]; }
+        if (tail_lanes > 2u) { a2 = a[a_row + k4 + 2u]; }
+        let a3: {{ aScalar }} = {{ "128u" if aUnsigned else "0" }};
+        let v = vec4<{{ aScalar }}>(a0, a1, a2, a3);
+        packed = pack_a(v);
+{% if needsASum %}
+        atomicAdd(&a_row_sum[tid], sum4_a(v));
+{% endif %}
+      }
+      a_tile[tid * KT_WORDS] = packed;
+    }
+    workgroupBarrier();
+    if (n_valid) {
+      let k4 = k_words_total * 4u;
+      let b_row = b_base + k4 * params.N + col;
+      let b0 = b[b_row];
+      var b1: {{ bScalar }} = {{ "128u" if bUnsigned else "0" }};
+      var b2: {{ bScalar }} = {{ "128u" if bUnsigned else "0" }};
+      if (tail_lanes > 1u) { b1 = b[b_row + params.N]; }
+      if (tail_lanes > 2u) { b2 = b[b_row + 2u * params.N]; }
+      let b3: {{ bScalar }} = {{ "128u" if bUnsigned else "0" }};
+{% if needsBSum %}
+      sum_b = sum_b + sum4_b(b0, b1, b2, b3);
+{% endif %}
+      let bp = pack_b(b0, b1, b2, b3);
+{% for m in range(tileM) %}
+      r{{ m }} = r{{ m }} + dot4_packed(a_tile[{{ m }}u * KT_WORDS], bp);
+{% endfor %}
+    }
+    workgroupBarrier();
+  }
+{% endif %}
+  if (!n_valid) {
+    return;
+  }
+  // Fold bias + zero-point corrections (exact in modular i32 arithmetic).
+  let cb: i32 = {{ "128" if bUnsigned else "0" }}{% if hasBZero %} - read_b_zero({% if bZeroPerColumn %}col{% endif %}){% endif %};
+{% if not aZeroPerRow %}
+  let ca: i32 = {{ "128" if aUnsigned else "0" }}{% if hasAZero %} - read_a_zero(){% endif %};
+  let kcc: i32 = i32(params.K) * ca * cb;
+{% endif %}
+{% if needsBSum and tileM == 1 %}
+  let sum_b_total = sum_b{% for u in range(1, GEMV_UNROLL) %} + sum_b_{{ u }}{% endfor %};
+{% endif %}
+{% for m in range(tileM) %}
+  {
+    let mr = m_block + {{ m }}u;
+    if (mr < params.M) {
+{% if aZeroPerRow %}
+      let ca: i32 = {{ "128" if aUnsigned else "0" }} - read_a_zero(mr);
+      let kcc: i32 = i32(params.K) * ca * cb;
+{% endif %}
+      var acc = r{{ m }}{% if tileM == 1 %}{% for u in range(1, GEMV_UNROLL) %} + r0_{{ u }}{% endfor %}{% endif %};
+{% if needsASum %}
+      acc = acc + cb * atomicLoad(&a_row_sum[{{ m }}u]);
+{% endif %}
+{% if needsBSum %}
+      acc = acc + ca * {% if tileM == 1 %}sum_b_total{% else %}sum_b{% endif %};
+{% endif %}
+      acc = acc + kcc;
+{% if hasBias %}
+      acc = acc + bias[mr];
+{% endif %}
+{% if requant == "none" %}
+      y[y_base + mr * params.N + col] = acc;
+{% else %}
+      let scaled = f32(acc) * f32(a_scale[0]) * f32(b_scale[{% if requant == "per_column" %}col{% elif requant == "per_row" %}mr{% else %}0u{% endif %}]) / f32(y_scale[0]);
+      // Clamp the rounded float into an i32-representable range before the cast
+      // so a huge-finite / Inf requantized value saturates (the final clamp pins
+      // it to [qMin,qMax]) instead of invoking undefined i32(huge_float).
+      let q = clamp(i32(clamp(round(scaled), -2.0e9, 2.0e9)) + read_y_zero(), {{ qMin }}, {{ qMax }});
+{% if yUnsigned %}
+      y[y_base + mr * params.N + col] = u32(q);
+{% else %}
+      y[y_base + mr * params.N + col] = q;
+{% endif %}
+{% endif %}
+    }
+  }
+{% endfor %}
+}

build/webgpu/quant-pad-k-matrices.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,50 @@

+// Pad the K dimension of widened quantized matrices for aligned DP4A. A is
+// [M, K], optional B is [batch, K, N], and outputs use K_PADDED. Padding stores
+// each tensor's raw zero point, so added lanes contribute exactly zero to
+// (a - a_zero) * (b - b_zero). This keeps odd-K matrices on the aligned DP4A
+// path instead of a scalar-tail path.
+{{ env.wgsl.resourceDeclarations }}
+fn read_a_zero({% if aZeroPerRow %}row: u32{% endif %}) -> i32 {
+{% if aZeroPerRow %}
+  return a_zero_point[row];
+{% else %}
+  return a_zero_point[0];
+{% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+) {
+  let flat = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  let a_count = params.M * params.KPadded;
+  if (flat < a_count) {
+    let row = flat / params.KPadded;
+    let k = flat - row * params.KPadded;
+    var value = read_a_zero({% if aZeroPerRow %}row{% endif %});
+    if (k < params.K) {
+      value = a[row * params.K + k];
+    }
+    a_padded[flat] = value;
+  }
+{% if padB %}
+  let b_count = params.batch * params.KPadded * params.N;
+  if (flat < b_count) {
+    let batch_stride_out = params.KPadded * params.N;
+    let batch = flat / batch_stride_out;
+    let rem = flat - batch * batch_stride_out;
+    let k = rem / params.N;
+    let col = rem - k * params.N;
+    var value = b_zero_point[0];
+    if (k < params.K) {
+      let src = (batch * params.K + k) * params.N + col;
+      value = b[src];
+    }
+    b_padded[flat] = value;
+  }
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1602 @@

+{
+  "op": "ai.onnx.QLinearConv",
+  "fixtureArrays": {
+    "same_upper_stride2_autopad_input_x": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16],
+    "backend_qlinearconv_pointwise_u8_input_x": [255, 174, 162, 25, 203, 168, 58, 15, 59, 237, 95, 129, 0, 64, 56, 242, 153, 221, 168, 12, 166, 232, 178, 186, 195, 237, 162, 237, 188, 39, 124, 77, 80, 102, 43, 127, 230, 21, 83, 41, 40, 134, 255, 154, 92, 141, 42, 148, 247],
+    "ort_depthwise_per_channel_weight_zero_points_s8s8_input_x": [-8, -4, 0, 4, 8, 12, -12, 16, -16, 3, -3, 6, -6, 9, -9, 12, -12, 15, -20, -10, 0, 10, 20, 30, -30, 40, -40],
+    "ort_depthwise_per_channel_weight_zero_points_s8s8_input_w": [-10, -8, -6, -4, -2, 0, 2, 4, 6, 3, 5, 7, 9, 11, 13, 15, 17, 19, -20, -15, -10, -5, 0, 5, 10, 15, 20],
+    "dp4a_pointwise_u8s8_c8_batched_input_w": [-77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90, -128, 127, 0, -1, 56],
+    "dp4a_pointwise_s8s8_c12_signed_output_input_x": [-128, 127, 0, -1, 56, -77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90]
+  },
+  "cases": [
+    {
+      "name": "same_upper_stride2_autopad",
+      "provenance": {
+        "source": "ONNX Runtime QLinearConv-10 CPUExecutionProvider",
+        "notes": "QLinearConv inherits Conv auto_pad semantics. This fixture exercises the exact SAME_UPPER spelling and derived asymmetric bottom/right padding; unit scales and zero points make the expected requantized values equal the integer convolution sums."
+      },
+      "attrs": { "auto_pad": "SAME_UPPER", "kernel_shape": [3, 3], "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/same_upper_stride2_autopad_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 3, 3], "data": { "kind": "constant", "value": 1 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [54, 45, 72, 54] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "dispatch_cliff_pointwise_u8",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 4096, 4097],
+          "data": { "kind": "cycle", "values": [120, 128, 136, 144] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.1] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [130] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.2] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 4096, 4097], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "uint8_padding",
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [128, 129, 130, 131, 132, 133, 134, 135, 136] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [127, 128, 129, 130] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 4, 4] } }
+    },
+    {
+      "name": "int8_output",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-4, -3, -2, -1, 0, 1, 2, 3, 4] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, -1, 2, -2] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 1, 2, 2] } },
+      "attrs": {}
+    },
+    {
+      "name": "int8_input_uint8_output_pairwise_route",
+      "provenance": {
+        "notes": "Pairs the signed input route with the independently typed uint8 output on a compact 1-D pointwise convolution."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1, 1, 1], "data": { "kind": "values", "values": [2] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 1], "data": { "kind": "values", "values": [3] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "requant_exact_half_ties_to_even",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Requantize_NoBias",
+        "notes": "ONNX quantized convolution requantization uses ORT's RoundHalfToEven helper; pointwise accumulators +/-1 and +/-5 with y_scale=2 produce exact +/-0.5 and +/-2.5 tie values."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [1] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "int8", "shape": [4, 1, 1, 1], "data": { "kind": "values", "values": [1, -1, 5, -5] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [128, 128, 130, 126] }
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "requant_huge_finite_saturates_before_i32_overflow",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Pointwise",
+        "notes": "Finite requantized outputs far outside uint8 must saturate; the scalar WebGPU requant pass currently converts to i32 before clamping."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [1] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "int8", "shape": [2, 1, 1, 1], "data": { "kind": "values", "values": [1, -1] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-20] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 2, 1, 1],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [255, 0] }
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "requant_subnormal_y_scale_saturates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Pointwise",
+        "notes": "A valid positive subnormal output scale can make a one-pixel convolution requantize beyond uint8; the result should saturate before any i32 conversion of an infinite f32."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [1] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "int8", "shape": [2, 1, 1, 1], "data": { "kind": "values", "values": [1, -1] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 2, 1, 1],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [255, 0] }
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_requant_subnormal_y_scale_saturates",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Pointwise",
+        "notes": "Fused DP4A pointwise route with a valid positive subnormal output scale should saturate uint8 results before any i32 conversion of an infinite f32."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 4, 1, 1], "data": { "kind": "values", "values": [1, 0, 0, 0] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [2, 4, 1, 1],
+          "data": { "kind": "values", "values": [1, 0, 0, 0, -1, 0, 0, 0] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 2, 1, 1],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [255, 0] }
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "uint8_pointwise_subnormal_scale_ratio_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Requant multiplier x_scale*w_scale/y_scale = 1e-40/1e-40 divides two denormals; Metal flushes denormals in floating-point division (-> NaN), while the CPU reference computes the ratio with denormal support. Subnormal scale-ratio cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2DTest",
+        "notes": "Valid positive subnormal input/output scales whose ratio is exactly meaningful for scalar quantization."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 4, 1, 1], "data": { "kind": "values", "values": [1, 0, 0, 0] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [4, 4, 1, 1],
+          "data": { "kind": "values", "values": [1, 0, 0, 0, 2, 0, 0, 0, 3, 0, 0, 0, 4, 0, 0, 0] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 4, 1, 1], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "backend_qlinearconv_pointwise_u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2DTest",
+        "notes": "Uses the same quantized tensors as ORT's handwritten pointwise Conv2D test."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_qlinearconv_pointwise_u8_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.00369204697] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [132] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [0] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.00172794575] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [255] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.00162681262] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [123] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 7, 7], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "group2_depthwise_scalar_quant_u8",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [128, 129, 130, 131, 120, 124, 128, 132] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.1] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [2, 1, 1, 1], "data": { "kind": "values", "values": [130, 126] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.2] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "dilated_stride2_padding_saturates_u8",
+      "attrs": { "strides": [2, 2], "dilations": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 5, 5],
+          "data": {
+            "kind": "values",
+            "values": [0, 255, 0, 255, 0, 255, 0, 255, 0, 255, 0, 255, 128, 255, 0, 255, 0, 255, 0, 255, 0, 255, 0, 255, 0]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [0, 255, 255, 0] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.04] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.001] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_style_multi_channel_four_outputs_pad1_u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.WithBias_2D",
+        "notes": "Adapts the handwritten 2D multi-channel/multi-output QLinearConv shape without optional bias."
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 4, 4],
+          "data": { "kind": "cycle", "values": [135, 140, 128, 120, 150, 110, 160, 100] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [4, 2, 3, 3],
+          "data": { "kind": "cycle", "values": [110, 115, 105, 130, 100, 125] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.03] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [110] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 4, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_style_group3_stride2_pad1_u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.WithGroup_2D",
+        "notes": "Covers grouped convolution with group=3 and stride=2 using symmetric padding supported by the WebGPU manifest."
+      },
+      "attrs": { "group": 3, "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 6, 5, 5],
+          "data": { "kind": "cycle", "values": [98, 166, 219, 195, 46, 97, 27, 211, 239, 1] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.01] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [135] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [6, 2, 3, 3],
+          "data": { "kind": "cycle", "values": [220, 111, 73, 254, 235, 151, 6, 156, 129, 204, 234, 198] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [110] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.04] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 6, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_with_group_2d_asymmetric_pad_no_bias_u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.WithGroup_2D",
+        "notes": "Uses ORT's grouped convolution tensors and asymmetric pads without optional bias."
+      },
+      "attrs": { "group": 3, "strides": [2, 2], "pads": [0, 0, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 6, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [98, 166, 219, 195, 46, 97, 27, 211, 239, 1, 28, 208, 143, 144, 215, 252, 79, 5, 154, 56, 122, 191, 94, 25, 221, 48, 37, 182, 68, 245, 210, 206, 183, 22, 163, 104, 242, 112, 161, 66, 181, 235, 117, 75, 236, 61, 115, 36, 120, 253, 165, 214, 159, 132, 11, 201, 30, 249, 89, 171, 186, 67, 225, 197, 135, 142, 241, 169, 170, 164, 178, 58, 50, 51, 200, 43, 199, 126, 222, 123, 227, 42, 3, 21, 124, 220, 24, 47, 63, 110]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.01] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [135] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [6, 2, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [220, 111, 73, 254, 235, 151, 6, 156, 129, 204, 234, 198, 44, 89, 202, 82, 118, 189, 71, 120, 123, 121, 110, 83, 173, 248, 108, 229, 124, 68, 85, 239, 133, 213, 112, 122, 170, 231, 225, 195, 192, 9, 232, 97, 160, 227, 67, 137]
+          }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.15] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [110] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.75] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 6, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "mixed_u8_i8_uint8_output_dilated_pad1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Dilations",
+        "notes": "Exercises the ORT-covered mixed uint8 activation / int8 weight quantization path with dilation."
+      },
+      "attrs": { "dilations": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 5, 5],
+          "data": { "kind": "cycle", "values": [0, 32, 96, 127, 128, 160, 224, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.04] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [-128, 127, -32, 96] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.01] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 5, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_u8s8_pointwise_multi_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Pointwise",
+        "notes": "Adapts QLinearConvTest.Conv2D_U8S8_Pointwise with multiple input and output channels."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [4, 15, 64, 128, 200, 255, 250, 180, 127, 80, 42, 0] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 2, 1, 1],
+          "data": { "kind": "values", "values": [0, 12, -7, 31, 63, -63] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.55] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [54] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 3, 2, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_u8s8_stride2_small",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Strides",
+        "notes": "Adapts QLinearConvTest.Conv2D_U8S8_Strides using stride=2 and scalar quantization."
+      },
+      "attrs": { "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 5, 5],
+          "data": {
+            "kind": "values",
+            "values": [4, 18, 32, 47, 63, 79, 95, 111, 127, 143, 159, 175, 191, 207, 223, 239, 255, 231, 199, 167, 135, 103, 71, 39, 7]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [-63, -9, 17, 63] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.31] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [30] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_u8s8_dilation2_small",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Dilations",
+        "notes": "Adapts QLinearConvTest.Conv2D_U8S8_Dilations using dilation=2."
+      },
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [4, 60, 116, 172, 228, 200, 144, 88, 32, 12, 68, 124, 180, 236, 252, 96]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [-3, 19, -41, 63] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.24] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [15] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_u8s8_group2_small",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Groups",
+        "notes": "Adapts QLinearConvTest.Conv2D_U8S8_Groups with two groups and scalar weight quantization."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 3, 3],
+          "data": { "kind": "cycle", "values": [4, 17, 33, 65, 97, 129, 161, 193, 225, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [4, 2, 2, 2],
+          "data": { "kind": "cycle", "values": [-63, -45, -18, -1, 0, 7, 23, 49, 63] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.76] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [88] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 4, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_s8s8_stride2_pad1_signed_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_S8S8_Strides",
+        "notes": "Adapts QLinearConvTest.Conv2D_S8S8_Strides with signed activations, signed weights, padding, and stride."
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 5, 5],
+          "data": {
+            "kind": "values",
+            "values": [-128, -96, -64, -32, 0, 32, 64, 96, 127, 95, 63, 31, -1, -33, -65, -97, -127, -80, -40, 8, 48, 88, 120, 72, 24]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-63, -31, 0, 17, 45, 63, -12, 29, -48] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.31] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-30] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 1, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_u8u8_pointwise_multi_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8U8_Pointwise",
+        "notes": "Adapts the ORT pointwise U8/U8 path using scalar quantization and multiple output channels."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [4, 15, 64, 128, 200, 255, 250, 180, 127, 80, 42, 0] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [3, 2, 1, 1],
+          "data": { "kind": "values", "values": [4, 12, 31, 63, 127, 255] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.03125] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.55] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [54] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 3, 2, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_s8s8_pointwise_signed_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_S8S8_Pointwise",
+        "notes": "Adapts the ORT signed-activation/signed-weight pointwise path with signed output saturation."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 2, 2, 3],
+          "data": { "kind": "values", "values": [-128, -96, -64, -32, 0, 31, 63, 95, 127, 64, 16, -16] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.04] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-3] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 2, 1, 1],
+          "data": { "kind": "values", "values": [-63, -31, -7, 0, 31, 63] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0275] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-5] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 3, 2, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_u8s8_depthwise_no_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_Depthwise_NoBias",
+        "notes": "Adapts ORT's depthwise no-bias coverage using scalar quantization and group=in_channels."
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 3, 3, 3],
+          "data": { "kind": "cycle", "values": [4, 31, 63, 95, 127, 159, 191, 223, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.04] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": { "dtype": "int8", "shape": [3, 1, 2, 2], "data": { "kind": "cycle", "values": [-63, -17, 0, 23, 63] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [80] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_s8s8_depthwise_no_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_S8S8_Depthwise_NoBias",
+        "notes": "Adapts ORT's signed depthwise no-bias coverage using scalar quantization and group=in_channels."
+      },
+      "attrs": { "group": 2, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "cycle", "values": [-128, -96, -64, -16, 0, 16, 64, 96, 127] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-1] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [2, 1, 2, 2],
+          "data": { "kind": "values", "values": [-63, -31, 17, 63, 45, -12, -48, 29] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.4] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-12] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 2, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_1d_u8s8_pad1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv1D_U8S8",
+        "notes": "Compact 1D convolution with symmetric padding."
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 5], "data": { "kind": "values", "values": [10, 11, 12, 13, 14] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 3], "data": { "kind": "values", "values": [1, 2, 3] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [100] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 5],
+          "data": { "kind": "values", "values": [103, 108, 114, 120, 111] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_1d_u8s8_stride2_dilation2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv1D_U8S8_Dilations",
+        "notes": "Compact 1D convolution combining padding, stride, and dilation."
+      },
+      "attrs": { "strides": [2], "dilations": [2], "pads": [1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7],
+          "data": { "kind": "values", "values": [10, 12, 14, 16, 18, 20, 22] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 3], "data": { "kind": "values", "values": [1, -2, 3] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [100] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [114, 120, 86] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_1d_s8s8_pointwise_signed_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv1D_S8S8_Pointwise",
+        "notes": "Compact signed 1D pointwise convolution."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1, 1, 4], "data": { "kind": "values", "values": [-3, -2, -1, 0] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-3] } },
+        "w": { "dtype": "int8", "shape": [1, 1, 1], "data": { "kind": "values", "values": [4] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [-10, -6, -2, 2] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "onnx_backend_qlinearconv",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_qlinearconv" },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/backend_qlinearconv_pointwise_u8_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.003692046971991658] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [132] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [0] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0017279457533732057] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [255] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.001626812620088458] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [123] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 7, 7], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_pointwise_2d_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.WithBias_2D",
+        "notes": "Compact ORT-valid pointwise adaptation that preserves optional int32 bias semantics."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [10, 11, 12, 13] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [14] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [100] } },
+        "B": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [98, 102, 106, 110] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_pointwise_3d_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv3DTest",
+        "notes": "Compact ORT-valid pointwise adaptation of the 3D QLinearConv path."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [10, 11, 12, 13, 14, 15, 16, 17] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [14] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [100] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [100, 104, 108, 112, 116, 120, 124, 128] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_3d_u8s8_compact",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv3D_U8S8",
+        "notes": "Compact deterministic adaptation of ORT's 3D uint8-input/int8-weight QLinearConv coverage."
+      },
+      "attrs": { "strides": [1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": [4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "w": { "dtype": "int8", "shape": [1, 2, 2, 1, 1], "data": { "kind": "values", "values": [2, -3, 4, 1] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.025] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [54] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 1, 2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [62, 63, 64, 65] }
+        }
+      }
+    },
+    {
+      "name": "ort_3d_s8s8_compact",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv3D_S8S8",
+        "notes": "Signed companion for compact 3D QLinearConv: int8 input, int8 weights, and int8 output."
+      },
+      "attrs": { "strides": [1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 2, 2, 2, 2],
+          "data": { "kind": "values", "values": [-4, -3, -2, -1, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-4] } },
+        "w": { "dtype": "int8", "shape": [1, 2, 2, 1, 1], "data": { "kind": "values", "values": [2, -3, 4, 1] } },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.125] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.025] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-54] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [1, 1, 1, 2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-46, -45, -44, -43] }
+        }
+      }
+    },
+    {
+      "name": "ort_per_channel_weight_zero_points_u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8U8_PerChannelZeroPoints",
+        "notes": "Compact ORT-validated per-output-channel weight scales and zero-points without optional bias."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [120, 128, 136, 144, 100, 110, 130, 150] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [2, 2, 1, 1], "data": { "kind": "values", "values": [100, 120, 140, 160] } },
+        "w_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 0.25] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [100, 150] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [0, 37, 137, 237, 102, 105, 119, 135] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_per_channel_weight_zero_points_u8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8S8_PerChannelZeroPoints",
+        "notes": "Compact mixed uint8 activation / int8 weight per-output-channel quantization case without optional bias."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [120, 128, 136, 144, 100, 110, 130, 150] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.2] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "int8", "shape": [2, 2, 1, 1], "data": { "kind": "values", "values": [-8, 12, 16, -4] } },
+        "w_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.25, 0.5] } },
+        "w_zero_point": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [-2, 6] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [59, 78, 124, 170, 208, 200, 152, 104] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_per_channel_weight_zero_points_s8s8_signed_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_S8S8_PerChannelZeroPoints",
+        "notes": "Compact signed activation / signed weight projection with per-output-channel weight scales and zero-points."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [-8, -4, 0, 4, 8, -12, 16, -16] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } },
+        "w": { "dtype": "int8", "shape": [2, 2, 1, 1], "data": { "kind": "values", "values": [-10, 5, 12, -6] } },
+        "w_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 0.25] } },
+        "w_zero_point": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [-8, 4] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-5] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [71, -63, 115, -97, -43, 15, -47, 41] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_grouped_per_channel_weight_zero_points_u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_U8U8_Groups_PerChannelZeroPoints",
+        "notes": "Compact grouped 1x1 projection with per-output-channel weight scales and zero-points."
+      },
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [120, 128, 136, 144, 100, 110, 130, 150, 126, 130, 134, 138, 90, 120, 160, 200]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [4, 2, 1, 1],
+          "data": { "kind": "values", "values": [100, 120, 140, 160, 110, 130, 170, 90] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 0.25, 0.4, 0.3] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [100, 150, 120, 80] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": [0, 37, 137, 237, 102, 105, 119, 135, 55, 107, 179, 251, 43, 142, 255, 255]
+          },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_depthwise_per_channel_weight_zero_points_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_S8S8_DepthwiseFallback_PerChannelZeroPoints",
+        "notes": "Compact depthwise projection with distinct per-output-channel weight scales and zero-points."
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_depthwise_per_channel_weight_zero_points_s8s8_input_x" }
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_depthwise_per_channel_weight_zero_points_s8s8_input_w" }
+          }
+        },
+        "w_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.1, 0.13, 0.16] } },
+        "w_zero_point": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-10, -4, 6] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-8] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 3, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_depthwise_uniform_per_channel_weight_zero_points_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
+        "test": "QLinearConvTest.Conv2D_S8S8_Depthwise_PerChannelUniformZeroPoints",
+        "notes": "Uses a per-channel weight zero-point tensor with identical values, while keeping per-output-channel scales."
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_depthwise_per_channel_weight_zero_points_s8s8_input_x" }
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 1, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_depthwise_per_channel_weight_zero_points_s8s8_input_w" }
+          }
+        },
+        "w_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.1, 0.13, 0.16] } },
+        "w_zero_point": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-7, -7, -7] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-8] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 3, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_pointwise_u8s8_c8_batched",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 8, 3, 5],
+          "data": {
+            "kind": "cycle",
+            "values": [0, 255, 128, 127, 56, 77, 3, 100, 200, 42, 5, 19, 64, 88, 33, 7, 125, 90, 211, 16]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [113] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [6, 8, 1, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_u8s8_c8_batched_input_w" } }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.4] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 6, 3, 5], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_s8s8_c12_signed_output",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 12, 7, 11],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_s8s8_c12_signed_output_input_x" }
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.025] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [9, 12, 1, 1],
+          "data": {
+            "kind": "cycle",
+            "values": [19, -64, 88, -33, 7, 125, -90, -128, 127, 0, -1, 56, -77, 3, 100, -100, 42, -5]
+          }
+        },
+        "w_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.018] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-3] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.6] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 9, 7, 11], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [0, 4, 2, 2], "data": { "kind": "values", "values": [] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [2, 4, 1, 1], "data": { "kind": "constant", "value": 128 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [0, 2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_direct_3x3_u8s8_padding_spatial_oc_tails",
+      "provenance": {
+        "notes": "Compact route lock for the direct tiled 3x3 DP4A kernel. Nonzero zero points exercise exact correction; 10x11 spatial and 12-channel output tails cover partial workgroups."
+      },
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 16, 10, 11],
+          "data": { "kind": "cycle", "values": [0, 255, 128, 127, 56, 77, 3, 100, 200, 42, 5, 19] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [113] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [12, 16, 3, 3],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_u8s8_c8_batched_input_w" } }
+        },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": {
+            "kind": "values",
+            "values": [0.011, 0.012, 0.013, 0.014, 0.015, 0.016, 0.017, 0.018, 0.019, 0.02, 0.021, 0.022]
+          }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [12],
+          "data": { "kind": "values", "values": [-7, -6, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.6] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 12, 10, 11], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_im2col_u8s8_oc64_c16_3x3_24x24",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 16, 24, 24],
+          "data": { "kind": "cycle", "values": [3, 1, 4, 1, 5, 9, 2, 6, 5, 3, 8] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [2] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [64, 16, 3, 3],
+          "data": { "kind": "cycle", "values": [2, -1, 3, -2, 1, 4, -3, 2, 0] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [1] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 64, 22, 22] } },
+      "attrs": {}
+    },
+    {
+      "name": "im2col_dp4a_requant_huge_finite_saturates_3x3",
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 16, 24, 24], "data": { "kind": "constant", "value": 255 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "int8", "shape": [16, 16, 3, 3], "data": { "kind": "constant", "value": 127 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-20] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 16, 22, 22], "tolerance": 0 } }
+    },
+    {
+      "name": "im2col_dp4a_zero_spatial_output_empty",
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [10, 20, 30, 40] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 3, 3], "data": { "kind": "constant", "value": 128 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 0, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "per_channel_grouped_requant_oc_indexing_g4_i8",
+      "attrs": { "group": 4, "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 8, 4, 4],
+          "data": { "kind": "cycle", "values": [-8, -4, 0, 4, 8, 12, -12, 16, -16, 3, -3, 6, -6, 9, -9, 12, 20, -20] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.05] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 2, 3, 3],
+          "data": { "kind": "cycle", "values": [-10, -8, -6, -4, -2, 0, 2, 4, 6, 3, 5, 7, 9, 11, 13, 15, 17, 19] }
+        },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.1, 0.13, 0.16, 0.09, 0.2, 0.07, 0.11, 0.14] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-10, -4, 6, 0, 3, -7, 5, -2] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-5] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 8, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_pointwise_u8_weight_zero_128_bias_shift_edge",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 8, 6, 6],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 254, 128, 127, 200, 55, 33, 222, 99, 156] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": {
+          "dtype": "uint8",
+          "shape": [4, 8, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 128, 64, 200, 30, 129, 127, 1, 254, 100, 156] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.4] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 4, 6, 6], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "conv3d_depth_dilation2_compact",
+      "provenance": {
+        "source": "ONNX QLinearConv-10 volumetric dilation semantics",
+        "notes": "Covers a non-default depth dilation. Existing 3D fixtures kept dilationD=1 and no 3D performance fixture existed."
+      },
+      "attrs": { "strides": [1, 1, 1], "dilations": [2, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 4, 2, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/same_upper_stride2_autopad_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [1, 2] } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [19, 22, 25, 28, 31, 34, 37, 40] }
+        }
+      }
+    },
+    {
+      "name": "dp4a_pointwise_per_channel_bias_compact",
+      "provenance": {
+        "source": "ONNX QLinearConv-10 per-output-channel quantization",
+        "notes": "Locks the common per-channel weight-scale/zero-point plus int32 bias combination on an aligned pointwise DP4A shape."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 3, 3],
+          "data": { "kind": "cycle", "values": [0, 17, 128, 255, 63, 191] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [4, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-127, -13, 0, 11, 63, 126] }
+        },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0.01, 0.015, 0.02, 0.025] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-3, 0, 2, 5] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "B": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [-17, 0, 31, 4096] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 4, 3, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "feature_independent_x4_odd_width_k5_compact",
+      "provenance": {
+        "source": "ONNX QLinearConv-10 scalar quantization",
+        "notes": "Locks four-column feature-independent accumulation with an odd output-width tail and K not divisible by four."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 5, 3, 5],
+          "data": { "kind": "cycle", "values": [0, 17, 63, 127, 128, 191, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 5, 1, 1],
+          "data": { "kind": "cycle", "values": [-127, -31, -1, 0, 7, 63, 126] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.015] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.4] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [129] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 3, 3, 5], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_padded_k66_baseline_16x16",
+      "provenance": {
+        "source": "ONNX QLinearConv-10 pointwise quantized convolution",
+        "notes": "Locks the aligned-padding DP4A fallback for a realistic odd channel count on devices without shader-f16. Padding uses raw zero points and must remain exact."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 66, 16, 16],
+          "data": { "kind": "cycle", "values": [0, 17, 63, 127, 128, 191, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-127, -31, -1, 0, 7, 63, 126] }
+        },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.01, 0.012, 0.014, 0.016, 0.018, 0.02, 0.022, 0.024] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-7, -3, -1, 0, 2, 4, 6, 9] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.4] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [129] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 8, 16, 16], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_im2col_padded_k81_baseline_exact",
+      "provenance": {
+        "notes": "Exactness and selector lock for no-feature 2-D im2col with odd K=81; im2col writes raw-x-zero-point padded rows and the shared pad pass extends per-row weight zero points."
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 9, 18, 18],
+          "data": { "kind": "cycle", "values": [0, 255, 127, 7, 200, 31, 129, 64, 99, 3, 254, 1, 88, 173, 240, 16, 5] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [32, 9, 3, 3],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_s8s8_c12_signed_output_input_x" }
+          }
+        },
+        "w_scale": { "dtype": "float32", "shape": [32], "data": { "kind": "constant", "value": 0.03 } },
+        "w_zero_point": { "dtype": "int8", "shape": [32], "data": { "kind": "constant", "value": -4 } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [8.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 32, 16, 16], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_im2col_bias_tail_k27_baseline_exact",
+      "provenance": {
+        "notes": "Exact selector lock for the portable DP4A bias path with odd K=27; the shared core packs its final three weight/input lanes without shader-f16."
+      },
+      "attrs": { "strides": [1, 1], "pads": [0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 3, 5, 5],
+          "data": { "kind": "cycle", "values": [0, 17, 63, 127, 128, 191, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 3, 3, 3],
+          "data": { "kind": "cycle", "values": [-128, -77, -5, 0, 7, 42, 88, 127] }
+        },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.03] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-4] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "B": {
+          "dtype": "int32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-4096, -31, -1, 0, 1, 29, 1024, 4096] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 8, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_im2col_3d_u8s8_b1c4m8_4x8x8_k3",
+      "provenance": {
+        "source": "ONNX QLinearConv-10 volumetric convolution",
+        "notes": "Locks the 3-D im2col-to-DP4A route, including zero-point-filled spatial padding and per-output-channel quantization."
+      },
+      "attrs": { "pads": [1, 1, 1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 4, 8, 8],
+          "data": { "kind": "cycle", "values": [0, 17, 63, 127, 128, 191, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 4, 3, 3, 3],
+          "data": { "kind": "cycle", "values": [-127, -31, -1, 0, 7, 63, 126] }
+        },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.01, 0.012, 0.014, 0.016, 0.018, 0.02, 0.022, 0.024] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-7, -3, -1, 0, 2, 4, 6, 9] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [129] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 8, 4, 8, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "feature_independent_x4_group2_octile4_odd_width_compact",
+      "provenance": {
+        "notes": "Compact selector and exactness lock for grouped 3x3 QLinearConv with four output channels per group. Odd output width also exercises the feature-independent x4 spatial tail."
+      },
+      "attrs": { "group": 2, "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 5, 5],
+          "data": { "kind": "cycle", "values": [0, 17, 63, 127, 128, 191, 255] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 2, 3, 3],
+          "data": { "kind": "cycle", "values": [-127, -31, -1, 0, 7, 63, 126] }
+        },
+        "w_scale": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.01, 0.012, 0.014, 0.016, 0.018, 0.02, 0.022, 0.024] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-7, -3, -1, 0, 2, 4, 6, 9] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [129] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 8, 5, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "same_lower_stride2_autopad",
+      "provenance": {
+        "source": "ONNX Runtime QLinearConv-10 CPUExecutionProvider",
+        "notes": "ONNX auto_pad SAME_LOWER puts the odd padding element at the START of the axis, so unlike its SAME_UPPER twin this case has a non-zero leading pad and fails if the kernel is handed the explicit pad attributes instead of the derived ones."
+      },
+      "attrs": { "auto_pad": "SAME_LOWER", "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/same_upper_stride2_autopad_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 3, 3], "data": { "kind": "constant", "value": 1 } },
+        "w_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 2, 2], "tolerance": 0 } }
+    }
+  ]
+}