sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +82 -0
build/webgpu/bench.json +328 -0
build/webgpu/manifest.json +839 -0
build/webgpu/metadata.json +20 -0
build/webgpu/quant-linear-blocked-axis.wgsl.jinja +156 -0
build/webgpu/quant-linear-scalar.wgsl.jinja +109 -0
build/webgpu/quant-linear-vec4.wgsl.jinja +137 -0
build/webgpu/test.json +1327 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.QuantizeLinear
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 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`.
+See the [ONNX `QuantizeLinear` spec](https://onnx.ai/onnx/operators/onnx__QuantizeLinear.html) for the reference semantics.
+## 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
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `1` | Axis of the quantization dimension in `x`, used for per-axis and blocked quantization; negative values count from the end. |
+| `block_size` | `0` | Number of elements along `axis` that share a single scale value for blocked quantization; 0 means blocked quantization is not used. |
+| `output_dtype` | `0` | ONNX TensorProto element-type code for `y`; 0 infers the type from `y_zero_point`, or uint8 when the zero point is omitted. |
+| `precision` | `0` | 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` | `1` | Controls out-of-range conversion for float8 outputs. The implemented int8/uint8 subset accepts the ONNX default `1`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `TX` | `float32`, `float16` |
+| `TS` | `float32`, `float16` |
+| `TQ` | `uint8`, `int8` |
+## 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
+- [`quant-linear-blocked-axis.wgsl.jinja`](build/webgpu/quant-linear-blocked-axis.wgsl.jinja)
+- [`quant-linear-scalar.wgsl.jinja`](build/webgpu/quant-linear-scalar.wgsl.jinja)
+- [`quant-linear-vec4.wgsl.jinja`](build/webgpu/quant-linear-vec4.wgsl.jinja)
+## 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.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.QuantizeLinear", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { y } = await kernel({ x: { data: xData, shape: [] }, y_scale: { data: y_scaleData, shape: [] } }, {
+  outputs: { y: { shape: [], dtype: "uint8" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,328 @@

+{
+  "op": "ai.onnx.QuantizeLinear",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "f32_to_u8_1m_scalar",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 1048576 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1024, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1024, 1024] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "f32_to_u8_1m_no_zero_point",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 1048576 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1024, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1024, 1024] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "f32_to_i8_1m_plus3_scalar",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "int8", "count": 1048579 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1048579],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "f32_to_u8_1m_plus3_no_zero_point",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 1048579 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1048579],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "f16_to_u8_1m_scalar",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float16", "outputDtype": "uint8", "count": 1048576 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1024, 1024], "data": { "kind": "constant", "value": 1.25 } },
+        "y_scale": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1024, 1024] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "blocked_axis1_4096x4096_block32_with_zp",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 16777216 },
+      "attrs": { "axis": 1, "block_size": 32 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4096, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [4096, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [4096, 128], "data": { "kind": "constant", "value": 128 } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4096, 4096] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "pertensor_4096x4096_scalar_healthy_sibling",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 16777216 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4096, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4096, 4096] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_axis0_inner2049_scalar_fallback_alignment_miss",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 1049088 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [512, 2049],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [512], "data": { "kind": "constant", "value": 128 } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [512, 2049] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_axis0_inner2048_vec4_healthy_sibling",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 1048576 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [512, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [512], "data": { "kind": "constant", "value": 128 } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [512, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "f16_per_axis0_inner2048_vec4_unbenched",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float16", "outputDtype": "uint8", "count": 1048576 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [512, 2048], "data": { "kind": "constant", "value": 1.25 } },
+        "y_scale": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [512], "data": { "kind": "constant", "value": 128 } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [512, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_axis1_lastdim_inner1_generic_scalar_cliff",
+      "preset": "stress",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 4194304 },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2048, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [2048], "data": { "kind": "constant", "value": 128 } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2048, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_axis1_lastdim_inner1_generic_scalar_no_zero_point_cliff",
+      "preset": "stress",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 4194304 },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2048, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2048, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "blocked-rank4-axis1-bs32-no-zp-scalar-pathology",
+      "preset": "stress",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 4194304 },
+      "attrs": { "axis": 1, "block_size": 32 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 512, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [4, 16, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.005, "offset": 0.02 }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4, 512, 32, 64], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "blocked-rank4-axis1-bs32-with-zp-scalar-pathology",
+      "preset": "stress",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint8", "count": 4194304 },
+      "attrs": { "axis": 1, "block_size": 32 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 512, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [4, 16, 32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.005, "offset": 0.02 }
+        },
+        "y_zero_point": {
+          "dtype": "uint8",
+          "shape": [4, 16, 32, 64],
+          "dist": "randint",
+          "seed": 204,
+          "min": 96,
+          "max": 160
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4, 512, 32, 64], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "dtypeBytes(args.inputDtype) * (numel(shapes.x) + numel(shapes.y_scale)) + dtypeBytes(args.outputDtype) * (numel(shapes.y_zero_point) + numel(shapes.y))"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,839 @@

+{
+  "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",
+    "quantDtypesOk": "(tensorDtypes.x != \"float16\" and tensorDtypes.y_scale != \"float16\" and not quantizeDivisionF16) or device.features.has(\"shader-f16\")",
+    "blockedScaleOk": "attrs.block_size > 0 and ranks.x >= 2 and ranks.y_scale == ranks.x",
+    "elementCount": "numel(shapes.y)",
+    "elementCount4": "floor(numel(shapes.y) / 4)",
+    "scaleSize": "1 if ranks.y_scale == 0 else dim(shapes.y_scale, 0)",
+    "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 ~ \">\"",
+    "yUnsigned": "tensorDtypes.y == \"uint8\"",
+    "qMin": "0 if tensorDtypes.y == \"uint8\" else 0 - 128",
+    "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" }] }
+      }
+    ]
+  }
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "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" }
+}

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

	@@ -0,0 +1,156 @@

+// 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 %}
+  return i32(y_zero_point[index]);
+{% else %}
+  return y_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+fn write_quantized(index: u32, value: i32) {
+{% if yUnsigned %}
+  y[index] = u32(value);
+{% else %}
+  y[index] = value;
+{% endif %}
+}
+{% endif %}
+{% if not lastAxisVectorized %}
+fn scale_index(index: u32) -> u32 {
+  let inner_index = index % params.inner;
+  let axis_index = (index / params.inner) % params.axisDim;
+  let outer_index = index / (params.axisDim * params.inner);
+  return (outer_index * params.scaleAxisDim + axis_index / params.blockSize) * params.inner + inner_index;
+}
+{% 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.
+{% if divisionF16 %}
+fn round_quotient_half_to_even(quotient: f32) -> i32 {
+  let v = clamp(quotient, -2.0e9, 2.0e9);
+  let fl = floor(v);
+  let hi = fl + 1.0;
+  let fraction = v - fl;
+  if (fraction < 0.5) { return i32(fl); }
+  if (fraction > 0.5) { return i32(hi); }
+  let half = floor(fl * 0.5);
+  let is_even = (fl - half * 2.0) == 0.0;
+  return i32(select(hi, fl, is_even));
+}
+fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
+  // ONNX precision=FLOAT16 (and an omitted precision with f16 y_scale) requires
+  // the division itself—not merely its operands—to round in f16.
+  return round_quotient_half_to_even(f32(f16(value) / f16(scale)));
+}
+{% else %}
+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;
+  let hi = fl + 1.0;
+  let lo_dist = abs(value - lo * scale);
+  let hi_dist = abs(hi * scale - value);
+  if (lo_dist < hi_dist) {
+    return i32(lo);
+  }
+  if (hi_dist < lo_dist) {
+    return i32(hi);
+  }
+  let half = floor(fl * 0.5);
+  let is_even = (fl - half * 2.0) == 0.0;
+  return i32(select(hi, lo, is_even));
+}
+{% 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];
+  let scale = f32(y_scale[scale_index]);
+{% if hasZero %}
+{% if yUnsigned %}
+  let zero = i32(y_zero_point[scale_index]);
+{% else %}
+  let zero = y_zero_point[scale_index];
+{% endif %}
+{% else %}
+  let zero = 0;
+{% endif %}
+  let quantized = vec4<i32>(
+    clamp(round_scaled_half_to_even(f32(values.x), scale) + zero, {{ qMin }}, {{ qMax }}),
+    clamp(round_scaled_half_to_even(f32(values.y), scale) + zero, {{ qMin }}, {{ qMax }}),
+    clamp(round_scaled_half_to_even(f32(values.z), scale) + zero, {{ qMin }}, {{ qMax }}),
+    clamp(round_scaled_half_to_even(f32(values.w), scale) + zero, {{ qMin }}, {{ qMax }})
+  );
+{% if yUnsigned %}
+  y[i4] = vec4<u32>(quantized);
+{% else %}
+  y[i4] = quantized;
+{% endif %}
+}
+{% 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];
+  let scales = y_scale[scale4];
+{% if hasZero %}
+{% if yUnsigned %}
+  let zero = vec4<i32>(y_zero_point[scale4]);
+{% else %}
+  let zero = y_zero_point[scale4];
+{% endif %}
+{% endif %}
+  let quantized = vec4<i32>(
+    clamp(round_scaled_half_to_even(f32(values.x), f32(scales.x)){% if hasZero %} + zero.x{% endif %}, {{ qMin }}, {{ qMax }}),
+    clamp(round_scaled_half_to_even(f32(values.y), f32(scales.y)){% if hasZero %} + zero.y{% endif %}, {{ qMin }}, {{ qMax }}),
+    clamp(round_scaled_half_to_even(f32(values.z), f32(scales.z)){% if hasZero %} + zero.z{% endif %}, {{ qMin }}, {{ qMax }}),
+    clamp(round_scaled_half_to_even(f32(values.w), f32(scales.w)){% if hasZero %} + zero.w{% endif %}, {{ qMin }}, {{ qMax }})
+  );
+{% if yUnsigned %}
+  y[i4] = vec4<u32>(quantized);
+{% else %}
+  y[i4] = quantized;
+{% endif %}
+}
+{% 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 %}

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

	@@ -0,0 +1,109 @@

+// 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 %}
+// Exact ONNX QuantizeLinear round-to-nearest-even, with identical handling of
+// infinities, saturation, and halfway values across every kernel route.
+{% if divisionF16 %}
+fn round_quotient_half_to_even(quotient: f32) -> i32 {
+  let v = clamp(quotient, -2.0e9, 2.0e9);
+  let fl = floor(v);
+  let hi = fl + 1.0;
+  let fraction = v - fl;
+  if (fraction < 0.5) { return i32(fl); }
+  if (fraction > 0.5) { return i32(hi); }
+  let half = floor(fl * 0.5);
+  let is_even = (fl - half * 2.0) == 0.0;
+  return i32(select(hi, fl, is_even));
+}
+fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
+  // ONNX precision=FLOAT16 (and an omitted precision with f16 y_scale) requires
+  // the division itself—not merely its operands—to round in f16.
+  return round_quotient_half_to_even(f32(f16(value) / f16(scale)));
+}
+{% else %}
+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;
+  let hi = fl + 1.0;
+  let lo_dist = abs(value - lo * scale);
+  let hi_dist = abs(hi * scale - value);
+  if (lo_dist < hi_dist) {
+    return i32(lo);
+  }
+  if (hi_dist < lo_dist) {
+    return i32(hi);
+  }
+  let half = floor(fl * 0.5);
+  let is_even = (fl - half * 2.0) == 0.0;
+  return i32(select(hi, lo, is_even));
+}
+{% endif %}
+fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
+{% if hasZero %}
+{% if yUnsigned %}
+  return i32(y_zero_point[index]);
+{% else %}
+  return y_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+fn write_value(index: u32, value: i32) {
+{% if yUnsigned %}
+  y[index] = u32(value);
+{% else %}
+  y[index] = value;
+{% endif %}
+}
+fn transform_one(index: u32, scale: f32, zero_point: i32) {
+  let rounded = round_scaled_half_to_even(f32(x[index]), scale) + zero_point;
+  write_value(index, clamp(rounded, {{ qMin }}, {{ qMax }}));
+}
+@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;
+  }
+  let base = params.count - params.count % 4u;
+  if (base >= params.count) {
+    return;
+  }
+  let scale = f32(y_scale[0]);
+  let zero_point = read_zero({% if hasZero %}0u{% endif %});
+  transform_one(base, scale, zero_point);
+  if (base + 1u < params.count) { transform_one(base + 1u, scale, zero_point); }
+  if (base + 2u < params.count) { transform_one(base + 2u, scale, zero_point); }
+  if (base + 3u < params.count) { transform_one(base + 3u, scale, zero_point); }
+{% else %}
+  if (invocation >= params.count) {
+    return;
+  }
+  var scale_index = 0u;
+{% if source.perAxis %}
+  scale_index = (invocation / params.inner) % params.scaleSize;
+{% endif %}
+  let scale = f32(y_scale[scale_index]);
+  transform_one(invocation, scale, read_zero({% if hasZero %}scale_index{% endif %}));
+{% endif %}
+}

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

	@@ -0,0 +1,137 @@

+// 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 %}
+// Exact ONNX QuantizeLinear round-to-nearest-even, with identical handling of
+// infinities, saturation, and halfway values across every kernel route.
+{% if divisionF16 %}
+fn round_quotient_half_to_even(quotient: f32) -> i32 {
+  let v = clamp(quotient, -2.0e9, 2.0e9);
+  let fl = floor(v);
+  let hi = fl + 1.0;
+  let fraction = v - fl;
+  if (fraction < 0.5) { return i32(fl); }
+  if (fraction > 0.5) { return i32(hi); }
+  let half = floor(fl * 0.5);
+  let is_even = (fl - half * 2.0) == 0.0;
+  return i32(select(hi, fl, is_even));
+}
+fn round_scaled_half_to_even(value: f32, scale: f32) -> i32 {
+  // ONNX precision=FLOAT16 (and an omitted precision with f16 y_scale) requires
+  // the division itself—not merely its operands—to round in f16.
+  return round_quotient_half_to_even(f32(f16(value) / f16(scale)));
+}
+{% else %}
+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;
+  let hi = fl + 1.0;
+  let lo_dist = abs(value - lo * scale);
+  let hi_dist = abs(hi * scale - value);
+  if (lo_dist < hi_dist) {
+    return i32(lo);
+  }
+  if (hi_dist < lo_dist) {
+    return i32(hi);
+  }
+  let half = floor(fl * 0.5);
+  let is_even = (fl - half * 2.0) == 0.0;
+  return i32(select(hi, lo, is_even));
+}
+{% endif %}
+{% if source.vectorParams is defined and source.vectorParams %}
+fn read_zero4({% if hasZero %}index: u32{% endif %}) -> vec4<i32> {
+{% if hasZero %}
+{% if yUnsigned %}
+  return vec4<i32>(y_zero_point[index]);
+{% else %}
+  return y_zero_point[index];
+{% endif %}
+{% else %}
+  return vec4<i32>(0);
+{% endif %}
+}
+{% else %}
+fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
+{% if hasZero %}
+{% if yUnsigned %}
+  return i32(y_zero_point[index]);
+{% else %}
+  return y_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+{% 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);
+  let zp3 = read_zero(scale_index3);
+{% 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 }});
+  let q3 = clamp(round_scaled_half_to_even(f32(xv.w), f32(y_scale[scale_index3])) + zp3, {{ qMin }}, {{ qMax }});
+{% else %}
+  let scale = f32(y_scale[scale_index]);
+  let q0 = clamp(round_scaled_half_to_even(f32(xv.x), scale) + zp, {{ qMin }}, {{ qMax }});
+  let q1 = clamp(round_scaled_half_to_even(f32(xv.y), scale) + zp, {{ qMin }}, {{ qMax }});
+  let q2 = clamp(round_scaled_half_to_even(f32(xv.z), scale) + zp, {{ qMin }}, {{ qMax }});
+  let q3 = clamp(round_scaled_half_to_even(f32(xv.w), scale) + zp, {{ qMin }}, {{ qMax }});
+{% endif %}
+{% if yUnsigned %}
+  y[i] = vec4<u32>(u32(q0), u32(q1), u32(q2), u32(q3));
+{% else %}
+  y[i] = vec4<i32>(q0, q1, q2, q3);
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1327 @@

+{
+  "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]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_vec4_no_zero_point",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 33554432], "data": { "kind": "linspace", "start": -64.0, "end": 64.0 } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 33554432], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_scalar_saturate_round_even",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "values", "values": [-100.0, -0.25, 0.25, 0.75, 1.25, 63.75, 200.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [7] } }
+    },
+    {
+      "name": "uint8_subnormal_scale_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8",
+        "notes": "Valid positive subnormal scale: quantization should preserve one-LSB steps instead of flushing the scale to zero."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_subnormal_scale_nonzero_zero_point_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8",
+        "notes": "Subnormal per-tensor scale with a nonzero zero point should preserve one-LSB signed offsets around the zero point."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 2e-40] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_subnormal_scale_preserves_signed_steps_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "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": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [-2e-40, -1e-40, 0.0, 1e-40, 2e-40] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-2, -1, 0, 1, 2] }
+        }
+      }
+    },
+    {
+      "name": "int8_subnormal_scale_vec4_preserves_signed_steps_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "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] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": { "dtype": "int8", "shape": [4], "tolerance": 0, "data": { "kind": "values", "values": [-1, 0, 1, 2] } }
+      }
+    },
+    {
+      "name": "uint8_axis1_per_channel_subnormal_scale_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "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": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.0, 1e-40, -1e-40, 2e-40, 2e-40, -2e-40, 0.0, 1e-40] }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, 1e-40, 1e-40, 1e-40] }
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [10, 20, 30, 40] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [10, 21, 29, 42, 12, 18, 30, 41] }
+        }
+      }
+    },
+    {
+      "name": "uint8_axis0_per_channel_subnormal_scale_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "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": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40, 2e-40, -2e-40, 0.0, 1e-40] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, 1e-40] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [10, 20] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [10, 11, 12, 13, 22, 18, 20, 21] }
+        }
+      }
+    },
+    {
+      "name": "uint8_axis1_per_channel_subnormal_scale_no_zero_point_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.QuantizeLinear_Without_Zero_Point_Opset13",
+        "notes": "Per-axis extension of omitted-zero-point coverage: default uint8 zero point is zero, but each channel still uses a valid subnormal scale."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40, 2e-40, 0.0, 1e-40, -1e-40] }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, 1e-40, 1e-40, 1e-40] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 2, 0, 1, 0] }
+        }
+      }
+    },
+    {
+      "name": "uint8_axis0_per_channel_subnormal_scale_no_zero_point_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: Metal flushes denormals in floating-point division (1e-40/1e-40 -> NaN); the CPU reference preserves the subnormal scale, so subnormal y_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.QuantizeLinear_Without_Zero_Point_Opset13",
+        "notes": "Vec4 per-axis extension of omitted-zero-point coverage with valid subnormal scales."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40, 2e-40, 0.0, -1e-40, 1e-40] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, 1e-40] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 2, 0, 0, 1] }
+        }
+      }
+    },
+    {
+      "name": "uint8_huge_finite_saturates_before_i32_overflow",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8",
+        "notes": "Finite values far outside the uint8 range should saturate; kernels must avoid converting an out-of-range rounded float to i32 before clamping."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "values", "values": [100000000000000000000.0, -100000000000000000000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2], "tolerance": 0, "data": { "kind": "values", "values": [255, 0] } }
+      }
+    },
+    {
+      "name": "uint8_infinities_saturate_before_i32_conversion",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8",
+        "notes": "Infinite inputs are valid float tensor values and should saturate to quantized bounds without first converting infinity to i32."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": ["Infinity", "-Infinity", 0.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [3], "tolerance": 0, "data": { "kind": "values", "values": [255, 0, 128] } }
+      }
+    },
+    {
+      "name": "uint8_infinities_saturate_vec4_before_i32_conversion",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8",
+        "notes": "Vec4 zero-point path: infinite inputs should saturate to uint8 bounds without converting infinity to i32 first."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": ["Infinity", "-Infinity", 0.0, 1.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [255, 0, 128, 129] }
+        }
+      }
+    },
+    {
+      "name": "int8_axis1_f16",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [-10.0, -0.5, 0.5, 1.5, 2.5, 4.0, 8.0, 64.0] }
+        },
+        "y_scale": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.0, 2.0, 0.25] } },
+        "y_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [0, -2, 1, 3] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [2, 4] } }
+    },
+    {
+      "name": "uint8_no_zero_negative_axis",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 0.5, 2.0, 3.0] } },
+        "y_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 2] } }
+    },
+    {
+      "name": "ort_int16_round_even_and_saturate_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The QuantizeLinear kernels only implement uint8/int8 output clamps; int16 needs round-to-even plus saturation to [-32768, 32767] before the standard route can be enabled."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int16"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": [0.0, -514.0, 3.0, -3.0, 2.9, -2.9, 3.1, -3.1, 65022.0, -66046.0, 65023.0, -66047.0, 65024.0, -66048.0, 70000.0, -70000.0]
+          }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "int16", "shape": [], "data": { "kind": "values", "values": [256] } }
+      },
+      "outputs": { "y": { "dtype": "int16", "shape": [16], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_scalar_onnx_backend",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 1000.0, -254.0, -1000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6] } }
+    },
+    {
+      "name": "ort_uint8_rank0_scale_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 1000.0, -254.0, -1000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6] } }
+    },
+    {
+      "name": "ort_int8_rank0_scale_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 5.0, -2.0, -5.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.039215686] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [6] } }
+    },
+    {
+      "name": "ort_scalar_input_rank0_with_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Scalar"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [] } }
+    },
+    {
+      "name": "ort_scalar_input_rank0_no_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.QuantizeLinear_Without_Zero_Point_Opset13"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [] } }
+    },
+    {
+      "name": "ort_scalar_input_rank0_zero_point_zero",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.QuantizeLinear_With_Zero_Point0"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.0] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [] } }
+    },
+    {
+      "name": "ort_singleton_input_no_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.QuantizeLinear_With_Zero_Dim1"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [3.0] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1] } }
+    },
+    {
+      "name": "scalar_x4_int8_with_zero_tail",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [17],
+          "data": {
+            "kind": "values",
+            "values": [-64.0, -32.0, -16.0, -8.0, -4.0, -2.0, -1.0, 0.0, 1.0, 2.0, 4.0, 8.0, 16.0, 32.0, 48.0, 64.0, 96.0]
+          }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [17] } }
+    },
+    {
+      "name": "scalar_x4_uint8_no_zero_tail",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [17],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.25, 0.5, 0.75, 1.0, 1.25, 2.0, 4.0, 8.0, 16.0, 32.0, 64.0, 128.0, 256.0, 384.0, 512.0, 1024.0]
+          }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [17] } }
+    },
+    {
+      "name": "vec4_tail_int8_with_zero_4097",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4097], "data": { "kind": "constant", "value": 1.25 } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [4097] } },
+      "provenance": { "notes": "Covers the packed bulk plus scalar tail with a scalar zero point." }
+    },
+    {
+      "name": "vec4_tail_uint8_no_zero_4097",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4097], "data": { "kind": "constant", "value": 4.0 } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4097] } },
+      "provenance": { "notes": "Covers the packed bulk plus scalar tail without a zero point." }
+    },
+    {
+      "name": "uint8_axis1_rank4_onnx_backend",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 3, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/uint8_axis1_rank4_onnx_backend_input_x" } }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 4.0, 5.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [84, 24, 196] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 3, 3, 2] } }
+    },
+    {
+      "name": "int8_scalar_zero_point_saturate_round_even",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [9],
+          "data": { "kind": "values", "values": [-100.0, -64.0, -63.5, -0.5, 0.0, 0.5, 63.5, 64.0, 100.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [9] } }
+    },
+    {
+      "name": "ort_int8_negative_zero_point_formulation",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8_NegativeZeroPoint"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 5.0, 6.0, -2.0, -5.0, -6.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.039215686] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-23] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [8] } }
+    },
+    {
+      "name": "ort_int8_positive_zero_point_formulation",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8_PositiveZeroPoint"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 5.0, 6.0, -2.0, -5.0, -6.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.039215686] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [23] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [8] } }
+    },
+    {
+      "name": "uint8_axis0_rank2_per_row",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [-2.0, 2.0, 20.0, 21.5, -20.0, 200.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, 1.0, 2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [10, 20, 30] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 2] } }
+    },
+    {
+      "name": "int8_negative_axis_rank3_per_last_dim",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [-4.0, -1.0, 0.0, 1.0, 2.5, 4.0, 8.0, -8.0, 0.5, -0.5, 63.0, -300.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 0.5, 2.0] } },
+        "y_zero_point": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-1, 0, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [2, 2, 3] } }
+    },
+    {
+      "name": "ort_f16_uint8_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpMLFloat16Test.Uint8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, 4.0, 1000.0, -254.0, -1000.0] }
+        },
+        "y_scale": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6] } }
+    },
+    {
+      "name": "ort_opset25_f16_int8_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpMLFloat16Test.Opset25_PerAxisInt8_Cuda"
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [-4.0, -2.0, 0.0, 2.0, 4.0, 6.0, 8.0, 10.0] }
+        },
+        "y_scale": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [2.0, 2.0, 4.0, 4.0] } },
+        "y_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_5d_per_tensor",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8_5D_DML_TypeMismatch"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [6, 1, 1, 1, 1] } }
+    },
+    {
+      "name": "ort_int8_5d_rank0_zero_point_opset21",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8_5D_WithZeroPoint_Opset21_DML"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "y_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [6, 1, 1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_5d_per_axis_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8_5D_PerAxis_DML_TypeMismatch"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0, 1.0, 1.0] }
+        },
+        "y_zero_point": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [6, 1, 1, 1, 1] } }
+    },
+    {
+      "name": "ort_uint8_5d_no_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8_5D_NoZeroPoint_Opset21_DML"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [0.0, 51.0, 102.0, 153.0, 204.0, 255.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6, 1, 1, 1, 1] } }
+    },
+    {
+      "name": "ort_uint8_2d_scalar_quantization",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.2D"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 1000.0, 0.0, 2.0, 3.0, 1000.0, 0.0, 2.0, 3.0, 1000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 4] } }
+    },
+    {
+      "name": "ort_uint8_per_channel_default_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Per_Channel_Axis_Default"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0.0, 2.0, 1.0, 1001.0, 1.0, 1.0, 2.0, 1100.0, 2.0, 4.2, 3.0, 1200.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 20.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [64, 100, 127, 127] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 4] } }
+    },
+    {
+      "name": "uint8_axis1_per_channel_innermost_vec4_no_zero_point",
+      "provenance": {
+        "notes": "Per-axis scale on the innermost axis with the zero point omitted, so the vec4-bound scale path runs with the implicit zero. Each channel uses a different scale and the first column ties on .5 in both directions, so a lane reading the wrong channel or rounding half-away-from-zero changes the output."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.25, 6.0, 20.0, 3.0, 3.75, 10.0, 36.0, 7.25, 50.0, 90.0, 172.0, 15.75]
+          }
+        },
+        "y_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 2.0, 4.0, 0.25] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_per_channel_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Per_Channel_Axis_0"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 1000.0, 0.0, 2.0, 3.0, 1000.0, 0.0, 2.0, 3.0, 1000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 4] } }
+    },
+    {
+      "name": "ort_uint8_per_channel_negative_axis_minus2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Per_Channel_Axis_neg"
+      },
+      "attrs": { "axis": -2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 1000.0, 0.0, 2.0, 3.0, 1000.0, 0.0, 2.0, 3.0, 1000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 4] } }
+    },
+    {
+      "name": "onnx_backend_quantizelinear",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_quantizelinear" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, 3.0, 1000.0, -254.0, -1000.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_quantizelinear_axis",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_quantizelinear_axis" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 3, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/uint8_axis1_rank4_onnx_backend_input_x" } }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 4.0, 5.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [84, 24, 196] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [1, 3, 3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_quantizelinear_blocked_asymmetric",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_quantizelinear_blocked_asymmetric"
+      },
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [6.0, 12.0, 50.0, 5.0, 1.0, 8.0, 4.0, 5.0, 0.0, 20.0, 10.0, 4.0] }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": {
+            "kind": "values",
+            "values": [1.5, 2.5, 3.0, 4.900000095367432, 5.099999904632568, 6.900000095367432]
+          }
+        },
+        "y_zero_point": {
+          "dtype": "uint8",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [0, 1, 1, 0, 2, 3] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_uint8_no_zero_point_4x8",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 8],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 2.0, 3.0, 6.0, 10.0, 14.0, 18.0, 22.0, 1.999, 2.001, 5.999, 6.001, 100.0, 250.0, 500.0, 750.0, 1000.0, 1016.0, 1018.0, 1020.0, 1022.0, 1024.0, 2000.0, 5000.0, -1.0, -100.0, 0.5, 1.5, 2.5, 3.5, 511.0, 513.0]
+          }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [4.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4, 8] } }
+    },
+    {
+      "name": "vec4_uint8_no_zero_point_per_axis0",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.25, 0.5, 0.75, 1.25, 63.75, 127.5, 200.0, 0.0, 2.0, 3.0, 6.0, 10.0, 250.0, 510.0, 1000.0]
+          }
+        },
+        "y_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 2, 4] } }
+    },
+    {
+      "name": "ort_uint8_f16_scalar_scale_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpMLFloat16Test.Uint8"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, 4.0, 1000.0, -254.0, -1000.0] }
+        },
+        "y_scale": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_5d_scalar_no_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Uint8_5D_NoZeroPoint_Opset21_DML"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [0.0, 51.0, 102.0, 153.0, 204.0, 255.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [6, 1, 1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_5d_per_axis0_with_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.Int8_5D_PerAxis_DML_TypeMismatch"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6, 1, 1, 1, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0, 1.0, 1.0] }
+        },
+        "y_zero_point": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [0, 0, 0, 0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [6, 1, 1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_blocked_uint8_with_zero_point_rank3_axis2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOp21BlockedTest.UnsignedInt_UseZeroPoint_LastAxis",
+        "notes": "Valid blocked quantization with per-block zero-points on the last axis."
+      },
+      "attrs": { "axis": 2, "block_size": 4 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 8],
+          "data": {
+            "kind": "values",
+            "values": [4.0, 2.0, 4.0, 2.0, -8.0, -12.0, -8.0, -12.0, 4.0, 2.0, 4.0, 2.0, -8.0, -12.0, -8.0, -12.0, 10.5, 14.0, 10.5, 14.0, -3.0, -2.0, -3.0, -2.0, 10.5, 14.0, 10.5, 14.0, -3.0, -2.0, -3.0, -2.0, -10.0, -8.0, -10.0, -8.0, 20.0, 24.0, 20.0, 24.0, -10.0, -8.0, -10.0, -8.0, 20.0, 24.0, 20.0, 24.0, -3.5, -7.0, -3.5, -7.0, -8.0, -9.0, -8.0, -9.0, -3.5, -7.0, -3.5, -7.0, -8.0, -9.0, -8.0, -9.0]
+          }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2, 4, 2],
+          "data": {
+            "kind": "values",
+            "values": [-2.0, -4.0, -2.0, -4.0, 3.5, 1.0, 3.5, 1.0, 2.0, 4.0, 2.0, 4.0, -3.5, -1.0, -3.5, -1.0]
+          }
+        },
+        "y_zero_point": {
+          "dtype": "uint8",
+          "shape": [2, 4, 2],
+          "data": { "kind": "values", "values": [2, 0, 2, 0, 1, 9, 1, 9, 13, 5, 13, 5, 11, 6, 11, 6] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_blocked_uint8_no_zero_point_rank3_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "QuantizeLinearOp21BlockedTest.UnsignedInt_NoZeroPoint_MiddleAxis"
+      },
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 4, 2],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 4.0, 5.0, 16.0, 18.0, 30.0, 33.0, 2.0, 3.0, 8.0, 10.0, 40.0, 42.0, 70.0, 72.0]
+          }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 0.5, 2.0, 3.0, 4.0, 5.0, 10.0, 16.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "blocked_uint8_saturates_extreme_finite_inputs",
+      "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": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [3.4028234663852886e+38, -3.4028234663852886e+38, 0.5, -0.5] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1, 2], "data": { "kind": "values", "values": [0.01, 0.01] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1, 2], "data": { "kind": "values", "values": [128, 128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [1, 4],
+          "data": { "kind": "values", "values": [255, 0, 178, 78] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "empty_zero_dim",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_dim_f16",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [0, 3], "data": { "kind": "values", "values": [] } },
+        "y_scale": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "blocked_uint8_no_zero_point_rank2_axis1",
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "values", "values": [0.0, 1.0, 4.0, 6.0, 16.0, 20.0, 2.0, 3.0, 10.0, 12.0, 40.0, 44.0] }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 0.5, 2.0, 4.0, 5.0, 8.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "blocked_uint8_partial_last_block_with_zp_axis1",
+      "attrs": { "axis": 1, "block_size": 4 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 6],
+          "data": { "kind": "values", "values": [0.0, 4.0, 8.0, 12.0, 20.0, 24.0, 2.0, 6.0, 10.0, 14.0, 30.0, 34.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [2.0, 4.0, 1.0, 3.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [2, 2], "data": { "kind": "values", "values": [1, 2, 0, 5] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_per_axis1_rank4_inner_not_mult4_scalar_fallback",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 3, 3, 1],
+          "data": { "kind": "values", "values": [-4.0, -2.0, 0.0, 2.0, 4.0, 6.0, -8.0, 8.0, 16.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 4.0, 8.0] } },
+        "y_zero_point": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [0, -1, 2] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 3, 3, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "per_axis1_scalar_kernel_dispatch_fold_over_16m",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [8388610, 3], "data": { "kind": "linspace", "start": -64.0, "end": 64.0 } },
+        "y_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, 1.0, 2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [10, 20, 30] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [8388610, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_zero_dim_per_axis_scale_nonempty",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [0, 4], "data": { "kind": "values", "values": [] } },
+        "y_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.0, 2.0, 4.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [128, 100, 64, 0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [0, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_huge_finite_tiny_scale_saturates_not_inf",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [60000.0, -60000.0, 0.0, 32.0] }
+        },
+        "y_scale": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.001] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [255, 0, 128, 255] }
+        }
+      }
+    },
+    {
+      "name": "int8_negative_zero_point_saturation_at_bounds",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-64.0, -62.5, 63.5, 64.0, 65.0, -100.0] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "y_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "blocked_vec4_int8_negative_axis_partial_block",
+      "provenance": {
+        "notes": "Route lock for blocked vec4 quantization: negative axis -3 normalizes to axis 1, block_size=2 leaves a partial fifth-axis entry, inner=8 keeps each four-lane x/scale access aligned, and signed int8 output exercises both saturation bounds."
+      },
+      "attrs": { "axis": -3, "block_size": 2, "output_dtype": 3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 5, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.17, "scale": 40.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 0.05, "offset": 0.2 }
+        }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 5, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "blocked_with_zp_int8_negative_axis_partial_block_vec4_candidate",
+      "provenance": {
+        "notes": "Asymmetric blocked-quantization lock adjacent to the no-zero-point vec4 route: negative axis -3, a partial final axis block, inner=8 alignment, lane-varying signed zero points, and saturating inputs distinguish signed ZP conversion from the symmetric path."
+      },
+      "attrs": { "axis": -3, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 5, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.17, "scale": 40.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [1, 3, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 0.05, "offset": 0.2 }
+        },
+        "y_zero_point": {
+          "dtype": "int8",
+          "shape": [1, 3, 2, 4],
+          "data": { "kind": "cycle", "values": [-128, -17, -1, 0, 1, 23, 64, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [1, 5, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "blocked_last_axis_vec4_no_zero_point_exact_blocks",
+      "provenance": {
+        "notes": "Route lock for the scalar-scale last-axis vec4 path: each aligned group of four values shares one scale, and signed output covers negative rounding and saturation semantics."
+      },
+      "attrs": { "axis": 1, "block_size": 4, "output_dtype": 3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 8.0, 12.0, 16.0, 20.0, 10.0, 20.0, 30.0, 40.0, -8.0, -4.0, 0.0, 4.0]
+          }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [1.0, 4.0, 10.0, 2.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [2, 8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 2, 3, 4, 5, 1, 2, 3, 4, -4, -2, 0, 2] }
+        }
+      }
+    },
+    {
+      "name": "blocked_last_axis_vec4_int8_with_zero_point_exact_blocks",
+      "provenance": {
+        "notes": "The blocked last-axis vec4 path uses an int8 zero point, checking the signed zero-point read independently of the uint8 form."
+      },
+      "attrs": { "axis": 1, "block_size": 4, "output_dtype": 3 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11, "scale": 2.0 }
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.25, 0.5, 0.125, 0.0625] }
+        },
+        "y_zero_point": { "dtype": "int8", "shape": [2, 2], "data": { "kind": "values", "values": [-8, 4, 0, -32] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [2, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_cross_axis_rows_width6_with_zero_point",
+      "provenance": {
+        "notes": "A vec4 crosses each six-element row boundary, so its four lanes may use two different per-axis scale and zero-point entries. Locks the lane-specific parameter-index path."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 6],
+          "data": { "kind": "cycle", "values": [-3.2, -1.1, 0.0, 0.9, 2.4, 7.8] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, 0.5, 1.0, 2.0] } },
+        "y_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [128, 120, 100, 80] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4, 6], "tolerance": 0 } }
+    },
+    {
+      "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": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4, 6],
+          "data": { "kind": "cycle", "values": [-3.2, -1.1, 0.0, 0.9, 2.4, 7.8, -5.6] }
+        },
+        "y_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, 0.5, 1.0, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [4, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_default_precision_rounds_division_in_f16",
+      "provenance": {
+        "notes": "Exact opset-25 precision witness: with f16 y_scale and omitted precision, x/y_scale is evaluated in f16. The f16 quotient rounds above 13.5 and then rounds to integer 14; an f32 division incorrectly produces 13."
+      },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [35.25] } },
+        "y_scale": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.611328125] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1], "tolerance": 0, "data": { "kind": "values", "values": [14] } }
+      }
+    },
+    {
+      "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": {
+        "x": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [35.25] } },
+        "y_scale": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.611328125] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1], "tolerance": 0, "data": { "kind": "values", "values": [13] } }
+      }
+    },
+    {
+      "name": "f32_x_f16_scale_default_precision",
+      "provenance": {
+        "notes": "Independent T1/T2 type-variable witness: float32 x and float16 y_scale are a standard mixed route. With precision omitted, the scale type selects f16 division and the quotient rounds to integer 14."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [35.25] } },
+        "y_scale": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.611328125] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1], "tolerance": 0, "data": { "kind": "values", "values": [14] } }
+      }
+    },
+    {
+      "name": "f16_x_f32_scale_default_precision",
+      "provenance": {
+        "notes": "Independent T1/T2 type-variable witness: float16 x and float32 y_scale are a standard mixed route. With precision omitted, the scale type selects f32 division and the quotient rounds to integer 13."
+      },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [35.25] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.611328125] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1], "tolerance": 0, "data": { "kind": "values", "values": [13] } }
+      }
+    },
+    {
+      "name": "f32_inputs_explicit_float16_precision",
+      "provenance": {
+        "notes": "Exact opset-25 precision=FLOAT16 witness with float32 x and y_scale. The division is explicitly evaluated in f16, where the quotient rounds above 13.5 and then rounds to integer 14; ignoring precision and dividing in f32 incorrectly produces 13."
+      },
+      "requires": { "features": ["shader-f16"] },
+      "attrs": { "precision": 10 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [35.25] } },
+        "y_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.611328125] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1], "tolerance": 0, "data": { "kind": "values", "values": [14] } }
+      }
+    }
+  ]
+}