sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +61 -0
build/webgpu/bench.json +177 -0
build/webgpu/dynamic-quantize-linear-quantize.wgsl.jinja +183 -0
build/webgpu/dynamic-quantize-linear-reduce.wgsl.jinja +132 -0
build/webgpu/dynamic-quantize-linear.wgsl.jinja +178 -0
build/webgpu/manifest.json +404 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +489 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.DynamicQuantizeLinear
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 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]`.
+See the [ONNX `DynamicQuantizeLinear` spec](https://onnx.ai/onnx/operators/onnx__DynamicQuantizeLinear.html) for the reference semantics.
+## 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
+| Variable | Allowed dtypes |
+| --- | --- |
+| `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
+- [`dynamic-quantize-linear-quantize.wgsl.jinja`](build/webgpu/dynamic-quantize-linear-quantize.wgsl.jinja)
+- [`dynamic-quantize-linear-reduce.wgsl.jinja`](build/webgpu/dynamic-quantize-linear-reduce.wgsl.jinja)
+- [`dynamic-quantize-linear.wgsl.jinja`](build/webgpu/dynamic-quantize-linear.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.DynamicQuantizeLinear", { version: 1 });
+const { y, y_scale, y_zero_point } = await kernel({ x: { data: xData, shape: [1] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,177 @@

+{
+  "op": "ai.onnx.DynamicQuantizeLinear",
+  "cases": [
+    {
+      "name": "f32_64k",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [65536],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [65536] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      }
+    },
+    {
+      "name": "f32_4m",
+      "preset": "smoke",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2048, 2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2048, 2048] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "name": "useful_onnx_payload", "value": "5 * numel(shapes.x) + 5" },
+          {
+            "type": "bandwidth",
+            "name": "physical_3pass_storage",
+            "value": "12 * numel(shapes.x) + 16 * min(256, ceilDiv(numel(shapes.x), 1024)) + 16"
+          }
+        ]
+      }
+    },
+    {
+      "name": "alignment_cliff_non_vec4_4194303",
+      "preset": "smoke",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4194303],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4194303] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "name": "useful_onnx_payload", "value": "5 * numel(shapes.x) + 5" },
+          {
+            "type": "bandwidth",
+            "name": "physical_3pass_storage",
+            "value": "12 * numel(shapes.x) + 16 * min(256, ceilDiv(numel(shapes.x), 1024)) + 16"
+          }
+        ]
+      }
+    },
+    {
+      "name": "alignment_healthy_vec4_4194304",
+      "preset": "smoke",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4194304],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4194304] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "name": "useful_onnx_payload", "value": "5 * numel(shapes.x) + 5" },
+          {
+            "type": "bandwidth",
+            "name": "physical_3pass_storage",
+            "value": "12 * numel(shapes.x) + 16 * min(256, ceilDiv(numel(shapes.x), 1024)) + 16"
+          }
+        ]
+      }
+    },
+    {
+      "name": "serial_finalize_uncapped_non_vec4_2m",
+      "preset": "smoke",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2097153],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2097153] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "name": "useful_onnx_payload", "value": "5 * numel(shapes.x) + 5" },
+          {
+            "type": "bandwidth",
+            "name": "physical_3pass_storage",
+            "value": "12 * numel(shapes.x) + 16 * min(256, ceilDiv(numel(shapes.x), 1024)) + 16"
+          }
+        ]
+      }
+    },
+    {
+      "name": "grid_stride_capped_finalize_2m",
+      "preset": "smoke",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2097152],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2097152] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "name": "useful_onnx_payload", "value": "5 * numel(shapes.x) + 5" },
+          {
+            "type": "bandwidth",
+            "name": "physical_3pass_storage",
+            "value": "12 * numel(shapes.x) + 16 * min(256, ceilDiv(numel(shapes.x), 1024)) + 16"
+          }
+        ]
+      }
+    },
+    {
+      "name": "below_gridcap_threshold_non_vec4_262143",
+      "preset": "smoke",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [262143],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [262143] },
+        "y_scale": { "dtype": "float32", "shape": [1] },
+        "y_zero_point": { "dtype": "uint8", "shape": [1] }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "name": "useful_onnx_payload", "value": "5 * numel(shapes.x) + 5" },
+          {
+            "type": "bandwidth",
+            "name": "physical_3pass_storage",
+            "value": "12 * numel(shapes.x) + 16 * min(256, ceilDiv(numel(shapes.x), 1024)) + 16"
+          }
+        ]
+      }
+    }
+  ]
+}

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

	@@ -0,0 +1,183 @@

+// Pass 3 of parallel DynamicQuantizeLinear: elementwise quantization using the
+// read-only y_scale/y_zero_point produced by the finalize pass. Each workgroup
+// covers the same WG * EPT contiguous chunk as the reduction pass, so their
+// dispatch counts match. Quantization applies round-to-even and saturates the
+// result to the uint8 range using the common scale and zero point.
+{{ env.wgsl.resourceDeclarations }}
+// ONNX DynamicQuantizeLinear uses correctly rounded f32 division followed by
+// round-half-to-even. Every execution path uses these helpers so their numerical
+// behavior cannot drift apart.
+// WGSL permits f32 division to differ from the correctly-rounded result by
+// 2.5 ULP, and fma() inherits separate multiply/add accuracy rather than
+// promising a fused residual. Reconstruct the correctly-rounded normal result
+// with integer significand division when the quotient can affect an integer
+// rounding boundary. This is backend-independent and uses only exact u32 ops.
+fn dynamic_quantize_exact_div_normal(numerator: f32, denominator: f32) -> f32 {
+  if (numerator == 0.0) {
+    return numerator;
+  }
+  let numerator_bits = bitcast<u32>(numerator);
+  let denominator_bits = bitcast<u32>(denominator);
+  let sign_bits = (numerator_bits ^ denominator_bits) & 0x80000000u;
+  let numerator_abs = numerator_bits & 0x7fffffffu;
+  let denominator_abs = denominator_bits & 0x7fffffffu;
+  if (denominator_abs == 0u
+      || (numerator_abs & 0x7f800000u) == 0x7f800000u
+      || (denominator_abs & 0x7f800000u) == 0x7f800000u) {
+    return numerator / denominator;
+  }
+  var numerator_mantissa = numerator_abs & 0x007fffffu;
+  var denominator_mantissa = denominator_abs & 0x007fffffu;
+  let numerator_biased_exponent = (numerator_abs >> 23u) & 0xffu;
+  let denominator_biased_exponent = (denominator_abs >> 23u) & 0xffu;
+  var numerator_exponent: i32;
+  var denominator_exponent: i32;
+  if (numerator_biased_exponent == 0u) {
+    numerator_exponent = -126;
+    // Zero returned above. A non-zero subnormal reaches the implicit-bit
+    // position in at most 23 exact shifts.
+    while ((numerator_mantissa & 0x00800000u) == 0u) {
+      numerator_mantissa = numerator_mantissa << 1u;
+      numerator_exponent = numerator_exponent - 1;
+    }
+  } else {
+    numerator_mantissa = numerator_mantissa | 0x00800000u;
+    numerator_exponent = i32(numerator_biased_exponent) - 127;
+  }
+  if (denominator_biased_exponent == 0u) {
+    denominator_exponent = -126;
+    while ((denominator_mantissa & 0x00800000u) == 0u) {
+      denominator_mantissa = denominator_mantissa << 1u;
+      denominator_exponent = denominator_exponent - 1;
+    }
+  } else {
+    denominator_mantissa = denominator_mantissa | 0x00800000u;
+    denominator_exponent = i32(denominator_biased_exponent) - 127;
+  }
+  var quotient_exponent = numerator_exponent - denominator_exponent;
+  var remainder = numerator_mantissa;
+  if (remainder < denominator_mantissa) {
+    remainder = remainder << 1u;
+    quotient_exponent = quotient_exponent - 1;
+  }
+  // The normalized ratio is now in [1, 2). Emit its implicit bit followed by
+  // all 23 stored significand bits using exact binary long division.
+  var quotient_mantissa = 0x00800000u;
+  remainder = remainder - denominator_mantissa;
+  for (var digit = 0u; digit < 23u; digit = digit + 1u) {
+    remainder = remainder << 1u;
+    if (remainder >= denominator_mantissa) {
+      remainder = remainder - denominator_mantissa;
+      quotient_mantissa = quotient_mantissa | (1u << (22u - digit));
+    }
+  }
+  // Round the 24-bit significand to nearest, ties to even. remainder and its
+  // doubled value are below 2^25, so no u32 overflow is possible.
+  let twice_remainder = remainder << 1u;
+  if (twice_remainder > denominator_mantissa
+      || (twice_remainder == denominator_mantissa && (quotient_mantissa & 1u) != 0u)) {
+    quotient_mantissa = quotient_mantissa + 1u;
+  }
+  if (quotient_mantissa == 0x01000000u) {
+    quotient_mantissa = quotient_mantissa >> 1u;
+    quotient_exponent = quotient_exponent + 1;
+  }
+  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
+    | (u32(biased_exponent) << 23u)
+    | (quotient_mantissa & 0x007fffffu);
+  return bitcast<f32>(result_bits);
+}
+fn dynamic_quantize_division_may_cross_half(estimate: f32) -> bool {
+  let lower = floor(estimate);
+  let fraction = estimate - lower;
+  let magnitude = abs(estimate);
+  let magnitude_bits = bitcast<u32>(magnitude);
+  let adjacent = bitcast<f32>(magnitude_bits + 1u);
+  let ulp = adjacent - magnitude;
+  // Division is allowed 2.5 ULP error. Eight ULP also covers the factor-of-two
+  // ULP change when an estimate straddles the 0.5 exponent boundary.
+  return abs(fraction - 0.5) <= ulp * 8.0;
+}
+fn round_dynamic_half_to_even(value: f32, scale: f32) -> i32 {
+  let estimate = value / scale;
+  var scaled = estimate;
+  // `select` evaluates both value operands in WGSL; use control flow so the
+  // 23-bit software divide remains a rare boundary fallback, not O(23) work
+  // for every quantized element.
+  if (dynamic_quantize_division_may_cross_half(estimate)) {
+    scaled = dynamic_quantize_exact_div_normal(value, scale);
+  }
+  let lower = floor(scaled);
+  let fraction = scaled - lower;
+  if (fraction < 0.5) {
+    return i32(lower);
+  }
+  if (fraction > 0.5) {
+    return i32(lower + 1.0);
+  }
+  let upper = lower + 1.0;
+  let half_lower = floor(lower * 0.5);
+  let lower_is_even = (lower - half_lower * 2.0) == 0.0;
+  return i32(select(upper, lower, lower_is_even));
+}
+const WG: u32 = {{ workgroupSize }}u;
+{% if not vec4 %}
+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 count4 = params.count / 4u;
+  let idx4 = blk * WG + tid;
+  if (idx4 < count4) {
+    let v = x[idx4];
+    let o = idx4 * 4u;
+    y[o + 0u] = u32(clamp(round_dynamic_half_to_even(v.x, scale) + zp_i32, 0, 255));
+    y[o + 1u] = u32(clamp(round_dynamic_half_to_even(v.y, scale) + zp_i32, 0, 255));
+    y[o + 2u] = u32(clamp(round_dynamic_half_to_even(v.z, scale) + zp_i32, 0, 255));
+    y[o + 3u] = u32(clamp(round_dynamic_half_to_even(v.w, scale) + zp_i32, 0, 255));
+  }
+{% else %}
+  let base = blk * WG * EPT;
+  for (var e = 0u; e < EPT; e = e + 1u) {
+    let idx = base + e * WG + tid;
+    if (idx < params.count) {
+      let q = clamp(round_dynamic_half_to_even(x[idx], scale) + zp_i32, 0, 255);
+      y[idx] = u32(q);
+    }
+  }
+{% endif %}
+}

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

	@@ -0,0 +1,132 @@

+// 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 %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+{% if not gridStride and not vec4 %}
+const EPT: u32 = {{ elemsPerThread }}u;
+{% endif %}
+var<workgroup> wgMin: array<f32, WG>;
+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;
+{% if gridStride %}
+  let blk = wg.x;
+  // The capped x-only dispatch is the full logical grid. Each workgroup writes
+  // exactly one partial, matching finalize's numPartials uniform.
+  let gridThreads = nwg.x * WG;
+{% else %}
+  // 2D-folded block index: wg.y carries the high bits past the device's
+  // 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 %}
+  let numBlocks = (params.count + WG * EPT - 1u) / (WG * EPT);
+{% endif %}
+  if (blk >= numBlocks) { return; }
+{% endif %}
+  var localMin = 0.0;
+  var localMax = 0.0;
+{% if gridStride %}
+{% if vec4 %}
+  // Adjacent lanes read adjacent vec4s on each grid-stride iteration.
+  let count4 = params.count / 4u;
+  for (var idx4 = blk * WG + tid; idx4 < count4; idx4 = idx4 + gridThreads) {
+    let v = x[idx4];
+    localMin = min(localMin, min(min(v.x, v.y), min(v.z, v.w)));
+    localMax = max(localMax, max(max(v.x, v.y), max(v.z, v.w)));
+  }
+{% else %}
+  for (var idx = blk * WG + tid; idx < params.count; idx = idx + gridThreads) {
+    let value = x[idx];
+    localMin = min(localMin, value);
+    localMax = max(localMax, value);
+  }
+{% endif %}
+{% else %}
+{% if vec4 %}
+  // Each thread reads one contiguous vec4; this path assumes params.count is
+  // divisible by 4.
+  let count4 = params.count / 4u;
+  let idx4 = blk * WG + tid;
+  if (idx4 < count4) {
+    let v = x[idx4];
+    localMin = min(localMin, min(min(v.x, v.y), min(v.z, v.w)));
+    localMax = max(localMax, max(max(v.x, v.y), max(v.z, v.w)));
+  }
+{% else %}
+  let base = blk * WG * EPT;
+  for (var e = 0u; e < EPT; e = e + 1u) {
+    let idx = base + e * WG + tid;
+    if (idx < params.count) {
+      let v = x[idx];
+      localMin = min(localMin, v);
+      localMax = max(localMax, v);
+    }
+  }
+{% endif %}
+{% endif %}
+{% 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;
+  }
+{% else %}
+  // No-subgroups fallback: full workgroup-tree min/max over per-thread partials.
+  wgMin[tid] = localMin;
+  wgMax[tid] = localMax;
+  workgroupBarrier();
+  for (var step = WG >> 1u; step > 0u; step = step >> 1u) {
+    if (tid < step) {
+      wgMin[tid] = min(wgMin[tid], wgMin[tid + step]);
+      wgMax[tid] = max(wgMax[tid], wgMax[tid + step]);
+    }
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+    partial_min[blk] = wgMin[0];
+    partial_max[blk] = wgMax[0];
+  }
+{% endif %}
+}

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

	@@ -0,0 +1,178 @@

+{% 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 %}
+// The serial path computes the complete range, scale, zero point, and quantized
+// output in one invocation.
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// ONNX DynamicQuantizeLinear uses correctly rounded f32 division followed by
+// round-half-to-even. Every execution path uses these helpers so their numerical
+// behavior cannot drift apart.
+// WGSL permits f32 division to differ from the correctly-rounded result by
+// 2.5 ULP, and fma() inherits separate multiply/add accuracy rather than
+// promising a fused residual. Reconstruct the correctly-rounded normal result
+// with integer significand division when the quotient can affect an integer
+// rounding boundary. This is backend-independent and uses only exact u32 ops.
+fn dynamic_quantize_exact_div_normal(numerator: f32, denominator: f32) -> f32 {
+  if (numerator == 0.0) {
+    return numerator;
+  }
+  let numerator_bits = bitcast<u32>(numerator);
+  let denominator_bits = bitcast<u32>(denominator);
+  let sign_bits = (numerator_bits ^ denominator_bits) & 0x80000000u;
+  let numerator_abs = numerator_bits & 0x7fffffffu;
+  let denominator_abs = denominator_bits & 0x7fffffffu;
+  if (denominator_abs == 0u
+      || (numerator_abs & 0x7f800000u) == 0x7f800000u
+      || (denominator_abs & 0x7f800000u) == 0x7f800000u) {
+    return numerator / denominator;
+  }
+  var numerator_mantissa = numerator_abs & 0x007fffffu;
+  var denominator_mantissa = denominator_abs & 0x007fffffu;
+  let numerator_biased_exponent = (numerator_abs >> 23u) & 0xffu;
+  let denominator_biased_exponent = (denominator_abs >> 23u) & 0xffu;
+  var numerator_exponent: i32;
+  var denominator_exponent: i32;
+  if (numerator_biased_exponent == 0u) {
+    numerator_exponent = -126;
+    // Zero returned above. A non-zero subnormal reaches the implicit-bit
+    // position in at most 23 exact shifts.
+    while ((numerator_mantissa & 0x00800000u) == 0u) {
+      numerator_mantissa = numerator_mantissa << 1u;
+      numerator_exponent = numerator_exponent - 1;
+    }
+  } else {
+    numerator_mantissa = numerator_mantissa | 0x00800000u;
+    numerator_exponent = i32(numerator_biased_exponent) - 127;
+  }
+  if (denominator_biased_exponent == 0u) {
+    denominator_exponent = -126;
+    while ((denominator_mantissa & 0x00800000u) == 0u) {
+      denominator_mantissa = denominator_mantissa << 1u;
+      denominator_exponent = denominator_exponent - 1;
+    }
+  } else {
+    denominator_mantissa = denominator_mantissa | 0x00800000u;
+    denominator_exponent = i32(denominator_biased_exponent) - 127;
+  }
+  var quotient_exponent = numerator_exponent - denominator_exponent;
+  var remainder = numerator_mantissa;
+  if (remainder < denominator_mantissa) {
+    remainder = remainder << 1u;
+    quotient_exponent = quotient_exponent - 1;
+  }
+  // The normalized ratio is now in [1, 2). Emit its implicit bit followed by
+  // all 23 stored significand bits using exact binary long division.
+  var quotient_mantissa = 0x00800000u;
+  remainder = remainder - denominator_mantissa;
+  for (var digit = 0u; digit < 23u; digit = digit + 1u) {
+    remainder = remainder << 1u;
+    if (remainder >= denominator_mantissa) {
+      remainder = remainder - denominator_mantissa;
+      quotient_mantissa = quotient_mantissa | (1u << (22u - digit));
+    }
+  }
+  // Round the 24-bit significand to nearest, ties to even. remainder and its
+  // doubled value are below 2^25, so no u32 overflow is possible.
+  let twice_remainder = remainder << 1u;
+  if (twice_remainder > denominator_mantissa
+      || (twice_remainder == denominator_mantissa && (quotient_mantissa & 1u) != 0u)) {
+    quotient_mantissa = quotient_mantissa + 1u;
+  }
+  if (quotient_mantissa == 0x01000000u) {
+    quotient_mantissa = quotient_mantissa >> 1u;
+    quotient_exponent = quotient_exponent + 1;
+  }
+  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
+    | (u32(biased_exponent) << 23u)
+    | (quotient_mantissa & 0x007fffffu);
+  return bitcast<f32>(result_bits);
+}
+fn dynamic_quantize_division_may_cross_half(estimate: f32) -> bool {
+  let lower = floor(estimate);
+  let fraction = estimate - lower;
+  let magnitude = abs(estimate);
+  let magnitude_bits = bitcast<u32>(magnitude);
+  let adjacent = bitcast<f32>(magnitude_bits + 1u);
+  let ulp = adjacent - magnitude;
+  // Division is allowed 2.5 ULP error. Eight ULP also covers the factor-of-two
+  // ULP change when an estimate straddles the 0.5 exponent boundary.
+  return abs(fraction - 0.5) <= ulp * 8.0;
+}
+fn round_dynamic_half_to_even(value: f32, scale: f32) -> i32 {
+  let estimate = value / scale;
+  var scaled = estimate;
+  // `select` evaluates both value operands in WGSL; use control flow so the
+  // 23-bit software divide remains a rare boundary fallback, not O(23) work
+  // for every quantized element.
+  if (dynamic_quantize_division_may_cross_half(estimate)) {
+    scaled = dynamic_quantize_exact_div_normal(value, scale);
+  }
+  let lower = floor(scaled);
+  let fraction = scaled - lower;
+  if (fraction < 0.5) {
+    return i32(lower);
+  }
+  if (fraction > 0.5) {
+    return i32(lower + 1.0);
+  }
+  let upper = lower + 1.0;
+  let half_lower = floor(lower * 0.5);
+  let lower_is_even = (lower - half_lower * 2.0) == 0.0;
+  return i32(select(upper, lower, lower_is_even));
+}
+@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) {
+    min_value = min(min_value, partial_min[i]);
+    max_value = max(max_value, partial_max[i]);
+  }
+{% else %}
+  var min_value = 0.0;
+  var max_value = 0.0;
+  for (var i = 0u; i < params.count; i = i + 1u) {
+    min_value = min(min_value, x[i]);
+    max_value = max(max_value, x[i]);
+  }
+{% endif %}
+  var scale = dynamic_quantize_exact_div_normal(max_value - min_value, 255.0);
+  if (scale == 0.0) {
+    scale = 1.0;
+  }
+  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);
+  }
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,404 @@

+{
+  "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)",
+    "storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
+    "inputCount": "numel(shapes.x)",
+    "elementsPerBlock": "tunables.WORKGROUP_SIZE * tunables.ELEMENTS_PER_THREAD",
+    "fullPartials": "ceilDiv(inputCount, elementsPerBlock)",
+    "gridPartials": "min(tunables.MAX_GRID_PARTIALS, fullPartials)",
+    "baseContract": "ranks.x >= 1 and ranks.y == ranks.x and inputCount == numel(shapes.y) and ranks.y_scale == 0 and ranks.y_zero_point == 0",
+    "serialContract": "ranks.x >= 0 and ranks.y == ranks.x and inputCount == numel(shapes.y) and ranks.y_scale == 0 and ranks.y_zero_point == 0",
+    "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,
+        "inputElement": "\"vec4<f32>\"",
+        "useSubgroups": "device.features.has(\"subgroups\")"
+      },
+      "intermediates": [
+        { "id": "partial_min", "dtype": "float32", "shape": "[fullPartials]" },
+        { "id": "partial_max", "dtype": "float32", "shape": "[fullPartials]" }
+      ],
+      "passes": [
+        {
+          "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,
+        "inputElement": "\"f32\"",
+        "useSubgroups": "device.features.has(\"subgroups\")"
+      },
+      "intermediates": [
+        { "id": "partial_min", "dtype": "float32", "shape": "[fullPartials]" },
+        { "id": "partial_max", "dtype": "float32", "shape": "[fullPartials]" }
+      ],
+      "passes": [
+        {
+          "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,
+        "inputElement": "\"vec4<f32>\"",
+        "useSubgroups": "device.features.has(\"subgroups\")"
+      },
+      "intermediates": [
+        { "id": "partial_min", "dtype": "float32", "shape": "[gridPartials]" },
+        { "id": "partial_max", "dtype": "float32", "shape": "[gridPartials]" }
+      ],
+      "passes": [
+        {
+          "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,
+        "inputElement": "\"f32\"",
+        "useSubgroups": "device.features.has(\"subgroups\")"
+      },
+      "intermediates": [
+        { "id": "partial_min", "dtype": "float32", "shape": "[gridPartials]" },
+        { "id": "partial_max", "dtype": "float32", "shape": "[gridPartials]" }
+      ],
+      "passes": [
+        {
+          "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"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

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

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,489 @@

+{
+  "op": "ai.onnx.DynamicQuantizeLinear",
+  "cases": [
+    {
+      "name": "mixed_sign_values",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-1.0, -0.5, 0.0, 0.5, 1.0, 2.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [6] },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "ort_mixed_sign_rounding_formulation",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.DynamicQuantizeLinear"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, -3.0, -2.5, 1.34, 0.5] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [6] },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "all_positive_includes_zero_range",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4] },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "constant_zero_scale_fallback",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4] },
+        "y_scale": { "dtype": "float32", "shape": [] },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "subnormal_positive_range_preserves_scale_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: DynamicQuantizeLinear computes scale=(max-min)/255 then x/scale; a subnormal input range yields a subnormal scale and Metal flushes denormals in floating-point division (-> NaN/0), so subnormal-range cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.DynamicQuantizeLinear",
+        "notes": "Tiny but valid positive range: y_scale is a float32 subnormal, so a backend must not flush it to zero and take the constant-input fallback."
+      },
+      "inputs": { "x": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1e-40] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "subnormal_negative_range_preserves_scale_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: DynamicQuantizeLinear computes scale=(max-min)/255 then x/scale; a subnormal input range yields a subnormal scale and Metal flushes denormals in floating-point division (-> NaN/0), so subnormal-range cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.DynamicQuantizeLinear_Max_Adjusted",
+        "notes": "Tiny but valid negative range: ONNX adjusts the max to zero, producing a float32 subnormal scale and a high zero point rather than the constant-input fallback."
+      },
+      "inputs": { "x": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-1e-40, 0.0] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "subnormal_mixed_sign_range_vec4_reduce_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal scale values required by this fixture. Backend evidence: DynamicQuantizeLinear computes scale=(max-min)/255 then x/scale; a subnormal input range yields a subnormal scale and Metal flushes denormals in floating-point division (-> NaN/0), so subnormal-range cases remain CPU-reference-only."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.DynamicQuantizeLinear",
+        "notes": "Tiny mixed-sign range with four elements exercises the vec4 reduce path: ONNX should keep a subnormal scale and nonzero zero point."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-1e-40, -5e-41, 5e-41, 1e-40] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "min_adjusted_positive_2d_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.DynamicQuantizeLinear_Min_Adjusted"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [1.0, 2.1, 1.3, 2.5, 3.34, 4.0, 1.5, 2.6, 3.9, 4.0, 3.0, 2.345] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [3, 4] },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "max_adjusted_negative_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/dynamic_quantize_linear_test.cc",
+        "test": "QuantizeLinearOpTest.DynamicQuantizeLinear_Max_Adjusted"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-1.0, -2.1, -1.3, -2.5, -3.34, -4.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [6] },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "singleton_positive_adjusts_min_to_zero",
+      "inputs": { "x": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [5.0] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1] },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "onnx_backend_dynamicquantizelinear",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dynamicquantizelinear" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.0, 2.0, -3.0, -2.5, 1.340000033378601, 0.5] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [6], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.00001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "onnx_backend_dynamicquantizelinear_max_adjusted",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dynamicquantizelinear_max_adjusted"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": {
+            "kind": "values",
+            "values": [-1.0, -2.0999999046325684, -1.2999999523162842, -2.5, -3.3399999141693115, -4.0]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [6], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.00001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "onnx_backend_dynamicquantizelinear_min_adjusted",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dynamicquantizelinear_min_adjusted"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0999999046325684, 1.2999999523162842, 2.5, 3.3399999141693115, 4.0, 1.5, 2.5999999046325684, 3.9000000953674316, 4.0, 3.0, 2.3450000286102295]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [3, 4], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.00001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "parallel_mixed_sign_5000",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [5000],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [5000], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "parallel_all_negative_1500_tail",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1500],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.0, -2.1, -1.3, -2.5, -3.34, -4.0, -0.25, -7.75, -0.001, -6.5, -2.75, -5.125, -3.5]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1500], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "parallel_const_zero_1027_tail",
+      "inputs": { "x": { "dtype": "float32", "shape": [1027], "data": { "kind": "cycle", "values": [0.0] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1027], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "scalar_rank0_positive",
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "scalar_rank0_negative",
+      "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": {
+        "y": { "dtype": "uint8", "shape": [], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "empty_zero_dim",
+      "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [0], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.0001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "ort_caseB_empty",
+      "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [0], "data": { "kind": "values", "values": [] }, "tolerance": 0.001 },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [1.0] },
+          "tolerance": 0.001
+        },
+        "y_zero_point": {
+          "dtype": "uint8",
+          "shape": [],
+          "data": { "kind": "values", "values": [0] },
+          "tolerance": 0.001
+        }
+      }
+    },
+    {
+      "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": {
+          "dtype": "float32",
+          "shape": [1024, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [1024, 1024], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "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 } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [67108864], "tolerance": 1 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [] }
+      }
+    },
+    {
+      "name": "grid_stride_reduce_clean_int_262145",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [262145],
+          "data": { "kind": "cycle", "values": [0.0, 32.0, 64.0, 96.0, 128.0, 160.0, 192.0, 224.0, 255.0] }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [262145], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "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": {
+          "dtype": "float32",
+          "shape": [4194305],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4194305], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "per_token_activation_2d_vec4",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.019, "scale": 2.5 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [3, 4096], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "non_vec4_odd_hidden_bridge",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [2049],
+          "data": {
+            "kind": "cycle",
+            "values": [-1.0, -2.1, -1.3, 2.5, 3.34, 4.0, -0.25, 7.75, -0.001, 6.5, -2.75, 5.125, -3.5]
+          }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [2049], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "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] } }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [4], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "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": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-2.0, -1.0, -0.5, 0.0, 0.5, 1.0, 1.5, 2.0] }
+        }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, 63, 95, 127, 159, 191, 223, 254] },
+          "tolerance": 0
+        },
+        "y_scale": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [0.01568627543747425] },
+          "tolerance": 0
+        },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [127] }, "tolerance": 0 }
+      }
+    },
+    {
+      "name": "grid_stride_threshold_boundary_vec4_262144",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [262144],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.017, "scale": 3.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [262144], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "multidim_non_vec4_parallel_reduce_2d",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [17, 120],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.023, "scale": 2.0 }
+        }
+      },
+      "outputs": {
+        "y": { "dtype": "uint8", "shape": [17, 120], "tolerance": 0 },
+        "y_scale": { "dtype": "float32", "shape": [], "tolerance": 0.000001 },
+        "y_zero_point": { "dtype": "uint8", "shape": [], "tolerance": 0 }
+      }
+    }
+  ]
+}