{% 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 %}
}