sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +25 -11
build/webgpu/bench.json +14 -15
build/webgpu/dynamic-quantize-linear-quantize.wgsl.jinja +6 -8
build/webgpu/dynamic-quantize-linear-reduce.wgsl.jinja +43 -22
build/webgpu/dynamic-quantize-linear.wgsl.jinja +6 -7
build/webgpu/manifest.json +98 -230
build/webgpu/metadata.json +18 -9
build/webgpu/test.json +6 -7

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX `DynamicQuantizeLinear` spec](https://onnx.ai/onnx/operators/onnx_
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `x` | `x` | `T` | — | — | Float32 input tensor to quantize. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `y` | `y` | `TQ` | same as `x` | same as `x` | Quantized output tensor; same shape as the input. | required |
-| `y_scale` | `y_scale` | `T` | `0` | `[]` | Per-tensor scale factor derived from the input min/max range; scalar. | required |
-| `y_zero_point` | `y_zero_point` | `TQ` | `0` | `[]` | Per-tensor zero point for the quantization; scalar. | required |
 ## Type constraints
@@ -37,9 +37,19 @@ See the [ONNX `DynamicQuantizeLinear` spec](https://onnx.ai/onnx/operators/onnx_
 | `T` | `float32` |
 | `TQ` | `uint8` |
 ## 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
@@ -49,10 +59,14 @@ See the [ONNX `DynamicQuantizeLinear` spec](https://onnx.ai/onnx/operators/onnx_
 ## 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.

 ## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `x` | `T` | — | — | Float32 input tensor to quantize. | required |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `y` | `TQ` | same as `x` | same as `x` | Quantized output tensor; same shape as the input. | required |
+| `y_scale` | `T` | `0` | `[]` | Per-tensor scale factor derived from the input min/max range; scalar. | required |
+| `y_zero_point` | `TQ` | `0` | `[]` | Per-tensor zero point for the quantization; scalar. | required |
 ## Type constraints
 | `T` | `float32` |
 | `TQ` | `uint8` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `single_invocation` — Uses one invocation to find the range and quantize the tensor, avoiding partial buffers for small inputs. It also provides the fallback when the parallel reduction cannot satisfy device limits.
+- `parallel_subgroup_reduce_vec4` — Reduces independent input blocks to min/max partials, combines them, and quantizes in a separate pass. The family uses packed reads when the input length is vec4-aligned.
+- `parallel_subgroup_reduce` — Reduces independent input blocks to min/max partials, combines them, and quantizes in a separate pass. The family uses packed reads when the input length is vec4-aligned.
+- `grid_stride_reduce_vec4` — Caps the number of min/max partials and grid-strides each workgroup across the input. This bounds scratch size and finalization work for large tensors.
+- `grid_stride_reduce` — Caps the number of min/max partials and grid-strides each workgroup across the input. This bounds scratch size and finalization work for large tensors.
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.DynamicQuantizeLinear",
   "cases": [
     {
       "name": "f32_64k",
@@ -12,8 +11,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [65536] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       }
     },
     {
@@ -28,8 +27,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [2048, 2048] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       },
       "bench": {
         "metrics": [
@@ -54,8 +53,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [4194303] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       },
       "bench": {
         "primary": true,
@@ -81,8 +80,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [4194304] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       },
       "bench": {
         "metrics": [
@@ -107,8 +106,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [2097153] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       },
       "bench": {
         "metrics": [
@@ -133,8 +132,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [2097152] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       },
       "bench": {
         "metrics": [
@@ -159,8 +158,8 @@
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [262143] },
-        "y_scale": { "dtype": "float32", "shape": [1] },
-        "y_zero_point": { "dtype": "uint8", "shape": [1] }
       },
       "bench": {
         "metrics": [

 {
   "cases": [
     {
       "name": "f32_64k",
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [65536] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       }
     },
     {
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [2048, 2048] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       },
       "bench": {
         "metrics": [
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [4194303] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       },
       "bench": {
         "primary": true,
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [4194304] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       },
       "bench": {
         "metrics": [
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [2097153] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       },
       "bench": {
         "metrics": [
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [2097152] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       },
       "bench": {
         "metrics": [
       },
       "outputs": {
         "y": { "dtype": "uint8", "shape": [262143] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
       },
       "bench": {
         "metrics": [

build/webgpu/dynamic-quantize-linear-quantize.wgsl.jinja CHANGED Viewed

@@ -93,10 +93,9 @@ fn dynamic_quantize_exact_div_normal(numerator: f32, denominator: f32) -> f32 {
   let biased_exponent = quotient_exponent + 127;
   if (biased_exponent <= 0 || biased_exponent >= 255) {
-    // DynamicQuantizeLinear's GPU-supported ranges are normal and finite.
-    // Preserve the implementation-defined subnormal/overflow behavior outside
-    // that contract; the integer quantizer never takes this branch near a
-    // half-integer boundary.
     return numerator / denominator;
   }
   let result_bits = sign_bits
@@ -149,14 +148,13 @@ const EPT: u32 = {{ elemsPerThread }}u;
 {% endif %}
 @compute @workgroup_size(WG, 1, 1)
 fn main(@builtin(workgroup_id) wg: vec3<u32>,
-        @builtin(local_invocation_id) lid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
   let tid = lid.x;
   let scale = y_scale[0];
   let zp_i32 = i32(y_zero_point[0]);
-  // Fold the block grid across x/y at the device's per-axis workgroup limit.
   // Per-element guards discard the over-dispatched tail.
-  let blk = wg.x + wg.y * nwg.x;
 {% if vec4 %}
   // The vec4 input load reads four scalars at once. Output storage still uses
   // one u32 element for each quantized value.

   let biased_exponent = quotient_exponent + 127;
   if (biased_exponent <= 0 || biased_exponent >= 255) {
+    // The exact normal-range reconstruction below cannot encode a subnormal or
+    // overflowed quotient. Use WGSL division for those exponent ranges; this
+    // branch is disjoint from the finite halfway cases handled below.
     return numerator / denominator;
   }
   let result_bits = sign_bits
 {% endif %}
 @compute @workgroup_size(WG, 1, 1)
 fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
   let tid = lid.x;
   let scale = y_scale[0];
   let zp_i32 = i32(y_zero_point[0]);
+  // Fold the block grid across x/y at a fixed per-axis workgroup width.
   // Per-element guards discard the over-dispatched tail.
+  let blk = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
 {% if vec4 %}
   // The vec4 input load reads four scalars at once. Output storage still uses
   // one u32 element for each quantized value.

build/webgpu/dynamic-quantize-linear-reduce.wgsl.jinja CHANGED Viewed

@@ -1,14 +1,17 @@
 // Pass 1 of parallel DynamicQuantizeLinear. Regular mode reduces one contiguous
 // WG * EPT block per workgroup. Grid-stride mode caps the partial count and has
 // every lane revisit the tensor at grid-sized strides. Both modes use the same
-// subgroup/tree combine and write one min/max pair per workgroup.
 //
 // The local min/max start at 0.0 because the ONNX DynamicQuantizeLinear range
 // always includes zero; out-of-range lanes contribute the same neutral value.
-// f32 min/max is order-independent, so the result is bit-identical to the serial
-// loop and between subgroup and workgroup-tree combines. Without subgroups, a
-// full workgroup tree combines the per-thread partials.
-{% set gridStride = source.gridStride if source.gridStride is defined else false %}
 {% if useSubgroups %}
 enable subgroups;
 {% endif %}
@@ -24,10 +27,11 @@ var<workgroup> wgMax: array<f32, WG>;
 @compute @workgroup_size(WG, 1, 1)
 fn main(@builtin(workgroup_id) wg: vec3<u32>,
-        @builtin(local_invocation_id) lid: vec3<u32>,
         @builtin(num_workgroups) nwg: vec3<u32>
 {%- if useSubgroups %},
-        @builtin(subgroup_invocation_id) sgLid: u32,
         @builtin(subgroup_size) sgSize: u32
 {%- endif %}) {
   let tid = lid.x;
@@ -41,7 +45,7 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>,
   // per-dimension dispatch limit. The whole over-dispatched workgroup returns
   // before barriers or scratch stores.
   // blk/numBlocks are workgroup-uniform, so the return can't split the barrier.
-  let blk = wg.x + wg.y * nwg.x;
 {% if vec4 %}
   let numBlocks = (params.count / 4u + WG - 1u) / WG;
 {% else %}
@@ -93,22 +97,39 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>,
 {% if useSubgroups %}
   let sgMin = subgroupMin(localMin);
   let sgMax = subgroupMax(localMax);
-  let safeSg = max(sgSize, 1u);
-  // Do not assume that the device subgroup size divides the workgroup size.
-  let slotCount = max(1u, (WG + safeSg - 1u) / safeSg);
-  if (sgLid == 0u) {
-    let slot = min(tid / safeSg, WG - 1u);
-    wgMin[slot] = sgMin;
-    wgMax[slot] = sgMax;
   }
-  workgroupBarrier();
   if (tid == 0u) {
-    var totalMin = wgMin[0];
-    var totalMax = wgMax[0];
-    for (var i = 1u; i < slotCount; i = i + 1u) {
-      totalMin = min(totalMin, wgMin[i]);
-      totalMax = max(totalMax, wgMax[i]);
-    }
     partial_min[blk] = totalMin;
     partial_max[blk] = totalMax;
   }

 // Pass 1 of parallel DynamicQuantizeLinear. Regular mode reduces one contiguous
 // WG * EPT block per workgroup. Grid-stride mode caps the partial count and has
 // every lane revisit the tensor at grid-sized strides. Both modes use the same
+// subgroup/tree combine and write one min/max pair per workgroup. The subgroup
+// combine publishes one shared-memory slot per invocation (its subgroup pair
+// from the elected lane, the neutral 0.0 from every other lane) and every
+// subgroup folds all of them.
 //
 // The local min/max start at 0.0 because the ONNX DynamicQuantizeLinear range
 // always includes zero; out-of-range lanes contribute the same neutral value.
+// f32 min/max is order-independent, so subgroup and workgroup-tree combinations
+// produce the same extrema. Without subgroups, a full workgroup tree combines
+// the per-invocation partials.
+{% set gridStride = gridStride if gridStride is defined else false %}
 {% if useSubgroups %}
 enable subgroups;
 {% endif %}
 @compute @workgroup_size(WG, 1, 1)
 fn main(@builtin(workgroup_id) wg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>
+{%- if gridStride %},
         @builtin(num_workgroups) nwg: vec3<u32>
+{%- endif %}
 {%- if useSubgroups %},
         @builtin(subgroup_size) sgSize: u32
 {%- endif %}) {
   let tid = lid.x;
   // per-dimension dispatch limit. The whole over-dispatched workgroup returns
   // before barriers or scratch stores.
   // blk/numBlocks are workgroup-uniform, so the return can't split the barrier.
+  let blk = wg.x + wg.y * {{ DISPATCH_FOLD_WIDTH }}u;
 {% if vec4 %}
   let numBlocks = (params.count / 4u + WG - 1u) / WG;
 {% else %}
 {% if useSubgroups %}
   let sgMin = subgroupMin(localMin);
   let sgMax = subgroupMax(localMax);
+  // Cross-subgroup fold that assumes nothing about which invocations share a
+  // subgroup, how many subgroups there are, or which of a subgroup's lanes are
+  // active: every invocation owns the slot at its own index, the elected lane
+  // publishes its subgroup pair there and every other lane publishes 0.0, which
+  // is an exact identity here because every lane's local range already includes
+  // zero (so every published minimum is <= 0 and every maximum >= 0). Each
+  // subgroup then folds all WG slots — lane `rank`, its dense position among
+  // the active lanes, walks slots rank, rank + count, ... — and one more
+  // collective merges the lane partials, so every slot is merged exactly once
+  // at any legal width and partition. min/max is commutative and associative,
+  // so the fold order does not change the result.
+  var totalMin = sgMin;
+  var totalMax = sgMax;
+  // A one-subgroup workgroup is already fully reduced by the collectives above.
+  // The test reads the `subgroup_size` builtin, which is uniform; a collective's
+  // result is not uniform to WGSL's analysis and may not guard a barrier.
+  if (sgSize != WG) {
+    let rank = subgroupExclusiveAdd(1u);
+    let count = subgroupAdd(1u);
+    let leader = rank == 0u;
+    wgMin[tid] = select(0.0, sgMin, leader);
+    wgMax[tid] = select(0.0, sgMax, leader);
+    workgroupBarrier();
+    var foldMin = 0.0;
+    var foldMax = 0.0;
+    for (var i = rank; i < WG; i = i + count) {
+      foldMin = min(foldMin, wgMin[i]);
+      foldMax = max(foldMax, wgMax[i]);
+    }
+    totalMin = subgroupMin(foldMin);
+    totalMax = subgroupMax(foldMax);
   }
   if (tid == 0u) {
     partial_min[blk] = totalMin;
     partial_max[blk] = totalMax;
   }

build/webgpu/dynamic-quantize-linear.wgsl.jinja CHANGED Viewed

@@ -1,4 +1,4 @@
-{% if source.fromPartials %}
 // Parallel range finalization folds the guarded min/max partials and computes
 // the output scale and zero point. Each partial already includes zero in its range.
 {% else %}
@@ -95,10 +95,9 @@ fn dynamic_quantize_exact_div_normal(numerator: f32, denominator: f32) -> f32 {
   let biased_exponent = quotient_exponent + 127;
   if (biased_exponent <= 0 || biased_exponent >= 255) {
-    // DynamicQuantizeLinear's GPU-supported ranges are normal and finite.
-    // Preserve the implementation-defined subnormal/overflow behavior outside
-    // that contract; the integer quantizer never takes this branch near a
-    // half-integer boundary.
     return numerator / denominator;
   }
   let result_bits = sign_bits
@@ -147,7 +146,7 @@ fn round_dynamic_half_to_even(value: f32, scale: f32) -> i32 {
 @compute @workgroup_size(1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   if (gid.x != 0u) { return; }
-{% if source.fromPartials %}
   var min_value = partial_min[0];
   var max_value = partial_max[0];
   for (var i = 1u; i < params.numPartials; i = i + 1u) {
@@ -169,7 +168,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   let zp_i32 = clamp(round_dynamic_half_to_even(-min_value, scale), 0, 255);
   y_scale[0] = scale;
   y_zero_point[0] = u32(zp_i32);
-{% if not source.fromPartials %}
   for (var i = 0u; i < params.count; i = i + 1u) {
     let q = clamp(round_dynamic_half_to_even(x[i], scale) + zp_i32, 0, 255);
     y[i] = u32(q);

+{% if fromPartials %}
 // Parallel range finalization folds the guarded min/max partials and computes
 // the output scale and zero point. Each partial already includes zero in its range.
 {% else %}
   let biased_exponent = quotient_exponent + 127;
   if (biased_exponent <= 0 || biased_exponent >= 255) {
+    // The exact normal-range reconstruction below cannot encode a subnormal or
+    // overflowed quotient. Use WGSL division for those exponent ranges; this
+    // branch is disjoint from the finite halfway cases handled below.
     return numerator / denominator;
   }
   let result_bits = sign_bits
 @compute @workgroup_size(1)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   if (gid.x != 0u) { return; }
+{% if fromPartials %}
   var min_value = partial_min[0];
   var max_value = partial_max[0];
   for (var i = 1u; i < params.numPartials; i = i + 1u) {
   let zp_i32 = clamp(round_dynamic_half_to_even(-min_value, scale), 0, 255);
   y_scale[0] = scale;
   y_zero_point[0] = u32(zp_i32);
+{% if not fromPartials %}
   for (var i = 0u; i < params.count; i = i + 1u) {
     let q = clamp(round_dynamic_half_to_even(x[i], scale) + zp_i32, 0, 255);
     y[i] = u32(q);

build/webgpu/manifest.json CHANGED Viewed

@@ -2,49 +2,19 @@
   "domain": "ai.onnx",
   "name": "DynamicQuantizeLinear",
   "sinceVersion": 11,
-  "description": "Computes a per-tensor scale and zero point from the range of floating-point input `x`, extending the range to include zero, then quantizes each value to `uint8` as `saturate(round(x / y_scale) + y_zero_point)`. Uses round-to-nearest-even and clamps results to `[0, 255]`.",
-  "inputs": [{ "role": "x", "dtype": "T", "description": "Float32 input tensor to quantize." }],
-  "outputs": [
-    {
-      "role": "y",
-      "dtype": "TQ",
-      "rank": "ranks.x",
-      "description": "Quantized output tensor; same shape as the input.",
-      "shape": "shapes.x"
-    },
-    {
-      "role": "y_scale",
-      "dtype": "T",
-      "rank": 0,
-      "description": "Per-tensor scale factor derived from the input min/max range; scalar.",
-      "shape": []
-    },
-    {
-      "role": "y_zero_point",
-      "dtype": "TQ",
-      "rank": 0,
-      "description": "Per-tensor zero point for the quantization; scalar.",
-      "shape": []
-    }
-  ],
-  "typeConstraints": { "T": ["float32"], "TQ": ["uint8"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "x", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "y", "role": "output" },
-    "y_scale": { "kind": "tensor", "semantic": "y_scale", "role": "output" },
-    "y_zero_point": { "kind": "tensor", "semantic": "y_zero_point", "role": "output" }
   },
   "tunables": {
-    "WORKGROUP_SIZE": 256,
-    "ELEMENTS_PER_THREAD": 4,
-    "GRID_STRIDE_MIN_ELEMENTS": 262144,
-    "MAX_GRID_PARTIALS": 256,
-    "SERIAL_MAX_ELEMENTS": 8192
-  },
-  "tunableDescriptions": {
-    "GRID_STRIDE_MIN_ELEMENTS": "Sets the minimum input element count admitted to the bounded grid-stride reduction.",
-    "MAX_GRID_PARTIALS": "Caps the number of min/max partials emitted by the grid-stride reduction, bounding scratch and finalization work.",
-    "SERIAL_MAX_ELEMENTS": "Sets the input-size ceiling that prefers the one-invocation route; larger inputs use it only when the parallel reduction cannot fit."
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
@@ -58,182 +28,45 @@
     "parallelDeviceOk": "tunables.WORKGROUP_SIZE <= deviceWorkgroupCap and 8 * tunables.WORKGROUP_SIZE <= device.limits.maxComputeWorkgroupStorageSize",
     "fullScratchFits": "4 * fullPartials <= storageBufferLimit",
     "gridScratchFits": "4 * gridPartials <= storageBufferLimit",
-    "fullDispatchFits": "ceilDiv(fullPartials, device.limits.maxComputeWorkgroupsPerDimension) <= device.limits.maxComputeWorkgroupsPerDimension",
-    "gridDispatchFits": "gridPartials <= device.limits.maxComputeWorkgroupsPerDimension",
     "parallelFullFits": "parallelDeviceOk and fullScratchFits and fullDispatchFits",
     "parallelGridFits": "parallelDeviceOk and gridScratchFits and gridDispatchFits and fullDispatchFits",
     "serialFallbackNeeded": "inputCount <= tunables.SERIAL_MAX_ELEMENTS or not parallelFullFits"
   },
-  "bindingSets": {
-    "reduce": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "x",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$inputElement"
-      },
-      { "name": "partial_min", "semantic": "partial_min", "buffer": { "type": "storage" }, "elementType": "f32" },
-      { "name": "partial_max", "semantic": "partial_max", "buffer": { "type": "storage" }, "elementType": "f32" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.x)" }] }
-      }
-    ],
-    "finalizeFull": [
-      {
-        "name": "partial_min",
-        "semantic": "partial_min",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "partial_max",
-        "semantic": "partial_max",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "storage" },
-        "elementType": "f32",
-        "length": 1
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "storage" },
-        "elementType": "u32",
-        "length": 1
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "numPartials", "type": "u32", "value": "fullPartials" }] }
-      }
-    ],
-    "finalizeGrid": [
-      {
-        "name": "partial_min",
-        "semantic": "partial_min",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "partial_max",
-        "semantic": "partial_max",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "storage" },
-        "elementType": "f32",
-        "length": 1
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "storage" },
-        "elementType": "u32",
-        "length": 1
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "numPartials", "type": "u32", "value": "gridPartials" }] }
-      }
-    ],
-    "quantize": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "x",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$inputElement"
-      },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "f32",
-        "length": 1
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "u32",
-        "length": 1
-      },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "u32" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.x)" }] }
-      }
-    ],
-    "serial": [
-      { "name": "x", "arg": "x", "semantic": "x", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "y", "arg": "y", "semantic": "y", "buffer": { "type": "storage" }, "elementType": "u32" },
-      {
-        "name": "y_scale",
-        "arg": "y_scale",
-        "semantic": "y_scale",
-        "buffer": { "type": "storage" },
-        "elementType": "f32",
-        "length": 1
-      },
-      {
-        "name": "y_zero_point",
-        "arg": "y_zero_point",
-        "semantic": "y_zero_point",
-        "buffer": { "type": "storage" },
-        "elementType": "u32",
-        "length": 1
-      },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.x)" }] }
-      }
-    ]
   },
   "variants": [
     {
       "id": "single_invocation",
-      "description": "Uses one invocation to find the range and quantize the tensor, avoiding partial buffers for small inputs. It also provides the fallback when the parallel reduction cannot satisfy device limits.",
       "when": ["serialContract", "serialFallbackNeeded"],
       "passes": [
         {
           "id": "main",
           "name": "DynamicQuantizeLinear",
-          "source": { "shader": "dynamic-quantize-linear.wgsl.jinja", "inputs": { "fromPartials": false } },
-          "bindings": "serial",
           "dispatch": { "x": 1 }
         }
       ]
     },
     {
       "id": "parallel_subgroup_reduce_vec4",
-      "description": "Reduces independent input blocks to min/max partials, combines them, and quantizes in a separate pass. The family uses packed reads when the input length is vec4-aligned.",
       "priority": 11,
       "when": ["baseContract", "inputCount > 0", "inputCount % 4 == 0", "parallelFullFits"],
-      "constants": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": true,
@@ -249,34 +82,42 @@
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
           "shader": "dynamic-quantize-linear-reduce.wgsl.jinja",
-          "bindings": "reduce",
-          "dispatch": { "workgroups": "fullPartials" }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
-          "source": { "shader": "dynamic-quantize-linear.wgsl.jinja", "inputs": { "fromPartials": true } },
-          "bindings": "finalizeFull",
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
-          "bindings": "quantize",
           "dispatch": {
-            "threads": "ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]
     },
     {
       "id": "parallel_subgroup_reduce",
-      "description": "Reduces independent input blocks to min/max partials, combines them, and quantizes in a separate pass. The family uses packed reads when the input length is vec4-aligned.",
       "priority": 10,
       "when": ["baseContract", "inputCount > 0", "true", "parallelFullFits"],
-      "constants": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": false,
@@ -292,34 +133,42 @@
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
           "shader": "dynamic-quantize-linear-reduce.wgsl.jinja",
-          "bindings": "reduce",
-          "dispatch": { "workgroups": "fullPartials" }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
-          "source": { "shader": "dynamic-quantize-linear.wgsl.jinja", "inputs": { "fromPartials": true } },
-          "bindings": "finalizeFull",
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
-          "bindings": "quantize",
           "dispatch": {
-            "threads": "ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]
     },
     {
       "id": "grid_stride_reduce_vec4",
-      "description": "Caps the number of min/max partials and grid-strides each workgroup across the input. This bounds scratch size and finalization work for large tensors.",
       "priority": 12,
       "when": ["baseContract", "inputCount > 0", "inputCount % 4 == 0", "inputCount >= tunables.GRID_STRIDE_MIN_ELEMENTS", "parallelGridFits"],
-      "constants": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": true,
@@ -334,35 +183,44 @@
         {
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
-          "source": { "shader": "dynamic-quantize-linear-reduce.wgsl.jinja", "inputs": { "gridStride": true } },
-          "bindings": "reduce",
           "dispatch": { "x": "gridPartials" }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
-          "source": { "shader": "dynamic-quantize-linear.wgsl.jinja", "inputs": { "fromPartials": true } },
-          "bindings": "finalizeGrid",
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
-          "bindings": "quantize",
           "dispatch": {
-            "threads": "ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]
     },
     {
       "id": "grid_stride_reduce",
-      "description": "Caps the number of min/max partials and grid-strides each workgroup across the input. This bounds scratch size and finalization work for large tensors.",
       "priority": 12,
       "when": ["baseContract", "inputCount > 0", "inputCount % 4 != 0", "inputCount >= tunables.GRID_STRIDE_MIN_ELEMENTS", "parallelGridFits"],
-      "constants": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": false,
@@ -377,25 +235,35 @@
         {
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
-          "source": { "shader": "dynamic-quantize-linear-reduce.wgsl.jinja", "inputs": { "gridStride": true } },
-          "bindings": "reduce",
           "dispatch": { "x": "gridPartials" }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
-          "source": { "shader": "dynamic-quantize-linear.wgsl.jinja", "inputs": { "fromPartials": true } },
-          "bindings": "finalizeGrid",
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
-          "bindings": "quantize",
           "dispatch": {
-            "threads": "ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]

   "domain": "ai.onnx",
   "name": "DynamicQuantizeLinear",
   "sinceVersion": 11,
+  "inputs": { "x": { "dtype": "T" } },
+  "outputs": {
+    "y": { "dtype": "TQ", "rank": "ranks.x", "shape": "shapes.x" },
+    "y_scale": { "dtype": "T", "rank": 0, "shape": [] },
+    "y_zero_point": { "dtype": "TQ", "rank": 0, "shape": [] }
   },
+  "typeConstraints": { "T": ["float32"], "TQ": ["uint8"] },
   "tunables": {
+    "WORKGROUP_SIZE": { "default": 256 },
+    "ELEMENTS_PER_THREAD": { "default": 4 },
+    "GRID_STRIDE_MIN_ELEMENTS": { "default": 262144 },
+    "MAX_GRID_PARTIALS": { "default": 256 },
+    "SERIAL_MAX_ELEMENTS": { "default": 8192 }
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "parallelDeviceOk": "tunables.WORKGROUP_SIZE <= deviceWorkgroupCap and 8 * tunables.WORKGROUP_SIZE <= device.limits.maxComputeWorkgroupStorageSize",
     "fullScratchFits": "4 * fullPartials <= storageBufferLimit",
     "gridScratchFits": "4 * gridPartials <= storageBufferLimit",
+    "fullDispatchFits": "ceilDiv(fullPartials, min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "gridDispatchFits": "gridPartials <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "parallelFullFits": "parallelDeviceOk and fullScratchFits and fullDispatchFits",
     "parallelGridFits": "parallelDeviceOk and gridScratchFits and gridDispatchFits and fullDispatchFits",
     "serialFallbackNeeded": "inputCount <= tunables.SERIAL_MAX_ELEMENTS or not parallelFullFits"
   },
+  "bindings": {
+    "y": { "buffer": "storage", "elementType": "u32" },
+    "y_scale": { "buffer": "storage", "elementType": "f32", "length": 1 },
+    "y_zero_point": { "buffer": "storage", "elementType": "u32", "length": 1 },
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.x)" }] },
+    "x_2": { "name": "x", "buffer": "read-only-storage", "elementType": "$inputElement" },
+    "partial_min": { "buffer": "storage", "elementType": "f32" },
+    "partial_max": { "buffer": "storage", "elementType": "f32" },
+    "partial_min_2": { "name": "partial_min", "buffer": "read-only-storage", "elementType": "f32" },
+    "partial_max_2": { "name": "partial_max", "buffer": "read-only-storage", "elementType": "f32" },
+    "y_scale_2": { "name": "y_scale", "buffer": "read-only-storage", "elementType": "f32", "length": 1 },
+    "y_zero_point_2": { "name": "y_zero_point", "buffer": "read-only-storage", "elementType": "u32", "length": 1 }
   },
   "variants": [
     {
       "id": "single_invocation",
       "when": ["serialContract", "serialFallbackNeeded"],
       "passes": [
         {
           "id": "main",
           "name": "DynamicQuantizeLinear",
+          "shader": "dynamic-quantize-linear.wgsl.jinja",
+          "derive": { "fromPartials": false },
+          "bindings": [{ "arg": "x", "elementType": "f32" }, "y", "y_scale", "y_zero_point", "params"],
           "dispatch": { "x": 1 }
         }
       ]
     },
     {
       "id": "parallel_subgroup_reduce_vec4",
       "priority": 11,
       "when": ["baseContract", "inputCount > 0", "inputCount % 4 == 0", "parallelFullFits"],
+      "derive": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": true,
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
           "shader": "dynamic-quantize-linear-reduce.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": ["x_2", "partial_min", "partial_max", "params"],
+          "dispatch": { "x": "min(fullPartials, 65535)", "y": "ceilDiv(fullPartials, 65535)", "z": 1 }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
+          "shader": "dynamic-quantize-linear.wgsl.jinja",
+          "derive": { "fromPartials": true },
+          "bindings": [
+            "partial_min_2",
+            "partial_max_2",
+            "y_scale",
+            "y_zero_point",
+            { "name": "params", "struct": [{ "name": "numPartials", "type": "u32", "value": "fullPartials" }] }
+          ],
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
+          "bindings": ["x_2", "y_scale_2", "y_zero_point_2", "y", "params"],
           "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
     },
     {
       "id": "parallel_subgroup_reduce",
       "priority": 10,
       "when": ["baseContract", "inputCount > 0", "true", "parallelFullFits"],
+      "derive": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": false,
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
           "shader": "dynamic-quantize-linear-reduce.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "bindings": ["x_2", "partial_min", "partial_max", "params"],
+          "dispatch": { "x": "min(fullPartials, 65535)", "y": "ceilDiv(fullPartials, 65535)", "z": 1 }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
+          "shader": "dynamic-quantize-linear.wgsl.jinja",
+          "derive": { "fromPartials": true },
+          "bindings": [
+            "partial_min_2",
+            "partial_max_2",
+            "y_scale",
+            "y_zero_point",
+            { "name": "params", "struct": [{ "name": "numPartials", "type": "u32", "value": "fullPartials" }] }
+          ],
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
+          "bindings": ["x_2", "y_scale_2", "y_zero_point_2", "y", "params"],
           "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
     },
     {
       "id": "grid_stride_reduce_vec4",
       "priority": 12,
       "when": ["baseContract", "inputCount > 0", "inputCount % 4 == 0", "inputCount >= tunables.GRID_STRIDE_MIN_ELEMENTS", "parallelGridFits"],
+      "derive": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": true,
         {
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
+          "shader": "dynamic-quantize-linear-reduce.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "derive": { "gridStride": true },
+          "bindings": ["x_2", "partial_min", "partial_max", "params"],
           "dispatch": { "x": "gridPartials" }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
+          "shader": "dynamic-quantize-linear.wgsl.jinja",
+          "derive": { "fromPartials": true },
+          "bindings": [
+            "partial_min_2",
+            "partial_max_2",
+            "y_scale",
+            "y_zero_point",
+            { "name": "params", "struct": [{ "name": "numPartials", "type": "u32", "value": "gridPartials" }] }
+          ],
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
+          "bindings": ["x_2", "y_scale_2", "y_zero_point_2", "y", "params"],
           "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
     },
     {
       "id": "grid_stride_reduce",
       "priority": 12,
       "when": ["baseContract", "inputCount > 0", "inputCount % 4 != 0", "inputCount >= tunables.GRID_STRIDE_MIN_ELEMENTS", "parallelGridFits"],
+      "derive": {
         "workgroupSize": "tunables.WORKGROUP_SIZE",
         "elemsPerThread": "tunables.ELEMENTS_PER_THREAD",
         "vec4": false,
         {
           "id": "reduce",
           "name": "DynamicQuantizeLinear.ReduceMinMax",
+          "shader": "dynamic-quantize-linear-reduce.wgsl.jinja",
+          "subgroupCollectivesWidth": "portable",
+          "derive": { "gridStride": true },
+          "bindings": ["x_2", "partial_min", "partial_max", "params"],
           "dispatch": { "x": "gridPartials" }
         },
         {
           "id": "finalize",
           "name": "DynamicQuantizeLinear.Finalize",
+          "shader": "dynamic-quantize-linear.wgsl.jinja",
+          "derive": { "fromPartials": true },
+          "bindings": [
+            "partial_min_2",
+            "partial_max_2",
+            "y_scale",
+            "y_zero_point",
+            { "name": "params", "struct": [{ "name": "numPartials", "type": "u32", "value": "gridPartials" }] }
+          ],
           "dispatch": { "x": 1 }
         },
         {
           "id": "quantize",
           "name": "DynamicQuantizeLinear.Quantize",
           "shader": "dynamic-quantize-linear-quantize.wgsl.jinja",
+          "bindings": ["x_2", "y_scale_2", "y_zero_point_2", "y", "params"],
           "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(inputCount, tunables.ELEMENTS_PER_THREAD)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,29 @@
 {
   "name": "ai.onnx.DynamicQuantizeLinear",
-  "id": "_ai_onnx_dynamicquantizelinear_webgpu_679d73b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "ic1d0BkyDcNr1B1KJ9oJN/5mnB04zIaCIW6y43kaEgY=",
-      "dynamic-quantize-linear-quantize.wgsl.jinja": "8J/MWQOE+FLfjP/5JCQvU7v0ncgxQcM79GMW6PtSPr8=",
-      "dynamic-quantize-linear-reduce.wgsl.jinja": "hbPwzMCANALDD58/USzBkSqY18jQS952IABcmsOy69w=",
-      "dynamic-quantize-linear.wgsl.jinja": "ccNGv53bI60K7evqWhmtmZi8ocQUp/zx9X9ICATi03M=",
-      "manifest.json": "i7HEkBE8BwjpTprLea26Sqfyb5DM1CVD7V8uRXE0DLI=",
-      "test.json": "I3ltf1DnGsVq5fWasOGlLYKfEDY9rtTLSV7xrLiU++M="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.DynamicQuantizeLinear" }
 }

 {
   "name": "ai.onnx.DynamicQuantizeLinear",
+  "id": "_ai_onnx_dynamicquantizelinear_webgpu_c781458",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "XDg7UAukA6OblWGm+ypY2zEJvMbW7c2WSl5jcuO5zxU=",
+      "dynamic-quantize-linear-quantize.wgsl.jinja": "ISRZ+RP/JwuhPFFtM/mSTJ+qGU8ATQz9B8b1bnZuS18=",
+      "dynamic-quantize-linear-reduce.wgsl.jinja": "MbIpHB8DKZZo8wVxryLnQkF9UeI5+GDuPb67Wq5CwhA=",
+      "dynamic-quantize-linear.wgsl.jinja": "tGwmn/Rl6x4gkOe+YNfxsBhJ6TrZgeNFvOXCNfOFeq4=",
+      "manifest.json": "bndT+hJeO/n9kaxatO81kxGcbg9bdnNYv9uEEOY8SHw=",
+      "test.json": "moR9c5IlMwac/omRuy88wGQgWfM/KKQRuM9+oU68p1M="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "single_invocation": ["dynamic-quantize-linear.wgsl.jinja"],
+      "parallel_subgroup_reduce_vec4": ["dynamic-quantize-linear-quantize.wgsl.jinja", "dynamic-quantize-linear-reduce.wgsl.jinja", "dynamic-quantize-linear.wgsl.jinja"],
+      "parallel_subgroup_reduce": ["dynamic-quantize-linear-quantize.wgsl.jinja", "dynamic-quantize-linear-reduce.wgsl.jinja", "dynamic-quantize-linear.wgsl.jinja"],
+      "grid_stride_reduce_vec4": ["dynamic-quantize-linear-quantize.wgsl.jinja", "dynamic-quantize-linear-reduce.wgsl.jinja", "dynamic-quantize-linear.wgsl.jinja"],
+      "grid_stride_reduce": ["dynamic-quantize-linear-quantize.wgsl.jinja", "dynamic-quantize-linear-reduce.wgsl.jinja", "dynamic-quantize-linear.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.DynamicQuantizeLinear",
   "cases": [
     {
       "name": "mixed_sign_values",
@@ -278,7 +277,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.DynamicQuantizeLinear_Max_Adjusted",
-        "notes": "Rank-0 scalar companion for the all-negative max-adjusted path: max is adjusted to include zero, producing zero_point=255."
       },
       "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } } },
       "outputs": {
@@ -318,7 +317,7 @@
     {
       "name": "grid_stride_reduce_1m_mixed_sign",
       "provenance": {
-        "notes": "numel = 1048576 >= 262144 selects grid_stride_reduce_vec4: the reduce dispatches a capped min(256, ceil(numel/1024))=256 workgroups, each grid-striding over the whole tensor and writing ONE partial, so the @workgroup_size(1) finalize folds only 256 partials (vs 4096 for the per-block reduce). Validates the grid-stride min/max is bit-identical to the serial reference (y_scale/y_zero_point exact) and that the capped 1D dispatch covers all elements."
       },
       "inputs": {
         "x": {
@@ -336,7 +335,7 @@
     {
       "name": "dispatch_cliff_blocks_65536",
       "provenance": {
-        "notes": "numel = 67108864 = 65536*1024 > 65535*1024 forces the folded block dispatch (x=65535, y=2) in the reduce + quantize passes. Validates blk = wg.x + wg.y*nwg.x AND the new `blk >= numBlocks` reduce-pass early-return that protects the partial_min/max[blk] scratch store from an over-dispatched OOB write. linspace puts the global min/max at the first/last block (the last block lives in the wg.y=1 folded region), so a wrong fold corrupts the scale. Before the 2D fold the plan-time dispatch-limit guard throws."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [67108864], "data": { "kind": "linspace", "start": -1.0, "end": 1.0 } }
@@ -365,7 +364,7 @@
     {
       "name": "non_vec4_large_finalize_fold_mixed_sign_corrected_division",
       "provenance": {
-        "notes": "Regression for cross-backend f32 division: raw and fma-residual quotients disagreed with correctly-rounded ORT at half-integer boundaries on Metal and shifted the zero point on Intel D3D. The backend-independent integer significand fallback must keep this odd-sized grid-stride path bit-exact."
       },
       "inputs": {
         "x": {
@@ -416,7 +415,7 @@
     {
       "name": "symmetric_half_step_tie_vec4_boundary",
       "provenance": {
-        "notes": "The scale 2/255 puts the zero-point quotient just below 127.5 after correctly-rounded f32 division. Intel D3D exposed an over-correction that changed zero_point 127 to 128."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1.0, -0.5, 0.5, 1.0] } }
@@ -430,7 +429,7 @@
     {
       "name": "intel_d3d_compensated_zero_point_regression",
       "provenance": {
-        "notes": "Explicit golden companion to symmetric_half_step_tie_vec4_boundary. It catches backend excess-precision or residual-division compensation shifting the correctly-rounded zero point from 127 to 128."
       },
       "inputs": {
         "x": {

 {
   "cases": [
     {
       "name": "mixed_sign_values",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
         "test": "QuantizeLinearOpTest.DynamicQuantizeLinear_Max_Adjusted",
+        "notes": "For an all-negative rank-0 input, the adjusted maximum includes zero and produces zero_point=255."
       },
       "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } } },
       "outputs": {
     {
       "name": "grid_stride_reduce_1m_mixed_sign",
       "provenance": {
+        "notes": "A 1,048,576-element input selects a grid-stride reduction capped at 256 workgroups, each producing one extrema partial. The final fold must include all elements and reproduce the exact global minimum, maximum, scale, and zero point."
       },
       "inputs": {
         "x": {
     {
       "name": "dispatch_cliff_blocks_65536",
       "provenance": {
+        "notes": "A 67,108,864-element input produces 65,536 reduction blocks and a folded 65,535-by-2 dispatch. A linspace places the extrema in the first and final blocks, making incorrect two-dimensional block reconstruction or a missing over-dispatch guard observable in the scale."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [67108864], "data": { "kind": "linspace", "start": -1.0, "end": 1.0 } }
     {
       "name": "non_vec4_large_finalize_fold_mixed_sign_corrected_division",
       "provenance": {
+        "notes": "Half-integer quotient boundaries can change zero-point rounding when f32 division precision differs across backends. The integer-significand fallback must keep this odd-sized grid-stride result bit-exact."
       },
       "inputs": {
         "x": {
     {
       "name": "symmetric_half_step_tie_vec4_boundary",
       "provenance": {
+        "notes": "The scale 2/255 puts the zero-point quotient just below 127.5 after correctly-rounded f32 division. The zero point must remain 127; rounding correction must not raise it to 128."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1.0, -0.5, 0.5, 1.0] } }
     {
       "name": "intel_d3d_compensated_zero_point_regression",
       "provenance": {
+        "notes": "An explicit expected output at a symmetric half-step boundary requires the correctly rounded zero point 127 rather than 128."
       },
       "inputs": {
         "x": {