sync 91d990483a17

Browse files

Files changed (12) hide show

README.md +23 -19
build/webgpu/bench.json +0 -1
build/webgpu/conv-int-accumulate-spatial.wgsl.jinja +34 -33
build/webgpu/conv-int-im2col-spatial.wgsl.jinja +18 -17
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +35 -13
build/webgpu/qlinear-conv-dp4a-direct-3x3.wgsl.jinja +23 -23
build/webgpu/qlinear-conv-nchw-x4.wgsl.jinja +15 -15
build/webgpu/qlinear-conv-requantize.wgsl.jinja +15 -15
build/webgpu/quant-dp4a-matmul.wgsl.jinja +29 -98
build/webgpu/quant-pad-k-matrices.wgsl.jinja +1 -2
build/webgpu/test.json +103 -10

README.md CHANGED Viewed

@@ -12,29 +12,29 @@ tags:
 ## 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
@@ -43,8 +43,8 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `auto_pad` | `"NOTSET"` | Automatic padding mode. `NOTSET` uses `pads`; `SAME_UPPER` and `SAME_LOWER` choose padding so each output spatial size is `ceil(input / stride)`; `VALID` uses no padding. |
-| `group` | `1` | Number of groups that input and output channels are split into; defaults to 1. |
 | `dilations` | — | Optional dilation factors, one positive integer per spatial axis. Omission means all ones. |
 | `kernel_shape` | — | Optional kernel shape, one positive integer per spatial axis. When present, it must match the spatial dimensions of the weight tensor; omission infers the shape from the weights. |
 | `pads` | — | Optional explicit padding in ONNX order `[begin_axis_0, ..., begin_axis_n, end_axis_0, ..., end_axis_n]`. Omission means all zeros; it cannot be combined with an automatic padding mode. |
 | `strides` | — | Optional stride factors, one positive integer per spatial axis. Omission means all ones. |
@@ -60,7 +60,7 @@ Attributes and default values (overridable per request):
 ## 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
@@ -74,10 +74,14 @@ Attributes and default values (overridable per request):
 ## 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.

 ## Description
+Performs convolution on a quantized input using quantized weights, producing a quantized output. Input and output quantization use per-tensor scales and zero points; weight quantization may be 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 | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `x` | `TX` | — | — | Quantized input data tensor shaped `(N x C x D1 x ... x Dn)`. | required |
+| `x_scale` | `TF` | — | — | Per-tensor scale for input `x`. | required |
+| `x_zero_point` | `TX` | — | — | Per-tensor zero point for input `x`. | required |
+| `w` | `TW` | — | — | Quantized weight tensor shaped `(M x C/group x k1 x ... x kn)`. | required |
+| `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` | `TW` | — | — | Zero point for weights `w`; scalar or 1-D of length M matching `w_scale`. | required |
+| `y_scale` | `TF` | — | — | Per-tensor scale for output `y`. | required |
+| `y_zero_point` | `TY` | — | — | Per-tensor zero point for output `y`. | required |
+| `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 | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `y` | `TY` | same as `x` | derived | Quantized output tensor; shape is determined by kernel size, strides, and padding. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `auto_pad` | `"NOTSET"` | Automatic padding mode. `NOTSET` uses `pads`; `SAME_UPPER` and `SAME_LOWER` choose padding so each output spatial size is `ceil(input / stride)`; `VALID` uses no padding. |
 | `dilations` | — | Optional dilation factors, one positive integer per spatial axis. Omission means all ones. |
+| `group` | `1` | Number of groups that input and output channels are split into; defaults to 1. |
 | `kernel_shape` | — | Optional kernel shape, one positive integer per spatial axis. When present, it must match the spatial dimensions of the weight tensor; omission infers the shape from the weights. |
 | `pads` | — | Optional explicit padding in ONNX order `[begin_axis_0, ..., begin_axis_n, end_axis_0, ..., end_axis_n]`. Omission means all zeros; it cannot be combined with an automatic padding mode. |
 | `strides` | — | Optional stride factors, one positive integer per spatial axis. Omission means all ones. |
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.QLinearConv",
   "tunableSpace": {
     "DP4A_DIRECT_SPATIAL_TILE": [8, 16],
     "DP4A_DIRECT_OC_TILE": [8, 16],

 {
   "tunableSpace": {
     "DP4A_DIRECT_SPATIAL_TILE": [8, 16],
     "DP4A_DIRECT_OC_TILE": [8, 16],

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

@@ -1,37 +1,36 @@
 {% 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 {
@@ -76,13 +75,13 @@ fn read_w_zero({% if wZeroPerChannel %}oc: u32{% endif %}) -> i32 {
 }
 @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 %}
@@ -94,28 +93,30 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   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 %}
   }

 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {%- endif %}
 {% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn read_x(index: u32) -> i32 {
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: 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 spatialRankSpec == 3 %}
   let od = t % params.outD;
   t = t / params.outD;
   {% endif %}
   var acc = 0i;
   for (var ic = 0u; ic < params.weightInChannels; ic = ic + 1u) {
     let input_channel = group * params.inChannelsPerGroup + ic;
+    {% if spatialRankSpec == 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 %}
+{% filter indent(2 if spatialRankSpec == 3 else 0, true) %}
+    for (var kh = 0u; kh < params.kernelH; kh = kh + 1u) {
+      let ih = i32(oh * params.strideH + kh * params.dilationH) - params.padH;
+      if (ih < 0 || ih >= i32(params.inH)) { continue; }
+      for (var kw = 0u; kw < params.kernelW; kw = kw + 1u) {
+        let iw = i32(ow * params.strideW + kw * params.dilationW) - params.padW;
+        if (iw < 0 || iw >= i32(params.inW)) { continue; }
+        {% if spatialRankSpec == 3 %}
+        let x_index = (((batch * params.inChannels + input_channel) * params.inD + u32(id)) * params.inH + u32(ih)) * params.inW + u32(iw);
+        let w_index = (((oc * params.weightInChannels + ic) * params.kernelD + kd) * params.kernelH + kh) * params.kernelW + kw;
         {% else %}
+        let x_index = ((batch * params.inChannels + input_channel) * params.inH + u32(ih)) * params.inW + u32(iw);
+        let w_index = ((oc * params.weightInChannels + ic) * params.kernelH + kh) * params.kernelW + kw;
         {% endif %}
+        acc = acc + (read_x(x_index) - xzp) * (read_w(w_index) - wzp);
+      }
+    }
+{% endfilter %}
+    {% if spatialRankSpec == 3 %}
     }
     {% endif %}
   }

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

@@ -5,27 +5,28 @@
  * 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>) {
@@ -35,7 +36,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   }
   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;
@@ -43,7 +44,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% 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;
@@ -63,7 +64,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   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)

  * Optional extra K rows contain the raw input zero point to align the DP4A
  * reduction. */
 {{ env.wgsl.resourceDeclarations }}
+{% set padK = padK is defined and padK %}
+{% if spatialRankSpec == 3 %}
+const KERNEL_D: u32 = {{ kernelDSpec }}u;
 {% endif %}
+const KERNEL_H: u32 = {{ kernelHSpec }}u;
+const KERNEL_W: u32 = {{ kernelWSpec }}u;
 const KERNEL_HW: u32 = KERNEL_H * KERNEL_W;
+{% if spatialRankSpec == 3 %}
 const KSIZE: u32 = KERNEL_D * KERNEL_HW;
+const STRIDE_D: u32 = {{ strideDSpec }}u;
+const DILATION_D: u32 = {{ dilationDSpec }}u;
+const PAD_FRONT: i32 = {{ padFrontSpec }};
 {% else %}
 const KSIZE: u32 = KERNEL_HW;
 {% endif %}
+const STRIDE_H: u32 = {{ strideHSpec }}u;
+const STRIDE_W: u32 = {{ strideWSpec }}u;
+const DILATION_H: u32 = {{ dilationHSpec }}u;
+const DILATION_W: u32 = {{ dilationWSpec }}u;
+const PAD_TOP: i32 = {{ padTopSpec }};
+const PAD_LEFT: i32 = {{ padLeftSpec }};
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   }
   let k = gid.y;
   let batch = gid.z;
+{% if 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 spatialRankSpec == 3 %}
   let kd = kq / KERNEL_HW;
   let khw = kq - kd * KERNEL_HW;
   let kh = khw / KERNEL_W;
   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 spatialRankSpec == 3 %}
   if (
     id >= 0 && id < i32(params.inD)
     && ih >= 0 && ih < i32(params.inH)

build/webgpu/manifest.json CHANGED Viewed

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

build/webgpu/metadata.json CHANGED Viewed

@@ -1,24 +1,46 @@
 {
   "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" }
 }

 {
   "name": "ai.onnx.QLinearConv",
+  "id": "_ai_onnx_qlinearconv_webgpu_5dc8940",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "spUQwS0Tr+3Vdfkxd9WepuT6g6ezyeV6aMao4r7OsdY=",
+      "conv-int-accumulate-spatial.wgsl.jinja": "TkacF0/fJ1yShdyPJWePK/+8XJFaV+ZDKWT6XvHOCI8=",
+      "conv-int-im2col-spatial.wgsl.jinja": "3vFeBqBVtyYtxV1yteSQ3O7skFur08E3vDoVtdwcYAo=",
+      "manifest.json": "VInAZzgFiEXRlHexx+MCXfvzy6hnwIGSsbdY6pw6Nas=",
+      "qlinear-conv-dp4a-direct-3x3.wgsl.jinja": "xt3DqIwhlzomIiLMuMZ+8F0+Xqg7AWfGCbeZnNT2rLU=",
+      "qlinear-conv-nchw-x4.wgsl.jinja": "XT1MIv1iIEOhj0keQUkU5cau6Nc/CTJCAgu+PardWaQ=",
+      "qlinear-conv-requantize.wgsl.jinja": "JCGp/e1wWm22EQlsYPNreGtdualHITd5dmvA9kGYZCk=",
+      "quant-dp4a-matmul.wgsl.jinja": "x1yY2y4Wb6vlzMVeRw+DbEmfvJdPBHkMF8ATG34ZGis=",
+      "quant-pad-k-matrices.wgsl.jinja": "hsLWgSTbpXBnsIOyrIzsmhQzVyIYALaFvx62VN46KyY=",
+      "test.json": "z2krhaHYtSozWBIW2J1ZKnlpd1S7dDdTQwYq83sf/7g="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "dp4a_pointwise_1x1_padded_k_baseline": ["quant-dp4a-matmul.wgsl.jinja", "quant-pad-k-matrices.wgsl.jinja"],
+      "dp4a_im2col_nchw_padded_k_baseline": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja", "quant-pad-k-matrices.wgsl.jinja"],
+      "dp4a_direct_3x3_nchw_u8s8": ["qlinear-conv-dp4a-direct-3x3.wgsl.jinja"],
+      "dp4a_pointwise_1x1": ["quant-dp4a-matmul.wgsl.jinja"],
+      "nchw2d_feature_independent_x4": ["qlinear-conv-nchw-x4.wgsl.jinja"],
+      "ncdhw3d_per_channel_quantization": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "ncw1d_scalar_quantization": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "ncdhw3d_scalar_quantization": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "nchw2d_scalar_quantization": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "ncw1d_scalar_quantization_bias": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "ncdhw3d_scalar_quantization_bias": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "ncw1d_per_channel_quantization": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "dp4a_im2col_ncdhw3d_scalar_quantization": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_im2col_nchw_bias_quantization": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_im2col_nchw_scalar_quantization": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
+      "nchw2d_per_channel_quantization": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "nchw2d_per_channel_quantization_bias": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"],
+      "nchw2d_scalar_quantization_bias": ["conv-int-accumulate-spatial.wgsl.jinja", "qlinear-conv-requantize.wgsl.jinja"]
+    }
+  }
 }

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

@@ -13,17 +13,17 @@ fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
 }
-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;
@@ -61,24 +61,24 @@ fn sum4_packed(v: u32) -> i32 {
     + 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,
@@ -87,7 +87,7 @@ fn requantize(acc: i32{% if source.wScalePerChannel %}, oc: u32{% endif %}) -> u
   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>
@@ -107,7 +107,7 @@ fn main(
   }
   workgroupBarrier();
-{% for oc in range(source.ocTile) %}
   var acc{{ oc }} = 0i;
 {% endfor %}
   var sumX = 0i;
@@ -167,7 +167,7 @@ fn main(
             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 %}
           }
@@ -181,17 +181,17 @@ fn main(
     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 %}

 }
+const IN_C: u32 = {{ inChannels }}u;
+const IN_H: u32 = {{ inHeight }}u;
+const IN_W: u32 = {{ inWidth }}u;
+const OUT_C: u32 = {{ outChannels }}u;
+const OUT_H: u32 = {{ outHeight }}u;
+const OUT_W: u32 = {{ outWidth }}u;
+const BATCH: u32 = {{ batch }}u;
+const SPATIAL_TILE: u32 = {{ spatialTile }}u;
+const OC_TILE: u32 = {{ ocTile }}u;
+const IC_TILE: u32 = {{ icTile }}u;
 const IC_WORDS: u32 = IC_TILE / 4u;
 const PATCH_EDGE: u32 = SPATIAL_TILE + 2u;
 const PATCH_AREA: u32 = PATCH_EDGE * PATCH_EDGE;
     + signed_byte(v, 16u) + signed_byte(v, 24u);
 }
+fn read_w_zero({% if wZeroPerChannel %}oc: u32{% endif %}) -> i32 {
+{% if wZeroPerChannel %}
   return w_zero_point[oc];
 {% else %}
   return w_zero_point[0];
 {% endif %}
 }
+fn read_w_scale({% if wScalePerChannel %}oc: u32{% endif %}) -> f32 {
+{% if wScalePerChannel %}
   return w_scale[oc];
 {% else %}
   return w_scale[0];
 {% endif %}
 }
+fn requantize(acc: i32{% if wScalePerChannel %}, oc: u32{% endif %}) -> u32 {
+  let scaled = f32(acc) * x_scale[0] * read_w_scale({% if wScalePerChannel %}oc{% endif %}) / y_scale[0];
   let q = clamp(
     i32(clamp(round(scaled), -2.0e9, 2.0e9)) + i32(y_zero_point[0]),
     0,
   return u32(q);
 }
+@compute @workgroup_size({{ spatialTile }}, {{ spatialTile }}, 1)
 fn main(
   @builtin(workgroup_id) wg: vec3<u32>,
   @builtin(local_invocation_id) lid: vec3<u32>
   }
   workgroupBarrier();
+{% for oc in range(ocTile) %}
   var acc{{ oc }} = 0i;
 {% endfor %}
   var sumX = 0i;
             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(ocTile) %}
             acc{{ oc }} += dot4_packed(wTile[({{ oc }}u * IC_WORDS + channelWord) * 9u + tap], xp);
 {% endfor %}
           }
     return;
   }
+{% for oc in range(ocTile) %}
   {
     let outOc = ocBase + {{ oc }}u;
     if (outOc < OUT_C) {
+      let ca = -read_w_zero({% if 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 wScalePerChannel %}, outOc{% endif %});
     }
   }
 {% endfor %}

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

@@ -1,30 +1,30 @@
 {% 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;
   }
@@ -97,7 +97,7 @@ fn write_y(i: u32, value: i32) {
 }
 @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;

 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: 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;

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

@@ -1,30 +1,30 @@
 {% 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;
   }
@@ -50,7 +50,7 @@ fn write_y(index: u32, value: i32) {
 }
 @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;

 {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
 {% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the output index.
 {% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width (the dispatch caps x and spills into y).
 {% elif note == "element-limit" %}
   // 2D-folded flat element index: gid.y carries the high bits past the
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d(guardInline=true) }}
 {% if hasBias or wScalePerChannel %}
   let oc = (i / params.outImageSize) % params.outChannels;

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

@@ -1,24 +1,36 @@
-{% 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 %}
@@ -28,7 +40,6 @@ fn pack_a(v: vec4<{{ aScalar }}>) -> u32 {
   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 {
@@ -39,7 +50,6 @@ fn sum4_a(v: vec4<{{ aScalar }}>) -> i32 {
 {% 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);
@@ -54,7 +64,6 @@ 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 %}
@@ -64,8 +73,6 @@ fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScala
 {% endif %}
 }
-{% endif %}
-{% if hasAZero %}
 fn read_a_zero({% if aZeroPerRow %}row: u32{% endif %}) -> i32 {
 {% if aZeroPerRow %}
 {% if aUnsigned %}
@@ -82,25 +89,14 @@ fn read_a_zero({% if aZeroPerRow %}row: u32{% endif %}) -> i32 {
 {% 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 %}
@@ -131,31 +127,15 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
   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);
@@ -179,49 +159,13 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
         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;
@@ -229,15 +173,12 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
         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();
   }
@@ -262,9 +203,7 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
         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;
     }
@@ -279,9 +218,7 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
       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);
@@ -296,14 +233,11 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
   }
   // 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) %}
   {
@@ -313,13 +247,9 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
       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];
@@ -327,7 +257,8 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
 {% 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).

+{% set yUnsigned = yUnsigned is defined and yUnsigned %}
+{% set qMin = qMin | default(0 if yUnsigned else 0 - 128) %}
+{% set qMax = qMax | default(255 if yUnsigned else 127) %}
+{% if not env.device.wgslLanguageFeatures.has("packed_4x8_integer_dot_product") %}
+// Runtime source-compatible fallbacks for the packed 4x8 integer built-ins used
+// by quantized kernels. Native implementations remain unshadowed when the
+// language feature is present. The user functions retain the built-ins'
+// @must_use contract, but cannot be @const; every call is runtime.
+// Inputs are function parameters, so replacing a built-in call with one of
+// these same-name functions preserves WGSL's single evaluation of each
+// argument. Byte extraction is unsigned. Signed lanes use the exact
+// i32(byte ^ 0x80u) - 128i sign extension, avoiding signed-left-shift overflow
+// while preserving every 8-bit two's-complement value.
+@must_use fn dot4I8Packed(e1: u32, e2: u32) -> i32 {
+  let shifts = vec4<u32>(0u, 8u, 16u, 24u);
+  let a_bytes = (vec4<u32>(e1) >> shifts) & vec4<u32>(0xffu);
+  let b_bytes = (vec4<u32>(e2) >> shifts) & vec4<u32>(0xffu);
+  let a = vec4<i32>(a_bytes ^ vec4<u32>(0x80u)) - vec4<i32>(128i);
+  let b = vec4<i32>(b_bytes ^ vec4<u32>(0x80u)) - vec4<i32>(128i);
+  return dot(a, b);
+}
 {% 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;
 var<workgroup> a_tile: array<u32, A_TILE_WORDS>;
 var<workgroup> a_row_sum: array<atomic<i32>, TILE_M>;
 // Pack 4 widened A values into one u32 (byte 0 = lowest k).
 fn pack_a(v: vec4<{{ aScalar }}>) -> u32 {
 {% if aUnsigned %}
   return (u32(v.x) & 0xFFu) | ((u32(v.y) & 0xFFu) << 8u) | ((u32(v.z) & 0xFFu) << 16u) | ((u32(v.w) & 0xFFu) << 24u);
 {% endif %}
 }
 // Sum of the 4 bias-shifted A values of one packed word.
 fn sum4_a(v: vec4<{{ aScalar }}>) -> i32 {
 {% 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);
   return dot4I8Packed(a_word, b_word);
 }
 fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> i32 {
 {% if bUnsigned %}
 {% endif %}
 }
 fn read_a_zero({% if aZeroPerRow %}row: u32{% endif %}) -> i32 {
 {% if aZeroPerRow %}
 {% if aUnsigned %}
 {% endif %}
 }
+fn read_b_zero() -> i32 {
 {% if bUnsigned %}
   return i32(b_zero_point[0]);
 {% else %}
   return b_zero_point[0];
 {% endif %}
 }
 {% if requant != "none" %}
 fn read_y_zero() -> i32 {
 {% if yUnsigned %}
   let k_words_total = params.K / 4u;
 {% endif %}
   if (tid < TILE_M) {
     atomicStore(&a_row_sum[tid], 0);
   }
   workgroupBarrier();
 {% for m in range(tileM) %}
   var r{{ m }}: i32 = 0;
 {% endfor %}
   var sum_b: i32 = 0;
   for (var kw0 = 0u; kw0 < k_words_total; kw0 = kw0 + KT_WORDS) {
     let words = min(KT_WORDS, k_words_total - kw0);
         let v = a[a_base4 + src_m * k_words_total + kw0 + ki];
 {% endif %}
         packed = pack_a(v);
         atomicAdd(&a_row_sum[mi], sum4_a(v));
       }
       a_tile[mi * KT_WORDS + ki] = packed;
     }
     workgroupBarrier();
     if (n_valid) {
       for (var kk = 0u; kk < words; kk = kk + 1u) {
         let k4 = (kw0 + kk) * 4u;
         let b_row = b_base + k4 * params.N + col;
         let b1 = b[b_row + params.N];
         let b2 = b[b_row + 2u * params.N];
         let b3 = b[b_row + 3u * params.N];
         sum_b = sum_b + sum4_b(b0, b1, b2, b3);
         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 %}
       }
     }
     workgroupBarrier();
   }
         let a3: {{ aScalar }} = {{ "128u" if aUnsigned else "0" }};
         let v = vec4<{{ aScalar }}>(a0, a1, a2, a3);
         packed = pack_a(v);
         atomicAdd(&a_row_sum[tid], sum4_a(v));
       }
       a_tile[tid * KT_WORDS] = packed;
     }
       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" }};
       sum_b = sum_b + sum4_b(b0, b1, b2, b3);
       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);
   }
   // Fold bias + zero-point corrections (exact in modular i32 arithmetic).
+  let cb: i32 = {{ "128" if bUnsigned else "0" }} - read_b_zero();
 {% if not aZeroPerRow %}
+  let ca: i32 = {{ "128" if aUnsigned else "0" }} - read_a_zero();
   let kcc: i32 = i32(params.K) * ca * cb;
 {% endif %}
 {% for m in range(tileM) %}
   {
       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 }};
       acc = acc + cb * atomicLoad(&a_row_sum[{{ m }}u]);
+      acc = acc + ca * sum_b;
       acc = acc + kcc;
 {% if hasBias %}
       acc = acc + bias[mr];
 {% if requant == "none" %}
       y[y_base + mr * params.N + col] = acc;
 {% else %}
+{% set scale_index = "col" if requant == "per_column" else ("mr" if requant == "per_row" else "0u") %}
+      let scaled = f32(acc) * f32(a_scale[0]) * f32(b_scale[{{ scale_index }}]) / 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).

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

@@ -16,9 +16,8 @@ fn read_a_zero({% if aZeroPerRow %}row: u32{% endif %}) -> i32 {
 @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) {

 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(
   @builtin(global_invocation_id) gid: vec3<u32>,
 ) {
+  let flat = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   let a_count = params.M * params.KPadded;
   if (flat < a_count) {

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "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],
@@ -83,7 +82,7 @@
         "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] } },
@@ -94,7 +93,10 @@
         "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",
@@ -881,7 +883,7 @@
       "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": {
@@ -1229,15 +1231,18 @@
         "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",
@@ -1327,7 +1332,7 @@
       "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": {
@@ -1579,7 +1584,7 @@
       "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": {
@@ -1597,6 +1602,94 @@
         "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
       },
       "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 2, 2], "tolerance": 0 } }
     }
   ]
 }

 {
   "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],
         "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] } },
         "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
       },
       "outputs": { "y": { "dtype": "int8", "shape": [1, 1, 2, 2] } },
+      "attrs": {},
+      "provenance": {
+        "notes": "Signed output storage with a nonzero output zero point. The input is not a ramp, so the 2x2 kernel gives each output position its own response."
+      }
     },
     {
       "name": "int8_input_uint8_output_pairwise_route",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/nn/qlinearconv_op_test.cc",
         "test": "QLinearConvTest.Conv3D_S8S8",
+        "notes": "A compact three-dimensional QLinearConv uses int8 input, weights, and output."
       },
       "attrs": { "strides": [1, 1, 1] },
       "inputs": {
         "w": {
           "dtype": "int8",
           "shape": [64, 16, 3, 3],
+          "data": { "kind": "cycle", "values": [2, -1, 3, -2, 1, 4, -3, 2, 0, -4] }
         },
         "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.002] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [128] } }
       },
       "outputs": { "y": { "dtype": "uint8", "shape": [1, 64, 22, 22] } },
+      "attrs": {},
+      "provenance": {
+        "notes": "The dot4 im2col route over 16 input channels. A ten-value weight cycle does not tile the nine taps of a 3x3 kernel, so every output channel carries its own filter, and the output scale resolves the accumulator instead of quantizing every position onto the zero point."
+      }
     },
     {
       "name": "im2col_dp4a_requant_huge_finite_saturates_3x3",
       "name": "conv3d_depth_dilation2_compact",
       "provenance": {
         "source": "ONNX QLinearConv-10 volumetric dilation semantics",
+        "notes": "A three-dimensional quantized convolution uses a nondefault depth dilation."
       },
       "attrs": { "strides": [1, 1, 1], "dilations": [2, 1, 1] },
       "inputs": {
       "name": "same_lower_stride2_autopad",
       "provenance": {
         "source": "ONNX Runtime QLinearConv-10 CPUExecutionProvider",
+        "notes": "ONNX auto_pad SAME_LOWER places the odd padding element at the start of the axis. The nonzero leading pad verifies that execution uses the derived padding rather than unrelated explicit pad attributes."
       },
       "attrs": { "auto_pad": "SAME_LOWER", "strides": [2, 2] },
       "inputs": {
         "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
       },
       "outputs": { "y": { "dtype": "uint8", "shape": [1, 1, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ncw1d_scalar_bias",
+      "provenance": {
+        "source": "ONNX Operators.md, QLinearConv",
+        "notes": "A one-dimensional quantized convolution carrying a bias, which is the ordinary audio form. The bias is non-zero so the case proves it is applied, not merely that a route exists."
+      },
+      "attrs": { "kernel_shape": [2], "pads": [0, 0] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 4], "data": { "kind": "values", "values": [1, 2, 3, 4] } },
+        "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": [1, 1, 2], "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": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "B": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [13, 15, 17] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ncdhw3d_scalar_bias",
+      "provenance": {
+        "source": "ONNX Operators.md, QLinearConv",
+        "notes": "A three-dimensional quantized convolution includes a nonzero bias."
+      },
+      "attrs": { "kernel_shape": [2, 2, 2], "pads": [0, 0, 0, 0, 0, 0] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8] }
+        },
+        "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": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1, 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": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "B": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [40] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ncw1d_per_channel_weight_quantization",
+      "provenance": {
+        "source": "ONNX Operators.md, QLinearConv",
+        "notes": "One-dimensional convolution uses weights quantized independently per output channel."
+      },
+      "attrs": { "kernel_shape": [2], "pads": [0, 0] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 4], "data": { "kind": "values", "values": [1, 2, 3, 4] } },
+        "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, 2], "data": { "kind": "values", "values": [1, 1, 2, -1] } },
+        "w_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 1.0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [0, 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, 2, 3],
+          "data": { "kind": "values", "values": [3, 5, 7, 0, 1, 2] },
+          "tolerance": 0
+        }
+      }
     }
   ]
 }