sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +74 -0
build/webgpu/bench.json +247 -0
build/webgpu/manifest.json +646 -0
build/webgpu/metadata.json +20 -0
build/webgpu/quant-linear-blocked-axis.wgsl.jinja +43 -0
build/webgpu/quant-linear-scalar.wgsl.jinja +63 -0
build/webgpu/quant-linear-vec4.wgsl.jinja +107 -0
build/webgpu/test.json +1020 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.DequantizeLinear
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 25
+## Description
+Dequantizes a quantized tensor back to full precision using the formula `y = (x - x_zero_point) * x_scale`. Scale and zero point determine quantization granularity: scalar for per-tensor, 1-D for per-axis, or same rank as input for blocked quantization. The output type matches `x_scale` unless overridden by `output_dtype`.
+See the [ONNX `DequantizeLinear` spec](https://onnx.ai/onnx/operators/onnx__DequantizeLinear.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `x` | `x` | `TQ` | — | — | N-D quantized input tensor to be dequantized. | required |
+| `x_scale` | `x_scale` | `TF` | — | — | Scale for `x`; scalar for per-tensor, 1-D for per-axis, or same-rank tensor for blocked dequantization. | required |
+| `x_zero_point` | `x_zero_point` | `TQ` | — | — | Zero point for `x` with shape matching `x_scale`; defaults to zero when absent. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `y` | `y` | `TF` | same as `x` | same as `x` | N-D full-precision output with the same shape as `x`, typed according to `x_scale` or `output_dtype`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `1` | The axis of the dequantizing dimension, used for per-axis and blocked quantization; negative values index from the end. |
+| `block_size` | `0` | Number of elements along `axis` that share each scale value for blocked quantization; 0 means not blocked. |
+| `output_dtype` | `0` | ONNX TensorProto element-type code for `y`; `0` uses the dtype of `x_scale`. The current float16/float32 routes require it to agree with `x_scale`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `TQ` | `uint8`, `int8`, `int16`, `int32` |
+| `TF` | `float32`, `float16` |
+## 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 derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.DequantizeLinear", { version: 1 });
+const { y } = await kernel({
+  x: { data: xData, shape: [4] },
+  x_scale: { data: x_scaleData, shape: [] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,247 @@

+{
+  "op": "ai.onnx.DequantizeLinear",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "uint8_1m_scalar",
+      "preset": "smoke",
+      "vars": { "inputDtype": "uint8", "outputDtype": "float32", "count": 1048576 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1024, 1024], "data": { "kind": "constant", "value": 127 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024, 1024] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "uint8_1m_no_zero_point",
+      "preset": "smoke",
+      "vars": { "inputDtype": "uint8", "outputDtype": "float32", "count": 1048576 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1024, 1024], "data": { "kind": "constant", "value": 127 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024, 1024] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "uint8_1m_plus3_scalar",
+      "preset": "smoke",
+      "vars": { "inputDtype": "uint8", "outputDtype": "float32", "count": 1048579 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1048579], "data": { "kind": "constant", "value": 127 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "int8_1m_plus3_no_zero_point",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 1048579 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1048579], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "int8_1m_f16",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float16", "count": 1048576 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1024, 1024], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1024, 1024] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "int8_per_axis0_512x2048",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 1048576 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [512, 2048], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [512], "data": { "kind": "constant", "value": 3 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [512, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_channel_innermost_axis1_scalar_fallback_2048x2048",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 4194304 },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [2048, 2048], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [2048], "data": { "kind": "constant", "value": 3 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2048, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_channel_axis0_vec4_healthy_2048x2048",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 4194304 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [2048, 2048], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [2048], "data": { "kind": "constant", "value": 3 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2048, 2048] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_channel_axis0_unaligned_inner2050_scalar_cliff",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 1049600 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [512, 2050], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "x_zero_point": { "dtype": "int8", "shape": [512], "data": { "kind": "constant", "value": 3 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [512, 2050] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "scalar_above_67M_vec4_ceiling_fallback",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Stress-only capacity case: uint8 elements use 4-byte WebGPU storage, bringing the declared footprint to 512 MiB."
+      },
+      "vars": { "inputDtype": "uint8", "outputDtype": "float32", "count": 67108864 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [8192, 8192], "data": { "kind": "constant", "value": 127 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.02] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8192, 8192] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "f16_tail_scalar_x4_1m_plus3",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float16", "count": 1048579 },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1048579], "data": { "kind": "constant", "value": -42 } },
+        "x_scale": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.125] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "blocked_rank4_axis1_bs32_stress_4m",
+      "preset": "stress",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 4194304 },
+      "attrs": { "axis": 1, "block_size": 32 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4, 512, 32, 64], "dist": "randint", "seed": 101, "min": -128, "max": 127 },
+        "x_scale": { "dtype": "float32", "shape": [4, 16, 32, 64], "dist": "normal", "seed": 102, "scale": 0.02 },
+        "x_zero_point": {
+          "dtype": "int8",
+          "shape": [4, 16, 32, 64],
+          "dist": "randint",
+          "seed": 103,
+          "min": -8,
+          "max": 8
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 512, 32, 64], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "per_axis_mid_axis1_rank3_inner_not_mult4_scalar_fallback_stress",
+      "preset": "stress",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 8396800 },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [8, 512, 2050], "dist": "randint", "seed": 201, "min": -128, "max": 127 },
+        "x_scale": { "dtype": "float32", "shape": [512], "dist": "normal", "seed": 202, "scale": 0.02 },
+        "x_zero_point": { "dtype": "int8", "shape": [512], "dist": "randint", "seed": 203, "min": -8, "max": 8 }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8, 512, 2050], "dist": "empty" } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,646 @@

+{
+  "domain": "ai.onnx",
+  "name": "DequantizeLinear",
+  "sinceVersion": 25,
+  "description": "Dequantizes a quantized tensor back to full precision using the formula `y = (x - x_zero_point) * x_scale`. Scale and zero point determine quantization granularity: scalar for per-tensor, 1-D for per-axis, or same rank as input for blocked quantization. The output type matches `x_scale` unless overridden by `output_dtype`.",
+  "inputs": [
+    { "role": "x", "dtype": "TQ", "description": "N-D quantized input tensor to be dequantized." },
+    {
+      "role": "x_scale",
+      "dtype": "TF",
+      "description": "Scale for `x`; scalar for per-tensor, 1-D for per-axis, or same-rank tensor for blocked dequantization."
+    },
+    {
+      "role": "x_zero_point",
+      "dtype": "TQ",
+      "optional": true,
+      "description": "Zero point for `x` with shape matching `x_scale`; defaults to zero when absent."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "y",
+      "dtype": "TF",
+      "rank": "ranks.x",
+      "description": "N-D full-precision output with the same shape as `x`, typed according to `x_scale` or `output_dtype`.",
+      "shape": "shapes.x"
+    }
+  ],
+  "attributes": { "axis": 1, "block_size": 0, "output_dtype": 0 },
+  "attributeDescriptions": {
+    "axis": "The axis of the dequantizing dimension, used for per-axis and blocked quantization; negative values index from the end.",
+    "block_size": "Number of elements along `axis` that share each scale value for blocked quantization; 0 means not blocked.",
+    "output_dtype": "ONNX TensorProto element-type code for `y`; `0` uses the dtype of `x_scale`. The current float16/float32 routes require it to agree with `x_scale`."
+  },
+  "typeConstraints": { "TQ": ["uint8", "int8", "int16", "int32"], "TF": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "x", "role": "input" },
+    "x_scale": { "kind": "tensor", "semantic": "x_scale", "role": "input" },
+    "x_zero_point": { "kind": "tensor", "semantic": "x_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",
+    "sameShapeOk": "ranks.y == ranks.x and numel(shapes.x) == numel(shapes.y) and outputDispatchFits",
+    "dequantDtypesOk": "(attrs.output_dtype == 0 or attrs.output_dtype == onnxDtypeCode(logicalDtypes.TF)) and f16Ok(dtypes.TF)",
+    "blockedScaleOk": "attrs.block_size > 0 and ranks.x >= 2 and ranks.x_scale == ranks.x and sameShapeOk and dim(shapes.x_scale, attrs.axis) == ceil(dim(shapes.x, attrs.axis) / attrs.block_size)",
+    "elementCount": "numel(shapes.y)",
+    "elementCount4": "floor(numel(shapes.y) / 4)",
+    "scaleSize": "1 if ranks.x_scale == 0 else dim(shapes.x_scale, 0)",
+    "scaleInner": "1 if ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1 else inner(shapes.x, attrs.axis)",
+    "blockedAxisDim": "dim(shapes.x, attrs.axis) if ranks.x >= 2 else 1",
+    "blockedScaleAxisDim": "dim(shapes.x_scale, attrs.axis) if ranks.x_scale >= 2 else 1",
+    "blockedInner": "inner(shapes.x, attrs.axis) if ranks.x >= 2 else 1"
+  },
+  "constants": {
+    "usesF16": "dtypes.TF == \"f16\"",
+    "xScalar": "dtypes.TQ",
+    "xVec4": "\"vec4<\" ~ dtypes.TQ ~ \">\"",
+    "scaleScalar": "dtypes.TF",
+    "scaleVec4": "\"vec4<\" ~ dtypes.TF ~ \">\"",
+    "yScalar": "dtypes.TF",
+    "yVec4": "\"vec4<\" ~ dtypes.TF ~ \">\"",
+    "xUnsigned": "tensorDtypes.x == \"uint8\""
+  },
+  "variants": [
+    {
+      "id": "blocked_with_zero_point",
+      "priority": 19,
+      "when": ["present.x_zero_point", "blockedScaleOk", "ranks.x_zero_point == ranks.x", "dim(shapes.x_zero_point, attrs.axis) == dim(shapes.x_scale, attrs.axis)", "dequantDtypesOk"],
+      "constants": { "hasZero": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.Blocked",
+          "source": { "shader": "quant-linear-blocked-axis.wgsl.jinja", "inputs": { "op": "\"dequantize\"" } },
+          "bindings": "blockedScalarWithZero",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "blocked_no_zero_point",
+      "priority": 19,
+      "when": ["not present.x_zero_point", "blockedScaleOk", "dequantDtypesOk"],
+      "constants": { "hasZero": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.BlockedNoZero",
+          "source": { "shader": "quant-linear-blocked-axis.wgsl.jinja", "inputs": { "op": "\"dequantize\"" } },
+          "bindings": "blockedScalarNoZero",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "innermost_axis_vec4_with_zero_point",
+      "priority": 20,
+      "when": ["present.x_zero_point", "ranks.x >= 1", "attrs.axis == -1 or attrs.axis == ranks.x - 1", "ranks.x_scale == 1", "dim(shapes.x_scale, 0) == dim(shapes.x, ranks.x - 1)", "dim(shapes.x_scale, 0) % 4 == 0", "ranks.x_zero_point == 1", "dim(shapes.x_zero_point, 0) == dim(shapes.x_scale, 0)", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "sameShapeOk", "dequantDtypesOk"],
+      "constants": { "hasZero": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.InnermostAxisVec4",
+          "source": {
+            "shader": "quant-linear-vec4.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "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.x_zero_point", "ranks.x >= 1", "attrs.axis == -1 or attrs.axis == ranks.x - 1", "ranks.x_scale == 1", "dim(shapes.x_scale, 0) == dim(shapes.x, ranks.x - 1)", "dim(shapes.x_scale, 0) % 4 == 0", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "sameShapeOk", "dequantDtypesOk"],
+      "constants": { "hasZero": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.InnermostAxisVec4NoZero",
+          "source": {
+            "shader": "quant-linear-vec4.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "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.x_zero_point", "attrs.block_size == 0", "ranks.x >= 1", "attrs.axis >= 0", "attrs.axis < ranks.x", "ranks.x_scale == 1", "dim(shapes.x_scale, 0) == dim(shapes.x, attrs.axis)", "inner(shapes.x, attrs.axis) % 4 != 0", "ranks.x_zero_point == 1", "dim(shapes.x_zero_point, 0) == dim(shapes.x_scale, 0)", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "sameShapeOk", "dequantDtypesOk"],
+      "constants": { "hasZero": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.Vec4CrossAxis",
+          "source": {
+            "shader": "quant-linear-vec4.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "perAxis": true, "crossingParams": true }
+          },
+          "bindings": "linearVec4WithZero",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_with_zero_point",
+      "priority": 15,
+      "when": ["present.x_zero_point", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "(ranks.x_scale == 0 or ranks.x_scale == 1)", "sameShapeOk", "(ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.x_scale, 0) == dim(shapes.x, attrs.axis) and inner(shapes.x, attrs.axis) % 4 == 0))", "(ranks.x_zero_point == 0 or ranks.x_zero_point == 1)", "(ranks.x_zero_point == 0 or dim(shapes.x_zero_point, 0) == 1 or (ranks.x >= 1 and dim(shapes.x_zero_point, 0) == dim(shapes.x, attrs.axis)))", "dequantDtypesOk"],
+      "constants": { "hasZero": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.Vec4",
+          "source": {
+            "shader": "quant-linear-vec4.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "perAxis": "not (ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1)" }
+          },
+          "bindings": "linearVec4WithZero",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_no_zero_point",
+      "priority": 15,
+      "when": ["not present.x_zero_point", "numel(shapes.y) > 0", "numel(shapes.y) % 4 == 0", "(ranks.x_scale == 0 or ranks.x_scale == 1)", "sameShapeOk", "(ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.x_scale, 0) == dim(shapes.x, attrs.axis) and inner(shapes.x, attrs.axis) % 4 == 0))", "dequantDtypesOk"],
+      "constants": { "hasZero": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear.Vec4",
+          "source": {
+            "shader": "quant-linear-vec4.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "perAxis": "not (ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1)" }
+          },
+          "bindings": "linearVec4NoZero",
+          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_tail_with_zero_point",
+      "priority": 10,
+      "when": ["present.x_zero_point", "numel(shapes.y) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.y) % 4 != 0", "(ranks.x_scale == 0 or (ranks.x_scale == 1 and dim(shapes.x_scale, 0) == 1))", "(ranks.x_zero_point == 0 or (ranks.x_zero_point == 1 and dim(shapes.x_zero_point, 0) == 1))", "sameShapeOk", "dequantDtypesOk"],
+      "constants": { "hasZero": true },
+      "passes": [
+        {
+          "id": "bulk",
+          "name": "DequantizeLinear.Vec4Bulk",
+          "source": { "shader": "quant-linear-vec4.wgsl.jinja", "inputs": { "op": "\"dequantize\"", "perAxis": false } },
+          "bindings": "linearVec4PerTensorWithZero",
+          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "tail",
+          "name": "DequantizeLinear.ScalarTail",
+          "source": {
+            "shader": "quant-linear-scalar.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "x4": true, "perAxis": false }
+          },
+          "bindings": "linearScalarPerTensorWithZero",
+          "dispatch": { "threads": 1, "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_tail_no_zero_point",
+      "priority": 10,
+      "when": ["not present.x_zero_point", "numel(shapes.y) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.y) % 4 != 0", "(ranks.x_scale == 0 or (ranks.x_scale == 1 and dim(shapes.x_scale, 0) == 1))", "sameShapeOk", "dequantDtypesOk"],
+      "constants": { "hasZero": false },
+      "passes": [
+        {
+          "id": "bulk",
+          "name": "DequantizeLinear.Vec4Bulk",
+          "source": { "shader": "quant-linear-vec4.wgsl.jinja", "inputs": { "op": "\"dequantize\"", "perAxis": false } },
+          "bindings": "linearVec4PerTensorNoZero",
+          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "tail",
+          "name": "DequantizeLinear.ScalarTail",
+          "source": {
+            "shader": "quant-linear-scalar.wgsl.jinja",
+            "inputs": { "op": "\"dequantize\"", "x4": true, "perAxis": false }
+          },
+          "bindings": "linearScalarPerTensorNoZero",
+          "dispatch": { "threads": 1, "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "linear_scalar_or_axis_zero",
+      "when": ["present.x_zero_point", "(ranks.x_scale == 0 or ranks.x_scale == 1)", "sameShapeOk", "(ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.x_scale, 0) == dim(shapes.x, attrs.axis)))", "(ranks.x_zero_point == 0 or ranks.x_zero_point == 1)", "(ranks.x_zero_point == 0 or dim(shapes.x_zero_point, 0) == 1 or (ranks.x >= 1 and dim(shapes.x_zero_point, 0) == dim(shapes.x, attrs.axis)))", "dequantDtypesOk"],
+      "constants": { "hasZero": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear",
+          "source": {
+            "shader": "quant-linear-scalar.wgsl.jinja",
+            "inputs": {
+              "op": "\"dequantize\"",
+              "x4": "false",
+              "perAxis": "not (ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1)"
+            }
+          },
+          "bindings": "linearScalarWithZero",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "linear_scalar_or_axis_no_zero",
+      "when": ["not present.x_zero_point", "(ranks.x_scale == 0 or ranks.x_scale == 1)", "sameShapeOk", "(ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1 or (ranks.x >= 1 and dim(shapes.x_scale, 0) == dim(shapes.x, attrs.axis)))", "dequantDtypesOk"],
+      "constants": { "hasZero": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "DequantizeLinear",
+          "source": {
+            "shader": "quant-linear-scalar.wgsl.jinja",
+            "inputs": {
+              "op": "\"dequantize\"",
+              "x4": "false",
+              "perAxis": "not (ranks.x_scale == 0 or dim(shapes.x_scale, 0) == 1)"
+            }
+          },
+          "bindings": "linearScalarNoZero",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ],
+  "bindingSets": {
+    "blockedScalarWithZero": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xScalar" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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" }
+          ]
+        }
+      }
+    ],
+    "linearInnermostVec4WithZero": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleVec4"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xVec4"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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" }
+          ]
+        }
+      }
+    ],
+    "linearVec4WithZero": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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" }
+          ]
+        }
+      }
+    ],
+    "linearVec4NoZero": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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" }
+          ]
+        }
+      }
+    ],
+    "linearVec4WithZeroIo": [
+      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$xVec4" },
+      {
+        "name": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_scale",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar"
+      },
+      {
+        "name": "x_zero_point",
+        "arg": "x_zero_point",
+        "semantic": "x_zero_point",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$xScalar"
+      },
+      { "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": "x_scale",
+        "arg": "x_scale",
+        "semantic": "x_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.DequantizeLinear",
+  "id": "_ai_onnx_dequantizelinear_webgpu_f3ea48b",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "2JjVTjr3ucOvyKXNWcKsP7v8YvivzR9BlRaf88Fb3ds=",
+      "manifest.json": "/hcoq4zBcnyUUUn2WGN+SXgx3E9rWCsw5mA5zEvo46k=",
+      "quant-linear-blocked-axis.wgsl.jinja": "haL3flgL3mhBuafYHdqoMWbtvP2znUTVbb/qxakgrxU=",
+      "quant-linear-scalar.wgsl.jinja": "9EKOoq47CEJJ0O/7QirueruVfkmZuOGcCKmov/VKacI=",
+      "quant-linear-vec4.wgsl.jinja": "yQkKYuSSLaGUTifXt1xBr54tkIjYD68ouw9UOLbPNW0=",
+      "test.json": "qEtiZ2PxB1uoK/fbGOrmRyt4bmOkki0eVGg92OA+6jI="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.DequantizeLinear" }
+}

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

	@@ -0,0 +1,43 @@

+// Blocked-axis scale indexing for {{ source.op | default("quantize") }}.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+fn read_x(index: u32) -> i32 {
+{% if xUnsigned %}
+  return i32(x[index]);
+{% else %}
+  return x[index];
+{% endif %}
+}
+fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
+{% if hasZero %}
+{% if xUnsigned %}
+  return i32(x_zero_point[index]);
+{% else %}
+  return x_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+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;
+}
+@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 result = f32(read_x(index) - read_zero({% if hasZero %}scaleIndex{% endif %})) * f32(x_scale[scaleIndex]);
+  y[index] = {{ yScalar }}(result);
+}

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

	@@ -0,0 +1,63 @@

+// 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 }}
+fn read_x(index: u32) -> i32 {
+{% if xUnsigned %}
+  return i32(x[index]);
+{% else %}
+  return x[index];
+{% endif %}
+}
+fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
+{% if hasZero %}
+{% if xUnsigned %}
+  return i32(x_zero_point[index]);
+{% else %}
+  return x_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+fn transform_one(index: u32, scale: f32, zero_point: i32) {
+  let result = f32(read_x(index) - zero_point) * scale;
+  y[index] = {{ yScalar }}(result);
+}
+@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(x_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(x_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,107 @@

+// 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 }}
+{% if source.vectorParams is defined and source.vectorParams %}
+fn read_zero4({% if hasZero %}index: u32{% endif %}) -> vec4<i32> {
+{% if hasZero %}
+{% if xUnsigned %}
+  return vec4<i32>(x_zero_point[index]);
+{% else %}
+  return x_zero_point[index];
+{% endif %}
+{% else %}
+  return vec4<i32>(0);
+{% endif %}
+}
+{% else %}
+fn read_zero({% if hasZero %}index: u32{% endif %}) -> i32 {
+{% if hasZero %}
+{% if xUnsigned %}
+  return i32(x_zero_point[index]);
+{% else %}
+  return x_zero_point[index];
+{% endif %}
+{% else %}
+  return 0;
+{% endif %}
+}
+{% 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>(x_scale[scale_index]);
+{% if xUnsigned %}
+  let r0 = f32(i32(xv.x) - zp4.x) * scale.x;
+  let r1 = f32(i32(xv.y) - zp4.y) * scale.y;
+  let r2 = f32(i32(xv.z) - zp4.z) * scale.z;
+  let r3 = f32(i32(xv.w) - zp4.w) * scale.w;
+{% else %}
+  let r0 = f32(xv.x - zp4.x) * scale.x;
+  let r1 = f32(xv.y - zp4.y) * scale.y;
+  let r2 = f32(xv.z - zp4.z) * scale.z;
+  let r3 = f32(xv.w - zp4.w) * scale.w;
+{% endif %}
+{% elif source.crossingParams is defined and source.crossingParams %}
+{% if xUnsigned %}
+  let r0 = f32(i32(xv.x) - zp0) * f32(x_scale[scale_index0]);
+  let r1 = f32(i32(xv.y) - zp1) * f32(x_scale[scale_index1]);
+  let r2 = f32(i32(xv.z) - zp2) * f32(x_scale[scale_index2]);
+  let r3 = f32(i32(xv.w) - zp3) * f32(x_scale[scale_index3]);
+{% else %}
+  let r0 = f32(xv.x - zp0) * f32(x_scale[scale_index0]);
+  let r1 = f32(xv.y - zp1) * f32(x_scale[scale_index1]);
+  let r2 = f32(xv.z - zp2) * f32(x_scale[scale_index2]);
+  let r3 = f32(xv.w - zp3) * f32(x_scale[scale_index3]);
+{% endif %}
+{% else %}
+  let scale = f32(x_scale[scale_index]);
+{% if xUnsigned %}
+  let r0 = f32(i32(xv.x) - zp) * scale;
+  let r1 = f32(i32(xv.y) - zp) * scale;
+  let r2 = f32(i32(xv.z) - zp) * scale;
+  let r3 = f32(i32(xv.w) - zp) * scale;
+{% else %}
+  let r0 = f32(xv.x - zp) * scale;
+  let r1 = f32(xv.y - zp) * scale;
+  let r2 = f32(xv.z - zp) * scale;
+  let r3 = f32(xv.w - zp) * scale;
+{% endif %}
+{% endif %}
+  y[i] = vec4<{{ yScalar }}>({{ yScalar }}(r0), {{ yScalar }}(r1), {{ yScalar }}(r2), {{ yScalar }}(r3));
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1020 @@

+{
+  "op": "ai.onnx.DequantizeLinear",
+  "fixtureArrays": {
+    "uint8_axis1_rank4_onnx_backend_input_x": [3, 89, 34, 200, 74, 59, 5, 24, 24, 87, 32, 13, 245, 99, 4, 142, 121, 102],
+    "blocked_no_zero_point_rank4_axis1_symmetric_input_x": [-8, -4, -2, -1, 0, 1, 2, 4, 8, 16, 32, 48, -16, -32, -48, -64]
+  },
+  "cases": [
+    {
+      "name": "int16_per_axis_vec4_boundaries",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "int16",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [-32768, -1, 0, 32767, 32767, 0, -1, -32768] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.0, 2.0, 0.25] } },
+        "x_zero_point": {
+          "dtype": "int16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-32768, -2, 1, 32767] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0.000001,
+          "data": { "kind": "values", "values": [0.0, 1.0, -2.0, 0.0, 32767.5, 2.0, -4.0, -16383.75] }
+        }
+      }
+    },
+    {
+      "name": "dispatch_cliff_scalar_scale_67M",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [67108864], "data": { "kind": "cycle", "values": [0, 1, 2, 3, 4, 5, 6, 7] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "constant", "value": 0.5 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [67108864], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "uint8_scalar_zero",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [2, 3], "data": { "kind": "values", "values": [0, 127, 128, 129, 255, 4] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3] } }
+    },
+    {
+      "name": "uint8_subnormal_scale_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Uint8",
+        "notes": "Valid positive subnormal scale: dequantized tiny values should not be flushed to zero."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 1, 2, 3] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Uint8",
+        "notes": "Subnormal scale with a nonzero zero point should preserve tiny signed dequantized offsets around zero."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [126, 127, 128, 129] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8",
+        "notes": "Signed-output companion: subnormal scale should preserve tiny signed dequantized offsets around zero."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [-2, -1, 0, 1, 2] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8",
+        "notes": "Vec4 signed-output companion for subnormal DequantizeLinear scale handling."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-1, 0, 1, 2] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Axis_Default",
+        "notes": "Per-axis companion for subnormal DequantizeLinear scale handling: each channel uses a valid subnormal scale and nonzero zero point."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [10, 21, 29, 42, 12, 18, 30, 41] }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, 1e-40, 1e-40, 1e-40] }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [10, 20, 30, 40] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 1e-40, -1e-40, 2e-40, 2e-40, -2e-40, 0.0, 1e-40] }
+        }
+      }
+    },
+    {
+      "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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Axis_0",
+        "notes": "Vec4 per-axis companion for subnormal DequantizeLinear 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": "uint8",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [10, 11, 12, 13, 22, 18, 20, 21] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, 1e-40] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [10, 20] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40, 2e-40, -2e-40, 0.0, 1e-40] }
+        }
+      }
+    },
+    {
+      "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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Without_Zero_Point",
+        "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": "uint8", "shape": [2, 4], "data": { "kind": "values", "values": [0, 1, 2, 3, 2, 0, 1, 0] } },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, 1e-40, 1e-40, 1e-40] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40, 2e-40, 0.0, 1e-40, 0.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 values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Without_Zero_Point",
+        "notes": "Vec4 per-axis extension of omitted-zero-point coverage with valid subnormal scales."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [2, 4], "data": { "kind": "values", "values": [0, 1, 2, 3, 2, 0, 0, 1] } },
+        "x_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, 1e-40] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 1e-40, 2e-40, 3e-40, 2e-40, 0.0, 0.0, 1e-40] }
+        }
+      }
+    },
+    {
+      "name": "int32_subnormal_scale_no_zero_point_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes denormal results of floating-point multiplication (e.g. 2 * 1e-40 -> 0); ORT computes (x - zp) * subnormal_scale on CPU with denormal support, so subnormal x_scale cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int32",
+        "notes": "Valid int32 dequantization with a subnormal scale and omitted zero point; outputs should remain finite subnormal values."
+      },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [-2, -1, 0, 1] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_axis1_f16",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [2, 3], "data": { "kind": "values", "values": [-128, -1, 0, 1, 2, 127] } },
+        "x_scale": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [0.25, 0.5, 2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [0, -1, 1] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3] } },
+      "tolerance": 0.002
+    },
+    {
+      "name": "uint8_no_zero_negative_axis",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [2, 2], "data": { "kind": "values", "values": [10, 20, 30, 40] } },
+        "x_scale": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.1, 0.25] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2] } }
+    },
+    {
+      "name": "ort_int16_negative_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int16"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int16", "shape": [4], "data": { "kind": "values", "values": [-300, -30, -1025, 1270] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int16", "shape": [], "data": { "kind": "values", "values": [-1024] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "uint8_scalar_onnx_backend",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Uint8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 3, 128, 255] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint8_rank0_scale_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Uint8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 3, 128, 255] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "uint8_axis1_rank4_onnx_backend",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 3, 3, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/uint8_axis1_rank4_onnx_backend_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 4.0, 5.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [84, 24, 196] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "int8_scalar_zero_point_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-30, -3, 100, 127] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_rank0_scale_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-30, -3, 100, 127] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint8_large_tail_1039",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Uint8_Large"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1, 1039], "data": { "kind": "constant", "value": 1 } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1039], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_large_tail_1039",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8_Large"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1, 1039], "data": { "kind": "constant", "value": 1 } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 1039], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint8_2d_scalar_scale_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.2D"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 0, 1, 2, 3, 0, 10, 20, 30] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_scalar_input_rank0_with_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Scalar"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [100] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_scalar_input_rank0_no_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Without_Zero_Point"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [100] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_scalar_input_rank0_f16_with_zero_point",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpMLFloat16Test.Scalar"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [100] } },
+        "x_scale": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [], "tolerance": 0.002 } }
+    },
+    {
+      "name": "ort_opset25_f16_int8_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpMLFloat16Test.Opset25_PerAxisInt8_Cuda"
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [2, 4], "data": { "kind": "values", "values": [-2, -1, 0, 1, 2, 3, 4, 5] } },
+        "x_scale": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [2.0, 2.0, 4.0, 4.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 4], "tolerance": 0.002 } }
+    },
+    {
+      "name": "ort_no_zero_point_int8_1x8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.No_Zero_Point_int8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [-10, 50, 100, 120, -9, 49, 99, 119] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_no_zero_point_uint8_1x8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.No_Zero_Point_uint8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [10, 50, 100, 180, 9, 49, 99, 179] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_point_int8_1x8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Zero_Point_int8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [-10, 50, 100, 120, -9, 49, 99, 119] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-10] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_point_uint8_1x8",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Zero_Point_uint8"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [10, 50, 100, 180, 9, 49, 99, 119] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [10] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "int8_no_zero_scalar_explicit_float32_output_dtype",
+      "provenance": {
+        "notes": "Explicit output_dtype=FLOAT is the non-default spelling of the same float32 route selected from x_scale when output_dtype is omitted."
+      },
+      "attrs": { "axis": 0, "output_dtype": 1 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [-128, -1, 0, 1, 127] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "scalar_x4_int8_no_zero_tail",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [17],
+          "data": { "kind": "values", "values": [-64, -32, -16, -8, -4, -2, -1, 0, 1, 2, 4, 8, 16, 32, 48, 64, 96] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [17], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "vec4_tail_int8_no_zero_4097",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4097], "data": { "kind": "constant", "value": -4 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4097], "tolerance": 0.000001 } },
+      "provenance": { "notes": "Covers the packed bulk plus scalar tail without a zero point." }
+    },
+    {
+      "name": "scalar_x4_uint8_with_zero_point_tail",
+      "attrs": { "axis": 0 },
+      "provenance": {
+        "notes": "Compact sibling of the large uint8 tail benchmark: scalar scale and scalar uint8 zero point with numel % 4 != 0 selects the scalar_x4_with_zero_point path and exercises tail handling."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [17],
+          "data": {
+            "kind": "values",
+            "values": [0, 3, 7, 11, 15, 31, 63, 95, 127, 128, 129, 160, 192, 224, 240, 252, 255]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [17], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "vec4_tail_uint8_with_zero_point_4097",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4097], "data": { "kind": "constant", "value": 7 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4097], "tolerance": 0.000001 } },
+      "provenance": { "notes": "Covers the packed bulk plus scalar tail with a scalar zero point." }
+    },
+    {
+      "name": "uint8_axis0_rank2_per_row",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [3, 2], "data": { "kind": "values", "values": [6, 14, 40, 42, 20, 130] } },
+        "x_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, 1.0, 2.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [10, 20, 30] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint8_per_channel_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Axis_0"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 0, 1, 2, 3, 0, 10, 20, 30] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 4.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_per_channel_axis1_nonzero_zero_points",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Axis_1_int8"
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 0, 2, 4, 6, 0, 10, 20, 30] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 4.0, 8.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [0, -10, -20, -30] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "int32_scalar_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int32",
+        "notes": "Uses a length-1 scale tensor for the framework's scalar-style fixture encoding."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [-30, -3, 100, 127] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "int32_nonzero_zero_point_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int32_Non_Zero_Zero_Point",
+        "notes": "Uses length-1 scale and zero-point tensors for the framework's scalar-style fixture encoding."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [-30, -3, 100, 127] } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_per_channel_default_axis_rank4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Axis_Default"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 3, 2, 4],
+          "data": {
+            "kind": "values",
+            "values": [7, 9, 10, 10, 5, 8, 9, 1, 8, 6, 7, 9, 10, 0, 7, 10, 8, 2, 6, 0, 5, 9, 8, 1, 2, 7, 5, 3, 2, 4, 1, 3, 8, 7, 4, 8, 10, 1, 5, 5, 7, 7, 0, 2, 4, 4, 0, 5]
+          }
+        },
+        "x_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 10.0, 7.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [10, 2, 1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int32_per_channel_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Axis_1_int32"
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 0, 2, 4, 6, 0, 10, 20, 30] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 4.0, 8.0] } },
+        "x_zero_point": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint8_per_channel_negative_axis_minus2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Per_Channel_Neg_2"
+      },
+      "attrs": { "axis": -2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 0, 1, 2, 3, 0, 10, 20, 30] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 4.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [0, 0, 0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_uint8_large_vectorized_zero",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Uint8_Large",
+        "notes": "Same non-vector-aligned element count as ORT, flattened from [1, 1039] to [1039]."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [1039], "data": { "kind": "constant", "value": 1 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1039], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_large_vectorized_zero",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8_Large",
+        "notes": "Same non-vector-aligned element count as ORT, flattened from [1, 1039] to [1039]."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1039], "data": { "kind": "constant", "value": 1 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1039], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_non_aligned_size_135",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOpTest.Int8_NonAlignedSize_Initializer",
+        "notes": "Same non-aligned length; constant data keeps the fixture compact."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [135], "data": { "kind": "constant", "value": 7 } },
+        "x_scale": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [135], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_dequantizelinear",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dequantizelinear" },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 3, 128, 255] } },
+        "x_scale": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_dequantizelinear_axis",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dequantizelinear_axis" },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 3, 3, 2],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/uint8_axis1_rank4_onnx_backend_input_x" } }
+        },
+        "x_scale": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 4.0, 5.0] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [84, 24, 196] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 3, 3, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_dequantizelinear_blocked",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dequantizelinear_blocked" },
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [1, 4, 3, 2],
+          "data": {
+            "kind": "values",
+            "values": [3, 89, 34, 200, 74, 59, 5, 24, 24, 87, 32, 13, 5, 12, 12, 33, 65, 42, 245, 99, 4, 142, 121, 102]
+          }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 2],
+          "data": { "kind": "values", "values": [3.0, 2.0, 4.0, 1.0, 2.0, 2.0, 5.0, 2.0, 4.0, 3.0, 5.0, 2.0] }
+        },
+        "x_zero_point": {
+          "dtype": "uint8",
+          "shape": [1, 2, 3, 2],
+          "data": { "kind": "values", "values": [1, 0, 0, 1, 2, 20, 3, 2, 4, 3, 15, 2] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 3, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_blocked_uint8_with_zero_point_rank3_axis2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOp21BlockedTest.UnsignedInt_UseZeroPoint_LastAxis",
+        "notes": "Valid blocked dequantization with per-block zero-points on the last axis."
+      },
+      "attrs": { "axis": 2, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [2, 2, 4],
+          "data": { "kind": "values", "values": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15] }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [-2.0, -4.0, 3.5, 1.0, 2.0, 4.0, -3.5, -1.0] }
+        },
+        "x_zero_point": {
+          "dtype": "uint8",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [2, 0, 1, 9, 13, 5, 11, 6] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_blocked_uint8_no_zero_point_rank3_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc",
+        "test": "DequantizeLinearOp21BlockedTest.UnsignedInt_NoZeroPoint_FirstAxis"
+      },
+      "attrs": { "axis": 0, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [4, 2, 2],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 20, 22, 24, 26, 40, 42, 44, 46] }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "x_scale": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "x_zero_point": { "dtype": "uint8", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_blocked_rank4_axis1_nonzero_zp",
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 3, 2],
+          "data": {
+            "kind": "values",
+            "values": [-50, 12, -7, 88, 30, -41, 5, -24, 24, 87, -32, 13, 5, -12, 12, -33, 65, 42, -120, 99, 4, -100, 121, -102]
+          }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 2],
+          "data": { "kind": "values", "values": [3.0, 2.0, 4.0, 1.0, 2.0, 2.0, 5.0, 2.0, 4.0, 3.0, 5.0, 2.0] }
+        },
+        "x_zero_point": {
+          "dtype": "int8",
+          "shape": [1, 2, 3, 2],
+          "data": { "kind": "values", "values": [1, 0, -2, 1, 2, -20, 3, 2, -4, 3, 15, 2] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "blocked_no_zero_point_rank4_axis1_symmetric",
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/blocked_no_zero_point_rank4_axis1_symmetric_input_x" }
+          }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [0.5, 1.0, 1.5, 2.0, 0.25, 0.75, 1.25, 1.75] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 4, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "empty_blocked_rank4_axis1_zero_channel",
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1, 0, 2, 2], "data": { "kind": "values", "values": [] } },
+        "x_scale": { "dtype": "float32", "shape": [1, 0, 2, 2], "data": { "kind": "values", "values": [] } },
+        "x_zero_point": { "dtype": "int8", "shape": [1, 0, 2, 2], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1, 0, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "per_axis0_dispatch_fold_17M_last_rows",
+      "requires": { "limits": { "maxBufferSize": 281018368, "maxStorageBufferBindingSize": 281018368 } },
+      "provenance": {
+        "notes": "Its largest tensor is 281018368 bytes, so the case needs an adapter whose maxBufferSize and maxStorageBufferBindingSize both reach it — declared, because the WebGPU guaranteed minimums (256 MiB / 128 MiB) do not, and a device at them must report the case inapplicable rather than fail allocating it."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": {
+          "dtype": "uint8",
+          "shape": [67, 1048576],
+          "data": { "kind": "cycle", "values": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9] }
+        },
+        "x_scale": {
+          "dtype": "float32",
+          "shape": [67],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.05, "scale": 0.01, "offset": 0.02 }
+        },
+        "x_zero_point": { "dtype": "uint8", "shape": [67], "data": { "kind": "constant", "value": 5 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [67, 1048576], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "innermost_axis_vec4_no_zero_two_rows",
+      "provenance": {
+        "notes": "Exercises vec4-bound innermost per-axis scales without a zero point. Expected values come from the TypeScript dequantization reference."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [-8, -4, 2, 7, 10, -12, 20, -24] }
+        },
+        "x_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, 0.5, 1.5, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "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": "int8", "shape": [4, 6], "data": { "kind": "cycle", "values": [-42, -8, 0, 7, 31, 63] } },
+        "x_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, 0.5, 1.0, 2.0] } },
+        "x_zero_point": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-3, 2, 5, -7] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "vec4_cross_axis_rows_width6_no_zero_point",
+      "provenance": {
+        "notes": "Symmetric per-channel dequantization over six-element rows with x_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; the seven-value input cycle is coprime with the six-wide row so no two rows repeat the same lane pattern. Expected values come from the TypeScript dequantization reference."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4, 6], "data": { "kind": "cycle", "values": [-42, -8, 0, 7, 31, 63, -19] } },
+        "x_scale": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, 0.5, 1.0, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "blocked_no_zero_point_rank4_axis1_symmetric_f16",
+      "provenance": {
+        "notes": "float16 output exercises half-precision storage on the blocked-axis route without a zero point."
+      },
+      "attrs": { "axis": 1, "block_size": 2 },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [1, 4, 2, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/blocked_no_zero_point_rank4_axis1_symmetric_input_x" }
+          }
+        },
+        "x_scale": {
+          "dtype": "float16",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [0.5, 1.0, 1.5, 2.0, 0.25, 0.75, 1.25, 1.75] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1, 4, 2, 2], "tolerance": 0 } }
+    }
+  ]
+}