[xpu-alignment] fake_quantize_per_channel_affine overflows to the top of the quantization range on XPU for an extreme finite float32 input

## Independent local verification

XPU rerun 2026-09-07 (Max 1100, torch 2.15.0.dev20260902+xpu): per-channel input -3.4028235e38, scale 1e-7, zero-point -128, quant range [-128,127] returned 2.55e-05 (0x37d5e8d5); CPU returned 0.0. The per-tensor control matched CPU.

The upstream result is still tracked as upstream-dependent; this issue records the concrete XPU defect and does not assert that upstream has finalized cross-backend semantics.

_{alignment scan `2026-08-30`, run `33350143612`}
### [xpu-alignment] fake_quantize_per_channel_affine overflows to the top of the quantization range on XPU for an extreme finite float32 input

## Summary

Upstream source: pytorch/pytorch#195343 (open) - `[CUDA] fake_quantize_per_channel_affine gives incorrect result for extreme finite float32 input`.
Alignment scan window: `[2026-08-30T00:00:00Z, 2026-08-31T00:00:00Z)` over `pytorch/pytorch`.

The upstream report is a CPU/CUDA divergence. XPU reproduces the wrong CUDA value bit for bit, and the defective code is the XPU SYCL functor in this repository, so the fix is owned here.

For a per-channel scale of `1e-7` and `quant_min = zero_point = -128`, the input `-3.4028235e38` must saturate to `quant_min` and dequantize to `(-128 - (-128)) * 1e-7 == 0.0`. CPU returns `0.0`. XPU returns `2.5499999537714757e-05`, which is the full quantization span `(127 - (-128)) * 1e-7`: the value saturated to the wrong end of the range.

## Observed XPU behavior

```
CPU: [[0.0, 0.0], [0.0, 0.0]]
XPU: [[2.5499999537714757e-05, 0.0], [0.0, 0.0]]
CPU bits: ['0x00000000', '0x00000000', '0x00000000', '0x00000000']
XPU bits: ['0x37d5e8d5', '0x00000000', '0x00000000', '0x00000000']
Exact equal: False
upstream cuda first element: 2.5499999537714757e-05
upstream cuda first element bits: 0x37d5e8d5
per-tensor control exact equal: True
```

The XPU result is bit-identical (`0x37d5e8d5`) to the CUDA value in the upstream report. The per-tensor variant of the same operation agrees with CPU bit for bit in the same run, which localizes the defect to the per-channel functor rather than to XPU float semantics in general.

## Reproducer

```python
import torch

x = torch.tensor(
[[-3.4028235e38, -1.0e-6], [1.0e6, 3.4028235e38]], dtype=torch.float32
)
scale = torch.tensor([1.0e-7, 1.0e20], dtype=torch.float32)
zero_point = torch.tensor([-128, 127], dtype=torch.int32)

cpu = torch.fake_quantize_per_channel_affine(x, scale, zero_point, 0, -128, 127)
xpu = torch.fake_quantize_per_channel_affine(
x.xpu(), scale.xpu(), zero_point.xpu(), 0, -128, 127
).cpu()

print("CPU:", cpu.tolist())
print("XPU:", xpu.tolist())
print("equal:", torch.equal(cpu, xpu))

# Control: the per-tensor variant agrees with CPU.
print(
"per-tensor equal:",
torch.equal(
torch.fake_quantize_per_tensor_affine(x, 1.0e-7, -128, -128, 127),
torch.fake_quantize_per_tensor_affine(x.xpu(), 1.0e-7, -128, -128, 127).cpu(),
),
)
```

## Target path and root cause

`fake_quantize_per_channel_affine_cachemask` dispatches `CPU, CUDA, XPU` to one shared host function, which calls `fake_quant_per_channel_cachemask_stub`. The XPU stub is registered in `src/ATen/native/quantized/FakeQuantizeCore.cpp` and implemented by `fake_quant_per_channel_cachemask_kernel` in `src/ATen/native/quantized/sycl/FakeQuantizeCoreKernels.cpp`.

`FakeQuantPerChannelCachemaskHelperDFunctor::operator()` computes:

```cpp
const float inv_scale = 1.0f / scale;
const auto qval =
static_cast(sycl::rint(input_val * inv_scale)) + zero_point;
const auto bounded_qval = sycl::clamp(qval, quant_min_, quant_max_);
return (bounded_qval - zero_point) * scale;
```

`-3.4028235e38 * 1e7` overflows float32 to `-inf`, `rint(-inf)` is `-inf`, and the `int64_t` cast of `-inf` is undefined; in practice it yields `INT64_MIN`. Adding `zero_point = -128` then wraps to a large positive value, so `sycl::clamp` saturates to `quant_max = 127` instead of `quant_min = -128`, and the dequantized result becomes `(127 - (-128)) * 1e-7 == 2.55e-05`. That arithmetic reproduces the observed value exactly.

The multiply-then-cast happens before any clamp, so the saturation the operator promises never gets a chance to run. The contrasting per-tensor functor in the same file adds `zero_point` inside the float expression and clamps with `fminf`/`fmaxf` on floats, which is why the control case survives.

Review note, not an observed result: `FakeQuantPerChannelCachemaskHelperCFunctor`, which produces the gradient mask, computes `qval` with the same expression. A fix should cover both functors consistently so the mask and the value stay in agreement.

Suggested direction: clamp in float before the integer cast (saturate `rint(input_val * inv_scale) + zero_point` to `[quant_min, quant_max]` while it is still floating point), which keeps `-inf` and `+inf` inputs saturating at the correct end.

## Ownership

Upstream pytorch/pytorch#195343 is open and scoped to CUDA; no upstream PR claims the fix, and #171777 (the ROCm-motivated fix for the same undefined behavior in the per-tensor CUDA kernel) was closed without merging. The XPU functor lives in `intel/torch-xpu-ops`, so an upstream CUDA fix cannot cover XPU. No existing `intel/torch-xpu-ops` issue or PR covers this operator. Independent XPU work is required here.

## Environment

```
torch 2.15.0.dev20260830+xpu
python 3.13.13 (/opt/conda/bin/python3.13)
torch path /opt/conda/lib/python3.13/site-packages/torch/__init__.py
xpu device Intel(R) Arc(TM) Pro B60 Graphics
```

Filed by the automated XPU alignment pipeline after independent review of the scan evidence. The reproducer above was executed by the deterministic runner, not by the reviewing agent.