sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +79 -0
build/webgpu/bench.json +172 -0
build/webgpu/conv-int-accumulate-spatial.wgsl.jinja +127 -0
build/webgpu/conv-int-im2col-spatial.wgsl.jinja +77 -0
build/webgpu/manifest.json +1221 -0
build/webgpu/metadata.json +20 -0
build/webgpu/quant-dp4a-matmul.wgsl.jinja +336 -0
build/webgpu/test.json +1282 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.ConvInteger
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 10
+## Description
+Performs integer convolution on quantized inputs `x` and filter `w`, each with an optional zero point, producing an `int32` output. Zero-point subtraction is applied before accumulation; the result must not overflow 32 bits during accumulation.
+See the [ONNX `ConvInteger` spec](https://onnx.ai/onnx/operators/onnx__ConvInteger.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `x` | `x` | `TX` | — | — | Input data tensor of shape `(N x C x D1 x ... x Dn)`. | required |
+| `w` | `w` | `TW` | — | — | Convolution weight tensor of shape `(M x C/group x k1 x ... x kn)`. | required |
+| `x_zero_point` | `x_zero_point` | `TX` | — | — | Optional scalar zero point for `x`; defaults to 0. | optional |
+| `w_zero_point` | `w_zero_point` | `TW` | — | — | Optional scalar or per-output-channel zero point for `w`; defaults to 0. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `y` | `y` | `TY` | same as `x` | derived; see description | Output tensor containing `int32` convolution results. | 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` | `int32` |
+## 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)
+- [`quant-dp4a-matmul.wgsl.jinja`](build/webgpu/quant-dp4a-matmul.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.ConvInteger", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [1, 1, 2, 1, 1] },
+  w: { data: wData, shape: [1, 1, 1, 1, 1] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,172 @@

+{
+  "op": "ai.onnx.ConvInteger",
+  "cases": [
+    {
+      "name": "u8_nchw_1x16x32x32",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 16, 32, 32], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "uint8", "shape": [16, 16, 3, 3], "data": { "kind": "constant", "value": 129 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 16, 30, 30] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "u8s8_pointwise_1x64x56x56_oc128",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 56, 56], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [128, 64, 1, 1], "data": { "kind": "constant", "value": -3 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 56, 56] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "u8s8_pointwise_unalignedC_c66_oc128_56x56_scalar",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 66, 56, 56], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [128, 66, 1, 1], "data": { "kind": "constant", "value": -3 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "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)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "u8s8_pointwise_rgb_stem_c3_oc32_112x112_scalar",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 3, 112, 112], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [32, 3, 1, 1], "data": { "kind": "constant", "value": -3 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 32, 112, 112] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "u8u8_depthwise_c128_3x3_group128_56x56_scalar",
+      "preset": "smoke",
+      "attrs": { "group": 128 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 128, 56, 56], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "uint8", "shape": [128, 1, 3, 3], "data": { "kind": "constant", "value": 129 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 54, 54] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      }
+    },
+    {
+      "name": "u8s8_pointwise_gemv_alignedN_n2048_c256_dp4a",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 256, 1, 1], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [2048, 256, 1, 1], "data": { "kind": "constant", "value": -3 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 2048, 1, 1] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "u8s8_pointwise_gemv_unalignedN_n2050_c256_dp4a",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 256, 1, 1], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [2050, 256, 1, 1], "data": { "kind": "constant", "value": -3 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 2050, 1, 1] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "u8s8_nchw_oc128_c64_3x3_32x32",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 32, 32], "dist": "randint", "seed": 21, "min": 0, "max": 255 },
+        "w": { "dtype": "int8", "shape": [128, 64, 3, 3], "dist": "randint", "seed": 22, "min": -127, "max": 127 },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 30, 30] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 128 * 30 * 30 * 64 * 3 * 3" }] },
+      "attrs": {}
+    },
+    {
+      "name": "s8s8_im2col_small_output_numel_scalar_fallback",
+      "preset": "stress",
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4, 64, 8, 8], "dist": "randint", "seed": 42, "min": -100, "max": 100 },
+        "w": { "dtype": "int8", "shape": [64, 64, 3, 3], "dist": "randint", "seed": 43, "min": -64, "max": 64 },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [4, 64, 6, 6], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" }
+        ]
+      },
+      "attrs": {}
+    },
+    {
+      "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": 7601, "min": 0, "max": 255 },
+        "w": { "dtype": "int8", "shape": [32, 16, 3, 3, 3], "dist": "randint", "seed": 7602, "min": -127, "max": 127 },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "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)"
+          }
+        ]
+      }
+    }
+  ]
+}

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

	@@ -0,0 +1,127 @@

+{% 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 source.xZeroOmitted is defined and source.xZeroOmitted %}
+  return 0;
+  {% elif 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 source.wZeroOmitted is defined and source.wZeroOmitted %}
+  return 0;
+  {% elif 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,77 @@

+/* 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;
+  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,1221 @@

+{
+  "domain": "ai.onnx",
+  "name": "ConvInteger",
+  "sinceVersion": 10,
+  "description": "Performs integer convolution on quantized inputs `x` and filter `w`, each with an optional zero point, producing an `int32` output. Zero-point subtraction is applied before accumulation; the result must not overflow 32 bits during accumulation.",
+  "inputs": [
+    { "role": "x", "dtype": "TX", "description": "Input data tensor of shape `(N x C x D1 x ... x Dn)`." },
+    { "role": "w", "dtype": "TW", "description": "Convolution weight tensor of shape `(M x C/group x k1 x ... x kn)`." },
+    {
+      "role": "x_zero_point",
+      "dtype": "TX",
+      "description": "Optional scalar zero point for `x`; defaults to 0.",
+      "optional": true
+    },
+    {
+      "role": "w_zero_point",
+      "dtype": "TW",
+      "description": "Optional scalar or per-output-channel zero point for `w`; defaults to 0.",
+      "optional": true
+    }
+  ],
+  "outputs": [
+    {
+      "role": "y",
+      "dtype": "TY",
+      "description": "Output tensor containing `int32` convolution results.",
+      "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": ["int32"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "x", "role": "input" },
+    "w": { "kind": "tensor", "semantic": "w", "role": "input" },
+    "x_zero_point": { "kind": "tensor", "semantic": "x_zero_point", "role": "input", "required": false },
+    "w_zero_point": { "kind": "tensor", "semantic": "w_zero_point", "role": "input", "required": false },
+    "y": { "kind": "tensor", "semantic": "y", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "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",
+    "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\")",
+    "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)",
+    "scalarXZeroOk": "ranks.x_zero_point == 0 or (ranks.x_zero_point == 1 and dim(shapes.x_zero_point, 0) == 1)",
+    "scalarWZeroOk": "ranks.w_zero_point == 0 or (ranks.w_zero_point == 1 and dim(shapes.w_zero_point, 0) == 1)",
+    "scalarXZeroRequiredOk": "present.x_zero_point and scalarXZeroOk",
+    "scalarWZeroRequiredOk": "present.w_zero_point and scalarWZeroOk",
+    "outputBatchChannelsOk": "dim(shapes.y, 0) == dim(shapes.x, 0) and dim(shapes.y, 1) == dim(shapes.w, 0)",
+    "ungroupedChannelsOk": "attrs.group == 1 and dim(shapes.w, 1) == dim(shapes.x, 1)",
+    "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",
+    "output3dShapeOk": "dim(shapes.y, 2) == expectedOutputDepth and dim(shapes.y, 3) == expectedOutputHeight and dim(shapes.y, 4) == expectedOutputWidth",
+    "rank3TensorOk": "spatialAttributesOk and (ranks.x == 3 and ranks.w == 3 and ranks.y == 3)",
+    "nchwTensorOk": "rank4TensorOk or rank3TensorOk",
+    "output1dShapeOk": "dim(shapes.y, 2) == expectedOutputWidth",
+    "outputSpatialShapeOk": "output2dShapeOk if rank4TensorOk else output1dShapeOk",
+    "zeroPointsOmitted": "not present.x_zero_point and not present.w_zero_point"
+  },
+  "bindingSets": {
+    "nchw2d_accumulate": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "inputHeight" },
+            { "name": "inW", "type": "u32", "value": "inputWidth" },
+            { "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": "kernelHeight" },
+            { "name": "kernelW", "type": "u32", "value": "kernelWidth" },
+            { "name": "outH", "type": "u32", "value": "outputHeight" },
+            { "name": "outW", "type": "u32", "value": "outputWidth" },
+            { "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)" }
+          ]
+        }
+      }
+    ],
+    "nchw2dAccumulateScalarWZero": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "inputHeight" },
+            { "name": "inW", "type": "u32", "value": "inputWidth" },
+            { "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": "kernelHeight" },
+            { "name": "kernelW", "type": "u32", "value": "kernelWidth" },
+            { "name": "outH", "type": "u32", "value": "outputHeight" },
+            { "name": "outW", "type": "u32", "value": "outputWidth" },
+            { "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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",
+        "length": 1
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "N", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4)" },
+            {
+              "name": "K",
+              "type": "u32",
+              "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3) * dim(shapes.w, 4)"
+            },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            {
+              "name": "bBatchStride",
+              "type": "u32",
+              "value": "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)"
+            },
+            {
+              "name": "yBatchStride",
+              "type": "u32",
+              "value": "dim(shapes.w, 0) * dim(shapes.y, 2) * dim(shapes.y, 3) * 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)" }
+          ]
+        }
+      }
+    ],
+    "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",
+        "length": 1
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "N", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
+            { "name": "K", "type": "u32", "value": "dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            {
+              "name": "bBatchStride",
+              "type": "u32",
+              "value": "(dim(shapes.w, 1) * dim(shapes.w, 2) * dim(shapes.w, 3)) * (dim(shapes.y, 2) * dim(shapes.y, 3))"
+            },
+            { "name": "yBatchStride", "type": "u32", "value": "dim(shapes.w, 0) * dim(shapes.y, 2) * dim(shapes.y, 3)" }
+          ]
+        }
+      }
+    ],
+    "dp4aPointwiseTail": [
+      { "name": "a", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$aScalar" },
+      { "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",
+        "length": 1
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "N", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
+            { "name": "K", "type": "u32", "value": "dim(shapes.x, 1)" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3)" },
+            { "name": "yBatchStride", "type": "u32", "value": "dim(shapes.w, 0) * dim(shapes.y, 2) * dim(shapes.y, 3)" }
+          ]
+        }
+      }
+    ],
+    "dp4aPointwise": [
+      { "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",
+        "length": 1
+      },
+      {
+        "name": "b_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "N", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
+            { "name": "K", "type": "u32", "value": "dim(shapes.x, 1)" },
+            { "name": "aBatchStride4", "type": "u32", "value": 0 },
+            { "name": "bBatchStride", "type": "u32", "value": "dim(shapes.x, 1) * dim(shapes.x, 2) * dim(shapes.x, 3)" },
+            { "name": "yBatchStride", "type": "u32", "value": "dim(shapes.w, 0) * dim(shapes.y, 2) * dim(shapes.y, 3)" }
+          ]
+        }
+      }
+    ],
+    "ncdhw3dAccumulate": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      {
+        "name": "w_zero_point",
+        "arg": "w_zero_point",
+        "semantic": "w_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$wScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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)" }
+          ]
+        }
+      }
+    ],
+    "nchwAccumulatePerChannelWZeroOnly": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "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", "arg": "y", "semantic": "y", "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": "inputHeight" },
+            { "name": "inW", "type": "u32", "value": "inputWidth" },
+            { "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": "kernelHeight" },
+            { "name": "kernelW", "type": "u32", "value": "kernelWidth" },
+            { "name": "outH", "type": "u32", "value": "outputHeight" },
+            { "name": "outW", "type": "u32", "value": "outputWidth" },
+            { "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)" }
+          ]
+        }
+      }
+    ],
+    "nchwAccumulateXZeroOnly": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "inputHeight" },
+            { "name": "inW", "type": "u32", "value": "inputWidth" },
+            { "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": "kernelHeight" },
+            { "name": "kernelW", "type": "u32", "value": "kernelWidth" },
+            { "name": "outH", "type": "u32", "value": "outputHeight" },
+            { "name": "outW", "type": "u32", "value": "outputWidth" },
+            { "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)" }
+          ]
+        }
+      }
+    ],
+    "nchwAccumulateWZeroOnly": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "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",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "inputHeight" },
+            { "name": "inW", "type": "u32", "value": "inputWidth" },
+            { "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": "kernelHeight" },
+            { "name": "kernelW", "type": "u32", "value": "kernelWidth" },
+            { "name": "outH", "type": "u32", "value": "outputHeight" },
+            { "name": "outW", "type": "u32", "value": "outputWidth" },
+            { "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)" }
+          ]
+        }
+      }
+    ],
+    "nchwAccumulateNoZeroPoints": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      { "name": "y", "arg": "y", "semantic": "y", "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": "inputHeight" },
+            { "name": "inW", "type": "u32", "value": "inputWidth" },
+            { "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": "kernelHeight" },
+            { "name": "kernelW", "type": "u32", "value": "kernelWidth" },
+            { "name": "outH", "type": "u32", "value": "outputHeight" },
+            { "name": "outW", "type": "u32", "value": "outputWidth" },
+            { "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)" }
+          ]
+        }
+      }
+    ],
+    "ncdhw3dAccumulateXZeroOnly": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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)" }
+          ]
+        }
+      }
+    ],
+    "ncdhw3dAccumulateWZeroOnly": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "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",
+        "length": 1
+      },
+      { "name": "y", "arg": "y", "semantic": "y", "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)" }
+          ]
+        }
+      }
+    ],
+    "ncdhw3dAccumulateNoZeroPoints": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      { "name": "w", "arg": "w", "semantic": "w", "buffer": { "type": "read-only-storage" }, "elementType": "$wScalar" },
+      { "name": "y", "arg": "y", "semantic": "y", "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)" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "dp4a_im2col_ncdhw3d_scalar_zero_points",
+      "priority": 20,
+      "when": ["packedFeature", "rank5TensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "ungroupedChannelsOk", "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", "outputBatchChannelsOk", "output3dShapeOk", "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) * 4 <= device.limits.maxStorageBufferBindingSize", "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>\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "aPacked": true,
+        "aZeroPerRow": false,
+        "bZeroPerColumn": false,
+        "requant": "\"none\"",
+        "scaleScalar": "\"f32\"",
+        "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": "ConvInteger.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": "ConvInteger.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_scalar_zero_points",
+      "priority": 20,
+      "when": ["packedFeature", "rank4TensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "ungroupedChannelsOk", "(dim(shapes.w, 2) > 1 or dim(shapes.w, 3) > 1)", "outputBatchChannelsOk", "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", "numel(shapes.y) >= 8192", "dim(shapes.y, 2) * dim(shapes.y, 3) <= 4194240", "dim(shapes.w, 0) <= 524280", "dim(shapes.y, 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>\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "bZeroPerColumn": false,
+        "requant": "\"none\"",
+        "scaleScalar": "\"f32\"",
+        "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": "ConvInteger.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": "ConvInteger.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": "dp4a_pointwise_1x1_tail",
+      "priority": 21,
+      "when": ["packedFeature", "rank4TensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "ungroupedChannelsOk", "dim(shapes.w, 2) == 1", "dim(shapes.w, 3) == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputBatchChannelsOk", "dim(shapes.y, 2) == dim(shapes.x, 2)", "dim(shapes.y, 3) == dim(shapes.x, 3)", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "dim(shapes.y, 2) * dim(shapes.y, 3) <= 4194240", "dim(shapes.w, 0) <= 524280", "dim(shapes.y, 0) <= device.limits.maxComputeWorkgroupsPerDimension"],
+      "constants": {
+        "aScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "bScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "aPacked": false,
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "bZeroPerColumn": false,
+        "requant": "\"none\"",
+        "scaleScalar": "\"f32\"",
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwiseTail",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aPointwiseTail",
+          "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_pointwise_1x1",
+      "priority": 20,
+      "when": ["packedFeature", "rank4TensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "ungroupedChannelsOk", "dim(shapes.w, 2) == 1", "dim(shapes.w, 3) == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputBatchChannelsOk", "dim(shapes.y, 2) == dim(shapes.x, 2)", "dim(shapes.y, 3) == dim(shapes.x, 3)", "dim(shapes.x, 1) % 4 == 0", "dim(shapes.y, 2) * dim(shapes.y, 3) <= 4194240", "dim(shapes.w, 0) <= 524280", "dim(shapes.y, 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>\"",
+        "aUnsigned": "tensorDtypes.w == \"uint8\"",
+        "bUnsigned": "tensorDtypes.x == \"uint8\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "bZeroPerColumn": false,
+        "requant": "\"none\"",
+        "scaleScalar": "\"f32\"",
+        "tileM": 8,
+        "tileN": 64,
+        "ktWords": 32
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwise",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": "dp4aPointwise",
+          "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_zero_point",
+      "priority": 5,
+      "when": ["nchwTensorOk", "scalarXZeroRequiredOk", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "groupChannelsOk", "outputBatchChannelsOk", "outputSpatialShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "wZeroPerChannel": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvIntegerPerChannel",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 2, "xZeroOmitted": "false", "wZeroOmitted": "false" }
+          },
+          "bindings": "nchw2d_accumulate",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_per_channel_w_zero_point_only",
+      "priority": 5,
+      "when": ["nchwTensorOk", "not present.x_zero_point", "present.w_zero_point", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "groupChannelsOk", "outputBatchChannelsOk", "outputSpatialShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\"",
+        "wZeroPerChannel": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvIntegerPerChannel",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 2, "xZeroOmitted": "true", "wZeroOmitted": "false" }
+          },
+          "bindings": "nchwAccumulatePerChannelWZeroOnly",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_scalar_zero_points",
+      "when": ["nchwTensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "groupChannelsOk", "outputBatchChannelsOk", "outputSpatialShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 2, "xZeroOmitted": "false", "wZeroOmitted": "false" }
+          },
+          "bindings": "nchw2dAccumulateScalarWZero",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_x_zero_point_only",
+      "when": ["nchwTensorOk", "present.x_zero_point", "not present.w_zero_point", "scalarXZeroRequiredOk", "groupChannelsOk", "outputBatchChannelsOk", "outputSpatialShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 2, "xZeroOmitted": "false", "wZeroOmitted": "true" }
+          },
+          "bindings": "nchwAccumulateXZeroOnly",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw2d_w_zero_point_only",
+      "when": ["nchwTensorOk", "not present.x_zero_point", "present.w_zero_point", "scalarWZeroRequiredOk", "groupChannelsOk", "outputBatchChannelsOk", "outputSpatialShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 2, "xZeroOmitted": "true", "wZeroOmitted": "false" }
+          },
+          "bindings": "nchwAccumulateWZeroOnly",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d_scalar_zero_points",
+      "when": ["rank5TensorOk", "present.x_zero_point", "present.w_zero_point", "scalarXZeroOk", "scalarWZeroOk", "groupChannelsOk", "outputBatchChannelsOk", "output3dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger3d",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 3, "xZeroOmitted": "false", "wZeroOmitted": "false" }
+          },
+          "bindings": "ncdhw3dAccumulate",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d_x_zero_point_only",
+      "when": ["rank5TensorOk", "present.x_zero_point", "not present.w_zero_point", "scalarXZeroOk", "groupChannelsOk", "outputBatchChannelsOk", "output3dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger3d",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 3, "xZeroOmitted": "false", "wZeroOmitted": "true" }
+          },
+          "bindings": "ncdhw3dAccumulateXZeroOnly",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d_w_zero_point_only",
+      "when": ["rank5TensorOk", "not present.x_zero_point", "present.w_zero_point", "scalarWZeroOk", "groupChannelsOk", "outputBatchChannelsOk", "output3dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger3d",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 3, "xZeroOmitted": "true", "wZeroOmitted": "false" }
+          },
+          "bindings": "ncdhw3dAccumulateWZeroOnly",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "ncdhw3d_implicit_zero_points",
+      "when": ["rank5TensorOk", "not present.x_zero_point", "not present.w_zero_point", "groupChannelsOk", "outputBatchChannelsOk", "output3dShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger3d",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 3, "xZeroOmitted": "true", "wZeroOmitted": "true" }
+          },
+          "bindings": "ncdhw3dAccumulateNoZeroPoints",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "nchw_implicit_zero_points",
+      "when": ["nchwTensorOk", "zeroPointsOmitted", "groupChannelsOk", "outputBatchChannelsOk", "outputSpatialShapeOk"],
+      "constants": {
+        "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+        "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+        "xUnsigned": "tensorDtypes.x == \"uint8\"",
+        "wUnsigned": "tensorDtypes.w == \"uint8\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger",
+          "source": {
+            "shader": "conv-int-accumulate-spatial.wgsl.jinja",
+            "inputs": { "spatialRank": 2, "xZeroOmitted": "true", "wZeroOmitted": "true" }
+          },
+          "bindings": "nchwAccumulateNoZeroPoints",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,20 @@

+{
+  "name": "ai.onnx.ConvInteger",
+  "id": "_ai_onnx_convinteger_webgpu_ac7ae27",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "FxmmRX4sXmysOymGqufcK/SEHqonPxx0mjgpdFtthw4=",
+      "conv-int-accumulate-spatial.wgsl.jinja": "n0fxsi3cnW+RmTawiPp6t8zgBzS/Sg+eKHARguziK7Q=",
+      "conv-int-im2col-spatial.wgsl.jinja": "yeNh5h+P9uP5JNa11SdeEd811/c9e6bJ1w/AqkzC0ck=",
+      "manifest.json": "GyuPs+3QC5Ok4y/pJ/R+ASsnICvEpfC6cbpOYA+Ppcs=",
+      "quant-dp4a-matmul.wgsl.jinja": "YO5g0dZL7JdLEzF04XK1cEtKp4P3i/IRnRUdFeuVsiw=",
+      "test.json": "5CyT/c0HcXkyJOl/Xem2F58Yi86+X6ZTltfq7ZUFBek="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.ConvInteger" }
+}

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

	@@ -0,0 +1,336 @@

+{% 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 {
+  return y_zero_point[0];
+}
+{% 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 }});
+      y[y_base + mr * params.N + col] = q;
+{% endif %}
+    }
+  }
+{% endfor %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1282 @@

+{
+  "op": "ai.onnx.ConvInteger",
+  "fixtureArrays": {
+    "ort_stride_padding_input_x": [10, 11, 12, 13, 14, 15, 16, 20, 21, 22, 23, 24, 25, 26, 30, 31, 32, 33, 34, 35, 36, 40, 41, 42, 43, 44, 45, 46, 50, 51, 52, 53, 54, 55, 56, 60, 61, 62, 63, 64, 65, 66, 70, 71, 72, 73, 74, 75, 76],
+    "ort_with_group_2d_u8u8_input_x": [2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28],
+    "ort_padded_3d_u8u8_output_y": [1, 3, 5, 3, 5, 12, 16, 9, 11, 24, 28, 15, 7, 15, 17, 9, 11, 24, 28, 15, 28, 60, 68, 36, 40, 84, 92, 48, 23, 48, 52, 27, 29, 60, 64, 33, 64, 132, 140, 72, 76, 156, 164, 84, 41, 84, 88, 45, 19, 39, 41, 21, 41, 84, 88, 45, 47, 96, 100, 51, 25, 51, 53, 27],
+    "dp4a_pointwise_u8s8_c8_batched_input_x": [0, 255, 128, 127, 56, 77, 3, 100, 200, 42, 5, 19, 64, 88, 33, 7, 125, 90, 211, 16],
+    "unaligned_kernel_small_numel_scalar_path_input_w": [11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30],
+    "same_upper_stride2_autopad_input_x": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_nchw2d_scalar_1x1",
+      "attrs": { "kernel_shape": [1, 1], "strides": [1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 64, 64], "data": { "kind": "cycle", "values": [1, 2, 3, 4, 5] } },
+        "w": { "dtype": "uint8", "shape": [4096, 1, 1, 1], "data": { "kind": "cycle", "values": [2, 1, 3] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 0 } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 0 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 4096, 64, 64], "tolerance": 0 } }
+    },
+    {
+      "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] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [127, 128, 129, 130] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4] } }
+    },
+    {
+      "name": "int8_grouped",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 2, 3, 3],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6, 7, 8, 10, 11, 12, 13, 14, 15, 16, 17, 18] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [2, 1, 2, 2],
+          "data": { "kind": "values", "values": [1, 0, 0, 1, -1, 1, 1, -1] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 2, 2, 2] } }
+    },
+    {
+      "name": "ort_without_padding_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithoutPadding_2D_u8u8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [2, 2, 2, 2] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 2, 2], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_with_padding_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_2D_u8u8"
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [2, 2, 2, 2] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_pointwise_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.Pointwise_2D_u8u8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 3, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_without_padding_u8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithoutPadding_2D_u8s8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [-9, -9, -9, -9] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 2, 2], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_with_padding_u8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_2D_u8s8"
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [-9, -9, -9, -9] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_without_padding_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithoutPadding_2D_s8s8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-1, 2, -3, 4, -5, 6, -7, 8, -9] }
+        },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, -2, 3, -4] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 2, 2], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_without_padding_s8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithoutPadding_2D_s8u8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-1, 2, -3, 4, -5, 6, -7, 8, -9] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 2, 2], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "ort_with_padding_s8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_2D_s8u8"
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-1, 2, -3, 4, -5, 6, -7, 8, -9] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_stride2_padding_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride2_2D_u8u8"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [11, 12, 11, 12, 13, 12, 11, 12, 11] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_with_group_2d_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithGroup_2D_u8u8"
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [3, 1, 2, 2],
+          "data": { "kind": "values", "values": [11, 12, 12, 11, 13, 14, 14, 13, 15, 16, 16, 15] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_with_group_2d_u8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithGroup_2D_u8s8"
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 1, 2, 2],
+          "data": { "kind": "values", "values": [-9, -8, -8, -9, -7, -6, -6, -7, -5, -4, -4, -5] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_stride2_padding_u8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride2_2D_u8s8"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-9, -8, -9, -8, -7, -8, -9, -8, -9] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_with_padding_2d_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_2D_s8s8"
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-1, 2, -3, 4, -5, 6, -7, 8, -9] }
+        },
+        "w": { "dtype": "int8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, -2, 3, -4] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_with_group_2d_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithGroup_2D_s8s8"
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [3, 1, 2, 2],
+          "data": { "kind": "values", "values": [-9, -8, -8, -9, -7, -6, -6, -7, -5, -4, -4, -5] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_with_group_2d_s8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithGroup_2D_s8u8"
+      },
+      "attrs": { "group": 3, "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [3, 1, 2, 2],
+          "data": { "kind": "values", "values": [11, 12, 12, 11, 13, 14, 14, 13, 15, 16, 16, 15] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_stride2_padding_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride2_2D_s8s8"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-9, -8, -9, -8, -7, -8, -9, -8, -9] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_stride2_padding_s8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride2_2D_s8u8"
+      },
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [11, 12, 11, 12, 13, 12, 11, 12, 11] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_no_x_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.NoXZeroPoint",
+        "notes": "Exercises the independent optional x_zero_point input; its omitted value defaults to zero while w_zero_point remains nonzero."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [2, 2, 2, 2] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [16, 20, 28, 32] }
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_no_w_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.NoWZeroPoint",
+        "notes": "Exercises the independent optional w_zero_point input; its omitted value defaults to zero while x_zero_point remains nonzero."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [2, 2, 2, 2] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [24, 32, 48, 56] }
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "ort_stride3_asymmetric_padding_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride3_2D_u8u8"
+      },
+      "attrs": { "strides": [3, 3], "pads": [2, 2, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [11, 12, 11, 12, 13, 12, 11, 12, 11] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 3, 3] } }
+    },
+    {
+      "name": "ort_stride3_asymmetric_padding_u8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride3_2D_u8s8"
+      },
+      "attrs": { "strides": [3, 3], "pads": [2, 2, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-9, -8, -9, -8, -7, -8, -9, -8, -9] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 3, 3] } }
+    },
+    {
+      "name": "ort_stride3_asymmetric_padding_s8s8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride3_2D_s8s8"
+      },
+      "attrs": { "strides": [3, 3], "pads": [2, 2, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [-9, -8, -9, -8, -7, -8, -9, -8, -9] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 3, 3] } }
+    },
+    {
+      "name": "ort_stride3_asymmetric_padding_s8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithStride3_2D_s8u8"
+      },
+      "attrs": { "strides": [3, 3], "pads": [2, 2, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 7, 7],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_stride_padding_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [11, 12, 11, 12, 13, 12, 11, 12, 11] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 3, 3] } }
+    },
+    {
+      "name": "ort_padded_3d_u8u8",
+      "attrs": { "pads": [1, 1, 1, 1, 1, 1] },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_3D_u8u8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [11, 11, 11, 11, 11, 11, 11, 11] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_padded_3d_u8u8_output_y" } },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_padded_3d_u8s8",
+      "attrs": { "pads": [1, 1, 1, 1, 1, 1] },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_3D_u8s8",
+        "notes": "ORT's mixed uint8/int8 3D padding case."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [-9, -9, -9, -9, -9, -9, -9, -9] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_padded_3d_u8u8_output_y" } },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_padded_3d_s8s8",
+      "attrs": { "pads": [1, 1, 1, 1, 1, 1] },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.WithPadding_3D_s8s8",
+        "notes": "ORT's signed 3D padding case."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [-9, -9, -9, -9, -9, -9, -9, -9] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 4, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_padded_3d_u8u8_output_y" } },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_pointwise_3d_u8u8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/nn/conv_integer_test.cc",
+        "test": "ConvIntegerTest.Pointwise_3D_u8u8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3, 3],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_with_group_2d_u8u8_input_x" } }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 3, 3, 3],
+          "data": {
+            "kind": "values",
+            "values": [4, 8, 12, 16, 20, 24, 28, 32, 36, 40, 44, 48, 52, 56, 60, 64, 68, 72, 76, 80, 84, 88, 92, 96, 100, 104, 108]
+          },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "onnx_backend_convinteger_with_padding",
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [2, 1, 2, 2],
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1, 1, 1] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 2, 4, 4] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convinteger_with_padding" }
+    },
+    {
+      "name": "onnx_backend_convinteger_without_padding",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [2, 3, 4, 5, 6, 7, 8, 9, 10] }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, 1, 1, 1] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 2, 2] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_convinteger_without_padding" },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_u8s8_c34_tail",
+      "provenance": {
+        "notes": "Pins the widened-storage DP4A tail route for a realistic pointwise projection whose channel count is not divisible by four."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 34, 2, 3],
+          "data": { "kind": "cycle", "values": [0, 255, 128, 127, 56, 77, 3, 100, 200, 42, 5, 19, 64] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [5, 34, 1, 1],
+          "data": { "kind": "cycle", "values": [-77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [113] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 5, 2, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_u8s8_c8_batched",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 8, 3, 5],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_u8s8_c8_batched_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [6, 8, 1, 1],
+          "data": {
+            "kind": "cycle",
+            "values": [-77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90, -128, 127, 0, -1, 56]
+          }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [113] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 6, 3, 5], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "dp4a_pointwise_s8u8_c12_tails",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 12, 7, 11],
+          "data": {
+            "kind": "cycle",
+            "values": [-128, 127, 0, -1, 56, -77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90]
+          }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [9, 12, 1, 1],
+          "data": {
+            "kind": "cycle",
+            "values": [19, 64, 88, 33, 7, 125, 90, 128, 127, 0, 1, 56, 77, 3, 100, 255, 42, 5]
+          }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [5] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [120] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 9, 7, 11], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [0, 1, 3, 3], "data": { "kind": "values", "values": [] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "constant", "value": 1 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 0 } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 0 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [0, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_caseB_empty",
+      "attrs": { "strides": [1, 1], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 0, 3], "data": { "kind": "values", "values": [] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [127, 128, 129, 130] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 1, 4],
+          "data": { "kind": "values", "values": [0, 0, 0, 0] },
+          "tolerance": 0.001
+        }
+      }
+    },
+    {
+      "name": "dilation2_2d_u8u8",
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 5, 5],
+          "data": {
+            "kind": "values",
+            "values": [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25]
+          }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 1, 3, 3],
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1, 1, 1, 1] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [2] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "dilation2_2d_u8s8_multichannel",
+      "attrs": { "dilations": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 8, 9, 9],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_u8s8_c8_batched_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [16, 8, 3, 3],
+          "data": { "kind": "cycle", "values": [-4, -3, -2, -1, 0, 1, 2, 3, 4, -2, 1, -1, 2, 0, 3, -3] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [100] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 16, 5, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "group2_in_channels_per_group2_u8s8",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 5, 5],
+          "data": {
+            "kind": "cycle",
+            "values": [3, 7, 12, 19, 25, 31, 44, 50, 8, 16, 22, 37, 41, 5, 18, 29, 33, 47, 11, 24]
+          }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [4, 2, 3, 3],
+          "data": { "kind": "cycle", "values": [-3, -2, -1, 0, 1, 2, 3, -1, 2, -2, 1, -3, 0, 3, -1, 1] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 4, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "dp4a_pointwise_c4_aligned_u8s8",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 8, 8],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_pointwise_u8s8_c8_batched_input_x" } }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 4, 1, 1],
+          "data": {
+            "kind": "cycle",
+            "values": [-77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90, -128, 127, 0]
+          }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [120] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 8, 8, 8], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "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] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [64, 16, 3, 3],
+          "data": { "kind": "cycle", "values": [2, -1, 3, -2, 1, 4, -3, 2, 0] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [2] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 64, 22, 22] } },
+      "attrs": {}
+    },
+    {
+      "name": "per_channel_w_zero_point_u8s8_3x3",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 5, 5],
+          "data": { "kind": "cycle", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90, 100, 110, 120, 130, 140, 150, 160] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [4, 4, 3, 3],
+          "data": {
+            "kind": "cycle",
+            "values": [-4, 3, -2, 1, 0, -1, 2, -3, 4, -1, 2, -2, 3, 0, -3, 1, 2, -4, 3, -1, 0, 1, -2, 4, -3, 2, 1, -1, 3, -4, 2, 0, -2, 4, -1, 3]
+          }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [0, 2, -3, 5] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 4, 3, 3], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "per_channel_w_zero_point_int8_pointwise",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 8, 4, 4],
+          "data": { "kind": "cycle", "values": [-50, 25, -10, 40, -30, 15, 50, -20, 35, -45, 5, -15, 20, -35, 45, -5] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [6, 8, 1, 1],
+          "data": {
+            "kind": "cycle",
+            "values": [3, -2, 5, -4, 1, -3, 6, -1, 4, -5, 2, -6, 7, -7, -2, 3, -4, 5, 1, -1, 3, -3, 2, -2]
+          }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [0, 1, -2, 3, -1, 2] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 6, 4, 4], "tolerance": 0 } },
+      "attrs": {}
+    },
+    {
+      "name": "unaligned_kernel_small_numel_scalar_path",
+      "attrs": { "strides": [2, 2], "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 6, 9, 9],
+          "data": {
+            "kind": "cycle",
+            "values": [10, 20, 30, 40, 50, 60, 70, 80, 90, 100, 110, 120, 130, 140, 150, 160, 170, 180]
+          }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [8, 6, 3, 3],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/unaligned_kernel_small_numel_scalar_path_input_w" }
+          }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [16] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 8, 5, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "group_conv_unequal_per_channel_zero_point_u8u8",
+      "attrs": { "group": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 5, 5],
+          "data": { "kind": "cycle", "values": [5, 10, 15, 20, 25, 30, 35, 40, 45, 50, 55, 60, 65, 70, 75, 80] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [4, 2, 3, 3],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/unaligned_kernel_small_numel_scalar_path_input_w" }
+          }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [10, 15, 12, 20] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 4, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "same_upper_stride2_autopad",
+      "provenance": {
+        "source": "ONNX Runtime ConvInteger-10 CPUExecutionProvider",
+        "notes": "ConvInteger inherits Conv auto_pad semantics. This fixture exercises the exact SAME_UPPER spelling and the derived asymmetric bottom/right padding for a stride-2 output."
+      },
+      "attrs": { "auto_pad": "SAME_UPPER", "strides": [2, 2] },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 1, 4, 4],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/same_upper_stride2_autopad_input_x" } }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 3, 3], "data": { "kind": "constant", "value": 1 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [54, 45, 72, 54] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "optional_zero_points_omitted",
+      "provenance": {
+        "source": "ONNX Runtime ConvInteger-10 CPUExecutionProvider",
+        "notes": "Both ConvInteger zero-point inputs are optional and independently default to zero."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 2],
+          "data": { "kind": "values", "values": [2, 4, 6, 8] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_u8u8_scalar_zero_points",
+      "provenance": {
+        "source": "ONNX Runtime ConvInteger-10 CPUExecutionProvider",
+        "notes": "ConvInteger is defined for N-dimensional convolution and rank-3 NCW is a realistic quantized audio shape."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 6], "data": { "kind": "values", "values": [10, 12, 13, 8, 15, 16] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 3], "data": { "kind": "values", "values": [3, 1, 4] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [4, -5, 15, 5] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_x_zero_point_only",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Exercises rank-3 NCW convolution with x_zero_point present and w_zero_point independently omitted."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 4], "data": { "kind": "values", "values": [10, 12, 13, 8] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2], "data": { "kind": "values", "values": [3, 1] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": {
+        "y": { "dtype": "int32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [2, 9, 7] }, "tolerance": 0 }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_w_zero_point_only",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Exercises rank-3 NCW convolution with w_zero_point present and x_zero_point independently omitted."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 4], "data": { "kind": "values", "values": [10, 12, 13, 8] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2], "data": { "kind": "values", "values": [3, 1] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [-2, -1, 5] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_zero_points_omitted",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Completes the rank-3 NCW optional-presence matrix with both zero points omitted and defaulting to zero."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 4], "data": { "kind": "values", "values": [10, 12, 13, 8] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [42, 49, 47] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "per_channel_w_zero_point_without_x_zero_point",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Exercises a per-output-channel w_zero_point while x_zero_point independently defaults to zero."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 2], "data": { "kind": "values", "values": [1, 2, 3, 4] } },
+        "w": { "dtype": "uint8", "shape": [2, 1, 1, 1], "data": { "kind": "values", "values": [3, 4] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [1, 2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [2, 4, 6, 8, 2, 4, 6, 8] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv3d_zero_points_omitted",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Exercises rank-5 NCDHW convolution with both independent zero-point inputs omitted."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [2, 5] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 1, 1],
+          "data": { "kind": "values", "values": [6, 15] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv3d_x_zero_point_only",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Exercises rank-5 NCDHW convolution with x_zero_point present and w_zero_point omitted."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [2, 5] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [3] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 1, 1],
+          "data": { "kind": "values", "values": [3, 12] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv3d_w_zero_point_only",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 optional-input contract",
+        "notes": "Exercises rank-5 NCDHW convolution with w_zero_point present and x_zero_point omitted."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [2, 5] } },
+        "w": { "dtype": "uint8", "shape": [1, 1, 1, 1, 1], "data": { "kind": "values", "values": [3] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 1, 1],
+          "data": { "kind": "values", "values": [2, 5] },
+          "tolerance": 0
+        }
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv3d_depth_dilation2_compact",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 volumetric dilation semantics",
+        "notes": "Covers a non-default depth dilation. Existing 3D fixtures varied data types and padding but kept dilationD=1."
+      },
+      "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" } }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 2, 1, 1], "data": { "kind": "values", "values": [1, 2] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int32",
+          "shape": [1, 1, 2, 2, 2],
+          "data": { "kind": "values", "values": [19, 22, 25, 28, 31, 34, 37, 40] }
+        }
+      }
+    },
+    {
+      "name": "dp4a_im2col_3d_u8s8_b1c4m8_4x8x8_k3",
+      "provenance": {
+        "source": "ONNX ConvInteger-10 volumetric convolution",
+        "notes": "Locks the shared 3-D im2col-to-DP4A route with zero-point-filled spatial padding."
+      },
+      "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] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [8, 4, 3, 3, 3],
+          "data": { "kind": "cycle", "values": [-127, -31, -1, 0, 7, 63, 126] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 8, 4, 8, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "same_lower_stride2_autopad",
+      "provenance": {
+        "source": "ONNX Runtime ConvInteger-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" } }
+        },
+        "w": { "dtype": "uint8", "shape": [1, 1, 3, 3], "data": { "kind": "constant", "value": 1 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 2, 2], "tolerance": 0 } }
+    }
+  ]
+}