sync 6fdf6301e2bb

Browse files

Files changed (8) hide show

README.md +19 -2
build/webgpu/bench.json +2187 -0
build/webgpu/conv-int-accumulate-spatial.wgsl.jinja +39 -58
build/webgpu/conv-int-im2col-spatial.wgsl.jinja +23 -21
build/webgpu/manifest.json +349 -164
build/webgpu/metadata.json +19 -9
build/webgpu/quant-dp4a-matmul.wgsl.jinja +103 -27
build/webgpu/test.json +0 -0

README.md CHANGED Viewed

@@ -52,12 +52,29 @@ Attributes and default values (overridable per request):
 | `TW` | `uint8`, `int8` |
 | `TY` | `int32` |
 ## Files
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
-- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`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)
@@ -65,7 +82,7 @@ Attributes and default values (overridable per request):
 ## Use with `@huggingface/kernels`
 ```sh
-npm install --save-exact @huggingface/kernels@0.0.1-preview.2
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

 | `TW` | `uint8`, `int8` |
 | `TY` | `int32` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `dp4a_pointwise_1x1_tail` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_tail_x_zero_point_only` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_tail_w_zero_point_only` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_tail_no_zero_points` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_x_zero_point_only` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_w_zero_point_only` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_no_zero_points` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_tail_per_channel` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_tail_per_channel_w_zero_point_only` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_per_channel` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
+- `dp4a_pointwise_1x1_per_channel_w_zero_point_only` — Exact tiled pointwise convolution for one or two spatial dimensions. Native packed dots or exact integer-vector dots share zero-point specialization; tile resources and dispatch eligibility follow device limits.
 ## Files
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark 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`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

build/webgpu/bench.json CHANGED Viewed

@@ -182,6 +182,2193 @@
         ]
       },
       "attrs": {}
     }
   ]
 }

         ]
       },
       "attrs": {}
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x1w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x1w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x1w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x1w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x1w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x1w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x1w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x1w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_int8_x1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_int8_x1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_int8_x0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_int8_x0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_int8_x1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_int8_x1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_int8_x0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_int8_x0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_uint8_x1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_uint8_x1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_uint8_x0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_int8_uint8_x0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_uint8_x1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_uint8_x1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_uint8_x0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 1, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 1, 3136] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "per_channel_uint8_uint8_x0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 1, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": {
+          "dtype": "uint8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 254] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 1, 65] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_uint8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_uint8_int8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_uint8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "pointwise_zeros_int8_int8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 1, 63] } },
+      "attrs": {},
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1)" }] }
+    },
+    {
+      "name": "conv1d_u8s8_pointwise_1x64x56x56_oc128",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 3136], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [128, 64, 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, 3136] } },
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_u8s8_pointwise_unalignedC_c66_oc128_56x56_scalar",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 66, 3136], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [128, 66, 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, 3136] } },
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_u8s8_pointwise_gemv_alignedN_n2048_c256_dp4a",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 256, 1], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [2048, 256, 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] } },
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_u8s8_pointwise_gemv_unalignedN_n2050_c256_dp4a",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 256, 1], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [2050, 256, 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] } },
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_u8s8_pointwise_1x64x56x56_oc128_wzp_omitted_pathology",
+      "preset": "smoke",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 64, 3136], "data": { "kind": "constant", "value": 127 } },
+        "w": { "dtype": "int8", "shape": [128, 64, 1], "data": { "kind": "constant", "value": -3 } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      },
+      "attrs": {}
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x1w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x1w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x0w0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x0w0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_per_channel_uint8_int8_x1_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_per_channel_uint8_int8_x1_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_per_channel_uint8_int8_x0_b1c64m128n3136",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 64, 3136],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [128, 64, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 128, 3136] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_per_channel_uint8_int8_x0_b2c66m129n65",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 66, 65],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [129, 66, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": {
+          "dtype": "int8",
+          "shape": [129],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [2, 129, 65] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_uint8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 251 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_uint8_int8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 128, 127, 254, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        },
+        "w_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "constant", "value": 253 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_uint8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "uint8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [1, 128, 127, 254, 63, 0, 255] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x1w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x1w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": -121 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x0w1_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x0w1_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        },
+        "w_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "constant", "value": 119 } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x0w0_b1c4m1n1",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [1, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
+    },
+    {
+      "name": "conv1d_pointwise_zeros_int8_int8_x0w0_b1c4m7n63",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 63],
+          "data": { "kind": "cycle", "values": [-128, 127, -1, 0, 1, -127, 63] }
+        },
+        "w": {
+          "dtype": "int8",
+          "shape": [7, 4, 1],
+          "data": { "kind": "cycle", "values": [-1, 0, 1, -127, 63, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 7, 63] } },
+      "attrs": {},
+      "bench": {
+        "metrics": [{ "type": "gflops", "value": "2 * numel(shapes.y) * dim(shapes.w, 1) * dim(shapes.w, 2)" }]
+      }
     }
   ]
 }

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

@@ -1,36 +1,41 @@
-{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
-{% if note == "dispatch-limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (outputs > 16.7M elements).
-{% elif note == "limit" %}
-  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width.
-{% elif note == "device-axis" %}
-  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
-  // width; gid.y carries the high portion of the output index.
-{% elif note == "vec4-limit" %}
-  // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
-  // per-axis workgroup fold width (the dispatch caps x and spills into y).
-{% elif note == "element-limit" %}
-  // 2D-folded flat element index: gid.y carries the high bits past the
-  // dispatch's per-axis workgroup fold width.
-{% elif note == "dispatch" %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
-{% endif %}
-{% if bound == "" %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-{%- elif guardInline %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) { return; }
-{%- else %}
-  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if ({{ name }} >= {{ bound }}) {
-    return;
-  }
-{%- endif %}
-{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn read_x(index: u32) -> i32 {
@@ -49,38 +54,14 @@ fn read_w(index: u32) -> i32 {
   {% endif %}
 }
-fn read_x_zero() -> i32 {
-  {% if xZeroOmitted is defined and 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 wZeroOmitted is defined and 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>) {
-{{ flat_index_2d("index", guardInline=true) }}
   let ow = index % params.outW;
   var t = index / params.outW;
   let oh = t % params.outH;

+{% macro quant_read_scalar_zero(fn_name, buffer_name, is_unsigned, present=true) %}
+fn {{ fn_name }}() -> i32 {
+  {% if present %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[0]);
+  {% else %}
+  return {{ buffer_name }}[0];
+  {% endif %}
+  {% else %}
+  return 0;
+  {% endif %}
+}{% endmacro %}
+{% macro quant_read_column_zero(fn_name, buffer_name, is_unsigned, present=true, per_column=false, param_name="col", return_type="i32") %}
+fn {{ fn_name }}({% if per_column %}{{ param_name }}: u32{% endif %}) -> {{ return_type }} {
+  {% if present %}
+  {% if per_column %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[{{ param_name }}]);
+  {% else %}
+  return {{ buffer_name }}[{{ param_name }}];
+  {% endif %}
+  {% else %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[0]);
+  {% else %}
+  return {{ buffer_name }}[0];
+  {% endif %}
+  {% endif %}
+  {% else %}
+  return 0;
+  {% endif %}
+}{% endmacro %}
+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
   // 2D-folded flat index: gid.y carries the high bits past the dispatch's
   // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) { return; }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 fn read_x(index: u32) -> i32 {
   {% endif %}
 }
+{{ quant_read_scalar_zero("read_x_zero", "x_zero_point", xUnsigned, not (xZeroOmitted is defined and xZeroOmitted)) }}
 // A per-output-channel weight zero point applies at every spatial rank.
+{{ quant_read_column_zero("read_w_zero", "w_zero_point", wUnsigned, not (wZeroOmitted is defined and wZeroOmitted), wZeroPerChannel, "oc") }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "index", guardInline=true) }}
   let ow = index % params.outW;
   var t = index / params.outW;
   let oh = t % params.outH;

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

@@ -5,7 +5,11 @@
  * Optional extra K rows contain the raw input zero point to align the DP4A
  * reduction. */
 {{ env.wgsl.resourceDeclarations }}
-{% set padK = padK is defined and padK %}
 {% if spatialRankSpec == 3 %}
 const KERNEL_D: u32 = {{ kernelDSpec }}u;
@@ -27,23 +31,10 @@ const DILATION_H: u32 = {{ dilationHSpec }}u;
 const DILATION_W: u32 = {{ dilationWSpec }}u;
 const PAD_TOP: i32 = {{ padTopSpec }};
 const PAD_LEFT: i32 = {{ padLeftSpec }};
-@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  let pos = gid.x;
-  if (pos >= params.outCount) {
-    return;
-  }
-  let k = gid.y;
-  let batch = gid.z;
-{% if padK %}
-  if (k >= params.sourceKRows) {
-    cols[(batch * params.kRows + k) * params.outCount + pos] = x_zero_point[0];
-    return;
-  }
-{% endif %}
-  let ic = k / KSIZE;
-  let kq = k - ic * KSIZE;
 {% if spatialRankSpec == 3 %}
   let kd = kq / KERNEL_HW;
   let khw = kq - kd * KERNEL_HW;
@@ -63,22 +54,33 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% 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 spatialRankSpec == 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;
 }

  * Optional extra K rows contain the raw input zero point to align the DP4A
  * reduction. */
 {{ env.wgsl.resourceDeclarations }}
+{% set kernelDSpec = kernelDSpec | default(1) %}
+{% set strideDSpec = strideDSpec | default(1) %}
+{% set dilationDSpec = dilationDSpec | default(1) %}
+{% set padFrontSpec = padFrontSpec | default(0) %}
+{% macro spatial_constants() %}
 {% if spatialRankSpec == 3 %}
 const KERNEL_D: u32 = {{ kernelDSpec }}u;
 const DILATION_W: u32 = {{ dilationWSpec }}u;
 const PAD_TOP: i32 = {{ padTopSpec }};
 const PAD_LEFT: i32 = {{ padLeftSpec }};
+{% endmacro %}
+{% macro spatial_sample(value="value", input="x", zero="x_zero_point", sampleK="k", sampleScalar="auto") %}
+  let ic = {{ sampleK }} / KSIZE;
+  let kq = {{ sampleK }} - ic * KSIZE;
 {% if spatialRankSpec == 3 %}
   let kd = kq / KERNEL_HW;
   let khw = kq - kd * KERNEL_HW;
 {% 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 if sampleScalar == "auto" else sampleScalar }} = {{ zero }}[0];
 {% if spatialRankSpec == 3 %}
   if (
     id >= 0 && id < i32(params.inD)
     && ih >= 0 && ih < i32(params.inH)
     && iw >= 0 && iw < i32(params.inW)
   ) {
+    {{ value }} = {{ input }}[
       (((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 }} = {{ input }}[((batch * params.inChannels + ic) * params.inH + u32(ih)) * params.inW + u32(iw)];
   }
 {% endif %}
+{% endmacro %}
+{{ spatial_constants() }}
+@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;
+{{ spatial_sample() }}
   cols[(batch * params.kRows + k) * params.outCount + pos] = value;
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -80,29 +80,43 @@
     "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"
   },
   "when": ["outputBatchChannelsOk"],
   "bindings": {
-    "x": { "buffer": "read-only-storage", "elementType": "$bScalar" },
-    "x_zero_point": { "buffer": "read-only-storage", "elementType": "$bScalar", "length": 1 },
-    "cols": { "scratch": "cols3d", "buffer": "storage", "elementType": "$bScalar" },
     "params": {
-      "buffer": "uniform",
       "struct": [
         { "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)" },
@@ -118,10 +132,9 @@
         { "name": "outW", "type": "u32", "value": "dim(shapes.y, 4)" }
       ]
     },
-    "cols_2": { "name": "cols", "buffer": "storage", "elementType": "$bScalar" },
-    "params_2": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "outCount", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
@@ -131,14 +144,13 @@
         { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" }
       ]
     },
-    "a": { "arg": "w", "buffer": "read-only-storage", "elementType": "$aVec4" },
     "b": { "scratch": "cols3d", "buffer": "read-only-storage", "elementType": "$bScalar" },
-    "a_zero_point": { "arg": "w_zero_point", "buffer": "read-only-storage", "elementType": "$aScalar", "length": 1 },
-    "b_zero_point": { "arg": "x_zero_point", "buffer": "read-only-storage", "elementType": "$bScalar", "length": 1 },
-    "y": { "buffer": "storage", "elementType": "i32" },
-    "params_3": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "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)" },
@@ -160,10 +172,9 @@
         }
       ]
     },
-    "b_2": { "scratch": "cols", "name": "b", "buffer": "read-only-storage", "elementType": "$bScalar" },
-    "params_4": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "dim(shapes.w, 0)" },
         { "name": "N", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
@@ -177,27 +188,25 @@
         { "name": "yBatchStride", "type": "u32", "value": "dim(shapes.w, 0) * dim(shapes.y, 2) * dim(shapes.y, 3)" }
       ]
     },
-    "a_2": { "arg": "w", "name": "a", "buffer": "read-only-storage", "elementType": "$aScalar" },
-    "b_3": { "arg": "x", "name": "b", "buffer": "read-only-storage", "elementType": "$bScalar" },
-    "params_5": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "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)" }
       ]
     },
-    "x_2": { "name": "x", "buffer": "read-only-storage", "elementType": "$xScalar" },
-    "w": { "buffer": "read-only-storage", "elementType": "$wScalar" },
-    "x_zero_point_2": { "name": "x_zero_point", "buffer": "read-only-storage", "elementType": "$xScalar", "length": 1 },
-    "w_zero_point": { "buffer": "read-only-storage", "elementType": "$wScalar" },
-    "params_6": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inH", "type": "u32", "value": "inputHeight" },
@@ -219,10 +228,9 @@
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
-    "w_zero_point_2": { "name": "w_zero_point", "buffer": "read-only-storage", "elementType": "$wScalar", "length": 1 },
-    "params_7": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inD", "type": "u32", "value": "dim(shapes.x, 2)" },
@@ -249,7 +257,8 @@
         { "name": "padW", "type": "i32", "value": "effectivePadLeft" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
-    }
   },
   "variants": [
     {
@@ -257,18 +266,10 @@
       "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", "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) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4), tunables.WORKGROUP_SIZE) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4), 64) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "ceilDiv(dim(shapes.w, 0), 8) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "dim(shapes.x, 0) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"],
       "derive": {
-        "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,
@@ -312,7 +313,7 @@
           "id": "main",
           "name": "ConvInteger.Im2Col3dDp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a", "b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4) / 64)",
             "y": "ceil(dim(shapes.w, 0) / 8)",
@@ -325,21 +326,7 @@
       "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)", "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) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"],
-      "derive": {
-        "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",
@@ -363,7 +350,7 @@
             "padTopSpec": "effectivePadTop",
             "padLeftSpec": "effectivePadLeft"
           },
-          "bindings": ["x", "x_zero_point", "cols_2", "params_2"],
           "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)",
@@ -374,7 +361,7 @@
           "id": "main",
           "name": "ConvInteger.Im2ColDp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a", "b_2", "a_zero_point", "b_zero_point", "y", "params_4"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
             "y": "ceil(dim(shapes.w, 0) / 8)",
@@ -386,31 +373,102 @@
     {
       "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", "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) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"],
       "derive": {
-        "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": ["a_2", "b_3", "a_zero_point", "b_zero_point", "y", "params_5"],
           "dispatch": {
-            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
-            "y": "ceil(dim(shapes.w, 0) / 8)",
             "z": "dim(shapes.y, 0)"
           }
         }
@@ -419,31 +477,210 @@
     {
       "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", "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) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"],
       "derive": {
-        "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": ["a", "b_3", "a_zero_point", "b_zero_point", "y", "params_5"],
           "dispatch": {
-            "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
-            "y": "ceil(dim(shapes.w, 0) / 8)",
             "z": "dim(shapes.y, 0)"
           }
         }
@@ -453,20 +690,14 @@
       "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", "outputSpatialShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "false", "wZeroOmitted": "false" },
-          "bindings": ["x_2", "w", "x_zero_point_2", "w_zero_point", "y", "params_6"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -479,20 +710,14 @@
       "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", "outputSpatialShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "true", "wZeroOmitted": "false" },
-          "bindings": ["x_2", "w", "w_zero_point", "y", "params_6"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -504,19 +729,14 @@
     {
       "id": "nchw2d_scalar_zero_points",
       "when": ["nchwTensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "groupChannelsOk", "outputSpatialShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "false", "wZeroOmitted": "false" },
-          "bindings": ["x_2", "w", "x_zero_point_2", "w_zero_point_2", "y", "params_6"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -528,19 +748,14 @@
     {
       "id": "nchw2d_x_zero_point_only",
       "when": ["nchwTensorOk", "present.x_zero_point", "not present.w_zero_point", "scalarXZeroRequiredOk", "groupChannelsOk", "outputSpatialShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "false", "wZeroOmitted": "true" },
-          "bindings": ["x_2", "w", "x_zero_point_2", "y", "params_6"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -552,19 +767,14 @@
     {
       "id": "nchw2d_w_zero_point_only",
       "when": ["nchwTensorOk", "not present.x_zero_point", "present.w_zero_point", "scalarWZeroRequiredOk", "groupChannelsOk", "outputSpatialShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "true", "wZeroOmitted": "false" },
-          "bindings": ["x_2", "w", "w_zero_point_2", "y", "params_6"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -576,19 +786,14 @@
     {
       "id": "ncdhw3d_scalar_zero_points",
       "when": ["rank5TensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "groupChannelsOk", "output3dShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "false", "wZeroOmitted": "false" },
-          "bindings": ["x_2", "w", "x_zero_point_2", "w_zero_point_2", "y", "params_7"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -600,19 +805,14 @@
     {
       "id": "ncdhw3d_x_zero_point_only",
       "when": ["rank5TensorOk", "scalarXZeroRequiredOk", "not present.w_zero_point", "groupChannelsOk", "output3dShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "false", "wZeroOmitted": "true" },
-          "bindings": ["x_2", "w", "x_zero_point_2", "y", "params_7"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -624,19 +824,14 @@
     {
       "id": "ncdhw3d_w_zero_point_only",
       "when": ["rank5TensorOk", "not present.x_zero_point", "scalarWZeroRequiredOk", "groupChannelsOk", "output3dShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "true", "wZeroOmitted": "false" },
-          "bindings": ["x_2", "w", "w_zero_point_2", "y", "params_7"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -648,19 +843,14 @@
     {
       "id": "ncdhw3d_implicit_zero_points",
       "when": ["rank5TensorOk", "zeroPointsOmitted", "groupChannelsOk", "output3dShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "true", "wZeroOmitted": "true" },
-          "bindings": ["x_2", "w", "y", "params_7"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -672,19 +862,14 @@
     {
       "id": "nchw_implicit_zero_points",
       "when": ["nchwTensorOk", "zeroPointsOmitted", "groupChannelsOk", "outputSpatialShapeOk"],
-      "derive": {
-        "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",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "true", "wZeroOmitted": "true" },
-          "bindings": ["x_2", "w", "y", "params_6"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",

     "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)",
+    "outputBatchChannelsOk": "dim(shapes.y, 0) == dim(shapes.x, 0) and dim(shapes.y, 1) == dim(shapes.w, 0)",
+    "groupChannelsOk": "attrs.group >= 1 and dim(shapes.w, 1) * attrs.group == dim(shapes.x, 1) and dim(shapes.w, 0) % attrs.group == 0",
+    "output2dShapeOk": "dim(shapes.y, 2) == expectedOutputHeight and dim(shapes.y, 3) == expectedOutputWidth",
+    "aScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+    "bScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+    "wScalar": "\"u32\" if tensorDtypes.w == \"uint8\" else \"i32\"",
+    "xScalar": "\"u32\" if tensorDtypes.x == \"uint8\" else \"i32\"",
+    "aUnsigned": "tensorDtypes.w == \"uint8\"",
+    "bUnsigned": "tensorDtypes.x == \"uint8\"",
+    "wUnsigned": "tensorDtypes.w == \"uint8\"",
+    "xUnsigned": "tensorDtypes.x == \"uint8\"",
+    "bZeroPerColumn": false,
     "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",
     "ungroupedChannelsOk": "attrs.group == 1 and dim(shapes.w, 1) == dim(shapes.x, 1)",
     "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",
+    "pointwiseTileM": 8,
+    "pointwiseTileN": 64,
+    "pointwiseKtWords": 32,
+    "pointwiseResourcesOk": "pointwiseTileN <= min(device.limits.maxComputeWorkgroupSizeX, device.limits.maxComputeInvocationsPerWorkgroup) and 4 * pointwiseTileM * (pointwiseKtWords + 1) <= device.limits.maxComputeWorkgroupStorageSize",
+    "pointwiseDispatchOk": "ceilDiv(outputHeight * outputWidth, pointwiseTileN) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(dim(shapes.w, 0), pointwiseTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and dim(shapes.y, 0) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\""
   },
   "when": ["outputBatchChannelsOk"],
   "bindings": {
+    "x": { "elementType": "$bScalar" },
+    "x_zero_point": { "elementType": "$bScalar", "length": 1 },
+    "cols": { "scratch": "cols3d", "elementType": "$bScalar" },
     "params": {
       "struct": [
         { "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": "outW", "type": "u32", "value": "dim(shapes.y, 4)" }
       ]
     },
+    "cols_im2col": { "name": "cols", "buffer": "storage", "elementType": "$bScalar" },
+    "params_im2col": {
       "name": "params",
       "struct": [
         { "name": "outCount", "type": "u32", "value": "dim(shapes.y, 2) * dim(shapes.y, 3)" },
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "outW", "type": "u32", "value": "dim(shapes.y, 3)" }
       ]
     },
+    "a": { "arg": "w", "elementType": "$aVec4" },
     "b": { "scratch": "cols3d", "buffer": "read-only-storage", "elementType": "$bScalar" },
+    "a_zero_point": { "arg": "w_zero_point", "elementType": "$aScalar", "length": 1 },
+    "b_zero_point": { "arg": "x_zero_point", "elementType": "$bScalar", "length": 1 },
+    "y": { "elementType": "i32" },
+    "params_main": {
       "name": "params",
       "struct": [
         { "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)" },
         }
       ]
     },
+    "b_main": { "scratch": "cols", "name": "b", "buffer": "read-only-storage", "elementType": "$bScalar" },
+    "params__uniform": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "dim(shapes.w, 0)" },
         { "name": "N", "type": "u32", "value": "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)" }
       ]
     },
+    "a_w": { "arg": "w", "name": "a", "elementType": "$aScalar" },
+    "b_x": { "arg": "x", "name": "b", "elementType": "$bScalar" },
+    "params_quant_dp4a_matmul": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "dim(shapes.w, 0)" },
+        { "name": "N", "type": "u32", "value": "outputHeight * outputWidth" },
         { "name": "K", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "aBatchStride4", "type": "u32", "value": 0 },
+        { "name": "bBatchStride", "type": "u32", "value": "dim(shapes.x, 1) * inputHeight * inputWidth" },
+        { "name": "yBatchStride", "type": "u32", "value": "dim(shapes.w, 0) * outputHeight * outputWidth" }
       ]
     },
+    "x_main": { "name": "x", "elementType": "$xScalar" },
+    "w": { "elementType": "$wScalar" },
+    "x_zero_point_main": { "name": "x_zero_point", "elementType": "$xScalar", "length": 1 },
+    "w_zero_point": { "elementType": "$wScalar" },
+    "params_nchw2d_x_zero_point_only": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inH", "type": "u32", "value": "inputHeight" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
     },
+    "w_zero_point_main": { "name": "w_zero_point", "elementType": "$wScalar", "length": 1 },
+    "params_ncdhw3d_x_zero_point_only": {
       "name": "params",
       "struct": [
         { "name": "inChannels", "type": "u32", "value": "dim(shapes.x, 1)" },
         { "name": "inD", "type": "u32", "value": "dim(shapes.x, 2)" },
         { "name": "padW", "type": "i32", "value": "effectivePadLeft" },
         { "name": "count", "type": "u32", "value": "numel(shapes.y)" }
       ]
+    },
+    "a_zero_point_per_row": { "arg": "w_zero_point", "name": "a_zero_point", "elementType": "$aScalar" }
   },
   "variants": [
     {
       "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", "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) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4), tunables.WORKGROUP_SIZE) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "ceilDiv(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4), 64) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "ceilDiv(dim(shapes.w, 0), 8) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)", "dim(shapes.x, 0) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"],
       "derive": {
         "hasAZero": true,
         "hasBZero": true,
         "aPacked": true,
         "aZeroPerRow": false,
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32,
           "id": "main",
           "name": "ConvInteger.Im2Col3dDp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) * dim(shapes.y, 4) / 64)",
             "y": "ceil(dim(shapes.w, 0) / 8)",
       "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)", "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) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"],
+      "derive": { "hasAZero": true, "hasBZero": true, "tileM": 8, "tileN": 64, "ktWords": 32 },
       "intermediates": [
         {
           "id": "cols",
             "padTopSpec": "effectivePadTop",
             "padLeftSpec": "effectivePadLeft"
           },
+          "bindings": ["x", "x_zero_point", "cols_im2col", "params_im2col"],
           "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)",
           "id": "main",
           "name": "ConvInteger.Im2ColDp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_main", "a_zero_point", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) * dim(shapes.y, 3) / 64)",
             "y": "ceil(dim(shapes.w, 0) / 8)",
     {
       "id": "dp4a_pointwise_1x1_tail",
       "priority": 21,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
       "derive": {
         "hasAZero": true,
         "hasBZero": true,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aPacked": false
       },
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger.Dp4aPointwiseTail",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_w", "b_x", "a_zero_point", "b_zero_point", "y", "params_quant_dp4a_matmul"],
           "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_tail_x_zero_point_only",
+      "priority": 21,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "scalarXZeroRequiredOk", "not present.w_zero_point", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "hasAZero": false,
+        "hasBZero": true,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aPacked": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwiseTail",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_w", "b_x", "b_zero_point", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_tail_w_zero_point_only",
+      "priority": 21,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "not present.x_zero_point", "scalarWZeroRequiredOk", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "hasAZero": true,
+        "hasBZero": false,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aPacked": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwiseTail",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_w", "b_x", "a_zero_point", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_tail_no_zero_points",
+      "priority": 21,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "not present.x_zero_point", "not present.w_zero_point", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "hasAZero": false,
+        "hasBZero": false,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aPacked": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwiseTail",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_w", "b_x", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
             "z": "dim(shapes.y, 0)"
           }
         }
     {
       "id": "dp4a_pointwise_1x1",
       "priority": 20,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "dim(shapes.x, 1) % 4 == 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
       "derive": {
         "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
         "hasAZero": true,
         "hasBZero": true,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords"
       },
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger.Dp4aPointwise",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_x", "a_zero_point", "b_zero_point", "y", "params_quant_dp4a_matmul"],
           "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_x_zero_point_only",
+      "priority": 20,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "scalarXZeroRequiredOk", "not present.w_zero_point", "dim(shapes.x, 1) % 4 == 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
+        "hasAZero": false,
+        "hasBZero": true,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwise",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_x", "b_zero_point", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_w_zero_point_only",
+      "priority": 20,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "not present.x_zero_point", "scalarWZeroRequiredOk", "dim(shapes.x, 1) % 4 == 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
+        "hasAZero": true,
+        "hasBZero": false,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwise",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_x", "a_zero_point", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_no_zero_points",
+      "priority": 20,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "not present.x_zero_point", "not present.w_zero_point", "dim(shapes.x, 1) % 4 == 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
+        "hasAZero": false,
+        "hasBZero": false,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwise",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_x", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_tail_per_channel",
+      "priority": 21,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "scalarXZeroRequiredOk", "present.w_zero_point", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "hasAZero": true,
+        "hasBZero": true,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aPacked": false,
+        "aZeroPerRow": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwiseTail",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_w", "b_x", "a_zero_point_per_row", "b_zero_point", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_tail_per_channel_w_zero_point_only",
+      "priority": 21,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "not present.x_zero_point", "present.w_zero_point", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "dim(shapes.x, 1) >= 32", "dim(shapes.x, 1) % 4 != 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "hasAZero": true,
+        "hasBZero": false,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aPacked": false,
+        "aZeroPerRow": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwiseTail",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_w", "b_x", "a_zero_point_per_row", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_per_channel",
+      "priority": 20,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "scalarXZeroRequiredOk", "present.w_zero_point", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "dim(shapes.x, 1) % 4 == 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
+        "hasAZero": true,
+        "hasBZero": true,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aZeroPerRow": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwise",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_x", "a_zero_point_per_row", "b_zero_point", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
+            "z": "dim(shapes.y, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "dp4a_pointwise_1x1_per_channel_w_zero_point_only",
+      "priority": 20,
+      "when": ["nchwTensorOk", "ungroupedChannelsOk", "kernelHeight == 1", "kernelWidth == 1", "strideH == 1", "strideW == 1", "padTop == 0 and padBottom == 0", "padLeft == 0 and padRight == 0", "outputHeight == inputHeight", "outputWidth == inputWidth", "not present.x_zero_point", "present.w_zero_point", "ranks.w_zero_point == 1", "dim(shapes.w_zero_point, 0) == dim(shapes.w, 0)", "dim(shapes.x, 1) % 4 == 0", "pointwiseResourcesOk", "pointwiseDispatchOk"],
+      "derive": {
+        "aVec4": "\"vec4<u32>\" if tensorDtypes.w == \"uint8\" else \"vec4<i32>\"",
+        "hasAZero": true,
+        "hasBZero": false,
+        "tileM": "pointwiseTileM",
+        "tileN": "pointwiseTileN",
+        "ktWords": "pointwiseKtWords",
+        "aZeroPerRow": true
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConvInteger.Dp4aPointwise",
+          "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a", "b_x", "a_zero_point_per_row", "y", "params_quant_dp4a_matmul"],
+          "dispatch": {
+            "x": "ceilDiv(outputHeight * outputWidth, pointwiseTileN)",
+            "y": "ceilDiv(dim(shapes.w, 0), pointwiseTileM)",
             "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", "outputSpatialShapeOk"],
+      "derive": { "wZeroPerChannel": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvIntegerPerChannel",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "false", "wZeroOmitted": "false" },
+          "bindings": ["x_main", "w", "x_zero_point_main", "w_zero_point", "y", "params_nchw2d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
       "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", "outputSpatialShapeOk"],
+      "derive": { "wZeroPerChannel": true },
       "passes": [
         {
           "id": "main",
           "name": "ConvIntegerPerChannel",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "true", "wZeroOmitted": "false" },
+          "bindings": ["x_main", "w", "w_zero_point", "y", "params_nchw2d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw2d_scalar_zero_points",
       "when": ["nchwTensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "groupChannelsOk", "outputSpatialShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "false", "wZeroOmitted": "false" },
+          "bindings": ["x_main", "w", "x_zero_point_main", "w_zero_point_main", "y", "params_nchw2d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw2d_x_zero_point_only",
       "when": ["nchwTensorOk", "present.x_zero_point", "not present.w_zero_point", "scalarXZeroRequiredOk", "groupChannelsOk", "outputSpatialShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "false", "wZeroOmitted": "true" },
+          "bindings": ["x_main", "w", "x_zero_point_main", "y", "params_nchw2d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw2d_w_zero_point_only",
       "when": ["nchwTensorOk", "not present.x_zero_point", "present.w_zero_point", "scalarWZeroRequiredOk", "groupChannelsOk", "outputSpatialShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "true", "wZeroOmitted": "false" },
+          "bindings": ["x_main", "w", "w_zero_point_main", "y", "params_nchw2d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncdhw3d_scalar_zero_points",
       "when": ["rank5TensorOk", "scalarXZeroRequiredOk", "scalarWZeroRequiredOk", "groupChannelsOk", "output3dShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger3d",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "false", "wZeroOmitted": "false" },
+          "bindings": ["x_main", "w", "x_zero_point_main", "w_zero_point_main", "y", "params_ncdhw3d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncdhw3d_x_zero_point_only",
       "when": ["rank5TensorOk", "scalarXZeroRequiredOk", "not present.w_zero_point", "groupChannelsOk", "output3dShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger3d",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "false", "wZeroOmitted": "true" },
+          "bindings": ["x_main", "w", "x_zero_point_main", "y", "params_ncdhw3d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncdhw3d_w_zero_point_only",
       "when": ["rank5TensorOk", "not present.x_zero_point", "scalarWZeroRequiredOk", "groupChannelsOk", "output3dShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger3d",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "true", "wZeroOmitted": "false" },
+          "bindings": ["x_main", "w", "w_zero_point_main", "y", "params_ncdhw3d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "ncdhw3d_implicit_zero_points",
       "when": ["rank5TensorOk", "zeroPointsOmitted", "groupChannelsOk", "output3dShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger3d",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 3, "xZeroOmitted": "true", "wZeroOmitted": "true" },
+          "bindings": ["x_main", "w", "y", "params_ncdhw3d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "nchw_implicit_zero_points",
       "when": ["nchwTensorOk", "zeroPointsOmitted", "groupChannelsOk", "outputSpatialShapeOk"],
+      "derive": {},
       "passes": [
         {
           "id": "main",
           "name": "ConvInteger",
           "shader": "conv-int-accumulate-spatial.wgsl.jinja",
           "derive": { "spatialRankSpec": 2, "xZeroOmitted": "true", "wZeroOmitted": "true" },
+          "bindings": ["x_main", "w", "y", "params_nchw2d_x_zero_point_only"],
           "dispatch": {
             "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",

build/webgpu/metadata.json CHANGED Viewed

@@ -1,28 +1,38 @@
 {
   "name": "ai.onnx.ConvInteger",
-  "id": "_ai_onnx_convinteger_webgpu_b3b1b7d",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "8qMXwF7LdNjySoXjQdUTVXPKqVAq0ugoDn2ViHhNbu0=",
-      "conv-int-accumulate-spatial.wgsl.jinja": "UdJ7LJuIFN2UaKUjcBJvtgzu0wm7HkD0nRTPZE1uMPU=",
-      "conv-int-im2col-spatial.wgsl.jinja": "3vFeBqBVtyYtxV1yteSQ3O7skFur08E3vDoVtdwcYAo=",
-      "manifest.json": "qqrXA/JviMtwmSxks7L2cI9lCxiVPDoEWotnWHEqlI4=",
-      "quant-dp4a-matmul.wgsl.jinja": "0D0fsGBbb4gNT5lh8UrHMsaGW5RdmB6mZ98DIrOCKso=",
-      "test.json": "mYyrEgbPAJwKlXwI3dnO0G6vgRAhDFPb0TuZiOG3Fn0="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "dp4a_im2col_ncdhw3d_scalar_zero_points": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
       "dp4a_im2col_nchw_scalar_zero_points": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
       "dp4a_pointwise_1x1_tail": ["quant-dp4a-matmul.wgsl.jinja"],
       "dp4a_pointwise_1x1": ["quant-dp4a-matmul.wgsl.jinja"],
       "nchw2d_per_channel_zero_point": ["conv-int-accumulate-spatial.wgsl.jinja"],
       "nchw2d_per_channel_w_zero_point_only": ["conv-int-accumulate-spatial.wgsl.jinja"],
       "nchw2d_scalar_zero_points": ["conv-int-accumulate-spatial.wgsl.jinja"],

 {
   "name": "ai.onnx.ConvInteger",
+  "id": "_ai_onnx_convinteger_webgpu_befbf8c",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "w1Y+YvLhm6rmHnI3I6fdQTErLMPyr5es6PymArrgGTA=",
+      "conv-int-accumulate-spatial.wgsl.jinja": "7393dtsdsogF8Zzus4+gZiwygpXHb15ZTWhE9VM83UA=",
+      "conv-int-im2col-spatial.wgsl.jinja": "rUkI3zRiXYyD3r5GvVRqfcLMijHGc2KA71kljasROJM=",
+      "manifest.json": "guLLSG4M/NsDkfMSVXE0qMhQ2L9fYa8oY2GLgeyOXME=",
+      "quant-dp4a-matmul.wgsl.jinja": "hfcRmXhc3rBZU5qOfHyf75KqjiAmJpqzTY7BlyvoyC8=",
+      "test.json": "T5Q9TkFCTbjYtbC//Z9DuWYARBX+XkneiZt+RCoNamI="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
       "dp4a_im2col_ncdhw3d_scalar_zero_points": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
       "dp4a_im2col_nchw_scalar_zero_points": ["conv-int-im2col-spatial.wgsl.jinja", "quant-dp4a-matmul.wgsl.jinja"],
       "dp4a_pointwise_1x1_tail": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_tail_x_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_tail_w_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_tail_no_zero_points": ["quant-dp4a-matmul.wgsl.jinja"],
       "dp4a_pointwise_1x1": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_x_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_w_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_no_zero_points": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_tail_per_channel": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_tail_per_channel_w_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_per_channel": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_pointwise_1x1_per_channel_w_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
       "nchw2d_per_channel_zero_point": ["conv-int-accumulate-spatial.wgsl.jinja"],
       "nchw2d_per_channel_w_zero_point_only": ["conv-int-accumulate-spatial.wgsl.jinja"],
       "nchw2d_scalar_zero_points": ["conv-int-accumulate-spatial.wgsl.jinja"],

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

@@ -1,3 +1,19 @@
 {% if not env.device.wgslLanguageFeatures.has("packed_4x8_integer_dot_product") %}
 // Runtime source-compatible fallbacks for the packed 4x8 integer built-ins used
 // by quantized kernels. Native implementations remain unshadowed when the
@@ -19,15 +35,22 @@
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const TILE_M: u32 = {{ tileM }}u;
 const TILE_N: u32 = {{ tileN }}u;
 const KT_WORDS: u32 = {{ ktWords }}u;
 const A_TILE_WORDS: u32 = TILE_M * KT_WORDS;
 var<workgroup> a_tile: array<u32, A_TILE_WORDS>;
 var<workgroup> a_row_sum: array<atomic<i32>, TILE_M>;
 // Pack 4 widened A values into one u32 (byte 0 = lowest k).
 fn pack_a(v: vec4<{{ aScalar }}>) -> u32 {
 {% if aUnsigned %}
@@ -37,6 +60,7 @@ fn pack_a(v: vec4<{{ aScalar }}>) -> u32 {
   return (u32(v.x) & 0xFFu) | ((u32(v.y) & 0xFFu) << 8u) | ((u32(v.z) & 0xFFu) << 16u) | ((u32(v.w) & 0xFFu) << 24u);
 {% endif %}
 }
 // Sum of the 4 bias-shifted A values of one packed word.
 fn sum4_a(v: vec4<{{ aScalar }}>) -> i32 {
@@ -47,6 +71,7 @@ fn sum4_a(v: vec4<{{ aScalar }}>) -> i32 {
 {% endif %}
 }
 fn pack_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> u32 {
 {% if bUnsigned %}
   return ((b0 ^ 0x80u) & 0xFFu) | (((b1 ^ 0x80u) & 0xFFu) << 8u) | (((b2 ^ 0x80u) & 0xFFu) << 16u) | (((b3 ^ 0x80u) & 0xFFu) << 24u);
@@ -60,7 +85,7 @@ fn pack_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScala
 fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
   return dot4I8Packed(a_word, b_word);
 }
 fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> i32 {
 {% if bUnsigned %}
@@ -70,30 +95,15 @@ fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScala
 {% endif %}
 }
-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 %}
-}
-fn read_b_zero() -> i32 {
-{% if bUnsigned %}
-  return i32(b_zero_point[0]);
-{% else %}
-  return b_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;
@@ -114,15 +124,31 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
   let k_words_total = params.K / 4u;
 {% endif %}
   if (tid < TILE_M) {
     atomicStore(&a_row_sum[tid], 0);
   }
   workgroupBarrier();
 {% for m in range(tileM) %}
   var r{{ m }}: i32 = 0;
 {% endfor %}
   var sum_b: i32 = 0;
   for (var kw0 = 0u; kw0 < k_words_total; kw0 = kw0 + KT_WORDS) {
     let words = min(KT_WORDS, k_words_total - kw0);
@@ -146,13 +172,49 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
         let v = a[a_base4 + src_m * k_words_total + kw0 + ki];
 {% endif %}
         packed = pack_a(v);
         atomicAdd(&a_row_sum[mi], sum4_a(v));
       }
       a_tile[mi * KT_WORDS + ki] = packed;
     }
     workgroupBarrier();
     if (n_valid) {
       for (var kk = 0u; kk < words; kk = kk + 1u) {
         let k4 = (kw0 + kk) * 4u;
         let b_row = b_base + k4 * params.N + col;
@@ -160,12 +222,15 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
         let b1 = b[b_row + params.N];
         let b2 = b[b_row + 2u * params.N];
         let b3 = b[b_row + 3u * params.N];
         sum_b = sum_b + sum4_b(b0, b1, b2, b3);
         let bp = pack_b(b0, b1, b2, b3);
 {% for m in range(tileM) %}
         r{{ m }} = r{{ m }} + dot4_packed(a_tile[{{ m }}u * KT_WORDS + kk], bp);
 {% endfor %}
       }
     }
     workgroupBarrier();
   }
@@ -190,7 +255,9 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
         let a3: {{ aScalar }} = {{ "128u" if aUnsigned else "0" }};
         let v = vec4<{{ aScalar }}>(a0, a1, a2, a3);
         packed = pack_a(v);
         atomicAdd(&a_row_sum[tid], sum4_a(v));
       }
       a_tile[tid * KT_WORDS] = packed;
     }
@@ -205,7 +272,9 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
       if (tail_lanes > 1u) { b1 = b[b_row + params.N]; }
       if (tail_lanes > 2u) { b2 = b[b_row + 2u * params.N]; }
       let b3: {{ bScalar }} = {{ "128u" if bUnsigned else "0" }};
       sum_b = sum_b + sum4_b(b0, b1, b2, b3);
       let bp = pack_b(b0, b1, b2, b3);
 {% for m in range(tileM) %}
       r{{ m }} = r{{ m }} + dot4_packed(a_tile[{{ m }}u * KT_WORDS], bp);
@@ -220,11 +289,14 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
   }
   // Fold bias + zero-point corrections (exact in modular i32 arithmetic).
-  let cb: i32 = {{ "128" if bUnsigned else "0" }} - read_b_zero();
 {% if not aZeroPerRow %}
-  let ca: i32 = {{ "128" if aUnsigned else "0" }} - read_a_zero();
   let kcc: i32 = i32(params.K) * ca * cb;
 {% endif %}
 {% for m in range(tileM) %}
   {
@@ -234,12 +306,16 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
       let ca: i32 = {{ "128" if aUnsigned else "0" }} - read_a_zero(mr);
       let kcc: i32 = i32(params.K) * ca * cb;
 {% endif %}
-      var acc = r{{ m }};
       acc = acc + cb * atomicLoad(&a_row_sum[{{ m }}u]);
-      acc = acc + ca * sum_b;
       acc = acc + kcc;
 {% if hasBias %}
-      acc = acc + bias[mr];
 {% endif %}
       y[y_base + mr * params.N + col] = acc;
     }

+{% macro quant_read_column_zero(fn_name, buffer_name, is_unsigned, present=true, per_column=false, param_name="col", return_type="i32") %}
+fn {{ fn_name }}({% if per_column %}{{ param_name }}: u32{% endif %}) -> {{ return_type }} {
+  {% if per_column %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[{{ param_name }}]);
+  {% else %}
+  return {{ buffer_name }}[{{ param_name }}];
+  {% endif %}
+  {% else %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[0]);
+  {% else %}
+  return {{ buffer_name }}[0];
+  {% endif %}
+  {% endif %}
+}{% endmacro %}
 {% if not env.device.wgslLanguageFeatures.has("packed_4x8_integer_dot_product") %}
 // Runtime source-compatible fallbacks for the packed 4x8 integer built-ins used
 // by quantized kernels. Native implementations remain unshadowed when the
 {% 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 = 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 %}
   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 {
 {% 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);
 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 %}
 {% endif %}
 }
 {% endif %}
+{% if hasAZero %}
+{{ quant_read_column_zero("read_a_zero", "a_zero_point", aUnsigned, per_column=aZeroPerRow, param_name="row") }}
 {% endif %}
+{% if hasBZero %}
+{{ quant_read_column_zero("read_b_zero", "b_zero_point", bUnsigned, per_column=bZeroPerColumn) }}
+{% 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 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);
         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 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();
   }
         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;
     }
       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);
   }
   // Fold bias + zero-point corrections (exact in modular i32 arithmetic).
+  let cb: i32 = {{ "128" if bUnsigned else "0" }}{% if hasBZero %} - read_b_zero(){% 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 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[{{ "col" if biasPerColumn is defined and biasPerColumn else "mr" }}];
 {% endif %}
       y[y_base + mr * params.N + col] = acc;
     }

build/webgpu/test.json CHANGED Viewed

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