sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +16 -13
build/webgpu/bench.json +0 -1
build/webgpu/manifest.json +256 -646
build/webgpu/metadata.json +28 -9
build/webgpu/quant-linear-blocked-axis.wgsl.jinja +47 -16
build/webgpu/quant-linear-scalar.wgsl.jinja +9 -8
build/webgpu/quant-linear-vec4.wgsl.jinja +17 -16
build/webgpu/test.json +7 -8

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX `QuantizeLinear` spec](https://onnx.ai/onnx/operators/onnx__Quanti
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `x` | `x` | `TX` | — | — | N-D full-precision input tensor to be quantized. | required |
-| `y_scale` | `y_scale` | `TS` | — | — | Scale factor; scalar for per-tensor, 1-D for per-axis, or same rank as `x` (with one axis blocked) for blocked quantization. | required |
-| `y_zero_point` | `y_zero_point` | `TQ` | — | — | Zero point for quantization; must have the same shape as `y_scale`. Defaults to zero if omitted. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `y` | `y` | `TQ` | same as `x` | same as `x` | N-D quantized output tensor with the same shape as `x`. | required |
 ## Attributes
@@ -52,7 +52,7 @@ Default values (overridable per request):
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
@@ -62,15 +62,18 @@ Default values (overridable per request):
 ## Use with `@huggingface/kernels`
-The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
-The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
-- `y`
-Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `x` | `TX` | — | — | N-D full-precision input tensor to be quantized. | required |
+| `y_scale` | `TS` | — | — | Scale factor; scalar for per-tensor, 1-D for per-axis, or same rank as `x` (with one axis blocked) for blocked quantization. | required |
+| `y_zero_point` | `TQ` | — | — | Zero point for quantization; must have the same shape as `y_scale`. Defaults to zero if omitted. | optional |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `y` | `TQ` | same as `x` | same as `x` | N-D quantized output tensor with the same shape as `x`. | required |
 ## Attributes
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
+This example supplies explicit metadata for:
+- `y`
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.QuantizeLinear",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

build/webgpu/manifest.json CHANGED Viewed

@@ -2,51 +2,26 @@
   "domain": "ai.onnx",
   "name": "QuantizeLinear",
   "sinceVersion": 25,
-  "description": "Linearly quantizes a high-precision tensor to a lower-precision integer type using the formula `y = saturate((x / y_scale) + y_zero_point)`, with rounding to nearest even. Supports per-tensor, per-axis, and blocked quantization granularities determined by the shape of `y_scale`.",
-  "inputs": [
-    { "role": "x", "dtype": "TX", "description": "N-D full-precision input tensor to be quantized." },
-    {
-      "role": "y_scale",
-      "dtype": "TS",
-      "description": "Scale factor; scalar for per-tensor, 1-D for per-axis, or same rank as `x` (with one axis blocked) for blocked quantization."
-    },
-    {
-      "role": "y_zero_point",
-      "dtype": "TQ",
-      "optional": true,
-      "description": "Zero point for quantization; must have the same shape as `y_scale`. Defaults to zero if omitted."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "y",
-      "dtype": "TQ",
-      "rank": "ranks.x",
-      "description": "N-D quantized output tensor with the same shape as `x`.",
-      "shape": "shapes.x"
-    }
-  ],
-  "attributes": { "axis": 1, "block_size": 0, "output_dtype": 0, "precision": 0, "saturate": 1 },
-  "attributeDescriptions": {
-    "axis": "Axis of the quantization dimension in `x`, used for per-axis and blocked quantization; negative values count from the end.",
-    "block_size": "Number of elements along `axis` that share a single scale value for blocked quantization; 0 means blocked quantization is not used.",
-    "output_dtype": "ONNX TensorProto element-type code for `y`; 0 infers the type from `y_zero_point`, or uint8 when the zero point is omitted.",
-    "precision": "ONNX TensorProto element-type code used for `x / y_scale`; `0` uses the dtype of `y_scale`, `1` selects FLOAT, and `10` selects FLOAT16.",
-    "saturate": "Controls out-of-range conversion for float8 outputs. The implemented int8/uint8 subset accepts the ONNX default `1`."
   },
   "attributeConstraints": { "precision": { "values": [0, 1, 10] }, "saturate": { "values": [1] } },
   "typeConstraints": { "TX": ["float32", "float16"], "TS": ["float32", "float16"], "TQ": ["uint8", "int8"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "x", "role": "input" },
-    "y_scale": { "kind": "tensor", "semantic": "y_scale", "role": "input" },
-    "y_zero_point": { "kind": "tensor", "semantic": "y_zero_point", "role": "input", "required": false },
-    "y": { "kind": "tensor", "semantic": "y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256, "VEC4_TAIL_MIN_ELEMENTS": 4096 },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "workgroupOk": "tunables.WORKGROUP_SIZE <= deviceWorkgroupCap",
-    "outputDispatchFits": "workgroupOk and ceilDiv(ceilDiv(numel(shapes.y), tunables.WORKGROUP_SIZE), device.limits.maxComputeWorkgroupsPerDimension) <= device.limits.maxComputeWorkgroupsPerDimension",
     "outputDtypeOk": "(attrs.output_dtype == 0 and (present.y_zero_point or tensorDtypes.y == \"uint8\")) or attrs.output_dtype == onnxDtypeCode(logicalDtypes.TQ)",
     "quantizeDivisionF16": "attrs.precision == onnxDtypeCode(\"float16\") or (attrs.precision == 0 and tensorDtypes.y_scale == \"float16\")",
     "sameShapeOk": "ranks.y == ranks.x and numel(shapes.x) == numel(shapes.y) and outputDispatchFits and outputDtypeOk",
@@ -58,13 +33,8 @@
     "scaleInner": "1 if ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 else inner(shapes.x, attrs.axis)",
     "blockedAxisDim": "dim(shapes.x, attrs.axis) if ranks.x >= 2 else 1",
     "blockedScaleAxisDim": "dim(shapes.y_scale, attrs.axis) if ranks.y_scale >= 2 else 1",
-    "blockedInner": "inner(shapes.x, attrs.axis) if ranks.x >= 2 else 1"
-  },
-  "constants": {
-    "usesF16": "dtypes.TX == \"f16\" or dtypes.TS == \"f16\" or quantizeDivisionF16",
-    "xScalar": "dtypes.TX",
     "xVec4": "\"vec4<\" ~ dtypes.TX ~ \">\"",
-    "scaleScalar": "dtypes.TS",
     "scaleVec4": "\"vec4<\" ~ dtypes.TS ~ \">\"",
     "yScalar": "dtypes.TQ",
     "yVec4": "\"vec4<\" ~ dtypes.TQ ~ \">\"",
@@ -73,767 +43,407 @@
     "qMax": "255 if tensorDtypes.y == \"uint8\" else 127",
     "divisionF16": "quantizeDivisionF16"
   },
   "variants": [
     {
       "id": "innermost_axis_vec4_with_zero_point",
       "priority": 20,
-      "when": ["present.y_zero_point", "ranks.x >= 1", "attrs.axis == -1 or attrs.axis == ranks.x - 1", "ranks.y_scale == 1", "dim(shapes.y_scale, 0) == dim(shapes.x, ranks.x - 1)", "dim(shapes.y_scale, 0) % 4 == 0", "ranks.y_zero_point == 1", "dim(shapes.y_zero_point, 0) == dim(shapes.y_scale, 0)", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "sameShapeOk", "quantDtypesOk"],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.InnermostAxisVec4",
-          "source": {
-            "shader": "quant-linear-vec4.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "perAxis": true, "vectorParams": true }
-          },
-          "bindings": "linearInnermostVec4WithZero",
-          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "innermost_axis_vec4_no_zero_point",
       "priority": 20,
-      "when": ["not present.y_zero_point", "ranks.x >= 1", "attrs.axis == -1 or attrs.axis == ranks.x - 1", "ranks.y_scale == 1", "dim(shapes.y_scale, 0) == dim(shapes.x, ranks.x - 1)", "dim(shapes.y_scale, 0) % 4 == 0", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "sameShapeOk", "quantDtypesOk"],
-      "constants": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.InnermostAxisVec4NoZero",
-          "source": {
-            "shader": "quant-linear-vec4.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "perAxis": true, "vectorParams": true }
-          },
-          "bindings": "linearInnermostVec4NoZero",
-          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "vec4_cross_axis_with_zero_point",
       "priority": 16,
-      "when": ["present.y_zero_point", "attrs.block_size == 0", "ranks.x >= 1", "attrs.axis >= 0", "attrs.axis < ranks.x", "ranks.y_scale == 1", "dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis)", "inner(shapes.x, attrs.axis) % 4 != 0", "ranks.y_zero_point == 1", "dim(shapes.y_zero_point, 0) == dim(shapes.y_scale, 0)", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "sameShapeOk", "quantDtypesOk"],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Vec4CrossAxis",
-          "source": {
-            "shader": "quant-linear-vec4.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "perAxis": true, "crossingParams": true }
-          },
-          "bindings": "linearVec4WithZero",
-          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "vec4_with_zero_point",
       "priority": 15,
-      "when": ["present.y_zero_point", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "sameShapeOk", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis) and inner(shapes.x, attrs.axis) % 4 == 0))", "(ranks.y_zero_point == 0 or ranks.y_zero_point == 1)", "(ranks.y_zero_point == 0 or dim(shapes.y_zero_point, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_zero_point, 0) == dim(shapes.x, attrs.axis)))", "quantDtypesOk"],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Vec4",
-          "source": {
-            "shader": "quant-linear-vec4.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)" }
-          },
-          "bindings": "linearVec4WithZero",
-          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "vec4_no_zero_point",
       "priority": 15,
-      "when": ["not present.y_zero_point", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "sameShapeOk", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis) and inner(shapes.x, attrs.axis) % 4 == 0))", "quantDtypesOk"],
-      "constants": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Vec4",
-          "source": {
-            "shader": "quant-linear-vec4.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)" }
-          },
-          "bindings": "linearVec4NoZero",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "vec4_tail_with_zero_point",
       "priority": 12,
-      "when": ["present.y_zero_point", "numel(shapes.y) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.y) % 4 != 0", "(ranks.y_scale == 0 or (ranks.y_scale == 1 and dim(shapes.y_scale, 0) == 1))", "(ranks.y_zero_point == 0 or (ranks.y_zero_point == 1 and dim(shapes.y_zero_point, 0) == 1))", "sameShapeOk", "quantDtypesOk"],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "bulk",
           "name": "QuantizeLinear.Vec4Bulk",
-          "source": { "shader": "quant-linear-vec4.wgsl.jinja", "inputs": { "op": "\"quantize\"", "perAxis": false } },
-          "bindings": "linearVec4PerTensorWithZero",
-          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         },
         {
           "id": "tail",
           "name": "QuantizeLinear.ScalarTail",
-          "source": {
-            "shader": "quant-linear-scalar.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "x4": true, "perAxis": false }
-          },
-          "bindings": "linearScalarPerTensorWithZero",
-          "dispatch": { "threads": 1, "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "vec4_tail_no_zero_point",
       "priority": 12,
-      "when": ["not present.y_zero_point", "numel(shapes.y) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.y) % 4 != 0", "(ranks.y_scale == 0 or (ranks.y_scale == 1 and dim(shapes.y_scale, 0) == 1))", "sameShapeOk", "quantDtypesOk"],
-      "constants": { "hasZero": false },
       "passes": [
         {
           "id": "bulk",
           "name": "QuantizeLinear.Vec4Bulk",
-          "source": { "shader": "quant-linear-vec4.wgsl.jinja", "inputs": { "op": "\"quantize\"", "perAxis": false } },
-          "bindings": "linearVec4PerTensorNoZero",
-          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         },
         {
           "id": "tail",
           "name": "QuantizeLinear.ScalarTail",
-          "source": {
-            "shader": "quant-linear-scalar.wgsl.jinja",
-            "inputs": { "op": "\"quantize\"", "x4": true, "perAxis": false }
-          },
-          "bindings": "linearScalarPerTensorNoZero",
-          "dispatch": { "threads": 1, "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "blocked_last_axis_vec4_with_zero_point",
       "priority": 24,
-      "when": ["present.y_zero_point", "blockedScaleOk", "attrs.block_size % 4 == 0", "ranks.y_zero_point == ranks.x", "sameShapeOk", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "inner(shapes.x, attrs.axis) == 1", "dim(shapes.x, attrs.axis) % attrs.block_size == 0", "dim(shapes.y_scale, attrs.axis) == dim(shapes.x, attrs.axis) / attrs.block_size", "numel(shapes.y_zero_point) == numel(shapes.y_scale)", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
-      "constants": { "hasZero": true, "blockVectors": "attrs.block_size / 4" },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedLastAxisVec4",
-          "source": { "shader": "quant-linear-blocked-axis.wgsl.jinja", "inputs": { "lastAxisVectorized": true } },
-          "bindings": "blockedLastVec4WithZero",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "blocked_last_axis_vec4_no_zero_point",
       "priority": 23,
-      "when": ["not present.y_zero_point", "blockedScaleOk", "attrs.block_size % 4 == 0", "sameShapeOk", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "inner(shapes.x, attrs.axis) == 1", "dim(shapes.x, attrs.axis) % attrs.block_size == 0", "dim(shapes.y_scale, attrs.axis) == dim(shapes.x, attrs.axis) / attrs.block_size", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
-      "constants": { "hasZero": false, "blockVectors": "attrs.block_size / 4" },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedLastAxisVec4",
-          "source": { "shader": "quant-linear-blocked-axis.wgsl.jinja", "inputs": { "lastAxisVectorized": true } },
-          "bindings": "blockedLastVec4NoZero",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "blocked_vec4_with_zero_point",
       "priority": 22,
-      "when": ["present.y_zero_point", "blockedScaleOk", "ranks.y_zero_point == ranks.x", "sameShapeOk", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "numel(shapes.y_scale) % 4 == 0", "numel(shapes.y_zero_point) == numel(shapes.y_scale)", "inner(shapes.x, attrs.axis) % 4 == 0", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "dim(shapes.y_zero_point, attrs.axis) == dim(shapes.y_scale, attrs.axis)", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedVec4WithZeroPoint",
-          "source": { "shader": "quant-linear-blocked-axis.wgsl.jinja", "inputs": { "vectorized": true } },
-          "bindings": "blockedVec4WithZero",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "blocked_vec4_no_zero_point",
       "priority": 21,
-      "when": ["not present.y_zero_point", "blockedScaleOk", "sameShapeOk", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "numel(shapes.y_scale) % 4 == 0", "inner(shapes.x, attrs.axis) % 4 == 0", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
-      "constants": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedVec4",
-          "source": { "shader": "quant-linear-blocked-axis.wgsl.jinja", "inputs": { "vectorized": true } },
-          "bindings": "blockedVec4NoZero",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "blocked_with_zero_point",
       "priority": 19,
-      "when": ["present.y_zero_point", "blockedScaleOk", "ranks.y_zero_point == ranks.x", "sameShapeOk", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "dim(shapes.y_zero_point, attrs.axis) == dim(shapes.y_scale, attrs.axis)", "quantDtypesOk"],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Blocked",
           "shader": "quant-linear-blocked-axis.wgsl.jinja",
-          "bindings": "blockedScalarWithZero",
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "blocked_no_zero_point",
       "priority": 19,
-      "when": ["not present.y_zero_point", "blockedScaleOk", "sameShapeOk", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "quantDtypesOk"],
-      "constants": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedNoZero",
           "shader": "quant-linear-blocked-axis.wgsl.jinja",
-          "bindings": "blockedScalarNoZero",
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "with_zero_point",
       "priority": 10,
-      "when": ["present.y_zero_point", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "sameShapeOk", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis)))", "(ranks.y_zero_point == 0 or ranks.y_zero_point == 1)", "(ranks.y_zero_point == 0 or dim(shapes.y_zero_point, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_zero_point, 0) == dim(shapes.x, attrs.axis)))", "quantDtypesOk"],
-      "constants": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear",
-          "source": {
-            "shader": "quant-linear-scalar.wgsl.jinja",
-            "inputs": {
-              "op": "\"quantize\"",
-              "x4": "false",
-              "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)"
-            }
           },
-          "bindings": "linearScalarWithZero",
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "no_zero_point",
-      "when": ["not present.y_zero_point", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "sameShapeOk", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis)))", "quantDtypesOk"],
-      "constants": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear",
-          "source": {
-            "shader": "quant-linear-scalar.wgsl.jinja",
-            "inputs": {
-              "op": "\"quantize\"",
-              "x4": "false",
-              "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)"
-            }
           },
-          "bindings": "linearScalarNoZero",
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }
-  ],
-  "bindingSets": {
-    "linearVec4WithZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count4", "type": "u32", "value": "elementCount4" },
-            { "name": "scaleSize", "type": "u32", "value": "scaleSize" },
-            { "name": "inner", "type": "u32", "value": "scaleInner" }
-          ]
-        }
-      }
-    ],
-    "linearInnermostVec4WithZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleVec4"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yVec4"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count4", "type": "u32", "value": "elementCount4" },
-            { "name": "scaleSize", "type": "u32", "value": "scaleSize" }
-          ]
-        }
-      }
-    ],
-    "linearInnermostVec4NoZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleVec4"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count4", "type": "u32", "value": "elementCount4" },
-            { "name": "scaleSize", "type": "u32", "value": "scaleSize" }
-          ]
-        }
-      }
-    ],
-    "linearVec4NoZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count4", "type": "u32", "value": "elementCount4" },
-            { "name": "scaleSize", "type": "u32", "value": "scaleSize" },
-            { "name": "inner", "type": "u32", "value": "scaleInner" }
-          ]
-        }
-      }
-    ],
-    "linearScalarWithZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "elementCount" },
-            { "name": "scaleSize", "type": "u32", "value": "scaleSize" },
-            { "name": "inner", "type": "u32", "value": "scaleInner" }
-          ]
-        }
-      }
-    ],
-    "linearScalarNoZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "elementCount" },
-            { "name": "scaleSize", "type": "u32", "value": "scaleSize" },
-            { "name": "inner", "type": "u32", "value": "scaleInner" }
-          ]
-        }
-      }
-    ],
-    "blockedLastVec4WithZero": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "x",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "vec4<f32>"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "elementCount4" }] }
-      }
-    ],
-    "blockedLastVec4NoZero": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "x",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "vec4<f32>"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "elementCount4" }] }
-      }
-    ],
-    "blockedVec4WithZero": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "x",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "vec4<f32>"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "vec4<f32>"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yVec4"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "elementCount4" },
-            { "name": "axisDim", "type": "u32", "value": "blockedAxisDim" },
-            { "name": "scaleAxisDim", "type": "u32", "value": "blockedScaleAxisDim" },
-            { "name": "inner", "type": "u32", "value": "blockedInner" },
-            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
-          ]
-        }
-      }
-    ],
-    "blockedVec4NoZero": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "x",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "vec4<f32>"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "vec4<f32>"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "elementCount4" },
-            { "name": "axisDim", "type": "u32", "value": "blockedAxisDim" },
-            { "name": "scaleAxisDim", "type": "u32", "value": "blockedScaleAxisDim" },
-            { "name": "inner", "type": "u32", "value": "blockedInner" },
-            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
-          ]
-        }
-      }
-    ],
-    "blockedScalarWithZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "elementCount" },
-            { "name": "axisDim", "type": "u32", "value": "blockedAxisDim" },
-            { "name": "scaleAxisDim", "type": "u32", "value": "blockedScaleAxisDim" },
-            { "name": "inner", "type": "u32", "value": "blockedInner" },
-            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
-          ]
-        }
-      }
-    ],
-    "blockedScalarNoZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "count", "type": "u32", "value": "elementCount" },
-            { "name": "axisDim", "type": "u32", "value": "blockedAxisDim" },
-            { "name": "scaleAxisDim", "type": "u32", "value": "blockedScaleAxisDim" },
-            { "name": "inner", "type": "u32", "value": "blockedInner" },
-            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
-          ]
-        }
-      }
-    ],
-    "linearVec4WithZeroIo": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" }
-    ],
-    "linearVec4NoZeroIo": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" }
-    ],
-    "linearScalarWithZeroIo": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" }
-    ],
-    "linearScalarNoZeroIo": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" }
-    ],
-    "linearVec4PerTensorWithZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count4", "type": "u32", "value": "elementCount4" }] }
-      }
-    ],
-    "linearVec4PerTensorNoZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yVec4" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count4", "type": "u32", "value": "elementCount4" }] }
-      }
-    ],
-    "linearScalarPerTensorWithZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$yScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "elementCount" }] }
-      }
-    ],
-    "linearScalarPerTensorNoZero": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scaleScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "$yScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "elementCount" }] }
-      }
-    ]
-  }
 }

   "domain": "ai.onnx",
   "name": "QuantizeLinear",
   "sinceVersion": 25,
+  "inputs": {
+    "x": { "dtype": "TX" },
+    "y_scale": { "dtype": "TS" },
+    "y_zero_point": { "dtype": "TQ", "optional": true }
+  },
+  "outputs": { "y": { "dtype": "TQ", "rank": "ranks.x", "shape": "shapes.x" } },
+  "attributes": {
+    "axis": { "default": 1 },
+    "block_size": { "default": 0 },
+    "output_dtype": { "default": 0 },
+    "precision": { "default": 0 },
+    "saturate": { "default": 1 }
   },
   "attributeConstraints": { "precision": { "values": [0, 1, 10] }, "saturate": { "values": [1] } },
   "typeConstraints": { "TX": ["float32", "float16"], "TS": ["float32", "float16"], "TQ": ["uint8", "int8"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 }, "VEC4_TAIL_MIN_ELEMENTS": { "default": 4096 } },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "workgroupOk": "tunables.WORKGROUP_SIZE <= deviceWorkgroupCap",
+    "outputDispatchFits": "workgroupOk and ceilDiv(ceilDiv(numel(shapes.y), tunables.WORKGROUP_SIZE), min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "outputDtypeOk": "(attrs.output_dtype == 0 and (present.y_zero_point or tensorDtypes.y == \"uint8\")) or attrs.output_dtype == onnxDtypeCode(logicalDtypes.TQ)",
     "quantizeDivisionF16": "attrs.precision == onnxDtypeCode(\"float16\") or (attrs.precision == 0 and tensorDtypes.y_scale == \"float16\")",
     "sameShapeOk": "ranks.y == ranks.x and numel(shapes.x) == numel(shapes.y) and outputDispatchFits and outputDtypeOk",
     "scaleInner": "1 if ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 else inner(shapes.x, attrs.axis)",
     "blockedAxisDim": "dim(shapes.x, attrs.axis) if ranks.x >= 2 else 1",
     "blockedScaleAxisDim": "dim(shapes.y_scale, attrs.axis) if ranks.y_scale >= 2 else 1",
+    "blockedInner": "inner(shapes.x, attrs.axis) if ranks.x >= 2 else 1",
     "xVec4": "\"vec4<\" ~ dtypes.TX ~ \">\"",
     "scaleVec4": "\"vec4<\" ~ dtypes.TS ~ \">\"",
     "yScalar": "dtypes.TQ",
     "yVec4": "\"vec4<\" ~ dtypes.TQ ~ \">\"",
     "qMax": "255 if tensorDtypes.y == \"uint8\" else 127",
     "divisionF16": "quantizeDivisionF16"
   },
+  "when": ["sameShapeOk"],
+  "bindings": {
+    "x": { "buffer": "read-only-storage", "elementType": "$xVec4" },
+    "y_scale": { "buffer": "read-only-storage", "elementType": "$scaleVec4" },
+    "y_zero_point": { "buffer": "read-only-storage", "elementType": "$yVec4" },
+    "y": { "buffer": "storage", "elementType": "$yVec4" },
+    "params": {
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count4", "type": "u32", "value": "elementCount4" },
+        { "name": "scaleSize", "type": "u32", "value": "scaleSize" }
+      ]
+    },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count4", "type": "u32", "value": "elementCount4" },
+        { "name": "scaleSize", "type": "u32", "value": "scaleSize" },
+        { "name": "inner", "type": "u32", "value": "scaleInner" }
+      ]
+    },
+    "params_3": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count4", "type": "u32", "value": "elementCount4" }]
+    },
+    "params_4": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "elementCount" }]
+    },
+    "x_3": { "name": "x", "buffer": "read-only-storage", "elementType": "vec4<f32>" },
+    "y_scale_3": { "name": "y_scale", "buffer": "read-only-storage", "elementType": "f32" },
+    "params_5": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "elementCount4" }]
+    },
+    "y_scale_4": { "name": "y_scale", "buffer": "read-only-storage", "elementType": "vec4<f32>" },
+    "params_6": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count", "type": "u32", "value": "elementCount4" },
+        { "name": "axisDim", "type": "u32", "value": "blockedAxisDim" },
+        { "name": "scaleAxisDim", "type": "u32", "value": "blockedScaleAxisDim" },
+        { "name": "inner", "type": "u32", "value": "blockedInner" },
+        { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
+      ]
+    },
+    "params_7": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count", "type": "u32", "value": "elementCount" },
+        { "name": "axisDim", "type": "u32", "value": "blockedAxisDim" },
+        { "name": "scaleAxisDim", "type": "u32", "value": "blockedScaleAxisDim" },
+        { "name": "inner", "type": "u32", "value": "blockedInner" },
+        { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
+      ]
+    },
+    "params_8": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        { "name": "count", "type": "u32", "value": "elementCount" },
+        { "name": "scaleSize", "type": "u32", "value": "scaleSize" },
+        { "name": "inner", "type": "u32", "value": "scaleInner" }
+      ]
+    }
+  },
   "variants": [
     {
       "id": "innermost_axis_vec4_with_zero_point",
       "priority": 20,
+      "when": ["present.y_zero_point", "ranks.x >= 1", "attrs.axis == -1 or attrs.axis == ranks.x - 1", "ranks.y_scale == 1", "dim(shapes.y_scale, 0) == dim(shapes.x, ranks.x - 1)", "dim(shapes.y_scale, 0) % 4 == 0", "ranks.y_zero_point == 1", "dim(shapes.y_zero_point, 0) == dim(shapes.y_scale, 0)", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "quantDtypesOk"],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.InnermostAxisVec4",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": true, "vectorParams": true },
+          "bindings": ["x", "y_scale", "y_zero_point", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "innermost_axis_vec4_no_zero_point",
       "priority": 20,
+      "when": ["not present.y_zero_point", "ranks.x >= 1", "attrs.axis == -1 or attrs.axis == ranks.x - 1", "ranks.y_scale == 1", "dim(shapes.y_scale, 0) == dim(shapes.x, ranks.x - 1)", "dim(shapes.y_scale, 0) % 4 == 0", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "quantDtypesOk"],
+      "derive": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.InnermostAxisVec4NoZero",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": true, "vectorParams": true },
+          "bindings": ["x", "y_scale", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_cross_axis_with_zero_point",
       "priority": 16,
+      "when": ["present.y_zero_point", "attrs.block_size == 0", "ranks.x >= 1", "attrs.axis >= 0", "attrs.axis < ranks.x", "ranks.y_scale == 1", "dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis)", "inner(shapes.x, attrs.axis) % 4 != 0", "ranks.y_zero_point == 1", "dim(shapes.y_zero_point, 0) == dim(shapes.y_scale, 0)", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "quantDtypesOk"],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Vec4CrossAxis",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": true, "crossingParams": true },
+          "bindings": ["x", { "arg": "y_scale" }, { "arg": "y_zero_point" }, "y", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_with_zero_point",
       "priority": 15,
+      "when": ["present.y_zero_point", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis) and inner(shapes.x, attrs.axis) % 4 == 0))", "(ranks.y_zero_point == 0 or ranks.y_zero_point == 1)", "(ranks.y_zero_point == 0 or dim(shapes.y_zero_point, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_zero_point, 0) == dim(shapes.x, attrs.axis)))", "quantDtypesOk"],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Vec4",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)" },
+          "bindings": ["x", { "arg": "y_scale" }, { "arg": "y_zero_point" }, "y", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_no_zero_point",
       "priority": 15,
+      "when": ["not present.y_zero_point", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis) and inner(shapes.x, attrs.axis) % 4 == 0))", "quantDtypesOk"],
+      "derive": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Vec4",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)" },
+          "bindings": ["x", { "arg": "y_scale" }, "y", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_tail_with_zero_point",
       "priority": 12,
+      "when": ["present.y_zero_point", "numel(shapes.y) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.y) % 4 != 0", "(ranks.y_scale == 0 or (ranks.y_scale == 1 and dim(shapes.y_scale, 0) == 1))", "(ranks.y_zero_point == 0 or (ranks.y_zero_point == 1 and dim(shapes.y_zero_point, 0) == 1))", "quantDtypesOk"],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "bulk",
           "name": "QuantizeLinear.Vec4Bulk",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": false },
+          "bindings": ["x", { "arg": "y_scale" }, { "arg": "y_zero_point" }, "y", "params_3"],
+          "dispatch": {
+            "x": "min(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         },
         {
           "id": "tail",
           "name": "QuantizeLinear.ScalarTail",
+          "shader": "quant-linear-scalar.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "x4": true, "perAxis": false },
+          "bindings": [{ "arg": "x" }, { "arg": "y_scale" }, { "arg": "y_zero_point" }, { "arg": "y" }, "params_4"],
+          "dispatch": {
+            "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_tail_no_zero_point",
       "priority": 12,
+      "when": ["not present.y_zero_point", "numel(shapes.y) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.y) % 4 != 0", "(ranks.y_scale == 0 or (ranks.y_scale == 1 and dim(shapes.y_scale, 0) == 1))", "quantDtypesOk"],
+      "derive": { "hasZero": false },
       "passes": [
         {
           "id": "bulk",
           "name": "QuantizeLinear.Vec4Bulk",
+          "shader": "quant-linear-vec4.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "perAxis": false },
+          "bindings": ["x", { "arg": "y_scale" }, "y", "params_3"],
+          "dispatch": {
+            "x": "min(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((floor(numel(shapes.y) / 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         },
         {
           "id": "tail",
           "name": "QuantizeLinear.ScalarTail",
+          "shader": "quant-linear-scalar.wgsl.jinja",
+          "derive": { "op": "\"quantize\"", "x4": true, "perAxis": false },
+          "bindings": [{ "arg": "x" }, { "arg": "y_scale" }, { "arg": "y" }, "params_4"],
+          "dispatch": {
+            "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "blocked_last_axis_vec4_with_zero_point",
       "priority": 24,
+      "when": ["present.y_zero_point", "blockedScaleOk", "attrs.block_size % 4 == 0", "ranks.y_zero_point == ranks.x", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "inner(shapes.x, attrs.axis) == 1", "dim(shapes.x, attrs.axis) % attrs.block_size == 0", "dim(shapes.y_scale, attrs.axis) == dim(shapes.x, attrs.axis) / attrs.block_size", "numel(shapes.y_zero_point) == numel(shapes.y_scale)", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
+      "derive": { "hasZero": true, "blockVectors": "attrs.block_size / 4" },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedLastAxisVec4",
+          "shader": "quant-linear-blocked-axis.wgsl.jinja",
+          "derive": { "lastAxisVectorized": true },
+          "bindings": ["x_3", "y_scale_3", { "arg": "y_zero_point" }, "y", "params_5"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "blocked_last_axis_vec4_no_zero_point",
       "priority": 23,
+      "when": ["not present.y_zero_point", "blockedScaleOk", "attrs.block_size % 4 == 0", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "inner(shapes.x, attrs.axis) == 1", "dim(shapes.x, attrs.axis) % attrs.block_size == 0", "dim(shapes.y_scale, attrs.axis) == dim(shapes.x, attrs.axis) / attrs.block_size", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
+      "derive": { "hasZero": false, "blockVectors": "attrs.block_size / 4" },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedLastAxisVec4",
+          "shader": "quant-linear-blocked-axis.wgsl.jinja",
+          "derive": { "lastAxisVectorized": true },
+          "bindings": ["x_3", "y_scale_3", "y", "params_5"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "blocked_vec4_with_zero_point",
       "priority": 22,
+      "when": ["present.y_zero_point", "blockedScaleOk", "ranks.y_zero_point == ranks.x", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "numel(shapes.y_scale) % 4 == 0", "numel(shapes.y_zero_point) == numel(shapes.y_scale)", "inner(shapes.x, attrs.axis) % 4 == 0", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "dim(shapes.y_zero_point, attrs.axis) == dim(shapes.y_scale, attrs.axis)", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedVec4WithZeroPoint",
+          "shader": "quant-linear-blocked-axis.wgsl.jinja",
+          "derive": { "vectorizedSpec": true },
+          "bindings": ["x_3", "y_scale_4", "y_zero_point", "y", "params_6"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "blocked_vec4_no_zero_point",
       "priority": 21,
+      "when": ["not present.y_zero_point", "blockedScaleOk", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "numel(shapes.y_scale) % 4 == 0", "inner(shapes.x, attrs.axis) % 4 == 0", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "tensorDtypes.x == \"float32\"", "tensorDtypes.y_scale == \"float32\""],
+      "derive": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedVec4",
+          "shader": "quant-linear-blocked-axis.wgsl.jinja",
+          "derive": { "vectorizedSpec": true },
+          "bindings": ["x_3", "y_scale_4", "y", "params_6"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "blocked_with_zero_point",
       "priority": 19,
+      "when": ["present.y_zero_point", "blockedScaleOk", "ranks.y_zero_point == ranks.x", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "dim(shapes.y_zero_point, attrs.axis) == dim(shapes.y_scale, attrs.axis)", "quantDtypesOk"],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.Blocked",
           "shader": "quant-linear-blocked-axis.wgsl.jinja",
+          "bindings": [{ "arg": "x" }, { "arg": "y_scale" }, { "arg": "y_zero_point" }, { "arg": "y" }, "params_7"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "blocked_no_zero_point",
       "priority": 19,
+      "when": ["not present.y_zero_point", "blockedScaleOk", "dim(shapes.y_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)", "quantDtypesOk"],
+      "derive": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear.BlockedNoZero",
           "shader": "quant-linear-blocked-axis.wgsl.jinja",
+          "bindings": [{ "arg": "x" }, { "arg": "y_scale" }, { "arg": "y" }, "params_7"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "with_zero_point",
       "priority": 10,
+      "when": ["present.y_zero_point", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis)))", "(ranks.y_zero_point == 0 or ranks.y_zero_point == 1)", "(ranks.y_zero_point == 0 or dim(shapes.y_zero_point, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_zero_point, 0) == dim(shapes.x, attrs.axis)))", "quantDtypesOk"],
+      "derive": { "hasZero": true },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear",
+          "shader": "quant-linear-scalar.wgsl.jinja",
+          "derive": {
+            "op": "\"quantize\"",
+            "x4": "false",
+            "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)"
           },
+          "bindings": [{ "arg": "x" }, { "arg": "y_scale" }, { "arg": "y_zero_point" }, { "arg": "y" }, "params_8"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "no_zero_point",
+      "when": ["not present.y_zero_point", "(ranks.y_scale == 0 or ranks.y_scale == 1)", "(ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.y_scale, 0) == dim(shapes.x, attrs.axis)))", "quantDtypesOk"],
+      "derive": { "hasZero": false },
       "passes": [
         {
           "id": "main",
           "name": "QuantizeLinear",
+          "shader": "quant-linear-scalar.wgsl.jinja",
+          "derive": {
+            "op": "\"quantize\"",
+            "x4": "false",
+            "perAxis": "not (ranks.y_scale == 0 or dim(shapes.y_scale, 0) == 1)"
           },
+          "bindings": [{ "arg": "x" }, { "arg": "y_scale" }, { "arg": "y" }, "params_8"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }
+  ]
 }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,39 @@
 {
   "name": "ai.onnx.QuantizeLinear",
-  "id": "_ai_onnx_quantizelinear_webgpu_a6a6d0c",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "dX5NKjv8Gz9VheT2QY/86ZiiTTguwW4MpDCtBw9kMHg=",
-      "manifest.json": "qDp2PlIs4E2AwjrcPIjItloZ2fWB900UI9LrGCs1PYg=",
-      "quant-linear-blocked-axis.wgsl.jinja": "964aieHZ1Jm6oeMgBxT1x9i3tEhFRUFJozYikdEIMwg=",
-      "quant-linear-scalar.wgsl.jinja": "191IHMqakZ9HaVXxPSwNl15B6ec70DLhx40LuSbeu30=",
-      "quant-linear-vec4.wgsl.jinja": "y48X/TVAaC+b/DD/AQT2ap35JVn7+KrdzlRrxrF8RDA=",
-      "test.json": "QV/it95sQgUj+dQu4OK/chmXHLuwBAsDJQc6f7RISfM="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.QuantizeLinear" }
 }

 {
   "name": "ai.onnx.QuantizeLinear",
+  "id": "_ai_onnx_quantizelinear_webgpu_bf60a95",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "B6CqLHGCbyXXUAAcatPURKRkCnRMOZCNu7ix44MZbZw=",
+      "manifest.json": "X5UX5Db4+s1nUDXg5jez+yaYuFnt+y5vPqCc1mWXy3I=",
+      "quant-linear-blocked-axis.wgsl.jinja": "fhTtN3vZ8jl4fTshwThxKYds9QrAZ9U2GzjPtWQs2OE=",
+      "quant-linear-scalar.wgsl.jinja": "/OPcCfkXBiT0dEm1wEnImi5EBY8l17n6jpS5KKCjvQ4=",
+      "quant-linear-vec4.wgsl.jinja": "qQrnnz6UMx0Pg/d7Cs2cjh1HBaurycYK0H24VIbq+Fs=",
+      "test.json": "cniKZLfeclIx3sXUNY1qFbwal5ui2JqFKfODpgZEoCA="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "innermost_axis_vec4_with_zero_point": ["quant-linear-vec4.wgsl.jinja"],
+      "innermost_axis_vec4_no_zero_point": ["quant-linear-vec4.wgsl.jinja"],
+      "vec4_cross_axis_with_zero_point": ["quant-linear-vec4.wgsl.jinja"],
+      "vec4_with_zero_point": ["quant-linear-vec4.wgsl.jinja"],
+      "vec4_no_zero_point": ["quant-linear-vec4.wgsl.jinja"],
+      "vec4_tail_with_zero_point": ["quant-linear-scalar.wgsl.jinja", "quant-linear-vec4.wgsl.jinja"],
+      "vec4_tail_no_zero_point": ["quant-linear-scalar.wgsl.jinja", "quant-linear-vec4.wgsl.jinja"],
+      "blocked_last_axis_vec4_with_zero_point": ["quant-linear-blocked-axis.wgsl.jinja"],
+      "blocked_last_axis_vec4_no_zero_point": ["quant-linear-blocked-axis.wgsl.jinja"],
+      "blocked_vec4_with_zero_point": ["quant-linear-blocked-axis.wgsl.jinja"],
+      "blocked_vec4_no_zero_point": ["quant-linear-blocked-axis.wgsl.jinja"],
+      "blocked_with_zero_point": ["quant-linear-blocked-axis.wgsl.jinja"],
+      "blocked_no_zero_point": ["quant-linear-blocked-axis.wgsl.jinja"],
+      "with_zero_point": ["quant-linear-scalar.wgsl.jinja"],
+      "no_zero_point": ["quant-linear-scalar.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/quant-linear-blocked-axis.wgsl.jinja CHANGED Viewed

@@ -1,12 +1,35 @@
-// Blocked-axis scale indexing for {{ source.op | default("quantize") }}.
-{% set vectorized = source.vectorized if source.vectorized is defined else false %}
-{% set lastAxisVectorized = source.lastAxisVectorized if source.lastAxisVectorized is defined else false %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% if not vectorized and not lastAxisVectorized %}
 fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
 {% if hasZero %}
 {% if yUnsigned %}
@@ -37,6 +60,7 @@ fn scale_index(index: u32) -> u32 {
 }
 {% endif %}
 {% set divisionF16 = divisionF16 is defined and divisionF16 %}
 // Exact ONNX QuantizeLinear round-to-nearest-even, with identical handling of
 // infinities, saturation, and halfway values across every kernel route.
@@ -63,7 +87,8 @@ fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
   // Clamp before the i32 cast so infinite and huge finite inputs saturate
   // instead of invoking undefined conversion behavior. Compare distances in
   // the input domain: doing the comparison on value / scale can move a value
-  // across a half-way boundary because GPU division is not correctly rounded.
   let v = clamp(value / scale, -2.0e9, 2.0e9);
   let fl = floor(v);
   let lo = fl;
@@ -83,12 +108,13 @@ fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
 {% endif %}
 {% if lastAxisVectorized %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // A vec4 cannot cross a block boundary on this route. Map vector indices
   // directly to their shared scale/zero-point entry.
-  let i4 = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (i4 >= params.count) { return; }
   let scale_index = i4 / {{ blockVectors }}u;
   let values = x[i4];
@@ -116,9 +142,9 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 }
 {% elif vectorized %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // Four adjacent inner elements map to four adjacent scale entries.
-  let i4 = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (i4 >= params.count) { return; }
   let scale4 = scale_index(i4 * 4u) / 4u;
   let values = x[i4];
@@ -144,13 +170,18 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 }
 {% else %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
-  // The flat dispatch is folded across x/y at the device's per-axis workgroup
-  // limit; gid.y carries the high portion of the element index.
-  let index = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (index >= params.count) { return; }
   let scaleIndex = scale_index(index);
   let rounded = round_scaled_half_to_even(f32(x[index]), f32(y_scale[scaleIndex])) + read_zero({% if hasZero %}scaleIndex{% endif %});
   write_quantized(index, clamp(rounded, {{ qMin }}, {{ qMax }}));
 }
 {% endif %}

+// Blocked-axis scale and zero-point indexing for linear quantization.
+{% set operation = op if op is defined else "quantize" %}
+{% set vectorized = vectorizedSpec if vectorizedSpec is defined else false %}
+{% set lastAxisVectorized = lastAxisVectorized if lastAxisVectorized is defined else false %}
+{% set xUnsigned = xUnsigned is defined and xUnsigned %}
+{% set yUnsigned = yUnsigned is defined and yUnsigned %}
+{% set qMin = qMin | default(0 if yUnsigned else 0 - 128) %}
+{% set qMax = qMax | default(255 if yUnsigned else 127) %}
+{% set blockVectors = blockVectors | default(1) %}
 {{ env.wgsl.resourceDeclarations }}
+{% if operation == "dequantize" %}
+fn read_x(index: u32) -> i32 {
+{% if xUnsigned %}
+  return i32(x[index]);
+{% else %}
+  return x[index];
+{% endif %}
+}
+fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
+{% if hasZero %}
+{% if xUnsigned %}
+  return i32(x_zero_point[index]);
+{% else %}
+  return x_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+{% elif not vectorized and not lastAxisVectorized %}
 fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
 {% if hasZero %}
 {% if yUnsigned %}
 }
 {% endif %}
+{% if operation == "quantize" %}
 {% set divisionF16 = divisionF16 is defined and divisionF16 %}
 // Exact ONNX QuantizeLinear round-to-nearest-even, with identical handling of
 // infinities, saturation, and halfway values across every kernel route.
   // Clamp before the i32 cast so infinite and huge finite inputs saturate
   // instead of invoking undefined conversion behavior. Compare distances in
   // the input domain: doing the comparison on value / scale can move a value
+  // across a halfway boundary because WGSL division is not required to be
+  // correctly rounded.
   let v = clamp(value / scale, -2.0e9, 2.0e9);
   let fl = floor(v);
   let lo = fl;
 {% endif %}
+{% endif %}
 {% if lastAxisVectorized %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // A vec4 cannot cross a block boundary on this route. Map vector indices
   // directly to their shared scale/zero-point entry.
+  let i4 = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i4 >= params.count) { return; }
   let scale_index = i4 / {{ blockVectors }}u;
   let values = x[i4];
 }
 {% elif vectorized %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // Four adjacent inner elements map to four adjacent scale entries.
+  let i4 = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i4 >= params.count) { return; }
   let scale4 = scale_index(i4 * 4u) / 4u;
   let values = x[i4];
 }
 {% else %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the element index.
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (index >= params.count) { return; }
   let scaleIndex = scale_index(index);
+{% if operation == "quantize" %}
   let rounded = round_scaled_half_to_even(f32(x[index]), f32(y_scale[scaleIndex])) + read_zero({% if hasZero %}scaleIndex{% endif %});
   write_quantized(index, clamp(rounded, {{ qMin }}, {{ qMax }}));
+{% else %}
+  let result = f32(read_x(index) - read_zero({% if hasZero %}scaleIndex{% endif %})) * f32(x_scale[scaleIndex]);
+  y[index] = {{ yScalar }}(result);
+{% endif %}
 }
 {% endif %}

build/webgpu/quant-linear-scalar.wgsl.jinja CHANGED Viewed

@@ -1,8 +1,8 @@
 // The x4 route handles widened int8/uint8 storage with scalar scale parameters;
 // the linear route supports both per-tensor and per-axis quantization.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% set divisionF16 = divisionF16 is defined and divisionF16 %}
@@ -31,7 +31,8 @@ fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
   // Clamp before the i32 cast so infinite and huge finite inputs saturate
   // instead of invoking undefined conversion behavior. Compare distances in
   // the input domain: doing the comparison on value / scale can move a value
-  // across a half-way boundary because GPU division is not correctly rounded.
   let v = clamp(value / scale, -2.0e9, 2.0e9);
   let fl = floor(v);
   let lo = fl;
@@ -77,11 +78,11 @@ fn transform_one(index: u32, scale: f32, zero_point: i32) {
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // Fold oversized flat dispatches into two dimensions; gid.y carries work
   // beyond the device-capped x dimension.
-  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
-{% if source.x4 %}
   if (invocation != 0u) {
     return;
   }
@@ -100,7 +101,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
     return;
   }
   var scale_index = 0u;
-{% if source.perAxis %}
   scale_index = (invocation / params.inner) % params.scaleSize;
 {% endif %}
   let scale = f32(y_scale[scale_index]);

+{% set yUnsigned = yUnsigned is defined and yUnsigned %}
+{% set qMin = qMin | default(0 if yUnsigned else 0 - 128) %}
+{% set qMax = qMax | default(255 if yUnsigned else 127) %}
 // The x4 route handles widened int8/uint8 storage with scalar scale parameters;
 // the linear route supports both per-tensor and per-axis quantization.
 {{ env.wgsl.resourceDeclarations }}
 {% set divisionF16 = divisionF16 is defined and divisionF16 %}
   // Clamp before the i32 cast so infinite and huge finite inputs saturate
   // instead of invoking undefined conversion behavior. Compare distances in
   // the input domain: doing the comparison on value / scale can move a value
+  // across a halfway boundary because WGSL division is not required to be
+  // correctly rounded.
   let v = clamp(value / scale, -2.0e9, 2.0e9);
   let fl = floor(v);
   let lo = fl;
 }
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // Fold oversized flat dispatches into two dimensions; gid.y carries work
   // beyond the device-capped x dimension.
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+{% if x4 %}
   if (invocation != 0u) {
     return;
   }
     return;
   }
   var scale_index = 0u;
+{% if perAxis %}
   scale_index = (invocation / params.inner) % params.scaleSize;
 {% endif %}
   let scale = f32(y_scale[scale_index]);

build/webgpu/quant-linear-vec4.wgsl.jinja CHANGED Viewed

@@ -1,9 +1,9 @@
 // int8/uint8 tensors are stored widened (one u32/i32 per element), so the vec4
 // binding gives 128-bit loads/stores of four elements. Per-component arithmetic
 // remains identical to the scalar quantize/dequantize paths.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% set divisionF16 = divisionF16 is defined and divisionF16 %}
@@ -32,7 +32,8 @@ fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
   // Clamp before the i32 cast so infinite and huge finite inputs saturate
   // instead of invoking undefined conversion behavior. Compare distances in
   // the input domain: doing the comparison on value / scale can move a value
-  // across a half-way boundary because GPU division is not correctly rounded.
   let v = clamp(value / scale, -2.0e9, 2.0e9);
   let fl = floor(v);
   let lo = fl;
@@ -52,7 +53,7 @@ fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
 {% endif %}
-{% if source.vectorParams is defined and source.vectorParams %}
 fn read_zero4({% if hasZero %}index: u32{% endif %}) -> vec4<i32> {
 {% if hasZero %}
 {% if yUnsigned %}
@@ -79,31 +80,31 @@ fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
-  // The flat dispatch is folded across x/y at the device's per-axis workgroup
-  // limit; gid.y carries the high portion of the vector index.
-  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count4) {
     return;
   }
-{% if source.vectorParams is defined and source.vectorParams %}
   let scale_index = i % (params.scaleSize / 4u);
-{% elif source.crossingParams is defined and source.crossingParams %}
   let base_index = i * 4u;
   let scale_index0 = (base_index / params.inner) % params.scaleSize;
   let scale_index1 = ((base_index + 1u) / params.inner) % params.scaleSize;
   let scale_index2 = ((base_index + 2u) / params.inner) % params.scaleSize;
   let scale_index3 = ((base_index + 3u) / params.inner) % params.scaleSize;
-{% elif source.perAxis %}
   // inner % 4 == 0, so all 4 lanes share one scale index.
   let scale_index = ((i * 4u) / params.inner) % params.scaleSize;
 {% else %}
   let scale_index = 0u;
 {% endif %}
   let xv = x[i];
-{% if source.vectorParams is defined and source.vectorParams %}
   let zp4 = read_zero4({% if hasZero %}scale_index{% endif %});
-{% elif source.crossingParams is defined and source.crossingParams %}
   let zp0 = read_zero(scale_index0);
   let zp1 = read_zero(scale_index1);
   let zp2 = read_zero(scale_index2);
@@ -111,13 +112,13 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 {% else %}
   let zp = read_zero({% if hasZero %}scale_index{% endif %});
 {% endif %}
-{% if source.vectorParams is defined and source.vectorParams %}
   let scale = vec4<f32>(y_scale[scale_index]);
   let q0 = clamp(round_scaled_half_to_even(f32(xv.x), scale.x) + zp4.x, {{ qMin }}, {{ qMax }});
   let q1 = clamp(round_scaled_half_to_even(f32(xv.y), scale.y) + zp4.y, {{ qMin }}, {{ qMax }});
   let q2 = clamp(round_scaled_half_to_even(f32(xv.z), scale.z) + zp4.z, {{ qMin }}, {{ qMax }});
   let q3 = clamp(round_scaled_half_to_even(f32(xv.w), scale.w) + zp4.w, {{ qMin }}, {{ qMax }});
-{% elif source.crossingParams is defined and source.crossingParams %}
   let q0 = clamp(round_scaled_half_to_even(f32(xv.x), f32(y_scale[scale_index0])) + zp0, {{ qMin }}, {{ qMax }});
   let q1 = clamp(round_scaled_half_to_even(f32(xv.y), f32(y_scale[scale_index1])) + zp1, {{ qMin }}, {{ qMax }});
   let q2 = clamp(round_scaled_half_to_even(f32(xv.z), f32(y_scale[scale_index2])) + zp2, {{ qMin }}, {{ qMax }});

+{% set yUnsigned = yUnsigned is defined and yUnsigned %}
+{% set qMin = qMin | default(0 if yUnsigned else 0 - 128) %}
+{% set qMax = qMax | default(255 if yUnsigned else 127) %}
 // int8/uint8 tensors are stored widened (one u32/i32 per element), so the vec4
 // binding gives 128-bit loads/stores of four elements. Per-component arithmetic
 // remains identical to the scalar quantize/dequantize paths.
 {{ env.wgsl.resourceDeclarations }}
 {% set divisionF16 = divisionF16 is defined and divisionF16 %}
   // Clamp before the i32 cast so infinite and huge finite inputs saturate
   // instead of invoking undefined conversion behavior. Compare distances in
   // the input domain: doing the comparison on value / scale can move a value
+  // across a halfway boundary because WGSL division is not required to be
+  // correctly rounded.
   let v = clamp(value / scale, -2.0e9, 2.0e9);
   let fl = floor(v);
   let lo = fl;
 {% endif %}
+{% if vectorParams is defined and vectorParams %}
 fn read_zero4({% if hasZero %}index: u32{% endif %}) -> vec4<i32> {
 {% if hasZero %}
 {% if yUnsigned %}
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; gid.y carries the high portion of the vector index.
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count4) {
     return;
   }
+{% if vectorParams is defined and vectorParams %}
   let scale_index = i % (params.scaleSize / 4u);
+{% elif crossingParams is defined and crossingParams %}
   let base_index = i * 4u;
   let scale_index0 = (base_index / params.inner) % params.scaleSize;
   let scale_index1 = ((base_index + 1u) / params.inner) % params.scaleSize;
   let scale_index2 = ((base_index + 2u) / params.inner) % params.scaleSize;
   let scale_index3 = ((base_index + 3u) / params.inner) % params.scaleSize;
+{% elif perAxis %}
   // inner % 4 == 0, so all 4 lanes share one scale index.
   let scale_index = ((i * 4u) / params.inner) % params.scaleSize;
 {% else %}
   let scale_index = 0u;
 {% endif %}
   let xv = x[i];
+{% if vectorParams is defined and vectorParams %}
   let zp4 = read_zero4({% if hasZero %}scale_index{% endif %});
+{% elif crossingParams is defined and crossingParams %}
   let zp0 = read_zero(scale_index0);
   let zp1 = read_zero(scale_index1);
   let zp2 = read_zero(scale_index2);
 {% else %}
   let zp = read_zero({% if hasZero %}scale_index{% endif %});
 {% endif %}
+{% if vectorParams is defined and vectorParams %}
   let scale = vec4<f32>(y_scale[scale_index]);
   let q0 = clamp(round_scaled_half_to_even(f32(xv.x), scale.x) + zp4.x, {{ qMin }}, {{ qMax }});
   let q1 = clamp(round_scaled_half_to_even(f32(xv.y), scale.y) + zp4.y, {{ qMin }}, {{ qMax }});
   let q2 = clamp(round_scaled_half_to_even(f32(xv.z), scale.z) + zp4.z, {{ qMin }}, {{ qMax }});
   let q3 = clamp(round_scaled_half_to_even(f32(xv.w), scale.w) + zp4.w, {{ qMin }}, {{ qMax }});
+{% elif crossingParams is defined and crossingParams %}
   let q0 = clamp(round_scaled_half_to_even(f32(xv.x), f32(y_scale[scale_index0])) + zp0, {{ qMin }}, {{ qMax }});
   let q1 = clamp(round_scaled_half_to_even(f32(xv.y), f32(y_scale[scale_index1])) + zp1, {{ qMin }}, {{ qMax }});
   let q2 = clamp(round_scaled_half_to_even(f32(xv.z), f32(y_scale[scale_index2])) + zp2, {{ qMin }}, {{ qMax }});

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.QuantizeLinear",
   "fixtureArrays": {
     "uint8_axis1_rank4_onnx_backend_input_x": [-162, 10, -100, 232, -20, -50, -76, 0, 0, 252, 32, -44, 245, -485, -960, -270, -375, -470]
   },
@@ -71,7 +70,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Int8",
-        "notes": "Signed-output subnormal scale companion: distinct one-LSB signed steps should survive instead of collapsing through zero-scale flushing."
       },
       "inputs": {
         "x": {
@@ -100,7 +99,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Int8",
-        "notes": "Vec4 signed-output companion for subnormal QuantizeLinear scale handling."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 2e-40] } },
@@ -120,7 +119,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Per_Channel_Axis_Default",
-        "notes": "Per-axis companion for subnormal QuantizeLinear scale handling: each channel uses a valid subnormal scale and nonzero zero point."
       },
       "attrs": { "axis": 1 },
       "inputs": {
@@ -154,7 +153,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Per_Channel_Axis_0",
-        "notes": "Vec4 per-axis companion for subnormal QuantizeLinear scale handling: axis=0 has inner size 4, so the vectorized channel path must preserve tiny one-LSB steps."
       },
       "attrs": { "axis": 0 },
       "inputs": {
@@ -989,7 +988,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOp21BlockedTest.UnsignedInt_UseZeroPoint_MiddleAxis",
-        "notes": "Blocked quantization saturates large finite scaled values after round-to-even, matching the non-blocked QuantizeLinear path (the blocked kernel now clamps value/scale into i32 range before the cast)."
       },
       "attrs": { "axis": 1, "block_size": 2 },
       "inputs": {
@@ -1242,7 +1241,7 @@
     {
       "name": "vec4_cross_axis_rows_width6_no_zero_point",
       "provenance": {
-        "notes": "Symmetric per-channel quantization over six-element rows with y_zero_point omitted (ONNX makes it optional and defaults it to 0), so a four-lane vector still crosses each row boundary and may need two different per-axis scale entries while no zero-point binding exists. Twin of vec4_cross_axis_rows_width6_with_zero_point; signed int8 output keeps the negative half of each row representable at zero point 0, and the seven-value input cycle is coprime with the six-wide row so no two rows repeat the same lane pattern. No quotient lands on a .5 tie, so the result is exact for either rounding of a tie."
       },
       "attrs": { "axis": 0, "output_dtype": 3 },
       "inputs": {
@@ -1271,7 +1270,7 @@
     {
       "name": "f16_inputs_explicit_float32_precision",
       "provenance": {
-        "notes": "Twin of the default-precision witness with precision=FLOAT (TensorProto code 1), proving that an explicit f32 division remains distinct and rounds to 13."
       },
       "attrs": { "precision": 1 },
       "inputs": {

 {
   "fixtureArrays": {
     "uint8_axis1_rank4_onnx_backend_input_x": [-162, 10, -100, 232, -20, -50, -76, 0, 0, 252, 32, -44, 245, -485, -960, -270, -375, -470]
   },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Int8",
+        "notes": "With signed output and a subnormal scale, distinct one-LSB steps must survive instead of collapsing through zero-scale flushing."
       },
       "inputs": {
         "x": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Int8",
+        "notes": "Vectorized signed-output quantization must preserve a valid subnormal scale."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 2e-40] } },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Per_Channel_Axis_Default",
+        "notes": "Each channel uses a valid subnormal scale and a nonzero zero point in per-axis quantization."
       },
       "attrs": { "axis": 1 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.Per_Channel_Axis_0",
+        "notes": "Axis 0 has inner size 4, so vectorized per-axis quantization must preserve tiny one-LSB steps under subnormal scales."
       },
       "attrs": { "axis": 0 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
         "test": "QuantizeLinearOp21BlockedTest.UnsignedInt_UseZeroPoint_MiddleAxis",
+        "notes": "Blocked quantization clamps the scaled value to int32 range before conversion, then applies round-to-even and output saturation."
       },
       "attrs": { "axis": 1, "block_size": 2 },
       "inputs": {
     {
       "name": "vec4_cross_axis_rows_width6_no_zero_point",
       "provenance": {
+        "notes": "Symmetric per-channel quantization over six-element rows omits y_zero_point, so it defaults to zero. Four-lane vectors cross row boundaries and may require two scale entries; signed int8 output preserves negative values, and a coprime seven-value cycle gives every row a distinct lane pattern. No quotient lands on a half-way tie."
       },
       "attrs": { "axis": 0, "output_dtype": 3 },
       "inputs": {
     {
       "name": "f16_inputs_explicit_float32_precision",
       "provenance": {
+        "notes": "With precision=FLOAT (TensorProto code 1), explicit f32 division produces a quotient that rounds to 13."
       },
       "attrs": { "precision": 1 },
       "inputs": {