ALIGNMENT_TRIAGE_ISSUE

XPU alignment triage

The daily XPU alignment workflow posts audit summaries and independently reviewed candidates here.

Every run for which the deterministic gate produces a decision posts one Run Summary, separate from any candidate draft:

XPU alignment run complete: collection and review completed without terminal warnings.
XPU alignment run completed with warnings: one or more units were isolated, while the remaining reviewed results are still valid and the workflow stays green.
XPU alignment run completed with partial collection: fetching stopped before every source completed; source progress and the terminal error are reported, trusted candidates may still be published or queued, and the workflow finishes red.
XPU alignment run failed: a malformed artifact, coverage gap, required producer failure, or other global blocker prevents all candidate publication and the workflow finishes red.

A Run Summary is a snapshot of the completed workflow attempt; each candidate draft comment carries the current filing state.

Only terminal conditions that affect the final outcome are shown in the Run Summary. Recovered retries, transient API failures, and internal tool errors remain in workflow logs and artifacts. A scheduled retry updates the existing Summary for the same run ID; duplicate scheduled Summary markers fail closed. Dry runs may post repeated summaries. If the publisher itself or the GitHub Issues API fails, the Actions result and gate artifact are the fallback because a comment cannot be guaranteed.

Formal scheduled drafts may be approved one at a time with:

@torchxpubot file <unit-id>

Run the command on this issue only.
Only repository OWNER or MEMBER accounts may approve filing.
A formal reviewed draft from either a complete scheduled scan or a structurally valid partial scheduled scan can be filed.
A scheduled run with exactly one reviewed payload may file it automatically, including when collection is structurally valid but partial.
Partial runs report source progress, last cursor, rate-limit reset, and terminal errors, notify maintainers, and finish red even when trusted candidates were published or queued.
[DRY RUN] drafts are test output, may repeat, never notify, and cannot be filed.
Malformed artifacts, incomplete stage coverage, core environment failure, or a failed producer job produce only a failed Run Summary and no formal draft.
A runner-backed blocker for one unit excludes that unit without invalidating other fully reviewed candidates.
Historical diagnostic comments remain historical evidence and cannot be filed because they do not carry a formal unit marker.
The bot copies the reviewed formal draft into a new [xpu-alignment] issue with exactly the ai_generated label.
A formal unit can be filed successfully only once; the draft is marked with the created issue number.
No agent runs between approval and issue creation.
Before approving, check the scan date, workflow run, reproducer, XPU result, independent reviewer verdict, and any incomplete-collection summary.

_{alignment scan 2026-08-24, run 32923438489}

[DRY RUN] [xpu-alignment] embedding_bag per_sample_weights backward indexes weight/grad with unchecked, int-truncated saved indices on XPU

Summary

aten::_embedding_bag_per_sample_weights_backward on XPU indexes grad and weight with the saved indices / offset2bag values without any bounds check, and truncates both to int. Upstream fixed the equivalent CUDA kernel in the 2026-08-24 window; that fix cannot reach this repository, and its regression tests are CUDA-only.

Upstream source

Commit: pytorch/pytorch@550fa84 -- "Fix CUDA illegal memory access in embedding_bag per_sample_weights backward (#193408)"
PR: Fix CUDA illegal memory access in embedding_bag per_sample_weights backward pytorch/pytorch#193408 (merged)
Issue: [CUDA] illegal memory read on _embedding_bag_per_sample_weights_backward_kernel pytorch/pytorch#192445 (closed)
Alignment scan window: [2026-08-24T00:00:00Z, 2026-08-25T00:00:00Z), repository pytorch/pytorch

The upstream change to aten/src/ATen/native/cuda/EmbeddingBag.cu does two things. It stops truncating offset2bag[i] and indices[i] with (int) and keeps both at index_t, and it adds CUDA_KERNEL_ASSERT guards for 0 <= bag_idx < num_bags and 0 <= embedding_idx < num_embeddings before those values are used to address grad and weight. Both regression tests it adds to test/nn/test_embedding.py -- test_embedding_bag_per_sample_weights_rejects_invalid_saved_indices and test_embedding_bag_per_sample_weights_large_index -- are decorated @onlyOn(["cuda"]), so nothing exercises XPU.

Target-path evidence

This op is implemented in this repository, so no upstream change reaches it. Line numbers are from intel/torch-xpu-ops main:

src/ATen/native/xpu/sycl/EmbeddingBag.cpp:912 -- _embedding_bag_per_sample_weights_backward_kernel
src/ATen/native/xpu/sycl/EmbeddingBag.cpp:722 -- _embedding_bag_per_sample_weights_backward_impl (called at :952)
src/ATen/native/xpu/sycl/EmbeddingBag.h:219 -- EmbeddingBagPerSampleWeightsBackwardKernelFunctor

The functor body is the exact pre-fix CUDA code:

// src/ATen/native/xpu/sycl/EmbeddingBag.h:233
const int bag_idx = (int)offset2bag_[sample_idx];
const int embedding_idx = (int)indices_[sample_idx];
if (embedding_idx != padding_idx_) {
  for (int feature_idx = sglid; feature_idx < embedding_features_;
       feature_idx += sgSize) {
    result +=
        grad_[grad_stride0_ * bag_idx + grad_stride1_ * feature_idx] *
        weight_[weight_stride0_ * embedding_idx + weight_stride1_ * feature_idx];
  }
}

Neither num_bags nor num_embeddings is passed into the functor, so no bound is even available to check against. The XPU forward kernel does guard its index -- SYCL_KERNEL_ASSERT(weight_index < num_row_) at src/ATen/native/xpu/sycl/EmbeddingBag.h:104, upper bound only -- which makes the missing backward guard an internal inconsistency as well as a CUDA divergence.

The (int) casts are a second, independent defect: an int64 index at or above 2^31 wraps and addresses the wrong row even when the index is perfectly valid. That is what upstream's test_embedding_bag_per_sample_weights_large_index covers, also CUDA-only.

Observed XPU behavior

Backward with an out-of-range saved indices value completes silently. weight has 5 rows, every saved index is set to 5 (== num_embeddings), and the per-sample-weights gradient is computed from a weight read one row past the allocation. Upstream's post-fix contract is that such an index is rejected by a device-side assertion.

Reproducer

# repro_embedding_bag_psw_backward.py
import subprocess
import sys

import torch

assert torch.xpu.is_available()
assert torch._C._dispatch_has_kernel_for_dispatch_key(
    "aten::_embedding_bag_per_sample_weights_backward", "XPU"
)

CHILD = '''
import torch

weight = torch.ones((5, 2), device="xpu")
indices = torch.zeros(1024, dtype=torch.int64, device="xpu")
offsets = torch.tensor([0], dtype=torch.int64, device="xpu")
per_sample_weights = torch.ones(1024, device="xpu", requires_grad=True)

output, offset2bag, _, _ = torch._embedding_bag(
    weight, indices, offsets, False, 0, False, per_sample_weights
)

# Alias through DLPack so the saved tensor is mutated without bumping its
# autograd version counter, exactly as the upstream CUDA test does.
torch.from_dlpack(indices).fill_(weight.size(0))  # 5 == num_embeddings

output.sum().backward()
torch.xpu.synchronize()
print("CLEAN_EXIT")
'''

proc = subprocess.run(
    [sys.executable, "-c", CHILD], capture_output=True, text=True, timeout=35
)
print("exit:", proc.returncode)
print((proc.stdout + proc.stderr).strip())
# Upstream post-fix contract: a device-side assertion naming an index bound.
# On XPU today: exit 0 and CLEAN_EXIT.

The child runs out of process because the expected upstream outcome is a device-side abort that cannot be caught in process, and it calls torch.xpu.synchronize() before printing so an asynchronous device fault would surface before exit.

Observed output

CASE: indices filled with its upper bound (int64)
  child exit=0
  child output: CLEAN_EXIT
REASON: indices: backward completed with an out-of-range saved index, so the XPU kernel indexed memory with it unguarded
VERDICT: SAME_FAILURE

No ALIAS_FALLBACK line appears, so the DLPack alias -- not .data -- mutated the saved tensor and autograd's version counter was not bumped: the value the kernel read really was the saved one.

Scope of the evidence

Only the indices upper-bound int64 case executed; the harness stops at its first failure. The offset2bag case, the negative-bound cases and the int32 cases in upstream's parametrization were not run. Their status here comes from source inspection: bag_idx is equally unguarded, and no lower bound is checked for either index.

One qualification on the signal. A clean exit demonstrates that the invalid index was used unguarded, not that memory outside the USM allocation was touched -- a two-element overread past a 5x2 fp32 buffer will usually stay inside the same page, so a clean run is expected even though the access is out of bounds. The defect is established by the source above; the run corroborates that no assertion exists.

Suggested fix

Mirror the upstream change:

Thread grad.size(0) (num_bags) and weight.size(0) (num_embeddings) from _embedding_bag_per_sample_weights_backward_impl into the functor.
Read offset2bag_[i] and indices_[i] at index_t instead of casting to int, and keep the addressing arithmetic at that width.
Add SYCL_KERNEL_ASSERT bounds checks for both indices before the grad_ / weight_ reads.
Add an XPU test covering {indices, offset2bag} x {negative, upper} x {int32, int64}, since upstream's tests are CUDA-only.

Environment

Python: /opt/conda/bin/python3.13, invoked with -I.
torch install: /opt/conda/lib/python3.13/site-packages/torch, built from /__w/pytorch/pytorch.
torch.xpu.is_available(): true, asserted by the script before the case ran.
torch version and XPU device name: not recorded by the run that produced this evidence. Stated rather than guessed; please confirm against your own build when reproducing.

Provenance

Produced by the xpu-alignment skill in automation mode, and reviewed independently of the scan that produced the local result. Reproducer alignment-run/scripts/repro_commit-550fa843275302768fc1fe758f308b2dccf543b2.py (sha256 6af1ae19088e7db7c2f8b19ef68c2dfb6771e524b41509a471665725c9431bc6), log alignment-run/runner/logs/commit-550fa843275302768fc1fe758f308b2dccf543b2.log.

_{alignment scan 2026-08-24, run 32923438489}

[DRY RUN] [xpu-alignment] XPU group_norm rejects fp32 affine params over an fp16 input, which the shared ATen mixed-dtype contract supports

Summary

XPU group_norm raises RuntimeError: expected scalar type Half but found Float for a reduced-precision input with fp32 affine parameters -- a combination the shared ATen layer explicitly validates and allocates buffers for, and that CPU computes. Root cause: both XPU dispatch stubs template the entire implementation on X.scalar_type(), so the fp32 buffers the shared layer allocated are accessed as the input type.

Upstream source

PR: Preserve channels_last in InstanceNorm and fix the channels-last group norm reductions at D == 1 pytorch/pytorch#194616 -- "Preserve channels_last in InstanceNorm and fix the channels-last group norm reductions at D == 1" (OPEN, unmerged at review time)
Alignment scan window: [2026-08-24T00:00:00Z, 2026-08-25T00:00:00Z), repository pytorch/pytorch

Among other changes, that PR threads a separate parameter type through the CUDA group norm kernels and allocates dgamma / dbeta from param_scalar_type(X, mixed_type) in shared aten/src/ATen/native/group_norm.cpp. Validating its mixed-dtype contract on XPU surfaced a different, pre-existing failure: XPU does not accept the mixed-dtype call at all.

Why this is a defect and not unsupported input

Shared aten/src/ATen/native/group_norm.cpp supports the combination explicitly:

bool mixed_type = is_mixed_type(X, gamma, beta);
if (mixed_type) {
  check_mixed_data_type(X, gamma, beta);
}
...
auto stat_options = X.options()
                        .memory_format(at::MemoryFormat::Contiguous)
                        .dtype(param_scalar_type(X, mixed_type));
Tensor mean = at::empty({N, group}, stat_options);
Tensor rstd = at::empty({N, group}, stat_options);

For an fp16 X with fp32 gamma / beta, the shared layer validates the combination and allocates mean / rstd as fp32 before calling the backend stub.

This is also the ordinary AMP shape rather than an exotic input: group_norm has no autocast cast policy outside MPS -- KERNEL_MPS(group_norm, fp32) is the only entry in aten/src/ATen/autocast_mode.cpp -- so under torch.autocast("xpu") a GroupNorm fed by a preceding fp16 op receives an fp16 activation together with its own fp32 parameters, unchanged.

Root cause and target-path evidence

Line numbers are from intel/torch-xpu-ops main, src/ATen/native/xpu/sycl/GroupNormKernels.cpp:

:1379 -- group_norm_kernel: AT_DISPATCH_FLOATING_TYPES_AND2(kHalf, kBFloat16, X.scalar_type(), ...)
:2445 -- group_norm_backward_kernel: the same, on X.scalar_type() alone
:1042 -- mean.mutable_data_ptr<T>(), the first throw site (mean is fp32 when mixed, T is bound to Half)
:1055, :1717, :2273 -- gamma.const_data_ptr<T>()
:1777, :2401 -- dgamma.mutable_data_ptr<T>()

Because T is bound to the input's scalar type for the whole implementation, mutable_data_ptr<Half>() on the fp32 mean throws, which is exactly the observed message. There is no separate parameter type anywhere in the file.

The target path is established rather than assumed: the reproducer asserts the XPU dispatch key for both aten::native_group_norm and aten::native_group_norm_backward before any call, and every tensor is allocated on xpu.

Reproducer

# repro_group_norm_mixed_dtype.py
import torch

N, C, H, W, G = 2, 6, 4, 4, 3
HXW, EPS = H * W, 1e-5

assert torch.xpu.is_available()
for op in ("aten::native_group_norm", "aten::native_group_norm_backward"):
    assert torch._C._dispatch_has_kernel_for_dispatch_key(op, "XPU"), op


def run(device):
    torch.manual_seed(7)
    x = torch.randn(N, C, H, W, dtype=torch.float32).to(device, torch.float16)
    torch.manual_seed(8)
    gamma = torch.randn(C, dtype=torch.float32).to(device)
    beta = torch.randn(C, dtype=torch.float32).to(device)
    torch.manual_seed(9)
    grad_y = torch.randn(N, C, H, W, dtype=torch.float32).to(device, torch.float16)
    x.requires_grad_(True)
    gamma.requires_grad_(True)
    beta.requires_grad_(True)

    y, mean, rstd = torch.native_group_norm(x, gamma, beta, N, C, HXW, G, EPS)
    dX, dgamma, dbeta = torch.autograd.grad(y, (x, gamma, beta), grad_y)
    if device == "xpu":
        torch.xpu.synchronize()
    print(device, "Y", y.dtype, "mean", mean.dtype, "dgamma", dgamma.dtype)


run("cpu")  # computes; mean / rstd / dgamma / dbeta are float32
run("xpu")  # RuntimeError: expected scalar type Half but found Float

The same dispatch path is reached from the user-facing module form below. This form was not executed in the run that produced the output quoted next; it is derived from source (nn.GroupNorm.forward -> F.group_norm -> at::group_norm -> native_group_norm) and is included because it is how the failure is likely to be met in practice:

import torch

m = torch.nn.GroupNorm(3, 6).to("xpu")  # fp32 weight and bias
x = torch.randn(2, 6, 4, 4, device="xpu", dtype=torch.float16)
m(x)

Observed output

Traceback (most recent call last):
  File ".../alignment-run/scripts/repro_pr-194616.py", line 142, in main
    xpu_vals, xpu_dtypes = run("xpu", with_gamma)
  File ".../alignment-run/scripts/repro_pr-194616.py", line 99, in run
    y, mean, rstd = torch.native_group_norm(
        x, gamma if with_gamma else None, beta, N, C, HXW, G, EPS
    )
RuntimeError: expected scalar type Half but found Float
CASE: affine (fp16 input, fp32 params)
REASON: affine: XPU rejected fp32 affine params over an fp16 input instead of computing them at the parameter type
VERDICT: RELATED_FAILURE

The cpu reference run happens immediately before the xpu run and outside the try, and it succeeded -- so the identical Python-level call computes on CPU and raises on XPU.

Relationship to CUDA and to PR #194616

Stated plainly so this is not mistaken for an XPU-only regression: CUDA on main has the same single-type templating (aten/src/ATen/native/cuda/group_norm_kernel.cu:844 mean.mutable_data_ptr<T>(), :899 and :962 gamma.const_data_ptr<T>()), so both GPU backends reject this call today. The divergence is from CPU, which computes it, and from the shared layer's own stated contract.

It still needs work here: PR #194616 fixes the CUDA half and touches no XPU file, so no upstream change can reach GroupNormKernels.cpp. #194616 is the reference design to mirror.

No upstream test would catch this on XPU either. test_groupnorm_nhwc is @onlyCPU; test_groupnorm_nhwc_cuda and test_groupnorm_unaligned_input are @onlyCUDA and convert the whole module to a single dtype; _test_GroupNorm_cpu_mixed_dtype runs only when device == "cpu".

Scope of the evidence

Only the affine case ran. The gamma-undefined case -- undefined gamma with an fp32 beta, where upstream's dgamma / dbeta allocation change matters -- never executed, because the harness stops at the first failure.
The channels_last half of PR #194616 was deliberately excluded from the oracle: shared group_norm_memory_format returns suggest_memory_format() only for CPU, CUDA and PrivateUse1, so channels_last is not part of XPU's group norm contract today and asserting it would not be a faithful oracle.
The PR's reduction rework at D == 1 is a throughput change and needs a performance harness; nothing here speaks to it.

Suggested fix

Introduce a separate parameter type in GroupNormKernels.cpp: use param_scalar_type(X, mixed_type) for the mean / rstd / gamma / beta / dgamma / dbeta accesses while keeping T for the activation reads and writes, and dispatch on both types, as upstream does for CUDA in #194616. Then add an XPU mixed-dtype test covering the affine and gamma-undefined cases for fp16 and bf16 inputs, since the existing upstream mixed-dtype test is CPU-gated.

Environment

Python: /opt/conda/bin/python3.13, invoked with -I.
torch install: /opt/conda/lib/python3.13/site-packages/torch, built from /__w/pytorch/pytorch.
torch.xpu.is_available(): true, asserted before the dispatch-key probes.
torch version and XPU device name: not recorded by the run that produced this evidence. Stated rather than guessed; please confirm against your own build when reproducing.

Provenance

Produced by the xpu-alignment skill in automation mode, and reviewed independently of the scan that produced the local result. Reproducer alignment-run/scripts/repro_pr-194616.py (sha256 f694e48f6f691c871d0bb11d89234ec15b94a0f22da283f8931c11251f28fc07), log alignment-run/runner/logs/pr-194616.log.

[DRY RUN] XPU alignment run

Dry run completed with triage and 2 review-approved candidate(s).

commit-550fa843275302768fc1fe758f308b2dccf543b2 — [xpu-alignment] embedding_bag per_sample_weights backward indexes weight/grad with unchecked, int-truncated saved indices on XPU
pr-194616 — [xpu-alignment] XPU group_norm rejects fp32 affine params over an fp16 input, which the shared ATen mixed-dtype contract supports

This run did not finish cleanly:

2 blocker(s): scan-blocked-result:commit-2509857d0e811629f1d34c3d187c172a7c8bc935:blocked-platform, scan-blocked-result:commit-e812f6fc74ce93d55284ca7534d84c7356c4ea94:blocked-env

Reviewed candidates were published or queued; blocked units were excluded.

_{alignment scan 2026-08-24, run 32923438489}

[Workflow] Add the daily XPU alignment scan, review and filing pipeli…

_{alignment scan 2026-08-28, run 33228378520}

[xpu-alignment] embedding_renorm_ derives the row width from stride(0) and faults on a zero-width embedding

Upstream source

Issue: [CUDA] illegal memory read on renorm_kernel pytorch/pytorch#195185 - [CUDA] illegal memory read on renorm_kernel (OPEN, created 2026-08-28).
Proposed upstream fix: Fix CUDA illegal memory access in embedding_renorm_ for zero-width embeddings pytorch/pytorch#195208 (OPEN, unmerged). It changes aten/src/ATen/native/cuda/Embedding.cu only, and its regression test is @onlyOn(["cuda"]).
Alignment scan window: [2026-08-28T00:00:00Z, 2026-08-29T00:00:00Z).

Observed XPU behavior

F.embedding(indices, weight, max_norm=1.0, norm_type=2.0) with a zero-width
weight fails on XPU:

weight: torch.Size([20, 0]) (1, 1) 0 xpu:0 | stride(0)= 1 size(1)= 0
cpu out: (1, 0)
RESULT: confirmed
reason: XPU embedding_renorm_ on a zero-width weight failed: level_zero backend failed with error: 40 (UR_RESULT_ERROR_OUT_OF_RESOURCES)

CPU, given the identical inputs, returns shape (1, 0) cleanly. The oracle here
is CPU plus the upstream signature, not a tolerance comparison.

Root cause and target-path evidence

src/ATen/native/xpu/sycl/Embedding.cpp:237 takes the same shortcut upstream
names as the CUDA root cause:

int dim = self.stride(0);

torch.empty(20, 0).stride() is (1, 1), so dim becomes 1 while the real
row width self.size(1) is 0. dim is passed to
embedding_renorm_template, and RenormKernelFunctor::operator()
(src/ATen/native/xpu/sycl/Embedding.cpp:100-141) then runs its accumulation
loop once and dereferences weights_[base_index]:

for (int i = tid; i < dim_; i += sgSize) {
  auto x = static_cast<accscalar_t>(weights_[base_index + i * weights_stride1_]);

The (20, 0) weight has zero elements, so that pointer has no backing
allocation - the same invalid read compute-sanitizer reports for CUDA. The fix
is the same one-line change: use self.size(1). The kernel already receives
weights_stride0/weights_stride1 separately for indexing, so it does not
depend on the tensor being contiguous, and dim == 0 makes both loops execute
zero times without skipping the launch.

Target-path proof that this is the XPU kernel and not setup data or an
unsupported input:

the log records the weight as resident on xpu:0 and in exactly the
stride(0)=1, size(1)=0 state the defect requires;
the CPU reference ran first on the same inputs and succeeded, so this is not
an input-validation rejection;
embedding_renorm_ is dispatched to XPU via embedding_renorm_xpu_ in
src/ATen/native/xpu/Embedding.cpp:39, so no fallback is involved;
the runner's environment probe was clean and five other reproducers did heavy
XPU work on the same device in the same run, so this is a kernel-level fault
rather than a global environment failure.

The UR error code alone does not distinguish an out-of-bounds device read from a
degenerate launch derived from dim; both trace to the same line, and either
way the oracle is violated on the XPU path.

Reproducer

import torch
import torch.nn.functional as F

weight = torch.empty((20, 0), dtype=torch.float32, device="xpu")
indices = torch.tensor([0], dtype=torch.int64, device="xpu")
print("cpu:", tuple(F.embedding(indices.cpu(), weight.cpu(), max_norm=1.0, norm_type=2.0).shape))
out = F.embedding(indices, weight, max_norm=1.0, norm_type=2.0)
torch.xpu.synchronize()
print("xpu:", tuple(out.shape))  # expected (1, 0)

Environment

torch 2.15.0.dev20260826+xpu, Python 3.13.13 (/opt/conda/bin/python3.13),
Intel(R) Arc(TM) Pro B60 Graphics, torch.xpu.is_available() == True, no
environment warnings. Neither upstream fix candidate is present in this build.

Ownership

The kernel is torch-xpu-ops code. pytorch/pytorch#195208 fixes only the CUDA
kernel and adds a CUDA-only test, so nothing upstream will fix or detect the XPU
copy. No pytorch/pytorch issue or PR claims the XPU work, and no
intel/torch-xpu-ops tracker covers it. This XPU fix is independent of whether
#195208 lands: XPU already diverges from CPU today, and no existing test
asserts the faulting behavior.

Suggested XPU regression test: the same reproducer, mirroring upstream's
test_embedding_max_norm_zero_embedding_dim.

_{alignment scan 2026-08-28, run 33228378520}

[xpu-alignment] exponential_ underflows to exactly 0 for float16 on XPU, outside the distribution support

Upstream source

PR: Fix exponential_ underflowing to exactly 0 for float16 pytorch/pytorch#195077 - Fix exponential_ underflowing to exactly 0 for float16 (OPEN, unmerged; fixes torch.multinomial selects zero-probability indices on CPU for float16, so exponential_ underflows to exactly 0 in the double to half cast pytorch/pytorch#192812).
It touches aten/src/ATen/native/cpu/DistributionTemplates.h,
aten/src/ATen/native/cuda/DistributionTemplates.h,
c10/util/safe_numerics.h (new c10::cast_no_underflow), and test/test_torch.py.
No XPU kernel is touched.
Alignment scan window: [2026-08-28T00:00:00Z, 2026-08-29T00:00:00Z).

Observed XPU behavior

The exponential distribution has support (0, inf), so Tensor.exponential_()
must never return exactly 0. Over 24e6 XPU draws per (dtype, lambd):

torch.float16    lambd=1.0        zeros=0        min_positive=4.883e-04  ok
torch.float16    lambd=1024.0     zeros=0        min_positive=4.768e-07  ok
torch.float16    lambd=16384.0    zeros=11655    min_positive=5.960e-08  WRONG
torch.float16    lambd=65504.0    zeros=47262    min_positive=5.960e-08  WRONG
torch.bfloat16   (all lambd)      zeros=0                                ok
torch.float32    (all lambd)      zeros=0                                ok

No sample was non-finite or negative. min_positive at the failing rows is
5.960e-08, exactly the smallest float16 subnormal, which is the signature of a
narrowing cast flushing a nonzero sample to zero.

Root cause and target-path evidence

Upstream states the root cause as: exponential_ samples in a wider type and
narrows with static_cast<scalar_t>; anything below float16's smallest
subnormal becomes an exact 0. XPU reproduces both halves of that:

ExponentialFunctor in src/ATen/native/xpu/sycl/DistributionTemplates.h:854
computes in accscalar_t and returns accscalar_t, applying the eps/2
floor in the accumulate type. For lambd=65504 that floor is
4.88e-4 / 2 / 65504 = 3.7e-9, already below the half subnormal.
DistributionElementwiseKernelFunctor::operator()
(src/ATen/native/xpu/sycl/DistributionTemplates.h:88 and :92) writes the
transform result straight into the output:
```
*out = transform_func_(static_cast<accscalar_t>((&rand.x)[i]));
```
out is scalar_t*, so this is the implicit accscalar_t -> scalar_t
narrowing with no floor in the output dtype - the same shape as the
static_cast<scalar_t> upstream replaces.

Target-path proof: the reproducer asserts t.device.type == "xpu" on every
draw and synchronizes before reading, so the samples were produced by the XPU
kernel, not by a host fallback or setup tensor.

Reproducer

import torch

torch.xpu.manual_seed(0)
for lambd in (1.0, 1024.0, 16384.0, 65504.0):
    zeros = 0
    for _ in range(3):
        t = torch.empty(8_000_000, dtype=torch.float16, device="xpu").exponential_(lambd)
        torch.xpu.synchronize()
        zeros += int((t.float() == 0).sum().item())
    print(f"lambd={lambd} zeros={zeros}")  # expected 0 for every lambd

Environment

torch 2.15.0.dev20260826+xpu, Python 3.13.13 (/opt/conda/bin/python3.13),
Intel(R) Arc(TM) Pro B60 Graphics, torch.xpu.is_available() == True, no
environment warnings. pytorch/pytorch#195077 is not in this build.

Ownership

The kernel is torch-xpu-ops code and the upstream PR does not touch it, so
landing #195077 leaves XPU broken. c10::cast_no_underflow will be available
from c10/util/safe_numerics.h once that PR lands, so the XPU change is a
single narrowing-site swap in DistributionTemplates.h. No pytorch/pytorch
issue or PR claims the XPU work, and no intel/torch-xpu-ops tracker covers
it - intel/torch-xpu-ops#4020 was considered and rejected, since it tracks
statistical-bias and moment assertion failures in test_distributions_xpu.py,
a different oracle and root cause.

CI impact: #195077 drops @onlyCUDA from test_exponential_no_zero in
test/test_torch.py and parametrizes it over lam in (1, 1024). That makes
it device-generic, so it will run on XPU through test/xpu/test_torch_xpu.py
and fail until this is fixed.

The XPU fix can land before or after #195077: nothing asserts that zeros are
produced, and preserving an exact input 0 keeps exponential_(inf) == 0
intact.

_{alignment scan 2026-08-28, run 33228378520}

[xpu-alignment] silu(-inf) returns NaN on XPU instead of the limit 0

Upstream source

PR: Fix silu(-inf) returning NaN instead of the limit 0 pytorch/pytorch#195197 - Fix silu(-inf) returning NaN instead of the limit 0 (OPEN, unmerged; fixes F.silu(torch.tensor(float('-inf'))) returns nan instead of 0 in PyTorch 2.7.0+cu128 pytorch/pytorch#160876, open since 2025-08-18).
It touches aten/src/ATen/native/cpu/Activation.cpp,
aten/src/ATen/native/cuda/ActivationSiluKernel.cu,
test/test_unary_ufuncs.py, and the nn.functional.silu OpInfo reference in
torch/testing/_internal/common_methods_invocations.py. No XPU kernel is
touched.
Alignment scan window: [2026-08-28T00:00:00Z, 2026-08-29T00:00:00Z).

Observed XPU behavior

silu(x) = x * sigmoid(x) tends to 0 as x -> -inf, so silu(-inf) must be
0. On XPU it is NaN for every floating dtype:

torch.float16    silu(-inf) xpu=nan cpu=nan NaN
torch.bfloat16   silu(-inf) xpu=nan cpu=nan NaN
torch.float32    silu(-inf) xpu=nan cpu=nan NaN
torch.float64    silu(-inf) xpu=nan cpu=nan NaN

Large finite negatives (-1e4, -1e3, -100) return -0.0 correctly, so this
is specifically the -inf / inf form and not early flushing. The backward at
-inf, whose limit is also 0, is NaN for all four dtypes as well.

The oracle is the mathematical limit, not CPU: CPU currently shares the same
defect, which is what #195197 fixes there.

Root cause and target-path evidence

SiluFunctor in src/ATen/native/xpu/sycl/ActivationSiluKernels.cpp:26-35
evaluates the algebraically equivalent but numerically unsafe form:

return x_acc / (opmath_t(1) + sycl::exp(-x_acc));

At x = -inf this is -inf / (1 + inf) = -inf / inf = NaN. Upstream's CUDA
fix adds an if constexpr (!c10::is_complex<scalar_t>::value) guard returning
scalar_t(0) at -inf; the XPU functor needs the same guard, and complex
silu must keep its current behavior since -inf has no complex ordering.

SiluBackwardFunctor (:38-46) has the same structure through
opmath_t(1) / (opmath_t(1) + sycl::exp(-x_acc)); upstream's PR does not
address the backward on any backend, so treat that as a related observation
rather than part of the parity change.

Target-path proof: the reproducer asserts the silu output is on XPU and
synchronizes before reading. Independently, the float64 row differs from CPU
in the low-order digits (-0.2689414322376251 on XPU vs -0.2689414213699951
on CPU), which shows the printed values came from the XPU kernel rather than a
host fallback.

Reproducer

import torch
import torch.nn.functional as F

for dtype in (torch.float16, torch.bfloat16, torch.float32, torch.float64):
    x = torch.tensor([float("-inf"), -1e4, -1.0, 0.0, 1.0], dtype=dtype, device="xpu")
    out = F.silu(x)
    torch.xpu.synchronize()
    print(dtype, out.float().cpu().tolist())  # first element expected 0.0

Environment

torch 2.15.0.dev20260826+xpu, Python 3.13.13 (/opt/conda/bin/python3.13),
Intel(R) Arc(TM) Pro B60 Graphics, torch.xpu.is_available() == True, no
environment warnings. pytorch/pytorch#195197 is not in this build.

Ownership

The kernel is torch-xpu-ops code and the upstream PR does not touch it, so
landing #195197 leaves XPU as the only backend still returning NaN. No
pytorch/pytorch issue or PR claims the XPU work, and no intel/torch-xpu-ops
tracker covers it.

CI impact once #195197 lands: test_silu in test/test_unary_ufuncs.py
gains a silu(-inf) == 0 assertion and the shared OpInfo reference becomes
np.where(x == -inf, 0.0, x / (1 + np.exp(-x))). Both are device-generic, so
test_silu and test_reference_numerics_extremal will fail on XPU through
test/xpu/test_unary_ufuncs_xpu.py.

Sequencing caveat: unlike the other units in this scan, the XPU change should
land with or after #195197, not before it. XPU's test_reference_numerics_*
compares against the shared OpInfo reference, so returning 0 while that
reference still evaluates the naive form would fail XPU CI in the opposite
direction. If #195197 is rejected upstream, this unit should be closed rather
than implemented.

@laifenxiawucha

3 XPU alignment candidates need triage

3 reviewed candidates are waiting for approval:

issue-195185 — [xpu-alignment] embedding_renorm_ derives the row width from stride(0) and faults on a zero-width embedding
pr-195077 — [xpu-alignment] exponential_ underflows to exactly 0 for float16 on XPU, outside the distribution support
pr-195197 — [xpu-alignment] silu(-inf) returns NaN on XPU instead of the limit 0

Approve them one at a time with @torchxpubot file <unit-id>.

_{alignment scan 2026-08-28, run 33228378520}

@laifenxiawucha @ZhaoqiongZ

@laifenxiawucha

XPU alignment run needs attention

This run did not finish cleanly:

1 blocker(s): scan-blocked-result:commit-1782e8ec673e736ffda3414fd921f3c1b1b62002:blocked-script-error

Nothing was published for the affected part; see the gate artifact.

_{alignment scan 2026-08-29, run 33287538032}

@laifenxiawucha @ZhaoqiongZ

@laifenxiawucha

XPU alignment run completed with warnings

Scan date: 2026-11-07
Collection: complete
Reviewed units: 2 (duplicate: 1, track-upstream: 1)
Review-approved tracker candidates: 0

No new XPU tracker was filed or drafted.

Excluded units:
- pr-203892 — blocked-script-error

_{workflow run 42437331929}

@laifenxiawucha @ZhaoqiongZ

@laifenxiawucha

XPU alignment run completed with warnings

Scan date: 2026-11-08
Collection: complete
Reviewed units: 5 (track-upstream: 5)
Review-approved tracker candidates: 0

No new XPU tracker was filed or drafted.

Excluded units:
- issue-203969 — blocked-platform
- issue-203979 — blocked-env

_{workflow run 42549395667}

@laifenxiawucha @ZhaoqiongZ

Filed as #6127
_{alignment scan 2026-11-09, workflow run 42740792617}

[xpu-alignment] [26-11-09] XPU segment_reduce accumulates fp16/bf16 in the input dtype: 4096 ones sum to 2048 (fp16) and 256 (bf16)

Summary

torch.segment_reduce on XPU keeps its running reduction value in the input dtype, so float16 and bfloat16 sum and mean stop growing once the partial sum passes the dtype's integer-resolution limit. A single segment of 4096 ones reduces to 2048 in float16 and 256 in bfloat16 instead of 4096, and mean returns 0.5 and 0.0625 instead of 1.0. Both expected values are exactly representable in either dtype, so this is accumulator loss, not representation error. The result is silently wrong: no warning, no error.

Upstream source

CPU segment_reduce undercounts FP16/BF16 sums and means pytorch/pytorch#204045 (open, filed 2026-11-09, labelled high priority and module: correctness (silent)) reports the same defect for the CPU kernel.
[CPU] Accumulate segment_reduce in opmath type pytorch/pytorch#204079 ("[CPU] Accumulate segment_reduce in opmath type", open, not landed) is the upstream fix. It changes aten/src/ATen/native/SegmentReduce.cpp and aten/src/ATen/native/cuda/SegmentReduce.cu to accumulate in at::opmath_type<scalar_t> and adds a device-generic regression test. It does not touch intel/torch-xpu-ops.
Found by an XPU alignment scan of the window [2026-11-09T00:00:00Z, 2026-11-10T00:00:00Z) over pytorch/pytorch, frozen default-branch head a68f117948121cdff578ff7813c5c906ff529b57.

Observed XPU behavior

Validated on XPU with torch 2.16.0.dev20261109+xpu:

torch.float16  sum   actual 2048.0   fp32_reference 4096.0
torch.float16  mean  actual 0.5      fp32_reference 1.0
torch.bfloat16 sum   actual 256.0    fp32_reference 4096.0
torch.bfloat16 mean  actual 0.0625   fp32_reference 1.0

All four cases reproduce the exact values upstream reports for CPU. The reference in each row is torch.segment_reduce(data.float(), ...) computed on the same XPU device and cast back, so the comparison is XPU against XPU and is not a host/device artifact.

Target path and root cause

segment_reduce_forward_kernel in src/ATen/native/xpu/sycl/SegmentReduceKernels.cpp:40. The accumulator is deduced from the scalar_t initial value and never widened:

// src/ATen/native/xpu/sycl/SegmentReduceKernels.cpp:49,61
    scalar_t initial_value_raw,
...
  auto initial_value = initial_value_raw;   // deduces scalar_t

// src/ATen/native/xpu/sycl/SegmentReduceKernels.cpp:88
      initial_value = initial_value + data;  // rounded to scalar_t every step

// src/ATen/native/xpu/sycl/SegmentReduceKernels.cpp:109
    initial_value = initial_value / lengths_data[lengths_idx];

For float16 every partial sum is rounded back to float16, so once the sum reaches 2^11 = 2048 adding 1.0 rounds to the same value and the sum stops advancing; for bfloat16 the limit is 2^8 = 256. mean inherits the truncated sum because the division happens after the loop, still in scalar_t. The initial value is scalar_t at the dispatch site too (src/ATen/native/xpu/sycl/SegmentReduceKernels.cpp:231).

Unlike CUDA, XPU has a single forward path (segment_reduce_forward_kernel_impl, line 117) behind both _segment_reduce_lengths_kernel (line 268) and _segment_reduce_offsets_kernel (line 278), so lengths- and offsets-based calls are both affected and one accumulator change covers both.

Target-path evidence: the lengths tensor and all four outputs are xpu:0, and torch.xpu.synchronize() is called before each value is read.

Reproducer

import torch

lengths = torch.tensor([4096], device="xpu")
for dtype in (torch.float16, torch.bfloat16):
    data = torch.ones(4096, dtype=dtype, device="xpu")
    for reduction in ("sum", "mean"):
        actual = torch.segment_reduce(data, reduction, lengths=lengths)
        reference = torch.segment_reduce(data.float(), reduction, lengths=lengths).to(dtype)
        print(dtype, reduction, "actual:", actual.item(), "fp32 reference:", reference.item())

Environment

torch 2.16.0.dev20261109+xpu
python 3.13.13
device Intel(R) Arc(TM) Pro B60 Graphics

Ownership

The XPU kernel carries its own copy of the accumulator. At the scanned head the same scalar_t accumulation is present upstream in aten/src/ATen/native/cuda/SegmentReduce.cu (initial_value = initial_value + data), so XPU is currently a faithful port rather than a divergence. The upstream fix PR #204079 rewrites only the CPU and CUDA accumulators; src/ATen/native/xpu/sycl/SegmentReduceKernels.cpp is untouched by it, so the XPU change is independent work in this repository and does not depend on #204079 landing.

Stage of the upstream fix: proposed, not landed (#204079 is open against a ghstack base as of this scan). Its approach, accumulate in at::opmath_type<scalar_t> and cast once when writing the output, is the one to mirror here. #204079 names two remaining upstream gaps that do not apply to the XPU kernel structure: the CUDA 1-D mean path storing the sum in the output dtype, and the CPU prod backward.

Related, not covering this finding:

[xpu-alignment] [26-10-21] XPU segment_reduce max/min backward gives a negative tied gradient to every tied element instead of splitting it #5808 (open) is a different defect in the same file: the max/min backward tied-gradient split. It does not change the forward accumulator.
[MPS] Add native segment_reduce forward pass pytorch/pytorch#197014 (open) adds an MPS segment_reduce forward and is cross-referenced from #204045; it has no XPU scope.

No open or closed intel/torch-xpu-ops tracker was found for the fp16/bf16 forward accumulation dtype (searched segment_reduce, opmath accumulation, fp16 bfloat16 accumulate sum, and PRs touching SegmentReduce*).

Regression test

test/xpu/test_segment_reductions_xpu.py runs the upstream suite with no skips (test/xpu/skip_list_common.py:229 maps it to None). The device-generic test #204079 adds, TestSegmentReductions.test_reduced_precision_accumulation, will therefore run on XPU as soon as that PR lands and will fail until the XPU accumulator is widened. No existing XPU test covers this: the current cases use segments short enough that scalar_t accumulation is exact.

Filed as #6128
_{alignment scan 2026-11-09, workflow run 42740792617}

[xpu-alignment] [26-11-09] XPU Tensor.renorm divides by norm + 1e-7, so renormalised rows land below max_norm (33% short at max_norm=1e-7)

Summary

Tensor.renorm / Tensor.renorm_ on XPU scale a too-long row by max_norm / (norm + 1e-7) instead of max_norm / norm, so the renormalised row never reaches max_norm, which is what the docstring promises. The 1e-7 is absolute, so the error is about 2e-8 relative at max_norm near 1 but grows without bound as max_norm shrinks: about 0.05% at 1e-4, 4.8% at 1e-6 and 33% at 1e-7. The guard is also unnecessary: the scaling branch only runs when norm > max_norm >= 0, so norm > 0 and the division is always safe.

Upstream source

F.embedding and nn.Embedding with max_norm renormalise to max_norm/(1+1e-7/norm) instead of max_norm pytorch/pytorch#204119 (open, filed 2026-11-09) reports the same +1e-7 for F.embedding(max_norm=...), nn.Embedding, embedding_bag and Tensor.renorm/renorm_ on CPU.
[nn] Preserve exact max norms during renormalization pytorch/pytorch#204132 ("[nn] Preserve exact max norms during renormalization", open, not landed) is the upstream fix. It divides by norm directly in aten/src/ATen/native/cpu/RenormKernel.cpp, aten/src/ATen/native/cuda/RenormKernel.cu, aten/src/ATen/native/mps/kernels/RenormKernel.metal, aten/src/ATen/native/Embedding.cpp and aten/src/ATen/native/cuda/Embedding.cu, and adds device-generic regressions. It does not touch intel/torch-xpu-ops.
Found by an XPU alignment scan of the window [2026-11-09T00:00:00Z, 2026-11-10T00:00:00Z) over pytorch/pytorch, frozen default-branch head a68f117948121cdff578ff7813c5c906ff529b57.

Observed XPU behavior

Validated on XPU with torch 2.16.0.dev20261109+xpu, float64:

renorm_row0_norm 0.9999999800000003    # x.renorm(2, 0, 1.0) on [[3.,4.],[.6,.8],[0.,0.]]

0.9999999800000003 is exactly 1.0 * 5.0 / (5.0 + 1e-7), the value upstream reports, so the XPU renorm scale factor carries the same absolute guard.

Scope of what was measured, stated explicitly:

Confirmed: Tensor.renorm in float64.
Not reproduced: F.embedding(..., max_norm=...) in float64 returned exactly 1.0 (scale 0.20000000000000004) with zero relative shortfall at max_norm 1e-4, 1e-6 and 1e-7. That is because the XPU embedding renorm uses a dtype-relative epsilon rather than upstream's absolute 1e-7 (see below), and numeric_limits<double>::epsilon() is about 2.2e-16.
Not measured: the float32 embedding path. There accscalar_t is float and numeric_limits<float>::epsilon() is 1.19e-7, the same order as upstream's 1e-7, so a comparable shortfall is expected but was not exercised by this reproducer. It should be checked as part of the fix.

Target path and root cause

Two separate XPU copies of the guard:

// src/ATen/native/xpu/sycl/RenormKernel.cpp:22-24  (Tensor.renorm / renorm_)
    const auto eps = static_cast<scalar_t>(1e-7);
    const auto one = static_cast<scalar_t>(1.0);
    return (norm > maxnorm_elm) ? maxnorm_elm / (norm + eps) : one;

// src/ATen/native/xpu/sycl/Embedding.cpp:146-148  (embedding / embedding_bag renorm)
  if (smem_[0] > max_norm) {
    auto factor = static_cast<scalar_t>(
        max_norm / (smem_[0] + std::numeric_limits<accscalar_t>::epsilon()));

RenormKernel.cpp is a faithful port of upstream's CUDA kernel at the scanned head (maxnorm_s / (norm + eps), eps = 1e-7). Embedding.cpp already deviates: upstream CUDA uses an absolute max_norm / (sdata[0] + 1e-7), XPU uses numeric_limits<accscalar_t>::epsilon(). After #204132 lands, upstream CPU/CUDA/MPS divide by norm exactly and both XPU copies are left adding an epsilon.

Target-path evidence: the weight, index, embedding output and renorm output are all xpu:0, and torch.xpu.synchronize() is called before each value is read.

Reproducer

import torch

x = torch.tensor([[3.0, 4.0], [0.6, 0.8], [0.0, 0.0]], dtype=torch.float64, device="xpu")
rn = x.renorm(2, 0, 1.0)
torch.xpu.synchronize()
print("row 0 norm after renorm:", float(rn[0].norm()))  # 0.9999999800000003, expected 1.0
print("1.0 * 5 / (5 + 1e-7)   :", 1.0 * 5.0 / (5.0 + 1e-7))

# float32, the dtype the upstream regression test uses
small = torch.tensor([[2e-6, 0.0]], device="xpu").renorm(2, 0, 1e-6)
torch.xpu.synchronize()
print("small max_norm:", float(small.norm()), "expected 1e-06")

Environment

torch 2.16.0.dev20261109+xpu
python 3.13.13
device Intel(R) Arc(TM) Pro B60 Graphics

Ownership

renorm_scale_factor_stub is registered by this repository (src/ATen/native/xpu/sycl/RenormKernel.cpp) and the embedding renorm kernel lives here too, so neither is reachable from the upstream fix. #204132 changes only CPU, CUDA and MPS; the XPU change is independent work in this repository and does not depend on #204132 landing.

Stage of the upstream fix: proposed, not landed (#204132 is open against main as of this scan, and its own test plan states that no local rebuild or device run was performed and that CI is required). Its approach, drop the guard and divide by norm, justified by norm > max_norm >= 0, applies unchanged to both XPU copies.

Related, not covering this finding:

[xpu-alignment] embedding_renorm_ derives the row width from stride(0) and faults on a zero-width embedding #5152 (closed, fixed by Fix embedding_renorm_ row width to use size(1) not stride(0) #5163) was a different embedding_renorm_ defect in the same file: the row width came from stride(0) instead of size(1). It did not touch the epsilon, and closing it did not fix the guard.
Align the semantic of accumulate type in the renorm operator with stock PyTorch #823 and Align the semantic of accumulate type in the renorm operator with stock PyTorch #822 (merged, 2024) aligned the renorm accumulate type with upstream, not the epsilon.

No open intel/torch-xpu-ops tracker was found for the renorm epsilon (searched renorm, max_norm, renorm 1e-7 epsilon norm, and PRs touching renorm).

Regression test

test/xpu/test_torch_xpu.py runs the upstream device-generic test_torch.py suite with no skips (test/xpu/skip_list_common.py:259 maps it to None). #204132 adds TestTorchDeviceType.test_renorm_exact_maxnorm, which asserts that torch.tensor([[2e-6, 0.0]], device=device).renorm(2, 0, 1e-6).norm() equals 1e-6 within atol=1e-8; in float32 the current XPU kernel returns about 9.524e-7, a 4.8e-8 miss, so that test will fail on XPU once the PR lands. #204132 also adds nn/test_embedding.py coverage for the functional, module and embedding_bag entry points on accelerator devices, which will exercise the Embedding.cpp copy. A fix here should additionally add a float32 F.embedding(max_norm=...) case, since the float64 path masks the dtype-relative epsilon.

@laifenxiawucha

XPU alignment run completed with warnings

Scan date: 2026-11-09
Collection: complete
Reviewed units: 7 (needs-xpu-fix: 2, track-upstream: 5)
Review-approved tracker candidates: 2

Automatically filed:

issue-204045 — [xpu-alignment] [26-11-09] XPU segment_reduce accumulates fp16/bf16 in the input dtype: 4096 ones sum to 2048 (fp16) and 256 (bf16) #6127
issue-204119 — [xpu-alignment] [26-11-09] XPU Tensor.renorm divides by norm + 1e-7, so renormalised rows land below max_norm (33% short at max_norm=1e-7) #6128
Excluded units:
- issue-204062 — blocked-script-error

_{workflow run 42740792617}

@laifenxiawucha @ZhaoqiongZ

label-issue: #5018

Root cause: not a defect; sink issue the alignment workflow posts to (.github/workflows/xpu_alignment.yml:28 ALIGNMENT_TRIAGE_ISSUE: '5018').

axis	value	reason
`type`	`Task`	`issue_type` empty; describes a standing workflow triage/posting channel, no failure reported.
`test`	`test: ut`	No e2e/oob/ut signal; default fallback surface.
`module`	`module: infra`	Human-set label `module: infra` preserved; matches CI workflow sink (xpu_alignment.yml:28).
`priority`	`Medium`	No tier evidence matches; JSON default tier.

Applied to this issue. Generated by a local label-issue run.

_{alignment scan 2026-11-10, workflow run 42943920909}

[xpu-alignment] [26-11-10] XPU multinomial(replacement=True) builds its CDF in the input dtype: bfloat16 resolves 847 of 10000 draws and a half row summing past 65504 returns index 0 for every draw

Upstream source

Commit pytorch/pytorch@500c660 -- Compute multinomial internals in float for half and bfloat16 (PR Compute multinomial internals in float for half and bfloat16 pytorch/pytorch#204161), committed 2026-11-10T09:09:13Z.
Reachable from the scanned pytorch/pytorch default-branch head 3aad3844fc331821235a833579a71323af3c3eaf (GitHub compare reports behind_by 0), so the change is landed upstream.
Scan window: 2026-11-10T00:00:00Z .. 2026-11-11T00:00:00Z (half-open, UTC).

The commit touched aten/src/ATen/native/Distributions.cpp, aten/src/ATen/native/cuda/MultinomialKernel.cu and test/distributions/test_distributions.py, and no XPU file. Its new regression test records the XPU gap rather than fixing it:

@skipXPUIf(True, "XPU's ported kernel overflows its half CDF (old CUDA bug)")
@dtypes(torch.half, torch.bfloat16)
def test_multinomial_reduced_precision_with_replacement(self, device, dtype):
    # Half row sums overflow past 65504, and a bfloat16 CDF is too coarse: the old
    # kernel's 10000 uniform draws hit about 830 distinct categories, not about 9630
    n = 1 << 17
    probs = torch.ones(n, device=device, dtype=dtype)
    samples = torch.multinomial(probs, 10000, replacement=True)
    upper_fraction = (samples >= n // 2).float().mean().item()
    self.assertEqual(upper_fraction, 0.5, atol=0.05, rtol=0)
    self.assertGreater(samples.unique().numel(), 9000)

Observed XPU behaviour

The XPU with-replacement kernel still does its renormalisation, prefix sum and binary search in the input dtype, so it reproduces both defects the commit describes for the pre-fix CUDA kernel:

bfloat16 CDF resolution. 2**17 uniform categories, 10000 draws, seed 0: XPU resolves 847 distinct categories with max index 65405, so the whole upper half of the category range (indices >= 65536) is unreachable and upper_fraction is 0 rather than 0.5. CPU, from the same seed, resolves 9622 distinct with max 131053. Upstream's figure for the broken CUDA kernel is about 830.
half L1 overflow. A row of 70000 ones has an L1 sum of 70000, past the half maximum 65504. XPU returns index 0 for all 64 draws (min 0 max 0 distinct 1) and exits cleanly; CPU returns min 312 max 69907 distinct 64. Note this signature differs from the CUDA one upstream describes: on XPU the overflow produces a silent wrong result, not a device-side assert.

Target-path evidence

torch.multinomial(w_xpu, n, replacement=True) -> multinomial_with_replacement_stub(XPU) -> at::native::xpu::multinomial_kernel in src/ATen/native/xpu/sycl/MultinomialKernel.cpp. The reproducer asserts out.device.type == "xpu" and prints the result device for every call; the log shows xpu:0 on both arms. In that file:

renormRowsL1<scalar_t> accumulates and divides the row in scalar_t (lines 31-74), called via renormRows (line 93);
the CDF is built with at::cumsum_out(prefixSum, normDist, 1) where prefixSum = native::empty_like(self_v), i.e. still scalar_t (lines 441-454);
binarySearchForMultinomial<scalar_t> compares scalar_t cumulative values (line 109).

The single-sample path sample_multinomial_once_kernel already uses acc_type_device<scalar_t, kXPU> (line 388) and is not affected; only the with-replacement path needs the float promotion.

Reproducer

import torch

torch.manual_seed(0)
w = torch.ones(1 << 17, dtype=torch.bfloat16, device="xpu")
out = torch.multinomial(w, 10000, replacement=True)
print("bf16", out.device, "distinct", int(out.unique().numel()), "max", int(out.max()))
# xpu:0 distinct 847  max 65405      (cpu: distinct 9622  max 131053)

torch.manual_seed(0)
w = torch.ones(70000, dtype=torch.half, device="xpu")   # L1 sum 70000 > 65504
out = torch.multinomial(w, 64, replacement=True)
print("half", out.device, "min", int(out.min()), "max", int(out.max()))
# xpu:0 min 0  max 0                 (cpu: min 312  max 69907)

Relevant output

torch 2.16.0.dev20261110+xpu
xpu_device Intel(R) Arc(TM) Pro B60 Graphics
--- case B: bfloat16 CDF resolution ---
  [cpu] result device=cpu
  cpu distinct=9622/10000 min=6 max=131053
  [xpu] result device=xpu:0
  xpu distinct=847/10000 min=1 max=65405
--- case A: half L1 overflow (sum > 65504) ---
  [cpu] stdout: device cpu min 312 max 69907 distinct 64
  [xpu] stdout: device xpu:0 min 0 max 0 distinct 1

Environment

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

Ownership

The defective arithmetic is in intel/torch-xpu-ops (src/ATen/native/xpu/sycl/MultinomialKernel.cpp). Upstream fixed only aten/src/ATen/native/cuda/MultinomialKernel.cu and deliberately skipped XPU in the new test, so the parity work is not owned upstream. The fix mirrors the CUDA one: promote the renorm, the prefix sum and the search comparison to opmath_type<scalar_t> / float for half and bfloat16.

Related trackers

Searched intel/torch-xpu-ops and pytorch/pytorch; no open or closed issue covers XPU multinomial reduced-precision sampling. Same class of defect in other XPU kernels, not duplicates:

[xpu-alignment] [26-11-09] XPU segment_reduce accumulates fp16/bf16 in the input dtype: 4096 ones sum to 2048 (fp16) and 256 (bf16) #6127 -- XPU segment_reduce accumulates fp16/bf16 in the input dtype.
[xpu-alignment] [26-09-24] XPU linspace/logspace compute the half and bfloat16 step in the output dtype, so results diverge from CPU #5523 -- XPU linspace/logspace compute the half/bfloat16 step in the output dtype.

_{alignment scan 2026-11-10, workflow run 42943920909}

[xpu-alignment] [26-11-10] XPU split_with_sizes_copy.out takes its byte-copy fast path for any dense input: wrong values for transposed and channels_last inputs, and a write past the output when the size-1 split dim has a large stride

Upstream source

Commit pytorch/pytorch@b0c56ca -- [CUDA] Only take the split_with_sizes_copy.out fast path for dense, matching slices (PR [CUDA] Only take the split_with_sizes_copy.out fast path for dense, matching slices pytorch/pytorch#203892), committed 2026-11-10T04:51:40Z.
Reachable from the scanned pytorch/pytorch default-branch head 3aad3844fc331821235a833579a71323af3c3eaf (GitHub compare reports behind_by 0), so the change is landed upstream.
Scan window: 2026-11-10T00:00:00Z .. 2026-11-11T00:00:00Z (half-open, UTC).

The commit touched aten/src/ATen/native/cuda/TensorShape.cu and test/test_torch.py, and no XPU file. Its own description of the bug applies verbatim to the XPU kernel: the fast path "accepted any dense input and outputs, so a transposed input, or a channels_last input split into contiguous outputs, produced wrong values. It also used the split dimension's stride as the chunk size, so an arbitrary stride on a size-1 split dimension could cause writes past the output."

Upstream recorded the XPU gap rather than fixing it; its new regression test is annotated:

@skipIfXpu(msg="XPU split_with_sizes_copy.out still uses stride(dim) as the chunk size")
def test_split_with_sizes_copy_out_non_canonical_layouts(self, device):

Observed XPU behaviour

Running that test's seven layout cases against the CPU composite result as the oracle (every case first shown to hold on CPU), 3 of 7 fail on XPU:

case	XPU result
transposed dense input, `torch.arange(12.).reshape(3, 4).t()`, split `[2, 2]` on dim 0	both outputs hold wrong values
contiguous input whose size-1 split dim has stride 16, `torch.zeros(64).as_strided((2, 1, 2), (2, 16, 1))`, split `[1]` on dim 1	writes 48 bytes past the output: the 12-element guard region after it goes from -1.0 to 0.0
channels_last input split into contiguous outputs, `(4, 3, 2, 5)`, split `[2, 2]` on dim 0	both outputs hold wrong values
channels_last -> channels_last, permuted leading dims, interleaved strides, single non-empty split	match CPU

The second case is an out-of-bounds device write, not just a wrong value.

Target-path evidence

torch.split_with_sizes_copy(x_xpu, sizes, dim, out=...) -> split_with_sizes_copy_out_xpu -> at::native::xpu::split_with_sizes_copy_out_xpu_kernel in src/ATen/native/xpu/sycl/TensorShapeKernels.cpp. The reproducer asserts that the input and every output tensor are on xpu and prints out[0].device per case; the log shows xpu:0 for all seven.

In that file:

split_with_sizes_copy_out_xpu_kernel (line 774) selects the byte-copy fast path on self.is_non_overlapping_and_dense() and the same predicate per output (lines 779-783) -- dense but non-canonical layouts are accepted, and nothing checks that each output stores its elements in the same order as the input slice;
get_split_chunk_sizes (line 161) computes split_size * tensor.stride(dim) * element_size, so a size-1 split dim with an arbitrary stride yields an oversized chunk and the copy runs past the destination.

Upstream's fix computes chunk sizes and offsets from dimension sizes instead of stride(dim), and restricts the fast path to dense, order-preserving slices; the equivalent change is needed in TensorShapeKernels.cpp.

Reproducer

import torch

# 1. transposed dense input: wrong values
x = torch.arange(12.0, device="xpu").reshape(3, 4).t()
out = [torch.empty(2, 3, device="xpu") for _ in range(2)]
torch.split_with_sizes_copy(x, [2, 2], dim=0, out=out)
print("case1 ok:", all(torch.equal(g, e) for g, e in zip(out, x.split_with_sizes([2, 2], dim=0))))

# 2. size-1 split dim with stride 16: writes past the output
x = torch.zeros(64, device="xpu").as_strided((2, 1, 2), (2, 16, 1))
x.copy_(torch.arange(4.0, device="xpu").view(2, 1, 2))
guard = torch.full((16,), -1.0, device="xpu")
torch.split_with_sizes_copy(x, [1], dim=1, out=[guard[:4].view(2, 1, 2)])
print("case2 guard[4:]:", guard[4:].tolist())   # must stay all -1.0

# 3. channels_last input into contiguous outputs: wrong values
x = torch.randn(4, 3, 2, 5, device="xpu").to(memory_format=torch.channels_last)
out = [torch.empty(2, 3, 2, 5, device="xpu") for _ in range(2)]
torch.split_with_sizes_copy(x, [2, 2], dim=0, out=out)
print("case3 ok:", all(torch.equal(g, e) for g, e in zip(out, x.split_with_sizes([2, 2], dim=0))))

Relevant output

torch 2.16.0.dev20261110+xpu
xpu_device Intel(R) Arc(TM) Pro B60 Graphics
--- transposed_dense_input ---
  cpu oracle holds
  [xpu] transposed_dense_input: out[0].device=xpu:0
  xpu WRONG: ['out[0] mismatch shape=(2, 3)/(2, 3)', 'out[1] mismatch shape=(2, 3)/(2, 3)']
--- size1_split_dim_stride ---
  cpu oracle holds
  [xpu] size1_split_dim_stride: out[0].device=xpu:0
  xpu WRONG: ['wrote past the output: guard[4:]=[0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0]']
--- channels_last_to_contiguous ---
  cpu oracle holds
  [xpu] channels_last_to_contiguous: out[0].device=xpu:0
  xpu WRONG: ['out[0] mismatch shape=(2, 3, 2, 5)/(2, 3, 2, 5)', 'out[1] mismatch shape=(2, 3, 2, 5)/(2, 3, 2, 5)']

Environment

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

Ownership

The defective fast path is in intel/torch-xpu-ops (src/ATen/native/xpu/sycl/TensorShapeKernels.cpp). Upstream fixed only the CUDA kernel and marked its regression test @skipIfXpu with a message naming this exact XPU gap, so the parity work is not owned upstream. Landing the fix here also lets that upstream skip be removed.

Related trackers

Searched intel/torch-xpu-ops and pytorch/pytorch; no open tracker covers the fast-path layout gate or the stride(dim) chunk size. Related but distinct, both about a zero-length split dividing by zero in the same kernel family rather than about layout:

test_split_with_sizes_copy_out_xpu crashes with Fatal Python error: Floating point exception (SIGFPE) #5661 -- open: test_split_with_sizes_copy_out_xpu crashes with SIGFPE, triggered by a zero-sized split. Fixing it would not fix the layout defect reported here, and vice versa.
[Bug Skip]: split_with_sizes_copy with a zero-length split size crashes the XPU worker process #6020 -- closed: split_with_sizes_copy with a zero-length split size crashes the XPU worker process.

_{alignment scan 2026-11-10, workflow run 42943920909}

[xpu-alignment] [26-11-10] Sparse COO mul kills the XPU process with SIGFPE: the torch-xpu-ops copy of the intersection heuristic divides by prod(shape[:sparse_dim]) accumulated in an int

Upstream source

Issue [Bug][Sparse] Multiplying empty sparse tensors with broadcastable shapes crashes due to integer division by zero pytorch/pytorch#204225 -- [Bug][Sparse] Multiplying empty sparse tensors with broadcastable shapes crashes due to integer division by zero, created 2026-11-10, still open.
Proposed fix [sparse] Avoid division by zero in the sparse intersection coalesce heuristic pytorch/pytorch#204266 -- [sparse] Avoid division ..., open, not merged (mergedAt null as of this review). It changes only aten/src/ATen/native/sparse/SparseBinaryOpIntersectionCommon.h (+3/-2) and test/test_sparse.py.
Scan window: 2026-11-10T00:00:00Z .. 2026-11-11T00:00:00Z (half-open, UTC).

Upstream roots the SIGFPE in the coalesce heuristic larger._nnz() / prod(larger.shape[:sparse_dim]) and names two ways the product reaches 0: a zero-size sparse dim, and -- not in the issue, only in the PR -- that std::accumulate(..., 1, std::multiplies<int64_t>()) keeps its accumulator in the type of the init value, int, so sparse sizes 65536 x 65536 truncate to 0.

Observed XPU behaviour

Both variants kill the XPU process with SIGFPE (child returncode -8), after a control multiplication proves the XPU intersection path works:

zero-size sparse dim. torch.zeros(0, 3, device='xpu').to_sparse(1) * torch.zeros(0, 1, device='xpu').to_sparse(1) dies with SIGFPE. Expected result, per the issue and dense torch.zeros(0, 3) * torch.zeros(0, 1), is an empty sparse tensor of shape (0, 3). CPU dies the same way, which is the open upstream bug.
int accumulator overflow. Two coalesced sparse_coo tensors of size (65536, 65536) with nnz 2 die with SIGFPE on XPU, while CPU returns the correct values [3.0, 8.0]. This asymmetry is XPU-specific: mul_out_sparse_cpu (aten/src/ATen/native/sparse/SparseTensorMath.cpp:1035) sends equal-size, coalesced sparse operands to a brute-force merge and never calls the intersection kernel, whereas mul_sparse_kernel in src/ATen/native/sparse/xpu/sycl/SparseTensorMathKernels.cpp:442 falls through to _mul_sparse_sparse_out for every sparse-sparse case except a zero-nnz short circuit. XPU therefore reaches the faulty division on inputs where CPU does not.

Target-path evidence

Tensor.mul on sparse COO XPU tensors -> mul_out_sparse_xpu -> xpu::mul_sparse_kernel -> _mul_sparse_sparse_out -> mul_sparse_sparse_out_stub(XPU), registered in src/ATen/native/sparse/xpu/SparseBinaryOpIntersection.cpp:16 to xpu::mul_sparse_sparse_kernel in src/ATen/native/sparse/xpu/sycl/SparseBinaryOpIntersectionKernels.cpp.

The control multiplication printed CONTROL device xpu:0 shape (2, 3) values [2.0, 4.0, 6.0, 12.0, 15.0, 18.0] before each crash, so the XPU intersection path executed; the crash then happened after the TARGET about to multiply line and before any result, in host code.

That torch-xpu-ops file carries its own copy of the upstream heuristic, byte-for-byte equivalent to the header at the scanned default-branch head 3aad3844fc331821235a833579a71323af3c3eaf:

// src/ATen/native/sparse/xpu/sycl/SparseBinaryOpIntersectionKernels.cpp:409
const auto sparse_dim_numel = std::accumulate(
    larger_sizes.begin(),
    larger_sizes.begin() + larger.sparse_dim(),
    1,                                  // int accumulator: 65536 * 65536 -> 0
    std::multiplies<int64_t>());
...
const auto max_count_lower_bound = larger._nnz() / sparse_dim_numel;   // SIGFPE

Because the XPU path never enters the upstream header, PR #204266 landing will not fix XPU. The same two-line change (accumulate in int64_t, skip the division when the product is 0) has to be applied to this copy.

Reproducer

Each variant aborts the process, so run them separately.

# variant 1: zero-size sparse dim (the upstream repro, on XPU)
import torch
x = torch.zeros(0, 3, device="xpu").to_sparse(sparse_dim=1)
y = torch.zeros(0, 1, device="xpu").to_sparse(sparse_dim=1)
z = x * y                      # SIGFPE; expected an empty sparse tensor of shape (0, 3)
print(z.device, tuple(z.shape), z._nnz())

# variant 2: prod(shape[:sparse_dim]) == 2**32 truncates to 0 in an int accumulator
import torch
idx = torch.tensor([[0, 1], [0, 1]], device="xpu")
x = torch.sparse_coo_tensor(idx, torch.tensor([1.0, 2.0], device="xpu"), (65536, 65536)).coalesce()
y = torch.sparse_coo_tensor(idx, torch.tensor([3.0, 4.0], device="xpu"), (65536, 65536)).coalesce()
z = (x * y).coalesce()         # SIGFPE on XPU; CPU returns values [3.0, 8.0]
print(z.device, tuple(z.shape), z._nnz(), z.values().cpu().tolist())

Relevant output

torch 2.16.0.dev20261110+xpu
xpu_device Intel(R) Arc(TM) Pro B60 Graphics
dense_reference_shape (0, 3)
=== variant zero_size_sparse_dim ===
  [cpu] returncode=-8
  [xpu] returncode=-8
  [xpu] stdout: CONTROL device xpu:0 shape (2, 3) values [2.0, 4.0, 6.0, 12.0, 15.0, 18.0]
                TARGET about to multiply (0, 3) (0, 1)
=== variant int32_overflow ===
  [cpu] returncode=0
  [cpu] stdout: TARGET ok device cpu shape (65536, 65536) nnz 2 values [3.0, 8.0]
  [xpu] returncode=-8
  [xpu] stdout: CONTROL device xpu:0 shape (2, 3) values [2.0, 4.0, 6.0, 12.0, 15.0, 18.0]
                TARGET about to multiply (65536, 65536) (65536, 65536)

(-8 is SIGFPE.)

Environment

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

Ownership

The faulting division is in intel/torch-xpu-ops. The upstream fix is still only proposed (PR #204266, open) and covers only the shared header, which the XPU dispatch path does not use, so this is independent XPU work whose fix does not depend on that PR landing. Worth fixing even ahead of it: the int-overflow variant crashes on XPU for inputs CPU handles correctly today.

Related trackers

Searched intel/torch-xpu-ops and pytorch/pytorch; no open or closed issue covers this XPU crash. Related but distinct:

[BMG] sparse mul hangs > 10 min for complex128/float64 (sub-issue of #4947) #5928 -- open: sparse mul hangs for complex128/float64 on BMG (performance, not a crash).
[xpu-alignment] [26-10-06] torch._masked_softmax kills the XPU process with SIGFPE on a zero-size softmax dim: masked_softmax_kernel has no numel() == 0 guard #5611 -- open: another missing zero-size guard causing SIGFPE on XPU (_masked_softmax), same class of defect in a different kernel.
[xpu-alignment] [26-09-13] torch.add(dense, dense, out=sparse_csr) aborts the process on XPU instead of raising #5343 -- open: torch.add(dense, dense, out=sparse_csr) aborts on XPU.

_{alignment scan 2026-11-10, workflow run 42943920909}

[xpu-alignment] [26-11-10] XPU deterministic reflection_pad2d backward drops the right-zone gradient for columns that are in both reflection zones: max|grad - CPU| = 3.97 for pad=(3, 2, 4, 3)

Upstream source

PR Fix deterministic CUDA reflection_pad2d backward for large pads pytorch/pytorch#204276 -- Fix deterministic CUDA reflection_pad2d backward for large pads, open, not merged (mergedAt null as of this review). Fixes use_deterministic_algorithms(True): CUDA reflection_pad2d backward returns wrong gradients for large pads pytorch/pytorch#203421 (open).
Scan window: 2026-11-10T00:00:00Z .. 2026-11-11T00:00:00Z (half-open, UTC).

The PR touches aten/src/ATen/native/cuda/ReflectionPad.cu (+5/-2) and test/test_nn.py, and no XPU file. Its root cause: the top-right and bottom-right corner terms sat in else if (is_right) branches after if (is_left), so they were skipped whenever a column is in both the left and the right reflection zone (pad_left + pad_right >= width - 1).

Observed XPU behaviour

Under torch.use_deterministic_algorithms(True), with the PR's own test shape -- a (2, 3, 5, 4) float64 input and pad = (3, 2, 4, 3) -- the XPU gradient disagrees with the CPU gradient computed from the identical input and the identical grad_output:

arm	non-overlapping control `pad=(1, 1, 1, 1)`	target `pad=(3, 2, 4, 3)`
float64	matches CPU, max abs diff 4.4e-16	mismatch, max abs diff 3.967
float32	matches CPU, max abs diff 2.4e-07	mismatch, max abs diff 3.450

The control isolates the defect to the overlapping zone: the kernel is correct when no column is in both zones. Gradients are silently wrong, with no error raised. For this shape input_dim_x = 4, pad_left = 3, pad_right = 2, so is_left holds for columns 1..3 and is_right for columns 1..2 -- columns 1 and 2 are in both and lose their right-zone contribution.

Target-path evidence

F.pad(x_xpu, pad, mode='reflect').backward(...) with determinism enabled -> reflection_pad2d_backward_det_template (selected at src/ATen/native/xpu/sycl/ReflectionPadKernels.cpp:842 by at::globalContext().deterministicAlgorithms()) -> reflection_pad2d_backward_det_kernel (line 310). The reproducer enables determinism, asserts the forward output and x.grad are on xpu, prints both devices (xpu:0), and calls torch.xpu.synchronize() before reading the gradient.

The kernel, under a comment at line 307 saying it "Mirrors the CUDA deterministic implementation", carries exactly the structure the PR replaces:

// src/ATen/native/xpu/sycl/ReflectionPadKernels.cpp:367
if (is_left) {   // top-left corner
  accum_grad_at(pad_top - inp_row, pad_left - inp_col);
} else if (is_right) {   // top-right corner
  accum_grad_at(pad_top - inp_row, pad_left + rightmost_col + dist_from_right);
}
...
// line 378
if (is_left) {   // bottom-left corner
  accum_grad_at(pad_top + bottom_row + dist_from_bottom, pad_left - inp_col);
} else if (is_right) {   // bottom-right corner
  accum_grad_at(pad_top + bottom_row + dist_from_bottom,
                pad_left + rightmost_col + dist_from_right);
}

The fix is the same as upstream's: make the two is_right blocks unconditional if (is_right) so both corner contributions accumulate. Note the edge blocks further down (lines 388-395) already use independent ifs, which is why only the corner terms are lost.

Reproducer

import torch
import torch.nn.functional as F

torch.use_deterministic_algorithms(True)

for pad in ((1, 1, 1, 1), (3, 2, 4, 3)):
    torch.manual_seed(0)
    base = torch.randn(2, 3, 5, 4, dtype=torch.float64)
    x = base.to("xpu").detach().requires_grad_()
    ref = base.clone().detach().requires_grad_()
    out = F.pad(x, pad, mode="reflect")
    g = torch.randn(*out.shape, dtype=torch.float64)      # one grad_output for both arms
    out.backward(g.to("xpu"))
    torch.xpu.synchronize()
    F.pad(ref, pad, mode="reflect").backward(g)
    print(pad, x.grad.device, "maxdiff", (x.grad.cpu() - ref.grad).abs().max().item())

# (1, 1, 1, 1) xpu:0 maxdiff 4.44e-16
# (3, 2, 4, 3) xpu:0 maxdiff 3.967

Relevant output

torch 2.16.0.dev20261110+xpu
xpu_device Intel(R) Arc(TM) Pro B60 Graphics
deterministic_algorithms_enabled True
--- float64 arm (upstream dtype) ---
  target out.device=xpu:0 grad.device=xpu:0
  control_non_overlapping: close=True maxdiff=4.441e-16
  target out.device=xpu:0 grad.device=xpu:0
  target_overlapping: close=False maxdiff=3.967e+00
--- float32 arm (secondary) ---
  control_non_overlapping: close=True maxdiff=2.384e-07
  target_overlapping: close=False maxdiff=3.450e+00

Environment

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

Ownership

The defect is in intel/torch-xpu-ops (src/ATen/native/xpu/sycl/ReflectionPadKernels.cpp) and is established here against the CPU gradient, independently of the upstream PR: the XPU deterministic backward is wrong today and the fix does not depend on PR #204276 landing. Upstream's fix is still at the proposed stage and touches only the CUDA kernel, so the XPU work is not owned upstream. The CUDA PR is worth tracking only as a reference for the shape of the change and for its regression test test_ReflectionPad2d_deterministic_overlapping_pads, which is device-generic and should pass on XPU once this is fixed.

Related trackers

Searched intel/torch-xpu-ops and pytorch/pytorch; no open or closed issue covers the XPU deterministic reflection_pad2d backward. Related but distinct:

[upstream_ut] XPU backward kernels missing input-validation error checks (replication_pad2d, grid_sample) #5508 -- open: missing input validation in XPU backward kernels (replication_pad2d, grid_sample), a validation gap rather than a wrong gradient.
[xpu-alignment] [26-11-02] XPU index_add_ on bool tensors raises under use_deterministic_algorithms(True): the deterministic branch promotes the source with source * alpha before the sort-based index_put_ #6002 -- open: another XPU defect specific to the use_deterministic_algorithms(True) branch (index_add_ on bool).

@laifenxiawucha

XPU alignment run completed with warnings

Scan date: 2026-11-10
Collection: complete
Reviewed units: 4 (needs-xpu-fix: 4)
Review-approved tracker candidates: 4

Formal candidate drafts:

commit-500c6606e46ae107a69deb006bde5e90d578660b — [xpu-alignment] [26-11-10] XPU multinomial(replacement=True) builds its CDF in the input dtype: bfloat16 resolves 847 of 10000 draws and a half row summing past 65504 returns index 0 for every draw
commit-b0c56ca7a17bd5dd2c52dd724f9c18ac46a3bde5 — [xpu-alignment] [26-11-10] XPU split_with_sizes_copy.out takes its byte-copy fast path for any dense input: wrong values for transposed and channels_last inputs, and a write past the output when the size-1 split dim has a large stride
issue-204225 — [xpu-alignment] [26-11-10] Sparse COO mul kills the XPU process with SIGFPE: the torch-xpu-ops copy of the intersection heuristic divides by prod(shape[:sparse_dim]) accumulated in an int
pr-204276 — [xpu-alignment] [26-11-10] XPU deterministic reflection_pad2d backward drops the right-zone gradient for columns that are in both reflection zones: max|grad - CPU| = 3.97 for pad=(3, 2, 4, 3)

Automatic filing is limited to 3 candidates; review these drafts and create trackers manually.

Excluded units:
- commit-fb6d70c154a6e8a58ff72e907c7f9959650c827d — blocked-env
- issue-204215 — blocked-env

_{workflow run 42943920909}

@laifenxiawucha @ZhaoqiongZ

Filed as #6162
_{alignment scan 2026-11-11, workflow run 43139444140}

[xpu-alignment] [26-11-11] XPU grid_sampler_2d backward accumulates the fp16/bf16 grid gradient in the input dtype: 2048 unit contributions sum to 1024 (fp16) and 128 (bf16)

Summary

grid_sampler_2d_backward on XPU keeps its grid-gradient accumulator in the input dtype, so for float16 and bfloat16 the per-channel sum stops growing once each tap falls below half a unit in the last place of the running sum. With 2048 channels that each contribute exactly 1 to the x-gradient and a coordinate scale of 1, XPU returns [1024.0, 0.0] for float16 and [128.0, 0.0] for bfloat16, where the exact answer [2048.0, 0.0] is representable in both dtypes. float32 on the same device is exact, so this is accumulator loss, not representation error, and it is silent: no warning, no error.

The practical effect is that half-precision F.grid_sample backward reports grid gradients that are too small by up to 16x for wide channel counts, which quietly stalls learning of the sampling grid (spatial transformers, optical-flow warps, deformable sampling).

Upstream source

[CUDA] grid_sampler_2d backward: FP16/BF16 grad_grid accumulates in the input dtype pytorch/pytorch#204496, "[CUDA] grid_sampler_2d backward: FP16/BF16 grad_grid accumulates in the input dtype", created 2026-11-11, open and labelled module: correctness (silent), module: numerical-stability, module: half, module: bfloat16, module: interpolation. The values were measured on ROCm (MI355X); the CUDA behaviour is derived from source and was not executed.
[ROCm] grid_sampler_2d backward: accumulate FP16/BF16 grid gradient in FP32 (changes values) pytorch/pytorch#204495, "[ROCm] grid_sampler_2d backward: accumulate FP16/BF16 grid gradient in FP32 (changes values)", is the proposed fix. It is open and not merged (re-checked 2026-11-12). It touches only aten/src/ATen/native/cuda/GridSampler.cu, aten/src/ATen/native/cuda/GridSampler.cuh, docs/source/notes/hip.rst and test/test_nn.py; it is deliberately scoped to the ROCm build and leaves the CUDA build byte-identical. It does not touch intel/torch-xpu-ops, so it cannot cover XPU, and all three tests it adds are @onlyCUDA, so XPU CI will not surface this gap either.
[CUDA][ROCm] grid_sampler_2d: keep 2D spatial offsets in the kernel's index type pytorch/pytorch#204493 and [ROCm] grid_sampler_2d backward: channel-parallel kernel, default on gfx1250 pytorch/pytorch#198032 are the other open members of the same grid-sampler series (32-bit offset overflow, channel-parallel ROCm kernel). Neither addresses the accumulator width and neither touches XPU.
Found by an XPU alignment scan of the half-open window [2026-11-11T00:00:00Z, 2026-11-12T00:00:00Z) over pytorch/pytorch, frozen default-branch head 125472c0cbd6259715738ed1a80cd0240f914ad1.

Observed XPU behavior

float16   mask=[False, True] device=xpu:0 grad_grid=[1024.0, 0.0] exact=False
float16   mask=[True, True]  device=xpu:0 grad_grid=[1024.0, 0.0] exact=False
bfloat16  mask=[False, True] device=xpu:0 grad_grid=[128.0, 0.0] exact=False
bfloat16  mask=[True, True]  device=xpu:0 grad_grid=[128.0, 0.0] exact=False
float32   mask=[False, True] device=xpu:0 grad_grid=[2048.0, 0.0] exact=True
float32   mask=[True, True]  device=xpu:0 grad_grid=[2048.0, 0.0] exact=True

The float16 and bfloat16 rows are the exact values upstream reports for ROCm. Both output masks are affected, which places the loss in the grid-gradient accumulation itself rather than in the grad_input scatter (the [False, True] mask skips the scatter entirely).

Why [2048.0, 0.0] is the exact answer: with align_corners=False and inp_H = inp_W = 2, a grid coordinate of (0, 0) unnormalizes to ix = iy = 0.5, so all four bilinear weights are 0.25 and gix_mult = inp_W / 2 = 1. Per channel the x-gradient is -nw_val * (iy_se - iy) + ne_val * (iy_sw - iy) - sw_val * (iy - iy_ne) + se_val * (iy - iy_nw), which for the input column pattern [0, 1] is -0 * 0.5 + 1 * 0.5 - 0 * 0.5 + 1 * 0.5 = 1; over 2048 channels that is exactly 2048. The y-gradient cancels to 0. Because each channel adds two 0.5 taps, float16 stalls at 1024 (there the ULP is 1.0, so +0.5 rounds back to even) and bfloat16 stalls at 128 for the same reason.

Target path and root cause

aten/src/ATen/native/native_functions.yaml at the frozen head dispatches grid_sampler_2d_backward to XPU: grid_sampler_2d_backward_xpu, which this repository implements, so the call reaches XPU code and not a CPU fallback. Target-path evidence in the run: every result row reports device=xpu:0, torch.xpu.synchronize() is called before each value is read, and the float32 control on the same device and same kernel is exact, which rules out a harness or host/device artifact.

The 2-D bilinear backward kernel declares the grid-gradient accumulator in scalar_t and never widens it:

// src/ATen/native/xpu/sycl/GridSampler.cpp:365
  scalar_t gix_mult, giy_mult;

// src/ATen/native/xpu/sycl/GridSampler.cpp:390
    scalar_t gix = static_cast<scalar_t>(0), giy = static_cast<scalar_t>(0);

// src/ATen/native/xpu/sycl/GridSampler.cpp:448 (inside the per-channel loop)
        gix -= nw_val * (iy_se - iy) * gOut;

// src/ATen/native/xpu/sycl/GridSampler.cpp:473-474
    gGrid_ptr_NHW[0] = gix_mult * gix;
    gGrid_ptr_NHW[1] = giy_mult * giy;

The loop at src/ATen/native/xpu/sycl/GridSampler.cpp:395 walks all C channels and rounds the running sum back to scalar_t on every tap, so once the sum reaches 2^10 (float16) or 2^7 (bfloat16) relative to the tap size it stops advancing. The coordinate scale gix_mult is applied only after the loop, so widening it alone would not help. The forward kernel in the same file already computes in at::opmath_type<scalar_t> (src/ATen/native/xpu/sycl/GridSampler.cpp:65, 104), which is why only the backward grid gradient is wrong.

Two sibling kernels in the same file have the same narrow accumulator and are not separately reproduced here:

2-D bicubic backward: src/ATen/native/xpu/sycl/GridSampler.cpp:531-532, stored at 583-584.
3-D bilinear backward: src/ATen/native/xpu/sycl/GridSampler.cpp:1299-1300, stored at 1485-1487.

The 3-D bicubic backward kernel already accumulates in opmath_t and casts once at the store (src/ATen/native/xpu/sycl/GridSampler.cpp:1619-1621, 1665-1667), so the intended style is already present in this file.

Suggested fix: accumulate gix/giy (and giz) in at::opmath_type<scalar_t>, apply the coordinate multiplier in that type, and cast once when writing gGrid_ptr_*, matching lines 1619-1667. Upstream #204495 additionally keeps a per-channel partial for reduced precision to limit intermediate overflow and cancellation; that refinement is worth reviewing, but the accumulator width is the part that produces the numbers above.

Reproducer

import torch

for dtype in (torch.float16, torch.bfloat16, torch.float32):
    inp = (
        torch.tensor([[0.0, 1.0], [0.0, 1.0]])
        .expand(1, 2048, 2, 2)
        .to("xpu", dtype)
        .contiguous()
    )
    grad_output = torch.ones(1, 2048, 1, 1, device="xpu", dtype=dtype)
    grid = torch.zeros(1, 1, 1, 2, device="xpu", dtype=dtype)
    for mask in ([False, True], [True, True]):
        grad_grid = torch.ops.aten.grid_sampler_2d_backward(
            grad_output, inp, grid, 0, 0, False, mask
        )[1]
        torch.xpu.synchronize()
        print(dtype, mask, grad_grid.float().flatten().tolist(), "expected [2048.0, 0.0]")

This is the upstream snippet with device="cuda" changed to "xpu"; shapes, dtypes, the expand/contiguous layout, both output masks and the direct torch.ops.aten.grid_sampler_2d_backward call are unchanged, and the float32 row is an added control.

Environment

torch_version:   2.16.0.dev20261111+xpu
torch_path:      /opt/conda/lib/python3.13/site-packages/torch/__init__.py
python:          3.13.13 (/opt/conda/bin/python3.13)
xpu_available:   True
xpu_device_name: Intel(R) Arc(TM) Pro B60 Graphics

The reproducer ran unmodified as python3.13 -I -u with no network access and exited 0 in 3.5 s.

Ownership

The XPU kernel is independent code in this repository: grid_sampler_2d_backward_xpu (src/ATen/native/xpu/sycl/GridSampler.cpp:676) is dispatched directly from native_functions.yaml and shares nothing with aten/src/ATen/native/cuda/GridSampler.cuh. The only upstream fix in flight (#204495) is scoped to the ROCm build, leaves CUDA unchanged, and does not modify this repository, so no upstream change will fix XPU. Searches of pytorch/pytorch for XPU-scoped grid-sampler issues or PRs returned nothing.

No open intel/torch-xpu-ops issue covers this work. The closest existing items are different defects:

[upstream_ut] XPU backward kernels missing input-validation error checks (replication_pad2d, grid_sample) #5508 (open) - XPU backward kernels missing input-validation error checks, including grid_sample; an error-checking gap, not accumulator width.
grid_sample: implement mode='bicubic' with 5-D input on XPU #5542 and [release/2.15] NotImplementedError: grid_sampler(): bicubic interpolation with 5D input is not implemented for xpu #5452 (both closed) - grid_sample bicubic with 5-D input on XPU; a feature gap, unrelated to precision.

Related accumulator-width parity work already tracked separately: #6127 (segment_reduce fp16/bf16 accumulation) and #5722 (nll_loss2d fp16 global atomics). This issue is the same class of defect in a different kernel and does not overlap their code.

Suggested verification once fixed: the reproducer above must print [2048.0, 0.0] for float16 and bfloat16; upstream's test_grid_sample_2d_backward_low_precision_grid_grad (added by #204495, currently @onlyCUDA) is the natural regression test to mirror for XPU.

@laifenxiawucha

XPU alignment run completed with warnings

Scan date: 2026-11-11
Collection: complete
Reviewed units: 6 (duplicate: 2, needs-xpu-fix: 1, track-upstream: 3)
Review-approved tracker candidates: 1

Automatically filed:

issue-204496 — [xpu-alignment] [26-11-11] XPU grid_sampler_2d backward accumulates the fp16/bf16 grid gradient in the input dtype: 2048 unit contributions sum to 1024 (fp16) and 128 (bf16) #6162
Excluded units:
- issue-204446 — blocked-script-error

_{workflow run 43139444140}

@laifenxiawucha @ZhaoqiongZ

ALIGNMENT_TRIAGE_ISSUE #5018

Description

XPU alignment triage

Activity

github-actions commented on Aug 26, 2026

[DRY RUN] [xpu-alignment] embedding_bag per_sample_weights backward indexes weight/grad with unchecked, int-truncated saved indices on XPU

Summary

Upstream source

Target-path evidence

Observed XPU behavior

Reproducer

Observed output

Scope of the evidence

Suggested fix

Environment

Provenance

github-actions commented on Aug 26, 2026

[DRY RUN] [xpu-alignment] XPU group_norm rejects fp32 affine params over an fp16 input, which the shared ATen mixed-dtype contract supports

Summary

Upstream source

Why this is a defect and not unsupported input

Root cause and target-path evidence

Reproducer

Observed output

Relationship to CUDA and to PR #194616

Scope of the evidence

Suggested fix

Environment

Provenance

github-actions commented on Aug 26, 2026

github-actions commented on Aug 27, 2026

github-actions commented on Aug 27, 2026

github-actions commented on Aug 29, 2026

[xpu-alignment] embedding_renorm_ derives the row width from stride(0) and faults on a zero-width embedding

Upstream source

Observed XPU behavior

Root cause and target-path evidence

Reproducer

Environment

Ownership

github-actions commented on Aug 29, 2026

[xpu-alignment] exponential_ underflows to exactly 0 for float16 on XPU, outside the distribution support

Upstream source

Observed XPU behavior

Root cause and target-path evidence

Reproducer

Environment

Ownership

github-actions commented on Aug 29, 2026

[xpu-alignment] silu(-inf) returns NaN on XPU instead of the limit 0

Upstream source

Observed XPU behavior

Root cause and target-path evidence

Reproducer

Environment

Ownership

github-actions commented on Aug 29, 2026

github-actions commented on Aug 30, 2026

121 remaining items

torchxpubot commented on Nov 8, 2026

torchxpubot commented on Nov 9, 2026

torchxpubot commented on Nov 10, 2026

[xpu-alignment] [26-11-09] XPU segment_reduce accumulates fp16/bf16 in the input dtype: 4096 ones sum to 2048 (fp16) and 256 (bf16)

Summary

Upstream source

Observed XPU behavior

Target path and root cause

Reproducer

Environment

Ownership

Regression test

torchxpubot commented on Nov 10, 2026

[xpu-alignment] [26-11-09] XPU Tensor.renorm divides by norm + 1e-7, so renormalised rows land below max_norm (33% short at max_norm=1e-7)

Summary

Upstream source

Observed XPU behavior

Target path and root cause

Reproducer

Environment

Ownership