[release/2.14] Build rocSHMEM for the requested gfx target, not gfx90a xnack variants by lucbruni-amd · Pull Request #3583 · ROCm/pytorch

lucbruni-amd · 2026-08-19T15:39:45Z

Drops the gfx90a→xnack± special case; builds torch_rocshmem for the requested target like every other arch. rocSHMEM ships xnack-agnostic device bitcode (librocshmem_device_gfx90a.bc, target-cpu=gfx90a), and the mapping matched per-xnack archives that ROCm/rocm-systems#4822 removed in April.

Fixes the split torch_gfx90a:xnack±.kpack stubs that break torch on MI250 (ROCm/TheRock#7081, ROCM-29326).

Scope: gfx90a only, minimal. The list(FIND ...) membership test is untouched; it needs bare-arch matching once ASan enables xnack uniformly (ROCm/TheRock#6624), tracked separately.

## Problem - PyTorch trunk ROCm CI still builds against ROCm 7.14. - ROCm 10.0 packages are published from the new `whl-next` index. - rocSHMEM must build for the requested gfx target instead of substituting gfx90a XNACK variants. ## Changes - Upgrade the shared `rocm-n` trunk image from ROCm 7.14 to 10.0. - Install ROCm 10.0 from `https://stable.repo.amd.com/rocm/whl-next/`. - Cherry-pick the rocSHMEM target fix from ROCm#3583. - Keep wheel, builder, and nightly changes independent in pytorch#194919. ## Validation - Bash syntax and stubbed trunk Docker argument routing pass. - The trunk image selects ROCm 10.0 and the `whl-next` index; the preview image retains its nightly index. - `git diff --check` and targeted IDE lint checks pass for the Docker and CMake changes. - Full image and rocSHMEM builds were not run locally; CI provides end-to-end validation. - [trunk workflow](https://github.com/pytorch/pytorch/actions/runs/33000092249/job/98385742860) ran successfully with ROCm10.0 - [rocm-mi200](https://github.com/pytorch/pytorch/actions/runs/33002840829/job/98340667528) ran successfully with ROCm10.0 Pull Request resolved: pytorch#194794 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily

…a xnack variants (#3583) Drops the gfx90a→xnack± special case; builds `torch_rocshmem` for the requested target like every other arch. rocSHMEM ships xnack-agnostic device bitcode (`librocshmem_device_gfx90a.bc`, `target-cpu=gfx90a`), and the mapping matched per-xnack archives that ROCm/rocm-systems#4822 removed in April. Fixes the split `torch_gfx90a:xnack±.kpack` stubs that break torch on MI250 (ROCm/TheRock#7081, ROCM-29326). Scope: gfx90a only, minimal. The `list(FIND ...)` membership test is untouched; it needs bare-arch matching once ASan enables xnack uniformly (ROCm/TheRock#6624), tracked separately.

* [Testcase Refactoring] Make test/test_module_tracker.py device-agnostic (#191838) ## Summary Register `TestModuleTracker` (test/test_module_tracker.py) with `instantiate_device_type_tests` instead of running as a plain CPU-only `unittest.TestCase`, so the suite exercises `torch.utils.module_tracker.ModuleTracker`'s hook-based forward/backward hierarchy tracking on every available device type — including MPS and XPU — and tags the class so it participates correctly in `--hw-classification` filtering (#186918). ## Changes **`test/test_module_tracker.py`**: - Added a `device` parameter to `test_module_hierarchy`, `test_confused_hierarchy`, and `test_bw_detection` - Threaded `device=device` through tensor construction and moved the `nn.Module` instances (`Mod`, `MyMod`, `nn.Linear`) onto `device` via `.to(device)` - Registered the class with `instantiate_device_type_tests(TestModuleTracker, globals(), allow_mps=True, allow_xpu=True)`, replacing implicit CPU-only execution - Added `hw_classification = HardwareClassification.GENERIC` to `TestModuleTracker` - No assertion values or control flow changed ## Motivation `ModuleTracker`'s hook-registration/ordering logic is pure Python bookkeeping — every assertion is on a `set`/`list`/`bool` value describing which modules ran and in what order, not on tensor numerics — so there's no reason it should be restricted to a single implicit device. The file had no device-generic test infrastructure at all before this change. `hw_classification = HardwareClassification.GENERIC` follows #186918's `HardwareClassification` enum (added to `common_utils.py` so test classes can opt into `--hw-classification` filtering). ## Test Plan python test/test_module_tracker.py -v python test/test_module_tracker.py --hw-classification GENERIC -v python test/test_module_tracker.py --hw-classification ACCELERATOR -v 6/6 tests pass (CPU + MPS). Pull Request resolved: https://github.com/pytorch/pytorch/pull/191838 Approved by: https://github.com/Niran814804102, https://github.com/fffrog, https://github.com/albanD * Assign the tensor type prototype instead of memcpy (#194590) py_initialize_tensor_type() copies a prototype PyTypeObject into the type it is initializing. Both sides are PyTypeObject, so this is a plain copy, not type punning, and assignment expresses it directly. It also drops the hand-written `sizeof`: with memcpy, changing the parameter's type would silently copy the wrong number of bytes, while assignment is a compile error. PyTypeObject is trivially copyable and copy-assignable, so the generated code is the same. Test Plan: ``` static_assert(std::is_copy_assignable_v<PyTypeObject>); static_assert(std::is_trivially_copyable_v<PyTypeObject>); ``` both hold, and the file compiles with -fsyntax-only against the CPython headers. This PR description was drafted with AI assistance. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194590 Approved by: https://github.com/albanD * Support torch.xpu.Event in torch.multiprocessing (#192514) # Motivation As title # Additional Context <img width="1349" height="298" alt="image" src="https://github.com/user-attachments/assets/729accca-951f-4e84-8ce8-aca64ee2f988" /> Pull Request resolved: https://github.com/pytorch/pytorch/pull/192514 Approved by: https://github.com/albanD ghstack dependencies: #192495 * Report a child's KeyboardInterrupt instead of exiting 0 in spawn (#194439) _wrap_ swallowed every `KeyboardInterrupt` (`pass` -> exit 0). That is only correct for pdeathsig (Linux parent death sends SIGINT); a `KeyboardInterrupt` raised by `fn` or sent to the child alone made a dead child look successful to `join()`. Now the child exits quietly only when the parent is gone (`getppid()` changed); otherwise the traceback goes to the error file like any other failure. Terminal-wide Ctrl+C is unaffected (the parent itself raises from `join()`). Test: `test_spawn_child_keyboard_interrupt`. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194439 Approved by: https://github.com/d4l3k * [inductor] Don't match pointless_cumsum_replacement on a symbolic fill (#193959) ## Issue Fixes #193876 ## Summary `pointless_cumsum_replacement` folds `full(shape, fill).cumsum(dim)` into `arange(1, n + 1) * fill`, which assumes `fill` is a Python constant. A recompile with a different closure constant, or an unbacked `.item()` fill, passes it in as a symbolic scalar, so `fill_value` binds to an fx `Node` and the fold dies: `int()` on a `Node` for an integral dtype, or `FakeTensor * Node` for a floating one. Eager is fine. The guard declines those matches so `cumsum` lowers normally, extending the scalar-shape check from #193422. A boolean `full` is exempt: the replacement collapses a bool fill with `int(bool(fill_value))` before any arithmetic, so bool still folds. Tested on `int64` and `float32`. Caveat: the periodic A100 `inductor_huggingface_unbacked_parity` job does not run on PRs, and this declines the fold when the fill is unbacked, so PR CI carries no signal there. ## Checklist - [x] Passes lint (`spin fixlint`) - [x] Added/updated tests - [ ] Updated documentation (if applicable) - [ ] Included benchmark results (for PRs impacting perf) ## BC-breaking? No. --- AI assistance was used in developing this change: reproduced the crash and test validation Pull Request resolved: https://github.com/pytorch/pytorch/pull/193959 Approved by: https://github.com/aorenste * Don't mutate the caller's annotation dict in mark_stream (#194851) Human note: backward annotation projection and cupti backend for getting annotations were merged approximately simultaneously, and cupti backend didn't correctly handle backward annotations - all the backward node stayed empty with cupti backend. mark_stream wrote the resolved stream id into the annotation dict it was handed. The annotation is then stored by reference in the kernel annotation registry, so the write had two effects the caller never asked for: the dict gained a "stream" key that outlived the call, and a dict reused across several mark_stream calls (a comms wrapper marking one lane per process group is the motivating case) ended up shared by every region recorded with it, so the last lane written overwrote all the earlier ones and every marked region rendered on a single lane. Build the tagged annotation as a copy instead, which is what the string input path already effectively did. Test Plan: ``` LD_LIBRARY_PATH=/usr/local/cuda-13.2/compat python test/test_cuda_graph_utils.py -k mark_stream ``` The new test marks two lanes with one shared dict; it fails on the parent commit (the caller's dict gains "stream", and both regions resolve to the second lane) and passes here. Authored with assistance from Claude Code. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194851 Approved by: https://github.com/drisspg * Fix grammatical errors in comments and docstrings (#194565) Fixes subject-verb agreement, missing articles, missing possessive apostrophes, and a few dropped or duplicated words across comments and docstrings in torch/. No functional changes. One fix is more than a pure wording change: the embedding sharding docstring in chunk_sharding_spec_ops/embedding.py contained a garbled sentence ("mask IDs whose embeddings are not stored in current rank will to an extra row will ensure ..."). The stray "will" is dropped so the sentence parses; the intended meaning is preserved. Test Plan: comment-only change, no tests run. Lint via: ``` lintrunner -a ``` Authored with the assistance of an AI agent. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194565 Approved by: https://github.com/albanD * Implement IValue::to without template specializations (#194769) Currently IValue::to is implemented using a combination of macros and template specializations. The implementation is not all in one place which makes it complicated and fragile. This commit is an attempt to reduce this complexity by replacing template specializations with if-constexpr. It also removes the need for _guarded_unsigned_long because the if-else chain is okay with multiple true conditions unlike template specializations. This change was motivated by: - https://github.com/pytorch/pytorch/pull/194659#pullrequestreview-5013557465 - https://github.com/pytorch/pytorch/pull/194659#issuecomment-5403190877 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194769 Approved by: https://github.com/cyyever, https://github.com/wenxin0319, https://github.com/eqy * [ROCm][CI] Upgrade trunk to ROCm 10.0 (#194794) ## Problem - PyTorch trunk ROCm CI still builds against ROCm 7.14. - ROCm 10.0 packages are published from the new `whl-next` index. - rocSHMEM must build for the requested gfx target instead of substituting gfx90a XNACK variants. ## Changes - Upgrade the shared `rocm-n` trunk image from ROCm 7.14 to 10.0. - Install ROCm 10.0 from `https://stable.repo.amd.com/rocm/whl-next/`. - Cherry-pick the rocSHMEM target fix from https://github.com/ROCm/pytorch/pull/3583. - Keep wheel, builder, and nightly changes independent in https://github.com/pytorch/pytorch/pull/194919. ## Validation - Bash syntax and stubbed trunk Docker argument routing pass. - The trunk image selects ROCm 10.0 and the `whl-next` index; the preview image retains its nightly index. - `git diff --check` and targeted IDE lint checks pass for the Docker and CMake changes. - Full image and rocSHMEM builds were not run locally; CI provides end-to-end validation. - [trunk workflow](https://github.com/pytorch/pytorch/actions/runs/33000092249/job/98385742860) ran successfully with ROCm10.0 - [rocm-mi200](https://github.com/pytorch/pytorch/actions/runs/33002840829/job/98340667528) ran successfully with ROCm10.0 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194794 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily * [torchcomms hash update] update the pinned torchcomms hash (#194671) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Update the pinned torchcomms hash. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194671 Approved by: https://github.com/pytorchbot * [ROCm][CI] Add ROCm 10.0 wheel and builder support (#194919) ## Motivation Add ROCm 10.0 support to nightly wheels and almalinux/manylinux docker images. ## Changes Binary matrix and image tags: - `ROCM_ARCHES` moves from `["7.2", "7.14"]` to `["7.14", "10.0"]`, with a matching `rocm[libraries,device-all]==10.0.*` entry in `PYTORCH_EXTRA_INSTALL_REQUIREMENTS`. - `build-manywheel-images.yml` and `build-almalinux-images.yml` build `rocm7.14` and `rocm10.0` instead of `rocm7.1`/`rocm7.2`. - `generated-linux-binary-manywheel-nightly.yml` regenerated to match. Single wheel-based builder path: - `manywheel/build.sh` installs every ROCm version from a TheRock wheel index. ROCm 7.14 keeps `repo.amd.com/rocm/whl-multi-arch/`; 10.0 and anything later default to `stable.repo.amd.com/rocm/whl-next/`, so new versions need no script change. The `rocm_final_legacy` target, the gfx900/gfx906 additions and the 7.1/7.2 patch-version fixups are gone. - `almalinux/build.sh` selects the index the same way. - `manywheel/Dockerfile_2_28` drops the `rocm_final_legacy` stage. - `almalinux/Dockerfile` collapses `rocm_base` and `rocm` into one `rocm` stage that installs the SDK with `RUN bash ./install_rocm.sh && rm install_rocm.sh`, consistent with `Dockerfile_2_28`. It sets `BASH_ENV=/etc/rocm_env.sh` because this image has no equivalent of manywheel's `build_env_setup.py` to read that file, and `install_rocm.sh` only hooks `/etc/bashrc`, which non-interactive bash skips. Legacy removals: - Deleted `install_miopen.sh`, `install_rocSHMEM.sh`, `install_rocm_drm.sh` and `install_rocm_magma.sh`. - Deleted the `.ci/magma-rocm/` tree and `build-magma-rocm-linux.yml`; ROCm MAGMA now comes from the SDK rather than standalone packaging. - Dropped the corresponding `.lintrunner.toml` exclusion and `CODEOWNERS` entry. ## Validation - [almalinux](https://github.com/pytorch/pytorch/actions/runs/33026534333/job/98369114589?pr=194919) and [manywheel](https://github.com/pytorch/pytorch/actions/runs/33026534333/job/98369114589?pr=194919) docker builds for ROCm10.0 succeeded. - [Wheel build jobs](https://github.com/pytorch/pytorch/actions/runs/33026859970/job/98372028809) succeeded (binary test jobs failed as expected, since they do not use ECR docker images; PR needs to be merged for public docker images to be available). - Regenerated all binary build workflows with `.github/scripts/generate_ci_workflows.py`; `git status` reports no drift, so the committed generated workflow matches the matrix change. Tracked in AIPYTORCH-1103. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194919 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily * Update Metamates merge rule (#194968) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Updates the Metamates merge rule with active team members. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194968 Approved by: https://github.com/pytorchbot * [dynamo, 3.15] Update test_decorators (#194424) binascii.b21_base64 changed its signature in 3.15 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194424 Approved by: https://github.com/hameerabbasi, https://github.com/guilhermeleobas * make sure `stream` is `torch.cuda.ExternalStream` or `torch.cuda.Stream` (#194727) To fix `AttributeError: 'cuda.bindings.driver.CUstream' object has no attribute 'device'` in some nv universal gemm tests. ref: https://hud.pytorch.org/pr/pytorch/pytorch/194364#96790367885 ``` TestNVUniversalGemm.test_scaled_gemm_mxfp8_layout_a_aligned_offset_m_128_n_256_k_512: File "/opt/pytorch/pytorch/test/inductor/test_nv_universal_gemm.py", line 428, in test_scaled_gemm_mxfp8 result = compiled_fn(a_fp8, b_fp8, scale_a, scale_b) ... File ".../torch/_inductor/codegen/nv_universal_gemm/nv_universal_gemm_kernel.py", line 869, in _nvgemm_run kernel.run( File ".../cutlass/operators/base.py", line 193, in run return self._run(args, compiled_artifact, stream, workspace) File ".../vendored_templates/cutedsl/wrappers/dense_blockscaled_gemm_kernel.py", line 355, in _run with torch.cuda.stream(stream): File ".../torch/cuda/__init__.py", line 869, in __enter__ if self.src_prev_stream.device != cur_stream.device: AttributeError: 'cuda.bindings.driver.CUstream' object has no attribute 'device' ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/194727 Approved by: https://github.com/ngimel * [vision hash update] update the pinned vision hash (#194669) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Update the pinned vision hash. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194669 Approved by: https://github.com/pytorchbot * Fix dispatch failure on a tensor subclass with no __torch_dispatch__ (#193968) ## Issue Fixes #193932 ## Summary `_AnalyzeCustomOpInputOutputMode.__torch_dispatch__` returns `NotImplemented` for any tensor whose exact type is outside `HANDLED_TYPES`, deferring to a subclass `__torch_dispatch__` below it. A bare subclass (`class Bare(torch.Tensor): pass`) has none, so dispatch runs out of candidates and raises: `aten.mm.out` from inductor's extern mm, `aten.t.default` under `aot_eager`. Pointwise ops inductor codegens never re-enter dispatch, so `add`/`sum`/`relu` on the same subclass already compile. The mode now defers only if the type really implements `__torch_dispatch__`, by the same `torch._C._disabled_torch_dispatch_impl` identity test `HigherOrderOperator.dispatch` uses in `torch/_ops.py`. Same shape as #174485, which handled the `nn.Parameter` case on this line. ## Checklist - [x] Passes lint (`spin fixlint`) - [x] Added/updated tests - [ ] Updated documentation (if applicable) - [ ] Included benchmark results (for PRs impacting perf) ## BC-breaking? No. --- AI assistance was used: Claude Code helped locate the root cause and draft the patch and test, which I reviewed and ran locally. Pull Request resolved: https://github.com/pytorch/pytorch/pull/193968 Approved by: https://github.com/aorenste * [3/N][Test] Enable `torch_ao_sparsity.py` for XPU (#191202) - rename `TestSaliencyPruner` to `TestSaliencyPrunerDevice` - enable XPU for `TestSaliencyPrunerDevice` - enable XPU for `TestBaseStructuredSparsifierDevice` - enable XPU for `TestFPGMPrunerDevice` - add hardware classification Test summary: Before: 88 passed After: 105 passed \---------------- Delta: +17 passed Newly passed tests: - TestSaliencyPrunerDeviceXPU::test_lstm_saliency_pruner_update_mask_xpu_float32 - TestSaliencyPrunerDeviceXPU::test_saliency_pruner_update_mask_xpu_float32 - TestBaseStructuredSparsifierDeviceXPU::test_complex_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_constructor_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prepare_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prepare_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_activation_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_bias_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_padding_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_pool_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_linear_activation_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_linear_bias_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_linear_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_step_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_step_linear_xpu - TestFPGMPrunerDeviceXPU::test_update_mask_xpu --- Stack of following PRs: - https://github.com/pytorch/pytorch/pull/191199 - https://github.com/pytorch/pytorch/pull/191200 - -> https://github.com/pytorch/pytorch/pull/191202 Pull Request resolved: https://github.com/pytorch/pytorch/pull/191202 Approved by: https://github.com/etaf, https://github.com/CuiYifeng, https://github.com/jansel * Drop unused includes in aten (#194960) ## Summary Split out of #194958. Removes unused `#include`s in aten, flagged by clang-tidy's misc-include-cleaner (not enabled by this repo's `.clang-tidy`/lintrunner config). Pull Request resolved: https://github.com/pytorch/pytorch/pull/194960 Approved by: https://github.com/eqy * [inductor] Fix test_comprehensive_addbmm_cpu_float16 precision mismatch (#187170) (#187171) ## Summary Fix `test_comprehensive_addbmm_cpu_float16` by comparing compiled f16 output against eager f16 output instead of an idealized f32 reference. Fixes #187170 ## Problem `addbmm` uses `make_fallback` in Inductor, meaning both compiled and eager paths call the same ATen kernel. The test was failing because it compared compiled f16 output against a float32 reference that avoids intermediate f16 rounding between the 5 batch `addmm_` iterations. The native f16 path rounds after each batch step, producing a deterministic ~0.019 absolute difference at one element — exceeding the 0.006 tolerance. ## Fix Set `reference_in_float=False` for `("addbmm", f16)` on CPU in `inductor_override_kwargs`, so the test compares compiled f16 against eager f16 directly. Since both paths execute the same ATen kernel with identical inputs, this is the correct comparison. ## Test plan - [x] `PYTORCH_OPINFO_SAMPLE_INPUT_INDEX=1 python test/inductor/test_torchinductor_opinfo.py TestInductorOpInfoCPU.test_comprehensive_addbmm_cpu_float16` passes - [ ] CI green Pull Request resolved: https://github.com/pytorch/pytorch/pull/187171 Approved by: https://github.com/jansel, https://github.com/fffrog, https://github.com/albanD, https://github.com/rogerchen838 * Label tests that should run on CuTeDSL or B200 appropriately (#193606) test/python_native/**, inductor/test_cutedsl_grouped_mm and inductor/test_cutedsl_template are CuTeDSL-gated but appear in no job's include list, so they never run anywhere; grouped_mm sits in the H100 smoke config, where the Blackwell gate skips it. Run them in smoke_b200, the one job with CuTeDSL, and label their paths and the templates they cover. The python_native list is discovered with find, like test_dtensor, so new DSL tests need no further edits here. Post-review cleanup: python_native/test_cutedsl_smoketest is picked up by finds, removing the explicit line so it doesn't run twice Test Plan: the job itself can only run in CI. Verified locally that the discovery pipeline resolves to the 14 python_native names run_test.py knows and that both inductor names are valid includes: ``` find test/python_native -name 'test_*.py' -printf '%P\n' | sed 's|\.py$||; s|^|python_native/|' | sort | tr '\n' ' ' python -c "from tools.testing.discover_tests import TESTS; print([t for t in TESTS if 'cutedsl' in t])" lintrunner -a --paths-cmd 'echo .ci/pytorch/test.sh' ``` Authored with assistance from an AI coding agent (Claude Code). Pull Request resolved: https://github.com/pytorch/pytorch/pull/193606 Approved by: https://github.com/drisspg, https://github.com/ngimel * [c10d] Add an opt-in uniform-rank simulation contract to FakeProcessGroup (#194545) `FakeProcessGroup` documents its collectives as "deterministic single-process approximations" whose "values are arbitrary; do not assert on them". That is unusable for any caller that reshapes collective output, and it is silently wrong for any caller that cares about the reduce op: - `allreduce` / `reduce` / `allreduce_coalesced` ignore `ReduceOp` entirely and are pure no-ops, so a SUM comes back unscaled - `reduce_scatter*` return the local chunk without applying the op - `all_to_all_single` fills the output with a prefix of the local input and zeros the remainder, disregarding the split structure The third one is not hypothetical. Running an internal TorchRec model against the stock backend fails in the first forward: ``` RuntimeError: shape '[16, 8]' is invalid for input of size 114 torchrec/sparse/jagged_tensor.py in KeyedJaggedTensor.dist_init stride_per_key_per_rank = stride_per_rank_per_key.view(...) ``` With a fake world of 16, TorchRec exchanges per-rank stride metadata over all-to-all and reshapes the result. The approximation does not preserve the split structure, so the element count does not line up. ## The contract `Options::simulate_uniform_ranks` (default `false`) makes the group model a world in which **every rank holds data identical to this one**. That single assumption makes every collective well defined from local inputs alone: | collective | under the contract | | --- | --- | | `broadcast` | identity, every rank already holds the value | | `all_gather*` | replicate the local input, unchanged from today | | `allreduce` / `reduce` | see the reduce-op table below | | `reduce_scatter*` | the local chunk, then the same op scaling | | `all_to_all_single` / `alltoall` | each output slot is filled from the segment this rank would send to its own position, tiled to fill the slot | The contract is **total** over reduce ops. Once the contributions are known to be equal, every op has a closed form, so none of them falls back to a silently wrong approximation: | op | result | | --- | --- | | `SUM` | `x * world_size` | | `AVG`, `MIN`, `MAX`, `BAND`, `BOR` | identity, all are idempotent on equal operands | | `PRODUCT` | `x ** world_size` | | `PREMUL_SUM(f)` | `x * f * world_size` | | `BXOR` | `x` for odd `world_size`, `0` for even, since `x ^ x == 0` | `FakeWork` now also resolves its future to the produced tensors, so `async_op=True` callers reading `get_future().value()` see the result instead of `None`. Default-constructed `FakeWork` is unchanged and still returns a valueless future. ## Type stubs `torch/_C/_distributed_c10d.pyi` did not declare `FakeProcessGroup.Options` at all, nor `_create_internal`'s `options` argument, so every Python caller of this flag was a type error. It also declared only the `(store, rank, size)` `ProcessGroup` constructor, though `init.cpp` has exposed a `(rank, size)` overload for as long as the store-less form has existed. Both are added here, which lets the callers in the diffs above drop their `cast(Any, ...)` workarounds. ## Scope Off by default: with the flag unset every code path and every existing test behaves exactly as before. The modeling cost when it is on is that ranks cannot diverge, so it cannot surface rank-divergence bugs; that is documented on the class in NOTE [FakeProcessGroup uniform-rank simulation]. @exported-using-ghexport Differential Revision: [D117139104](https://our.internmc.facebook.com/intern/diff/D117139104/) Differential Revision: [D117139104](https://our.internmc.facebook.com/intern/diff/D117139104) Pull Request resolved: https://github.com/pytorch/pytorch/pull/194545 Approved by: https://github.com/aorenste * [CUDA][ROCm] Support pruned lookups in quantized embedding_bag_byte_rowwise_offsets (#192573) (#192573) Summary: `quantized::embedding_bag_byte_rowwise_offsets` takes an optional `compressed_indices_mapping`: a table mapping an index in the original, uncompressed row space onto a row of a pruned weight table, with `-1` marking a pruned row. The CUDA/ROCm path never supported it — the kernel accepts the argument and ignores it, and `embedding_bag_byte_impl` bails out with "Compressed indices mapping not yet implemented", so pruned quantized embedding tables were CPU-only. ## How Translate the indices into the compressed row space, then re-enter this entry point as an ordinary dense lookup, so the kernel is reused unchanged. This mirrors the CPU op in `native/quantized/cpu/qembeddingbag.cpp`. - Rows marked pruned (`-1`) must contribute nothing. CPU skips them; the dense kernel cannot, so the index is clamped to a valid row and its contribution cancelled with a zero per-sample weight. Branchless: testing for pruned rows on the host would synchronize the device on every call. - A fully pruned table has no row to index, so the all-zero result is returned directly, which is what CPU produces. - The caller's tensors are screened in one block at the top of the entry point. The remap re-enters with rewritten tensors and returns early for a fully pruned table, so a check left further down would inspect the wrong tensors or be skipped entirely. - The mapping must be `Int` — CPU enforces that implicitly by reading it through `const_data_ptr<int32_t>()`, whereas here it is `index_select`'s `self`, which takes any dtype and would silently truncate — and must be on the weight's device, since a per-call host-to-device copy of a table with one entry per uncompressed row is not a cost to hide inside an operator. - `indices` must be 1-D when a mapping is given. CPU accepts 2-D and synthesizes offsets, but that path is already dead on CUDA: `at::arange` builds those offsets on the host and then trips `TORCH_CHECK(offsets.is_cuda())`. Only the 8-bit op is covered; `embedding_bag_4bit_rowwise_offsets` still rejects a mapping. ## Deliberate divergences from CPU All three follow from keeping a device synchronization out of the forward pass. 1. **Any single-entry mapping is read as the "not pruned" sentinel.** Two spellings circulate: `{0}`, which CPU honours, and `{-1}`, which at least one out-of-tree producer emits for "no pruning was performed". CPU reads any non-`{0}` single value as a genuine one-row mapping. Telling them apart needs a device read, so for such a mapping this does a dense lookup where CPU would index through it, or reject ids `>= 1` against its length. Pinned by tests on both sides. 2. **Out-of-range indices fault device-side** rather than raising CPU's "... is out of bounds: ...". Not pinned by a test: on ROCm the failure is a bare `abort()` that kills the process instead of raising. 3. **A pruned entry yields NaN when the data is not finite.** Cancelling with a zero weight is exact only for finite values, and the clamped row is still read, so `0 * inf = NaN` where CPU skips the entry outright. ## Cost The remap keeps two per-index tensors alive across the launch: the translated ids, and a float32 weight vector, because cancelling pruned rows turns an unweighted call into a weighted one. Peak device memory grows by 8 bytes per index for `Int` indices, 16 for `Long`. Doing the lookup inside the kernel would remove that, but its `compressed_indices_mapping` parameter arrives as a host `std::optional<Tensor>&` and would have to become a device accessor first, so that is left as a follow-up. ## Also in this PR Small fixes in the same file, each surfaced by the change or by reading around it: - Both `embedding_bag_byte_impl` and `embedding_bag_4bit_impl` built their `offsets` accessor from the index type deduced from `indices`, so the mismatched dtype pairs their own dispatch tables enumerate threw from the accessor. Both now cast `offsets` first. - The byte path's unreachable "not yet implemented" check and the device check beside it are removed; the 4-bit copy of that message named the 8-bit op and is corrected. - `#include <ATen/ops/zeros.h>`, which the per-operator-header build needs for the new `at::zeros`. - An `offsets.dim() == 1` check, and a `weight.size(1) >= 8` bound that the existing `(weight.size(1) - 8) % 4 == 0` does not imply, since `(4 - 8) % 4` is also 0. `IndexType`/`OffsetType` are unused in both impls, so the four-way dispatch could collapse to one call each. That is a pure deletion, but ~90 lines across two functions, so it is left as a follow-up rather than folded in here. Test Plan: Eight new CUDA tests in `TestQuantizedEmbeddingOps` (`test/quantization/core/test_quantized_op.py`), covering: CUDA against CPU on a pruned table, across per-sample weights and `include_last_offset`; that the remap adds no numeric error of its own, asserted bitwise on-device; the `{0}` and `{-1}` sentinels, with CPU's divergent rejection pinned; a fully pruned table; the input guards via `assertRaisesRegex`, over both a populated and a fully pruned table, so the early return cannot start letting input through; a mapping passed with `pruned_weights=False`; mixed `indices`/`offsets` dtypes for the 8-bit and the 4-bit op; and that the 4-bit op still rejects a mapping. CPU comparisons on a populated table use `rtol=0, atol=1e-6`. They are not bitwise, because the two kernels factor the dequant and the per-sample-weight multiply differently, so the remaining bitwise assertions are on-device -- the exception is the fully pruned table, where both sides are exactly zero. The worst measured CPU-vs-CUDA difference is one float32 ulp: 8.9e-08 on AMD, 1.2e-07 on NVIDIA. These are device-pair tests for a single operator, so they are deliberately not `OpInfo`-based or device-generic; the adjacent pre-existing CUDA test in this class uses the same style. Axes such as `include_last_offset` and the per-sample weights run as `subTest` rather than `parametrize`, so the method names stay stable: parametrization renames them, which breaks name-based test selection and disable lists. python test/test_quantization.py TestQuantizedEmbeddingOps Passes on both an AMD GPU (ROCm) and an NVIDIA H100 (sm_90, CUDA 12.4): 15 tests, none skipped. Reviewed By: sacdroid Differential Revision: D115070363 Pull Request resolved: https://github.com/pytorch/pytorch/pull/192573 Approved by: https://github.com/vkuzo, https://github.com/liangel-02 * [Eager] Add FlyDSL native RMSNorm forward override (#191447) Add a FlyDSL-based fused RMSNorm forward kernel as a native op override, on top of the backend in the PR below. ### Supported cases fp16, bf16 and fp32 tensors on gfx950, contiguous input and weight, 1-D `normalized_shape`, and `weight` of shape `(N,)` with matching dtype and device. N-D inputs are logically flattened to `(M, N)`. `eps` is passed through at runtime and does not participate in the cache key. The kernel returns the forward result and the FP32 `rstd` that ATen's backward consumes. ### When FlyDSL acceleration kicks in gfx950, non-copy-on-write contiguous input and weight, one normalized dimension, non-negative `eps`, and: | N | minimum rows M | |---|---:| | `4096 <= N < 8192` | 8192 | | `8192 <= N < 16384` | 4096 | | `16384 <= N <= 114688` | 2048 | ### Test plan: - lintrunner on all changed files - python test/python_native/test_flydsl_rmsnorm_fwd.py - OPINFO_RESTRICT_TO_DSL=flydsl python test/test_ops.py -k rms_norm ### Performance methodology Latency is measured with CUDA events after 10 warmup iterations and over 50 timed iterations. The ATen baseline runs with FlyDSL disabled; the FlyDSL measurement invokes the same operator with the override enabled. Every row below was confirmed to actually dispatch, by checking the compile cache call count advanced, and its output and `rstd` were compared against the ATen baseline. Each row is a single run. Repeating the smallest shape ten times spread its speedup over 1.16x-1.26x. ```python def bench(fn, iters=50, warmup=10): for _ in range(warmup): fn() torch.cuda.synchronize() start = torch.cuda.Event(enable_timing=True) end = torch.cuda.Event(enable_timing=True) start.record() for _ in range(iters): fn() end.record() torch.cuda.synchronize() return start.elapsed_time(end) / iters fwd = lambda: torch.ops.aten._fused_rms_norm(x, [N], weight, eps) with pn.flydsl.disabled(): aten_ms = bench(fwd) flydsl_ms = bench(fwd) speedup = aten_ms / flydsl_ms ``` ### Performance GPU: AMD 256-CU MI355X (gfx950). FlyDSL 0.3.0. | dtype | M | N | ATen (us) | FlyDSL (us) | speedup | |:---|---:|---:|---:|---:|---:| | fp16 | 8192 | 4096 | 26.7 | 22.9 | 1.16x | | fp16 | 4096 | 8192 | 30.1 | 22.6 | 1.33x | | fp16 | 2048 | 16384 | 31.0 | 23.1 | 1.34x | | fp16 | 2048 | 65536 | 173.2 | 112.3 | 1.54x | | fp16 | 2048 | 114688 | 303.2 | 258.8 | 1.17x | | fp16 | 8192 | 4097 | 96.7 | 26.4 | 3.66x | | fp16 | 4096 | 8193 | 82.7 | 26.0 | 3.18x | | fp16 | 2048 | 16385 | 91.9 | 25.9 | 3.55x | | bf16 | 8192 | 4096 | 26.7 | 22.2 | 1.20x | | bf16 | 2048 | 16384 | 30.9 | 22.4 | 1.38x | | bf16 | 8192 | 4097 | 96.4 | 26.4 | 3.65x | | bf16 | 2048 | 16385 | 73.7 | 25.8 | 2.86x | | fp32 | 8192 | 4096 | 56.3 | 40.1 | 1.41x | | fp32 | 4096 | 8192 | 58.5 | 40.7 | 1.44x | | fp32 | 2048 | 16384 | 57.7 | 44.1 | 1.31x | | fp32 | 2048 | 32768 | 154.1 | 105.5 | 1.46x | | fp32 | 2048 | 65536 | 321.8 | 210.3 | 1.53x | | fp32 | 2048 | 114688 | 564.7 | 483.8 | 1.17x | | fp32 | 8192 | 4097 | 91.8 | 41.6 | 2.20x | | fp32 | 4096 | 8193 | 82.5 | 42.3 | 1.95x | | fp32 | 2048 | 16385 | 78.0 | 43.9 | 1.78x | | fp32 | 2048 | 32769 | 212.7 | 105.0 | 2.03x | Aligned shapes improve by 1.2x to 1.5x, with both ends of the band -- `N` = 4096 and `N` = 114688 -- sitting near 1.2x. Off-by-one `N` improves by up to 3.7x. Unsupported shapes retain ATen behavior. Co-authored-by: lirui927 <ruili@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Pull Request resolved: https://github.com/pytorch/pytorch/pull/191447 Approved by: https://github.com/jgong5, https://github.com/slayton58 Co-authored-by: lirui927 <Rui.Li@amd.com> * Extend accelerator graph-capture health check beyond CUDA/XPU (#194362) ## Summary `_accelerator_graph_capture_health_check()` gated its whole body behind `if accelerator and accelerator.type in {"cuda", "xpu"}:`, but everything inside that block is already backend-agnostic: `torch.accelerator.current_stream().is_capturing()` and `accelerator.type.upper()` in the warning/error messages both work for any accelerator already. That gate traces back to #180453, which modernized this function's internals but preserved the previous code's explicit `torch.cuda.is_available()`/`torch.xpu.is_available()` restriction as a faithful port rather than a deliberate decision tied to `is_capturing()`'s semantics. `c10::Stream::is_capturing()` calls `DeviceGuardImplInterface::isStreamCapturing()`, whose base implementation is `return false;` — only `CUDAGuardImpl` and `XPUGuardImpl` override it with real logic (`c10/core/impl/DeviceGuardImplInterface.h`). Every other backend (MPS, MTIA, HPU, PrivateUse1) already safely returns `False` through the same call, so restricting the check to `{"cuda", "xpu"}` served no functional purpose — it just meant this health check (and its `capturable=True`/`False` warnings) silently never ran on any other backend. Replaced the gate with `if accelerator is not None:`, which behaves identically for CUDA/XPU and correctly extends the same already-safe check to every other accelerator. Part of the device-decoupling initiative proposed in #194355. ## Test Plan Verified using: python -c " import torch import torch.optim as optim p = torch.nn.Parameter(torch.randn(3, 3)) opt = optim.Adam([p], capturable=False) opt._accelerator_graph_capture_health_check() print('ok') " `is_capturing()` correctly returns `False` on MPS and the health check runs without error. Ran `lintrunner -a` on `torch/optim/optimizer.py` — no issues. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194362 Approved by: https://github.com/janeyx99, https://github.com/guangyey * Fix positional arg overflow in TritonBenchmarkRequest when template inputs are deduplicated (#188268) # Fix positional arg overflow in `TritonBenchmarkRequest` when template inputs are deduplicated Fixes #188069 When bmm template codegen deduplicates kernel inputs (both operands alias the same buffer, e.g. `einsum("...ab,...Ab->...aA", grad, grad)`), the compiled kernel expects fewer positional parameters than `TritonBenchmarkRequest.make_run_fn()` provides. The excess tensor overflows into the `stream` parameter slot, producing `TypeError: too many positional arguments`. This happens when two compilations run in sequence: the first caches a deduplicated bmm kernel, the second's autotuner reuses that cached candidate but passes more input tensors than the kernel accepts. The CUTLASS path handles a similar case via `dict.fromkeys` on `input_tensor_meta` names, but that doesn't work here because the meta reflects the *current* graph's inputs (2 unique names), not the *cached* kernel's actual param count (1 input after dedup). Tested it and it still crashes 4/5 on 2.12.0. This PR reads `triton_meta["signature"]` from the loaded kernel to get its actual parameter count and trims `input_tensors` to match. The check is a no-op when counts already agree. Tested on RTX 3090, 5 trials per configuration with isolated caches: | | Baseline | Patched | |---|---|---| | 2.12.0+cu130 | 1-2/5 crash | 0/5 | | Nightly 2.14.0.dev20260625 | 0-2/5 crash | 0/5 | Pull Request resolved: https://github.com/pytorch/pytorch/pull/188268 Approved by: https://github.com/jananisriram, https://github.com/jansel, https://github.com/mlazos * Reenable test_addcmul_alpha_one_fma_parity dtypes F32 and F64 on XPU. (#182811) To be merged after https://github.com/intel/torch-xpu-ops/pull/3275 is merged and reference updated. Pull Request resolved: https://github.com/pytorch/pytorch/pull/182811 Approved by: https://github.com/astachowiczhabana, https://github.com/pbielak, https://github.com/janeyx99, https://github.com/guangyey * Revert "Label tests that should run on CuTeDSL or B200 appropriately (#193606)" This reverts commit 926a5451a8eb984d35c4e6b17543aa8e9ad76d7f. Reverted https://github.com/pytorch/pytorch/pull/193606 on behalf of https://github.com/jansel due to Broke smoke_b200 on trunk ([comment](https://github.com/pytorch/pytorch/pull/193606#issuecomment-5443173771)) * [inductor] Fix inductor dropping ordering dep between effectful ops with different kernel types (#188301) `with_effects` records an ordering edge between consecutive ORDERED-effectful ops so the scheduler preserves their program order. That edge was both dropped and, when it did apply, too strong. Fixing the strength then exposed a latent conflation in how pinned buffers are freed. **The edge was dropped.** It was recorded in `V.graph.additional_star_deps` keyed by `new_op.get_name()` (a buffer name), but the scheduler reads it back by operation name. Ops lowering to the same kernel type gave names that coincided, masking the mismatch; different kernel types (e.g. an `ExternKernel` followed by a `_CollectiveKernel`) made the key miss, so the dep vanished and the scheduler could reorder `op_b` ahead of `op_a`. Fixed by keying on `get_operation_name()`. **The edge was too strong.** `additional_star_deps` is installed as a `StarDep`, which asserts ordering *and* liveness, so an effect edge counted as a real read of the previous op's output buffer and held it live past its last true use, blocking reuse and deallocation. An effect edge carries no data, so it belongs in `additional_buffer_deps`, installed as `WeakDep(is_fake=True)`, filtered out of `used_or_aliased_buffer_names`, so it orders without extending lifetimes. This is the pattern `OrderingBarrier` already documents. `additional_star_deps` had no other users and is removed. **Retention is a separate concern, and also needed.** The above is about the *previous* op's output buffer. Independently, *this* op's inputs may be retained in state inductor cannot see (`tq.push(x)` stashes a tensor in a torchbind queue), and no dependency edge can express "the callee kept a reference", so those inputs must not have their storage recycled. The axes are independent: without the pin `test_torchbind_queue` gives silently wrong numerics under either dep strength. **Pinning must block recycling without blocking deallocation,** since freeing only drops inductor's own reference and storage retained behind its back survives. That distinction did not exist: `never_reuse_buffers` both withheld a buffer from the reuse pool and suppressed its free. Putting retention pins there would therefore keep every effectful op's input allocated for the whole graph, and peak memory grows with the number of effectful ops — measured locally, though the memory test is not part of this PR. But the suppression half is load-bearing elsewhere: `ShallowCopyDataKernel` pins its own output, which in the cpp wrapper has no output variable at all (the call is emitted bare), so emitting a free for it generates a `reset()` on an undeclared name and fails to compile. Rather than redefine the existing set, retention pins go in a new `never_reuse_but_free_buffers`, and the split is expressed as `can_free` (body unchanged from the old `can_reuse`, so every pre-existing pin site behaves exactly as before) plus `can_reuse` (`can_free` and not retention-pinned). Review `lowering.py` first for the keying fix, the weakened dep, and the pin; then `scheduler.py` for the dep the switch now goes through, and `graph.py` / `wrapper.py` for the two sets and the `can_free`/`can_reuse` split. ## Test plan `test_inductor_preserves_effect_order_across_kernel_types` covers the keying fix: two ORDERED ops lowering to different kernel types (`op_a` -> `FallbackKernel`, `op_b` -> `_CollectiveKernel`), asserting `op_a` is emitted first. It fails without the fix because the dropped dep lets the scheduler reorder them. `test_inductor_effect_order_does_not_extend_lifetimes` covers the dep strength. To discriminate it needs an effectful op with a real output buffer for a strong dep to over-retain, plus a later effectful op consuming a value derived from that buffer so it cannot be hoisted above the buffer's last use; it then asserts the buffer is freed before that later op. Verified to fail with `is_fake=False` and pass with `is_fake=True`. Pull Request resolved: https://github.com/pytorch/pytorch/pull/188301 Approved by: https://github.com/izaitsevfb * Create a @periodic Decorator (#191946) Solves https://github.com/pytorch/test-infra/issues/8296 Adds a `@periodic` decorator for tests that should run only in periodic CI, including interoperability with `@slowTest`. Periodic Linux, ROCm, Inductor, and s390x workflows now explicitly enable these tests. For now, the only periodic tests are the ones I'm introducing to test the decorator. _Authored with Codex_ Pull Request resolved: https://github.com/pytorch/pytorch/pull/191946 Approved by: https://github.com/izaitsevfb, https://github.com/huydhn * [ROCm] Fixed test_cuda preferred_blas_library_settings for gfx1250 arch (#194785) Fixed test_cuda.py::TestCuda::test_preferred_blas_library_settings Pull Request resolved: https://github.com/pytorch/pytorch/pull/194785 Approved by: https://github.com/jeffdaily Co-authored-by: Jeff Daily <jeff.daily@amd.com> * [c10d] Decouple completion hooks from Work lifetime (#194830) Completion hooks correlated operations through Work addresses, so consumers had to retain each Work to prevent address reuse. Give Work an opaque completion key, use a monotonic key for NCCL2, and pass the key through completion hooks so the flight recorder can release Work references safely. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_fr_hook.py Nccl2FlightRecorderHookTest.test_records_and_retires_collectives Nccl2FlightRecorderHookTest.test_completion_arriving_before_registration_still_retires Nccl2FlightRecorderHookTest.test_duration_ms_from_backend_timing ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194830 Approved by: https://github.com/fduwjj, https://github.com/kapilsh * [c10d] Move NCCL2 completion data into shared state (#194831) NCCL2 completion tracking is owned directly by WorkNCCL, preventing backend tracking from having a narrower lifetime than the caller-visible Work. Move status, events, timing, timeout, and completion notification into a shared State while retaining the existing Work queue behavior. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2Test.test_watchdog_does_not_release_python_backed_tensor ProcessGroupNCCL2Test.test_ephemeral_timeout ``` ```bash env CUDA_VISIBLE_DEVICES=2 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2UninitializedCudaTest.test_barrier_after_init ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194831 Approved by: https://github.com/kapilsh ghstack dependencies: #194830 * [c10d] Share NCCL2 input tensor ownership (#194832) NCCL2 stores input tensors directly on each Work and keeps child Work objects for coalesced operations. Replace both forms with a synchronized shared shelf so input lifetime can later be managed independently of the caller-visible Work. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2Test.test_watchdog_does_not_release_python_backed_tensor ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194832 Approved by: https://github.com/kapilsh ghstack dependencies: #194830, #194831 * [c10d] Release unused functional Work results (#193239) NCCL2 queues caller-visible Work objects for watchdog and CUDA graph tracking. Those references retain collective tensors after wait_tensor removes the functional registry entry. Queue only shared completion state and an input tensor shelf, and retain only completion state during graph capture. Caller waits clear the shared shelf, including every child of a coalesced operation, while direct Work results remain available. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2Test.test_wait_tensor_releases_work_tensors ProcessGroupNCCL2Test.test_wait_tensors_releases_coalesced_work_tensors ProcessGroupNCCL2Test.test_watchdog_does_not_release_python_backed_tensor ``` ```bash env CUDA_VISIBLE_DEVICES=0,1 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_inductor_collectives.py TestCollectivesMultiProc.test_allreduce_inductor_cudagraph_trees ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_fr_hook.py Nccl2FlightRecorderHookTest.test_records_and_retires_collectives Nccl2FlightRecorderHookTest.test_completion_arriving_before_registration_still_retires Nccl2FlightRecorderHookTest.test_duration_ms_from_backend_timing ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/193239 Approved by: https://github.com/kapilsh ghstack dependencies: #194830, #194831, #194832 * [aoti] Prevent exceptions from crossing the AOTI C ABI (#194613) (#194613) Summary: Prevent exceptions from unwinding out of generated AOTI DSOs without changing the public C API function types. Each exported runtime entrypoint uses a function-try-block that records an escaping exception and returns `AOTI_RUNTIME_FAILURE`; secondary failures while recording the error are contained. Consumers retain potentially throwing function pointers for compatibility with older artifacts. Test Plan: `buck2 build fbcode//caffe2/test/inductor/fb:test_aot_inductor_model_runner_pybind fbcode//sigmoid/core/executor:aoti_model_impl_cpu` `buck2 run fbcode//caffe2/test/inductor:aot_inductor_arrayref_cpu -- caffe2.test.inductor.test_aot_inductor_arrayref.AOTInductorTestABICompatibleCpuWithStackAllocationAndMinimalArrayRefInterface.test_runtime_input_count_error_does_not_cross_c_abi_cpu_with_stack_allocation_and_minimal_arrayref_interface` `buck2 run fbcode//caffe2/test/inductor:aot_inductor_arrayref_cpu_local -- caffe2.test.inductor.test_aot_inductor_arrayref.AOTInductorTestABICompatibleCpuWithStackAllocationAndMinimalArrayRefInterface.test_runtime_input_count_error_does_not_cross_c_abi_cpu_with_stack_allocation_and_minimal_arrayref_interface` Reviewed By: desertfire Differential Revision: D116951104 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194613 Approved by: https://github.com/desertfire * Revert "Enable torch.compile for RNN modules by default (#185902)" This reverts commit a2ed3ae5b69bb2ec3427f3ca18a27ae1c100b229. Reverted https://github.com/pytorch/pytorch/pull/185902 on behalf of https://github.com/jansel due to Causes perf regressions in dashboard ([comment](https://github.com/pytorch/pytorch/pull/185902#issuecomment-5443506876)) * Revert "[CUDA][ROCm] Support pruned lookups in quantized embedding_bag_byte_rowwise_offsets (#192573) (#192573)" This reverts commit a0edfca08d1316d2ce990fc6b892366bb03e1040. Reverted https://github.com/pytorch/pytorch/pull/192573 on behalf of https://github.com/georgehong due to Reverting due to OSS and internal sync issues (ERROR: does NOT match the GitHub patch), will be succeeded by https://github.com/pytorch/pytorch/pull/194934 ([comment](https://github.com/pytorch/pytorch/pull/192573#issuecomment-5443720585)) * [Test] Enable XPU for `test_hop.py` and add hardware classification (#192772) This PR enables XPU coverage for `test_hop.py` by opting `TestHOP` into XPU instantiation with `allow_xpu=True` and classifying it as `ACCELERATOR`. Previously, the file collected zero XPU tests under the XPU CI device filter because `XPUTestBase` was never added. This change ensures the file is exercised on XPU instead of silently being dropped. The tests are already device-generic, so no logic changes were needed beyond the XPU registration/classification. On a single-XPU machine: Before: 44 tests (35 passed, 9 xfailed), all `TestHOPCPU` After: 88 tests (70 passed, 18 xfailed) Test plan: ```bash python -m pytest test/export/test_hop.py -q PYTORCH_TESTING_DEVICE_ONLY_FOR=xpu python -m pytest test/export/test_hop.py -q ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/192772 Approved by: https://github.com/pbielak, https://github.com/etaf, https://github.com/jansel * Reland #188184: derive ROCM_VERSION from torch.version.rocm (#194985) Summary: Relands #188184 with the import-time failure fixed. #188184 changed `ROCM_VERSION` in `torch/utils/cpp_extension.py` to come from `torch.version.rocm` rather than `torch.version.hip`, so extensions can compare against the ROCm release version rather than the HIP runtime version. It was reverted out of Meta's internal monorepo because it raised at module import. `torch.version.rocm` was added by #168097, which wired it through `tools/generate_torch_version.py` and CMake only. A build that produces `torch/version.py` some other way -- Meta's Buck build substitutes a hand-written module -- leaves `rocm` absent or None while `torch.version.hip` is set. #188184 raised `AssertionError` in that state; where the attribute was absent entirely it raised `AttributeError` instead, because `if torch.version.rocm is None:` dereferences the very attribute whose absence it reports, so the assertion message was unreachable. Either way, `torch/testing/_internal/common_utils.py` imports `cpp_extension` at module scope, so this took down test collection rather than a single test. Read `torch/utils/cpp_extension.py` first. `_derive_rocm_version()` prefers `torch.version.rocm` and falls back to the HIP version with a warning instead of raising, which keeps `ROCM_VERSION` set whenever `torch.version.hip` is -- `_get_hipcc_path` compares it with no None guard. That function also moves to the newly added `HIP_VERSION`: it selects between `hipcc.exe` and `hipcc.bat`, which is a HIP SDK layout property, and #188184 left it reading `ROCM_VERSION` after repointing that name at a different quantity. The remaining files put the other readers of the ROCm release version on the same footing. `common_cuda.ROCM_VERSION`, `common_cuda._get_torch_rocm_version()` and the `skip_if_rocm_ver_*` decorators in `common_distributed.py` all derived it from HIP even though their thresholds mirror the C++ `ROCM_VERSION` macro. `blaslt_supported_device()` now uses the shared `common_cuda.ROCM_VERSION` instead of recomputing the same tuple, and keeps its `torch.version.hip` guard; switching that guard, as the original did, makes a build with a falsy `rocm` fall through to `return True` and report hipBLASLt support on every AMD GPU without checking `gcnArchName`. `torch/_inductor/heuristics/template/triton.py` reads the attribute through `getattr`, since its `torch.version.rocm is not None` test is itself an attribute access and raised at import on any build lacking the attribute, CUDA builds included. On shipped ROCm the HIP and ROCm major and minor components agree, so this is behaviour-preserving there. It changes results only on builds where the two diverge, which is what the original change was for. Test Plan: Internal context: D116640565 was the import of #188184 and was backed out by D117079854 after S699381. The internal producer that defines `rocm` in `caffe2/fb/version.py` is already on trunk via D117070955, so this diff needs no `caffe2/fb/**` change -- the consumer-side fallback makes a None `rocm` harmless, which is also why the previously proposed `rocm = hip` producer fallback is not needed. Done locally: - `arc f` and `arc lint` clean. - fbcode and xplat copies verified byte-identical for all six files. - The derivation was exercised over every `(hip, rocm)` state: both set and equal, both set and different, `rocm` None, `rocm` attribute absent, `rocm` empty string, `hip` None, `hip` absent, and a `-rc1` third component. No state raises; a differing `rocm` is preserved rather than overwritten by `hip`. To run before landing (needs AMD hardware): - `buck2 test fbcode//caffe2/test/inductor:max_autotune_amd -- --list-only` - `buck2 test fbcode//deeplearning/aot_inductor/inductor_passes/post_grad/tests:test_move_cat_before_batch_expand_in_bmm_amd -- --list-only` - `buck2 test fbcode//mode/opt-amd-gpu fbcode//dper_lib/silvertorch/core/experimental/mrs_publish/utils/tests:li_utils_test` - the APS checkpoint-compatibility target, testx 844425290139603 - the owning targets for `caffe2/test/test_utils.py` and `caffe2/test/test_linalg.py` The first two are the land-blocking AMD test-listing signals that were red on D116640565 V1 and were landed through anyway. Differential Revision: D117630130 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194985 Approved by: https://github.com/pytorchgreenlight, https://github.com/jeffdaily * Fix nn.functional.threshold OpInfo reference for strict NumPy uint8 casting (#189273) ## Summary Fixes #189267 by updating the `nn.functional.threshold` OpInfo reference inputs for unsigned dtypes. ## Problem NumPy ≥ 2.5 deliberately stopped truncating out-of-range Python ints in `np.where` (the reference path). The threshold OpInfo used `value=-9` against `uint8`, which now raises `OverflowError`. ## Fix Use a non-negative fill (`9`) when `dtype == torch.uint8`; keep `value=-9` for signed and float dtypes to preserve negative-value coverage. This is the code-side fix; #189566 proposes capping numpy `<2.5` in CI requirements as an alternative. Pull Request resolved: https://github.com/pytorch/pytorch/pull/189273 Approved by: https://github.com/jeffdaily * [MPS] enable masked_fill fast path for scalar boolean indexing (#195020) Keeps MPS indexing scalars on CPU, allowing `x[bool_mask] = scalar` to use the existing `masked_fill_` fast path instead of generic `index_put_`. Perf: | dtype | Shape | Layout | Before (us) | After (us) | Speedup | |---|---|---|---:|---:|---:| | bf16 | 256 x 1024 | contiguous | 205.0 | 18.4 | 11.14x | | bf16 | 8 x 512 x 1024 | contiguous | 847.6 | 51.3 | 16.52x | | bf16 | 4096 x 4096 | contiguous | 2291.5 | 535.4 | 4.28x | | bf16 | 2048 x 4096 | transposed dense | 1656.8 | 508.2 | 3.26x | | bf16 | 16 x 128 x 64 x 64 | channels-last | 2602.0 | 962.4 | 2.70x | | bf16 | 4 x 32 x 32 x 32 x 32 | channels-last-3d | 1656.0 | 626.1 | 2.64x | | bf16 | 2048 x 2048 | sliced non-dense | 751.7 | 292.9 | 2.57x | | fp16 | 4096 x 4096 | contiguous | 2375.4 | 544.7 | 4.36x | | fp32 | 4096 x 4096 | contiguous | 2841.0 | 1068.5 | 2.66x | | fp32 | 16 x 128 x 64 x 64 | channels-last | 2639.9 | 989.3 | 2.67x | Smaller shapes | dtype | Shape | Layout | Before (us) | After (us) | Speedup | |---|---|---|---:|---:|---:| | bf16 | 128 | contiguous | 168.7 | 7.0 | 24.24x | | bf16 | 256 | contiguous | 168.4 | 7.5 | 22.61x | | bf16 | 512 | contiguous | 168.3 | 6.7 | 25.19x | | bf16 | 1024 | contiguous | 169.2 | 6.8 | 25.00x | | bf16 | 32 x 32 | contiguous | 168.1 | 5.8 | 28.94x | | bf16 | 64 x 64 | contiguous | 169.7 | 4.0 | 42.95x | | bf16 | 128 x 128 | contiguous | 169.0 | 5.2 | 32.63x | | bf16 | 256 x 256 | contiguous | 173.1 | 6.1 | 28.19x | | fp16 | 128 | contiguous | 168.9 | 2.9 | 57.86x | | fp16 | 256 | contiguous | 169.3 | 3.8 | 45.02x | | fp16 | 512 | contiguous | 168.6 | 2.9 | 57.93x | | fp16 | 1024 | contiguous | 169.3 | 3.0 | 56.05x | | fp16 | 32 x 32 | contiguous | 169.3 | 3.0 | 57.00x | | fp16 | 64 x 64 | contiguous | 169.0 | 2.6 | 64.02x | | fp16 | 128 x 128 | contiguous | 169.3 | 2.7 | 63.42x | | fp16 | 256 x 256 | contiguous | 173.5 | 3.1 | 56.14x | | fp32 | 128 | contiguous | 169.8 | 2.3 | 74.82x | | fp32 | 256 | contiguous | 168.3 | 2.3 | 73.51x | | fp32 | 512 | contiguous | 169.3 | 2.4 | 71.72x | | fp32 | 1024 | contiguous | 169.0 | 2.5 | 68.97x | | fp32 | 32 x 32 | contiguous | 169.0 | 2.5 | 68.15x | | fp32 | 64 x 64 | contiguous | 169.9 | 2.5 | 68.24x | | fp32 | 128 x 128 | contiguous | 168.7 | 2.5 | 66.69x | | fp32 | 256 x 256 | contiguous | 175.1 | 3.2 | 54.71x | Pull Request resolved: https://github.com/pytorch/pytorch/pull/195020 Approved by: https://github.com/malfet * Drop unused includes in c10 (#194959) ## Summary Split out of #194958. Removes unused `#include`s in c10, flagged by clang-tidy's misc-include-cleaner (not enabled by this repo's `.clang-tidy`/lintrunner config). Pull Request resolved: https://github.com/pytorch/pytorch/pull/194959 Approved by: https://github.com/albanD * [Inductor] Reuse DeviceInterface for out-of-tree C++ device options (#190614) Reuse the existing DeviceInterface registration path so out-of-tree backends (e.g. NPU) can provide device-specific C++ compile options without monkey-patching `get_cpp_torch_device_options()`. Changes: - Add an optional `get_cpp_device_options()` method to `DeviceInterface` - `get_cpp_torch_device_options()` queries the registered device interface for device-specific C++ options - Preserve existing behavior when the device interface does not provide C++ options @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo @azahed98 Pull Request resolved: https://github.com/pytorch/pytorch/pull/190614 Approved by: https://github.com/JiasenTian, https://github.com/jansel * Sync requires_grad after _set_tensor_requires_grad under torch.func (#193724) ## Issue Fixes #193279 ## Summary `_create_differentiable` flips `requires_grad` via `_set_tensor_requires_grad`, an in-graph function whose mutation never reached its `TensorVariable`. `call_apply` gates autograd customization on that stale `False` (`misc.py:1118`), so it inlined `forward` and dropped the custom `backward`. The fix re-syncs metadata after fake propagation; the graph node is unchanged. Any `x.requires_grad` read in a transform was stale too - second test. Behaviour change: a custom `vjp` or `jvp` now hits its pre-existing graph break under `func.grad` (`misc.py:1143`, `:1158`); both were previously inlined, so a `fullgraph=True` compile that happened to be correct now errors…

* Assign the tensor type prototype instead of memcpy (#194590) py_initialize_tensor_type() copies a prototype PyTypeObject into the type it is initializing. Both sides are PyTypeObject, so this is a plain copy, not type punning, and assignment expresses it directly. It also drops the hand-written `sizeof`: with memcpy, changing the parameter's type would silently copy the wrong number of bytes, while assignment is a compile error. PyTypeObject is trivially copyable and copy-assignable, so the generated code is the same. Test Plan: ``` static_assert(std::is_copy_assignable_v<PyTypeObject>); static_assert(std::is_trivially_copyable_v<PyTypeObject>); ``` both hold, and the file compiles with -fsyntax-only against the CPython headers. This PR description was drafted with AI assistance. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194590 Approved by: https://github.com/albanD * Support torch.xpu.Event in torch.multiprocessing (#192514) # Motivation As title # Additional Context <img width="1349" height="298" alt="image" src="https://github.com/user-attachments/assets/729accca-951f-4e84-8ce8-aca64ee2f988" /> Pull Request resolved: https://github.com/pytorch/pytorch/pull/192514 Approved by: https://github.com/albanD ghstack dependencies: #192495 * Report a child's KeyboardInterrupt instead of exiting 0 in spawn (#194439) _wrap_ swallowed every `KeyboardInterrupt` (`pass` -> exit 0). That is only correct for pdeathsig (Linux parent death sends SIGINT); a `KeyboardInterrupt` raised by `fn` or sent to the child alone made a dead child look successful to `join()`. Now the child exits quietly only when the parent is gone (`getppid()` changed); otherwise the traceback goes to the error file like any other failure. Terminal-wide Ctrl+C is unaffected (the parent itself raises from `join()`). Test: `test_spawn_child_keyboard_interrupt`. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194439 Approved by: https://github.com/d4l3k * [inductor] Don't match pointless_cumsum_replacement on a symbolic fill (#193959) ## Issue Fixes #193876 ## Summary `pointless_cumsum_replacement` folds `full(shape, fill).cumsum(dim)` into `arange(1, n + 1) * fill`, which assumes `fill` is a Python constant. A recompile with a different closure constant, or an unbacked `.item()` fill, passes it in as a symbolic scalar, so `fill_value` binds to an fx `Node` and the fold dies: `int()` on a `Node` for an integral dtype, or `FakeTensor * Node` for a floating one. Eager is fine. The guard declines those matches so `cumsum` lowers normally, extending the scalar-shape check from #193422. A boolean `full` is exempt: the replacement collapses a bool fill with `int(bool(fill_value))` before any arithmetic, so bool still folds. Tested on `int64` and `float32`. Caveat: the periodic A100 `inductor_huggingface_unbacked_parity` job does not run on PRs, and this declines the fold when the fill is unbacked, so PR CI carries no signal there. ## Checklist - [x] Passes lint (`spin fixlint`) - [x] Added/updated tests - [ ] Updated documentation (if applicable) - [ ] Included benchmark results (for PRs impacting perf) ## BC-breaking? No. --- AI assistance was used in developing this change: reproduced the crash and test validation Pull Request resolved: https://github.com/pytorch/pytorch/pull/193959 Approved by: https://github.com/aorenste * Don't mutate the caller's annotation dict in mark_stream (#194851) Human note: backward annotation projection and cupti backend for getting annotations were merged approximately simultaneously, and cupti backend didn't correctly handle backward annotations - all the backward node stayed empty with cupti backend. mark_stream wrote the resolved stream id into the annotation dict it was handed. The annotation is then stored by reference in the kernel annotation registry, so the write had two effects the caller never asked for: the dict gained a "stream" key that outlived the call, and a dict reused across several mark_stream calls (a comms wrapper marking one lane per process group is the motivating case) ended up shared by every region recorded with it, so the last lane written overwrote all the earlier ones and every marked region rendered on a single lane. Build the tagged annotation as a copy instead, which is what the string input path already effectively did. Test Plan: ``` LD_LIBRARY_PATH=/usr/local/cuda-13.2/compat python test/test_cuda_graph_utils.py -k mark_stream ``` The new test marks two lanes with one shared dict; it fails on the parent commit (the caller's dict gains "stream", and both regions resolve to the second lane) and passes here. Authored with assistance from Claude Code. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194851 Approved by: https://github.com/drisspg * Fix grammatical errors in comments and docstrings (#194565) Fixes subject-verb agreement, missing articles, missing possessive apostrophes, and a few dropped or duplicated words across comments and docstrings in torch/. No functional changes. One fix is more than a pure wording change: the embedding sharding docstring in chunk_sharding_spec_ops/embedding.py contained a garbled sentence ("mask IDs whose embeddings are not stored in current rank will to an extra row will ensure ..."). The stray "will" is dropped so the sentence parses; the intended meaning is preserved. Test Plan: comment-only change, no tests run. Lint via: ``` lintrunner -a ``` Authored with the assistance of an AI agent. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194565 Approved by: https://github.com/albanD * Implement IValue::to without template specializations (#194769) Currently IValue::to is implemented using a combination of macros and template specializations. The implementation is not all in one place which makes it complicated and fragile. This commit is an attempt to reduce this complexity by replacing template specializations with if-constexpr. It also removes the need for _guarded_unsigned_long because the if-else chain is okay with multiple true conditions unlike template specializations. This change was motivated by: - https://github.com/pytorch/pytorch/pull/194659#pullrequestreview-5013557465 - https://github.com/pytorch/pytorch/pull/194659#issuecomment-5403190877 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194769 Approved by: https://github.com/cyyever, https://github.com/wenxin0319, https://github.com/eqy * [ROCm][CI] Upgrade trunk to ROCm 10.0 (#194794) ## Problem - PyTorch trunk ROCm CI still builds against ROCm 7.14. - ROCm 10.0 packages are published from the new `whl-next` index. - rocSHMEM must build for the requested gfx target instead of substituting gfx90a XNACK variants. ## Changes - Upgrade the shared `rocm-n` trunk image from ROCm 7.14 to 10.0. - Install ROCm 10.0 from `https://stable.repo.amd.com/rocm/whl-next/`. - Cherry-pick the rocSHMEM target fix from https://github.com/ROCm/pytorch/pull/3583. - Keep wheel, builder, and nightly changes independent in https://github.com/pytorch/pytorch/pull/194919. ## Validation - Bash syntax and stubbed trunk Docker argument routing pass. - The trunk image selects ROCm 10.0 and the `whl-next` index; the preview image retains its nightly index. - `git diff --check` and targeted IDE lint checks pass for the Docker and CMake changes. - Full image and rocSHMEM builds were not run locally; CI provides end-to-end validation. - [trunk workflow](https://github.com/pytorch/pytorch/actions/runs/33000092249/job/98385742860) ran successfully with ROCm10.0 - [rocm-mi200](https://github.com/pytorch/pytorch/actions/runs/33002840829/job/98340667528) ran successfully with ROCm10.0 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194794 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily * [torchcomms hash update] update the pinned torchcomms hash (#194671) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Update the pinned torchcomms hash. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194671 Approved by: https://github.com/pytorchbot * [ROCm][CI] Add ROCm 10.0 wheel and builder support (#194919) ## Motivation Add ROCm 10.0 support to nightly wheels and almalinux/manylinux docker images. ## Changes Binary matrix and image tags: - `ROCM_ARCHES` moves from `["7.2", "7.14"]` to `["7.14", "10.0"]`, with a matching `rocm[libraries,device-all]==10.0.*` entry in `PYTORCH_EXTRA_INSTALL_REQUIREMENTS`. - `build-manywheel-images.yml` and `build-almalinux-images.yml` build `rocm7.14` and `rocm10.0` instead of `rocm7.1`/`rocm7.2`. - `generated-linux-binary-manywheel-nightly.yml` regenerated to match. Single wheel-based builder path: - `manywheel/build.sh` installs every ROCm version from a TheRock wheel index. ROCm 7.14 keeps `repo.amd.com/rocm/whl-multi-arch/`; 10.0 and anything later default to `stable.repo.amd.com/rocm/whl-next/`, so new versions need no script change. The `rocm_final_legacy` target, the gfx900/gfx906 additions and the 7.1/7.2 patch-version fixups are gone. - `almalinux/build.sh` selects the index the same way. - `manywheel/Dockerfile_2_28` drops the `rocm_final_legacy` stage. - `almalinux/Dockerfile` collapses `rocm_base` and `rocm` into one `rocm` stage that installs the SDK with `RUN bash ./install_rocm.sh && rm install_rocm.sh`, consistent with `Dockerfile_2_28`. It sets `BASH_ENV=/etc/rocm_env.sh` because this image has no equivalent of manywheel's `build_env_setup.py` to read that file, and `install_rocm.sh` only hooks `/etc/bashrc`, which non-interactive bash skips. Legacy removals: - Deleted `install_miopen.sh`, `install_rocSHMEM.sh`, `install_rocm_drm.sh` and `install_rocm_magma.sh`. - Deleted the `.ci/magma-rocm/` tree and `build-magma-rocm-linux.yml`; ROCm MAGMA now comes from the SDK rather than standalone packaging. - Dropped the corresponding `.lintrunner.toml` exclusion and `CODEOWNERS` entry. ## Validation - [almalinux](https://github.com/pytorch/pytorch/actions/runs/33026534333/job/98369114589?pr=194919) and [manywheel](https://github.com/pytorch/pytorch/actions/runs/33026534333/job/98369114589?pr=194919) docker builds for ROCm10.0 succeeded. - [Wheel build jobs](https://github.com/pytorch/pytorch/actions/runs/33026859970/job/98372028809) succeeded (binary test jobs failed as expected, since they do not use ECR docker images; PR needs to be merged for public docker images to be available). - Regenerated all binary build workflows with `.github/scripts/generate_ci_workflows.py`; `git status` reports no drift, so the committed generated workflow matches the matrix change. Tracked in AIPYTORCH-1103. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194919 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily * Update Metamates merge rule (#194968) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Updates the Metamates merge rule with active team members. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194968 Approved by: https://github.com/pytorchbot * [dynamo, 3.15] Update test_decorators (#194424) binascii.b21_base64 changed its signature in 3.15 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194424 Approved by: https://github.com/hameerabbasi, https://github.com/guilhermeleobas * make sure `stream` is `torch.cuda.ExternalStream` or `torch.cuda.Stream` (#194727) To fix `AttributeError: 'cuda.bindings.driver.CUstream' object has no attribute 'device'` in some nv universal gemm tests. ref: https://hud.pytorch.org/pr/pytorch/pytorch/194364#96790367885 ``` TestNVUniversalGemm.test_scaled_gemm_mxfp8_layout_a_aligned_offset_m_128_n_256_k_512: File "/opt/pytorch/pytorch/test/inductor/test_nv_universal_gemm.py", line 428, in test_scaled_gemm_mxfp8 result = compiled_fn(a_fp8, b_fp8, scale_a, scale_b) ... File ".../torch/_inductor/codegen/nv_universal_gemm/nv_universal_gemm_kernel.py", line 869, in _nvgemm_run kernel.run( File ".../cutlass/operators/base.py", line 193, in run return self._run(args, compiled_artifact, stream, workspace) File ".../vendored_templates/cutedsl/wrappers/dense_blockscaled_gemm_kernel.py", line 355, in _run with torch.cuda.stream(stream): File ".../torch/cuda/__init__.py", line 869, in __enter__ if self.src_prev_stream.device != cur_stream.device: AttributeError: 'cuda.bindings.driver.CUstream' object has no attribute 'device' ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/194727 Approved by: https://github.com/ngimel * [vision hash update] update the pinned vision hash (#194669) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Update the pinned vision hash. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194669 Approved by: https://github.com/pytorchbot * Fix dispatch failure on a tensor subclass with no __torch_dispatch__ (#193968) ## Issue Fixes #193932 ## Summary `_AnalyzeCustomOpInputOutputMode.__torch_dispatch__` returns `NotImplemented` for any tensor whose exact type is outside `HANDLED_TYPES`, deferring to a subclass `__torch_dispatch__` below it. A bare subclass (`class Bare(torch.Tensor): pass`) has none, so dispatch runs out of candidates and raises: `aten.mm.out` from inductor's extern mm, `aten.t.default` under `aot_eager`. Pointwise ops inductor codegens never re-enter dispatch, so `add`/`sum`/`relu` on the same subclass already compile. The mode now defers only if the type really implements `__torch_dispatch__`, by the same `torch._C._disabled_torch_dispatch_impl` identity test `HigherOrderOperator.dispatch` uses in `torch/_ops.py`. Same shape as #174485, which handled the `nn.Parameter` case on this line. ## Checklist - [x] Passes lint (`spin fixlint`) - [x] Added/updated tests - [ ] Updated documentation (if applicable) - [ ] Included benchmark results (for PRs impacting perf) ## BC-breaking? No. --- AI assistance was used: Claude Code helped locate the root cause and draft the patch and test, which I reviewed and ran locally. Pull Request resolved: https://github.com/pytorch/pytorch/pull/193968 Approved by: https://github.com/aorenste * [3/N][Test] Enable `torch_ao_sparsity.py` for XPU (#191202) - rename `TestSaliencyPruner` to `TestSaliencyPrunerDevice` - enable XPU for `TestSaliencyPrunerDevice` - enable XPU for `TestBaseStructuredSparsifierDevice` - enable XPU for `TestFPGMPrunerDevice` - add hardware classification Test summary: Before: 88 passed After: 105 passed \---------------- Delta: +17 passed Newly passed tests: - TestSaliencyPrunerDeviceXPU::test_lstm_saliency_pruner_update_mask_xpu_float32 - TestSaliencyPrunerDeviceXPU::test_saliency_pruner_update_mask_xpu_float32 - TestBaseStructuredSparsifierDeviceXPU::test_complex_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_constructor_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prepare_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prepare_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_activation_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_bias_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_padding_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_conv2d_pool_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_linear_activation_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_linear_bias_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_prune_linear_linear_xpu - TestBaseStructuredSparsifierDeviceXPU::test_step_conv2d_xpu - TestBaseStructuredSparsifierDeviceXPU::test_step_linear_xpu - TestFPGMPrunerDeviceXPU::test_update_mask_xpu --- Stack of following PRs: - https://github.com/pytorch/pytorch/pull/191199 - https://github.com/pytorch/pytorch/pull/191200 - -> https://github.com/pytorch/pytorch/pull/191202 Pull Request resolved: https://github.com/pytorch/pytorch/pull/191202 Approved by: https://github.com/etaf, https://github.com/CuiYifeng, https://github.com/jansel * Drop unused includes in aten (#194960) ## Summary Split out of #194958. Removes unused `#include`s in aten, flagged by clang-tidy's misc-include-cleaner (not enabled by this repo's `.clang-tidy`/lintrunner config). Pull Request resolved: https://github.com/pytorch/pytorch/pull/194960 Approved by: https://github.com/eqy * [inductor] Fix test_comprehensive_addbmm_cpu_float16 precision mismatch (#187170) (#187171) ## Summary Fix `test_comprehensive_addbmm_cpu_float16` by comparing compiled f16 output against eager f16 output instead of an idealized f32 reference. Fixes #187170 ## Problem `addbmm` uses `make_fallback` in Inductor, meaning both compiled and eager paths call the same ATen kernel. The test was failing because it compared compiled f16 output against a float32 reference that avoids intermediate f16 rounding between the 5 batch `addmm_` iterations. The native f16 path rounds after each batch step, producing a deterministic ~0.019 absolute difference at one element — exceeding the 0.006 tolerance. ## Fix Set `reference_in_float=False` for `("addbmm", f16)` on CPU in `inductor_override_kwargs`, so the test compares compiled f16 against eager f16 directly. Since both paths execute the same ATen kernel with identical inputs, this is the correct comparison. ## Test plan - [x] `PYTORCH_OPINFO_SAMPLE_INPUT_INDEX=1 python test/inductor/test_torchinductor_opinfo.py TestInductorOpInfoCPU.test_comprehensive_addbmm_cpu_float16` passes - [ ] CI green Pull Request resolved: https://github.com/pytorch/pytorch/pull/187171 Approved by: https://github.com/jansel, https://github.com/fffrog, https://github.com/albanD, https://github.com/rogerchen838 * Label tests that should run on CuTeDSL or B200 appropriately (#193606) test/python_native/**, inductor/test_cutedsl_grouped_mm and inductor/test_cutedsl_template are CuTeDSL-gated but appear in no job's include list, so they never run anywhere; grouped_mm sits in the H100 smoke config, where the Blackwell gate skips it. Run them in smoke_b200, the one job with CuTeDSL, and label their paths and the templates they cover. The python_native list is discovered with find, like test_dtensor, so new DSL tests need no further edits here. Post-review cleanup: python_native/test_cutedsl_smoketest is picked up by finds, removing the explicit line so it doesn't run twice Test Plan: the job itself can only run in CI. Verified locally that the discovery pipeline resolves to the 14 python_native names run_test.py knows and that both inductor names are valid includes: ``` find test/python_native -name 'test_*.py' -printf '%P\n' | sed 's|\.py$||; s|^|python_native/|' | sort | tr '\n' ' ' python -c "from tools.testing.discover_tests import TESTS; print([t for t in TESTS if 'cutedsl' in t])" lintrunner -a --paths-cmd 'echo .ci/pytorch/test.sh' ``` Authored with assistance from an AI coding agent (Claude Code). Pull Request resolved: https://github.com/pytorch/pytorch/pull/193606 Approved by: https://github.com/drisspg, https://github.com/ngimel * [c10d] Add an opt-in uniform-rank simulation contract to FakeProcessGroup (#194545) `FakeProcessGroup` documents its collectives as "deterministic single-process approximations" whose "values are arbitrary; do not assert on them". That is unusable for any caller that reshapes collective output, and it is silently wrong for any caller that cares about the reduce op: - `allreduce` / `reduce` / `allreduce_coalesced` ignore `ReduceOp` entirely and are pure no-ops, so a SUM comes back unscaled - `reduce_scatter*` return the local chunk without applying the op - `all_to_all_single` fills the output with a prefix of the local input and zeros the remainder, disregarding the split structure The third one is not hypothetical. Running an internal TorchRec model against the stock backend fails in the first forward: ``` RuntimeError: shape '[16, 8]' is invalid for input of size 114 torchrec/sparse/jagged_tensor.py in KeyedJaggedTensor.dist_init stride_per_key_per_rank = stride_per_rank_per_key.view(...) ``` With a fake world of 16, TorchRec exchanges per-rank stride metadata over all-to-all and reshapes the result. The approximation does not preserve the split structure, so the element count does not line up. ## The contract `Options::simulate_uniform_ranks` (default `false`) makes the group model a world in which **every rank holds data identical to this one**. That single assumption makes every collective well defined from local inputs alone: | collective | under the contract | | --- | --- | | `broadcast` | identity, every rank already holds the value | | `all_gather*` | replicate the local input, unchanged from today | | `allreduce` / `reduce` | see the reduce-op table below | | `reduce_scatter*` | the local chunk, then the same op scaling | | `all_to_all_single` / `alltoall` | each output slot is filled from the segment this rank would send to its own position, tiled to fill the slot | The contract is **total** over reduce ops. Once the contributions are known to be equal, every op has a closed form, so none of them falls back to a silently wrong approximation: | op | result | | --- | --- | | `SUM` | `x * world_size` | | `AVG`, `MIN`, `MAX`, `BAND`, `BOR` | identity, all are idempotent on equal operands | | `PRODUCT` | `x ** world_size` | | `PREMUL_SUM(f)` | `x * f * world_size` | | `BXOR` | `x` for odd `world_size`, `0` for even, since `x ^ x == 0` | `FakeWork` now also resolves its future to the produced tensors, so `async_op=True` callers reading `get_future().value()` see the result instead of `None`. Default-constructed `FakeWork` is unchanged and still returns a valueless future. ## Type stubs `torch/_C/_distributed_c10d.pyi` did not declare `FakeProcessGroup.Options` at all, nor `_create_internal`'s `options` argument, so every Python caller of this flag was a type error. It also declared only the `(store, rank, size)` `ProcessGroup` constructor, though `init.cpp` has exposed a `(rank, size)` overload for as long as the store-less form has existed. Both are added here, which lets the callers in the diffs above drop their `cast(Any, ...)` workarounds. ## Scope Off by default: with the flag unset every code path and every existing test behaves exactly as before. The modeling cost when it is on is that ranks cannot diverge, so it cannot surface rank-divergence bugs; that is documented on the class in NOTE [FakeProcessGroup uniform-rank simulation]. @exported-using-ghexport Differential Revision: [D117139104](https://our.internmc.facebook.com/intern/diff/D117139104/) Differential Revision: [D117139104](https://our.internmc.facebook.com/intern/diff/D117139104) Pull Request resolved: https://github.com/pytorch/pytorch/pull/194545 Approved by: https://github.com/aorenste * [CUDA][ROCm] Support pruned lookups in quantized embedding_bag_byte_rowwise_offsets (#192573) (#192573) Summary: `quantized::embedding_bag_byte_rowwise_offsets` takes an optional `compressed_indices_mapping`: a table mapping an index in the original, uncompressed row space onto a row of a pruned weight table, with `-1` marking a pruned row. The CUDA/ROCm path never supported it — the kernel accepts the argument and ignores it, and `embedding_bag_byte_impl` bails out with "Compressed indices mapping not yet implemented", so pruned quantized embedding tables were CPU-only. ## How Translate the indices into the compressed row space, then re-enter this entry point as an ordinary dense lookup, so the kernel is reused unchanged. This mirrors the CPU op in `native/quantized/cpu/qembeddingbag.cpp`. - Rows marked pruned (`-1`) must contribute nothing. CPU skips them; the dense kernel cannot, so the index is clamped to a valid row and its contribution cancelled with a zero per-sample weight. Branchless: testing for pruned rows on the host would synchronize the device on every call. - A fully pruned table has no row to index, so the all-zero result is returned directly, which is what CPU produces. - The caller's tensors are screened in one block at the top of the entry point. The remap re-enters with rewritten tensors and returns early for a fully pruned table, so a check left further down would inspect the wrong tensors or be skipped entirely. - The mapping must be `Int` — CPU enforces that implicitly by reading it through `const_data_ptr<int32_t>()`, whereas here it is `index_select`'s `self`, which takes any dtype and would silently truncate — and must be on the weight's device, since a per-call host-to-device copy of a table with one entry per uncompressed row is not a cost to hide inside an operator. - `indices` must be 1-D when a mapping is given. CPU accepts 2-D and synthesizes offsets, but that path is already dead on CUDA: `at::arange` builds those offsets on the host and then trips `TORCH_CHECK(offsets.is_cuda())`. Only the 8-bit op is covered; `embedding_bag_4bit_rowwise_offsets` still rejects a mapping. ## Deliberate divergences from CPU All three follow from keeping a device synchronization out of the forward pass. 1. **Any single-entry mapping is read as the "not pruned" sentinel.** Two spellings circulate: `{0}`, which CPU honours, and `{-1}`, which at least one out-of-tree producer emits for "no pruning was performed". CPU reads any non-`{0}` single value as a genuine one-row mapping. Telling them apart needs a device read, so for such a mapping this does a dense lookup where CPU would index through it, or reject ids `>= 1` against its length. Pinned by tests on both sides. 2. **Out-of-range indices fault device-side** rather than raising CPU's "... is out of bounds: ...". Not pinned by a test: on ROCm the failure is a bare `abort()` that kills the process instead of raising. 3. **A pruned entry yields NaN when the data is not finite.** Cancelling with a zero weight is exact only for finite values, and the clamped row is still read, so `0 * inf = NaN` where CPU skips the entry outright. ## Cost The remap keeps two per-index tensors alive across the launch: the translated ids, and a float32 weight vector, because cancelling pruned rows turns an unweighted call into a weighted one. Peak device memory grows by 8 bytes per index for `Int` indices, 16 for `Long`. Doing the lookup inside the kernel would remove that, but its `compressed_indices_mapping` parameter arrives as a host `std::optional<Tensor>&` and would have to become a device accessor first, so that is left as a follow-up. ## Also in this PR Small fixes in the same file, each surfaced by the change or by reading around it: - Both `embedding_bag_byte_impl` and `embedding_bag_4bit_impl` built their `offsets` accessor from the index type deduced from `indices`, so the mismatched dtype pairs their own dispatch tables enumerate threw from the accessor. Both now cast `offsets` first. - The byte path's unreachable "not yet implemented" check and the device check beside it are removed; the 4-bit copy of that message named the 8-bit op and is corrected. - `#include <ATen/ops/zeros.h>`, which the per-operator-header build needs for the new `at::zeros`. - An `offsets.dim() == 1` check, and a `weight.size(1) >= 8` bound that the existing `(weight.size(1) - 8) % 4 == 0` does not imply, since `(4 - 8) % 4` is also 0. `IndexType`/`OffsetType` are unused in both impls, so the four-way dispatch could collapse to one call each. That is a pure deletion, but ~90 lines across two functions, so it is left as a follow-up rather than folded in here. Test Plan: Eight new CUDA tests in `TestQuantizedEmbeddingOps` (`test/quantization/core/test_quantized_op.py`), covering: CUDA against CPU on a pruned table, across per-sample weights and `include_last_offset`; that the remap adds no numeric error of its own, asserted bitwise on-device; the `{0}` and `{-1}` sentinels, with CPU's divergent rejection pinned; a fully pruned table; the input guards via `assertRaisesRegex`, over both a populated and a fully pruned table, so the early return cannot start letting input through; a mapping passed with `pruned_weights=False`; mixed `indices`/`offsets` dtypes for the 8-bit and the 4-bit op; and that the 4-bit op still rejects a mapping. CPU comparisons on a populated table use `rtol=0, atol=1e-6`. They are not bitwise, because the two kernels factor the dequant and the per-sample-weight multiply differently, so the remaining bitwise assertions are on-device -- the exception is the fully pruned table, where both sides are exactly zero. The worst measured CPU-vs-CUDA difference is one float32 ulp: 8.9e-08 on AMD, 1.2e-07 on NVIDIA. These are device-pair tests for a single operator, so they are deliberately not `OpInfo`-based or device-generic; the adjacent pre-existing CUDA test in this class uses the same style. Axes such as `include_last_offset` and the per-sample weights run as `subTest` rather than `parametrize`, so the method names stay stable: parametrization renames them, which breaks name-based test selection and disable lists. python test/test_quantization.py TestQuantizedEmbeddingOps Passes on both an AMD GPU (ROCm) and an NVIDIA H100 (sm_90, CUDA 12.4): 15 tests, none skipped. Reviewed By: sacdroid Differential Revision: D115070363 Pull Request resolved: https://github.com/pytorch/pytorch/pull/192573 Approved by: https://github.com/vkuzo, https://github.com/liangel-02 * [Eager] Add FlyDSL native RMSNorm forward override (#191447) Add a FlyDSL-based fused RMSNorm forward kernel as a native op override, on top of the backend in the PR below. ### Supported cases fp16, bf16 and fp32 tensors on gfx950, contiguous input and weight, 1-D `normalized_shape`, and `weight` of shape `(N,)` with matching dtype and device. N-D inputs are logically flattened to `(M, N)`. `eps` is passed through at runtime and does not participate in the cache key. The kernel returns the forward result and the FP32 `rstd` that ATen's backward consumes. ### When FlyDSL acceleration kicks in gfx950, non-copy-on-write contiguous input and weight, one normalized dimension, non-negative `eps`, and: | N | minimum rows M | |---|---:| | `4096 <= N < 8192` | 8192 | | `8192 <= N < 16384` | 4096 | | `16384 <= N <= 114688` | 2048 | ### Test plan: - lintrunner on all changed files - python test/python_native/test_flydsl_rmsnorm_fwd.py - OPINFO_RESTRICT_TO_DSL=flydsl python test/test_ops.py -k rms_norm ### Performance methodology Latency is measured with CUDA events after 10 warmup iterations and over 50 timed iterations. The ATen baseline runs with FlyDSL disabled; the FlyDSL measurement invokes the same operator with the override enabled. Every row below was confirmed to actually dispatch, by checking the compile cache call count advanced, and its output and `rstd` were compared against the ATen baseline. Each row is a single run. Repeating the smallest shape ten times spread its speedup over 1.16x-1.26x. ```python def bench(fn, iters=50, warmup=10): for _ in range(warmup): fn() torch.cuda.synchronize() start = torch.cuda.Event(enable_timing=True) end = torch.cuda.Event(enable_timing=True) start.record() for _ in range(iters): fn() end.record() torch.cuda.synchronize() return start.elapsed_time(end) / iters fwd = lambda: torch.ops.aten._fused_rms_norm(x, [N], weight, eps) with pn.flydsl.disabled(): aten_ms = bench(fwd) flydsl_ms = bench(fwd) speedup = aten_ms / flydsl_ms ``` ### Performance GPU: AMD 256-CU MI355X (gfx950). FlyDSL 0.3.0. | dtype | M | N | ATen (us) | FlyDSL (us) | speedup | |:---|---:|---:|---:|---:|---:| | fp16 | 8192 | 4096 | 26.7 | 22.9 | 1.16x | | fp16 | 4096 | 8192 | 30.1 | 22.6 | 1.33x | | fp16 | 2048 | 16384 | 31.0 | 23.1 | 1.34x | | fp16 | 2048 | 65536 | 173.2 | 112.3 | 1.54x | | fp16 | 2048 | 114688 | 303.2 | 258.8 | 1.17x | | fp16 | 8192 | 4097 | 96.7 | 26.4 | 3.66x | | fp16 | 4096 | 8193 | 82.7 | 26.0 | 3.18x | | fp16 | 2048 | 16385 | 91.9 | 25.9 | 3.55x | | bf16 | 8192 | 4096 | 26.7 | 22.2 | 1.20x | | bf16 | 2048 | 16384 | 30.9 | 22.4 | 1.38x | | bf16 | 8192 | 4097 | 96.4 | 26.4 | 3.65x | | bf16 | 2048 | 16385 | 73.7 | 25.8 | 2.86x | | fp32 | 8192 | 4096 | 56.3 | 40.1 | 1.41x | | fp32 | 4096 | 8192 | 58.5 | 40.7 | 1.44x | | fp32 | 2048 | 16384 | 57.7 | 44.1 | 1.31x | | fp32 | 2048 | 32768 | 154.1 | 105.5 | 1.46x | | fp32 | 2048 | 65536 | 321.8 | 210.3 | 1.53x | | fp32 | 2048 | 114688 | 564.7 | 483.8 | 1.17x | | fp32 | 8192 | 4097 | 91.8 | 41.6 | 2.20x | | fp32 | 4096 | 8193 | 82.5 | 42.3 | 1.95x | | fp32 | 2048 | 16385 | 78.0 | 43.9 | 1.78x | | fp32 | 2048 | 32769 | 212.7 | 105.0 | 2.03x | Aligned shapes improve by 1.2x to 1.5x, with both ends of the band -- `N` = 4096 and `N` = 114688 -- sitting near 1.2x. Off-by-one `N` improves by up to 3.7x. Unsupported shapes retain ATen behavior. Co-authored-by: lirui927 <ruili@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Pull Request resolved: https://github.com/pytorch/pytorch/pull/191447 Approved by: https://github.com/jgong5, https://github.com/slayton58 Co-authored-by: lirui927 <Rui.Li@amd.com> * Extend accelerator graph-capture health check beyond CUDA/XPU (#194362) ## Summary `_accelerator_graph_capture_health_check()` gated its whole body behind `if accelerator and accelerator.type in {"cuda", "xpu"}:`, but everything inside that block is already backend-agnostic: `torch.accelerator.current_stream().is_capturing()` and `accelerator.type.upper()` in the warning/error messages both work for any accelerator already. That gate traces back to #180453, which modernized this function's internals but preserved the previous code's explicit `torch.cuda.is_available()`/`torch.xpu.is_available()` restriction as a faithful port rather than a deliberate decision tied to `is_capturing()`'s semantics. `c10::Stream::is_capturing()` calls `DeviceGuardImplInterface::isStreamCapturing()`, whose base implementation is `return false;` — only `CUDAGuardImpl` and `XPUGuardImpl` override it with real logic (`c10/core/impl/DeviceGuardImplInterface.h`). Every other backend (MPS, MTIA, HPU, PrivateUse1) already safely returns `False` through the same call, so restricting the check to `{"cuda", "xpu"}` served no functional purpose — it just meant this health check (and its `capturable=True`/`False` warnings) silently never ran on any other backend. Replaced the gate with `if accelerator is not None:`, which behaves identically for CUDA/XPU and correctly extends the same already-safe check to every other accelerator. Part of the device-decoupling initiative proposed in #194355. ## Test Plan Verified using: python -c " import torch import torch.optim as optim p = torch.nn.Parameter(torch.randn(3, 3)) opt = optim.Adam([p], capturable=False) opt._accelerator_graph_capture_health_check() print('ok') " `is_capturing()` correctly returns `False` on MPS and the health check runs without error. Ran `lintrunner -a` on `torch/optim/optimizer.py` — no issues. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194362 Approved by: https://github.com/janeyx99, https://github.com/guangyey * Fix positional arg overflow in TritonBenchmarkRequest when template inputs are deduplicated (#188268) # Fix positional arg overflow in `TritonBenchmarkRequest` when template inputs are deduplicated Fixes #188069 When bmm template codegen deduplicates kernel inputs (both operands alias the same buffer, e.g. `einsum("...ab,...Ab->...aA", grad, grad)`), the compiled kernel expects fewer positional parameters than `TritonBenchmarkRequest.make_run_fn()` provides. The excess tensor overflows into the `stream` parameter slot, producing `TypeError: too many positional arguments`. This happens when two compilations run in sequence: the first caches a deduplicated bmm kernel, the second's autotuner reuses that cached candidate but passes more input tensors than the kernel accepts. The CUTLASS path handles a similar case via `dict.fromkeys` on `input_tensor_meta` names, but that doesn't work here because the meta reflects the *current* graph's inputs (2 unique names), not the *cached* kernel's actual param count (1 input after dedup). Tested it and it still crashes 4/5 on 2.12.0. This PR reads `triton_meta["signature"]` from the loaded kernel to get its actual parameter count and trims `input_tensors` to match. The check is a no-op when counts already agree. Tested on RTX 3090, 5 trials per configuration with isolated caches: | | Baseline | Patched | |---|---|---| | 2.12.0+cu130 | 1-2/5 crash | 0/5 | | Nightly 2.14.0.dev20260625 | 0-2/5 crash | 0/5 | Pull Request resolved: https://github.com/pytorch/pytorch/pull/188268 Approved by: https://github.com/jananisriram, https://github.com/jansel, https://github.com/mlazos * Reenable test_addcmul_alpha_one_fma_parity dtypes F32 and F64 on XPU. (#182811) To be merged after https://github.com/intel/torch-xpu-ops/pull/3275 is merged and reference updated. Pull Request resolved: https://github.com/pytorch/pytorch/pull/182811 Approved by: https://github.com/astachowiczhabana, https://github.com/pbielak, https://github.com/janeyx99, https://github.com/guangyey * Revert "Label tests that should run on CuTeDSL or B200 appropriately (#193606)" This reverts commit 926a5451a8eb984d35c4e6b17543aa8e9ad76d7f. Reverted https://github.com/pytorch/pytorch/pull/193606 on behalf of https://github.com/jansel due to Broke smoke_b200 on trunk ([comment](https://github.com/pytorch/pytorch/pull/193606#issuecomment-5443173771)) * [inductor] Fix inductor dropping ordering dep between effectful ops with different kernel types (#188301) `with_effects` records an ordering edge between consecutive ORDERED-effectful ops so the scheduler preserves their program order. That edge was both dropped and, when it did apply, too strong. Fixing the strength then exposed a latent conflation in how pinned buffers are freed. **The edge was dropped.** It was recorded in `V.graph.additional_star_deps` keyed by `new_op.get_name()` (a buffer name), but the scheduler reads it back by operation name. Ops lowering to the same kernel type gave names that coincided, masking the mismatch; different kernel types (e.g. an `ExternKernel` followed by a `_CollectiveKernel`) made the key miss, so the dep vanished and the scheduler could reorder `op_b` ahead of `op_a`. Fixed by keying on `get_operation_name()`. **The edge was too strong.** `additional_star_deps` is installed as a `StarDep`, which asserts ordering *and* liveness, so an effect edge counted as a real read of the previous op's output buffer and held it live past its last true use, blocking reuse and deallocation. An effect edge carries no data, so it belongs in `additional_buffer_deps`, installed as `WeakDep(is_fake=True)`, filtered out of `used_or_aliased_buffer_names`, so it orders without extending lifetimes. This is the pattern `OrderingBarrier` already documents. `additional_star_deps` had no other users and is removed. **Retention is a separate concern, and also needed.** The above is about the *previous* op's output buffer. Independently, *this* op's inputs may be retained in state inductor cannot see (`tq.push(x)` stashes a tensor in a torchbind queue), and no dependency edge can express "the callee kept a reference", so those inputs must not have their storage recycled. The axes are independent: without the pin `test_torchbind_queue` gives silently wrong numerics under either dep strength. **Pinning must block recycling without blocking deallocation,** since freeing only drops inductor's own reference and storage retained behind its back survives. That distinction did not exist: `never_reuse_buffers` both withheld a buffer from the reuse pool and suppressed its free. Putting retention pins there would therefore keep every effectful op's input allocated for the whole graph, and peak memory grows with the number of effectful ops — measured locally, though the memory test is not part of this PR. But the suppression half is load-bearing elsewhere: `ShallowCopyDataKernel` pins its own output, which in the cpp wrapper has no output variable at all (the call is emitted bare), so emitting a free for it generates a `reset()` on an undeclared name and fails to compile. Rather than redefine the existing set, retention pins go in a new `never_reuse_but_free_buffers`, and the split is expressed as `can_free` (body unchanged from the old `can_reuse`, so every pre-existing pin site behaves exactly as before) plus `can_reuse` (`can_free` and not retention-pinned). Review `lowering.py` first for the keying fix, the weakened dep, and the pin; then `scheduler.py` for the dep the switch now goes through, and `graph.py` / `wrapper.py` for the two sets and the `can_free`/`can_reuse` split. ## Test plan `test_inductor_preserves_effect_order_across_kernel_types` covers the keying fix: two ORDERED ops lowering to different kernel types (`op_a` -> `FallbackKernel`, `op_b` -> `_CollectiveKernel`), asserting `op_a` is emitted first. It fails without the fix because the dropped dep lets the scheduler reorder them. `test_inductor_effect_order_does_not_extend_lifetimes` covers the dep strength. To discriminate it needs an effectful op with a real output buffer for a strong dep to over-retain, plus a later effectful op consuming a value derived from that buffer so it cannot be hoisted above the buffer's last use; it then asserts the buffer is freed before that later op. Verified to fail with `is_fake=False` and pass with `is_fake=True`. Pull Request resolved: https://github.com/pytorch/pytorch/pull/188301 Approved by: https://github.com/izaitsevfb * Create a @periodic Decorator (#191946) Solves https://github.com/pytorch/test-infra/issues/8296 Adds a `@periodic` decorator for tests that should run only in periodic CI, including interoperability with `@slowTest`. Periodic Linux, ROCm, Inductor, and s390x workflows now explicitly enable these tests. For now, the only periodic tests are the ones I'm introducing to test the decorator. _Authored with Codex_ Pull Request resolved: https://github.com/pytorch/pytorch/pull/191946 Approved by: https://github.com/izaitsevfb, https://github.com/huydhn * [ROCm] Fixed test_cuda preferred_blas_library_settings for gfx1250 arch (#194785) Fixed test_cuda.py::TestCuda::test_preferred_blas_library_settings Pull Request resolved: https://github.com/pytorch/pytorch/pull/194785 Approved by: https://github.com/jeffdaily Co-authored-by: Jeff Daily <jeff.daily@amd.com> * [c10d] Decouple completion hooks from Work lifetime (#194830) Completion hooks correlated operations through Work addresses, so consumers had to retain each Work to prevent address reuse. Give Work an opaque completion key, use a monotonic key for NCCL2, and pass the key through completion hooks so the flight recorder can release Work references safely. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_fr_hook.py Nccl2FlightRecorderHookTest.test_records_and_retires_collectives Nccl2FlightRecorderHookTest.test_completion_arriving_before_registration_still_retires Nccl2FlightRecorderHookTest.test_duration_ms_from_backend_timing ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194830 Approved by: https://github.com/fduwjj, https://github.com/kapilsh * [c10d] Move NCCL2 completion data into shared state (#194831) NCCL2 completion tracking is owned directly by WorkNCCL, preventing backend tracking from having a narrower lifetime than the caller-visible Work. Move status, events, timing, timeout, and completion notification into a shared State while retaining the existing Work queue behavior. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2Test.test_watchdog_does_not_release_python_backed_tensor ProcessGroupNCCL2Test.test_ephemeral_timeout ``` ```bash env CUDA_VISIBLE_DEVICES=2 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2UninitializedCudaTest.test_barrier_after_init ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194831 Approved by: https://github.com/kapilsh ghstack dependencies: #194830 * [c10d] Share NCCL2 input tensor ownership (#194832) NCCL2 stores input tensors directly on each Work and keeps child Work objects for coalesced operations. Replace both forms with a synchronized shared shelf so input lifetime can later be managed independently of the caller-visible Work. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2Test.test_watchdog_does_not_release_python_backed_tensor ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/194832 Approved by: https://github.com/kapilsh ghstack dependencies: #194830, #194831 * [c10d] Release unused functional Work results (#193239) NCCL2 queues caller-visible Work objects for watchdog and CUDA graph tracking. Those references retain collective tensors after wait_tensor removes the functional registry entry. Queue only shared completion state and an input tensor shelf, and retain only completion state during graph capture. Caller waits clear the shared shelf, including every child of a coalesced operation, while direct Work results remain available. Test Plan: ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv pip install -e . -v --no-build-isolation ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_nccl2.py ProcessGroupNCCL2Test.test_wait_tensor_releases_work_tensors ProcessGroupNCCL2Test.test_wait_tensors_releases_coalesced_work_tensors ProcessGroupNCCL2Test.test_watchdog_does_not_release_python_backed_tensor ``` ```bash env CUDA_VISIBLE_DEVICES=0,1 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_inductor_collectives.py TestCollectivesMultiProc.test_allreduce_inductor_cudagraph_trees ``` ```bash env CUDA_VISIBLE_DEVICES=2,3 NCCL_DEBUG=WARN VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin /home/tristanr/.local/bin/uv run --no-project test/distributed/test_c10d_fr_hook.py Nccl2FlightRecorderHookTest.test_records_and_retires_collectives Nccl2FlightRecorderHookTest.test_completion_arriving_before_registration_still_retires Nccl2FlightRecorderHookTest.test_duration_ms_from_backend_timing ``` ```bash env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/spin quicklint env VIRTUAL_ENV=/home/tristanr/pytorch/.venv PATH=/home/tristanr/pytorch/.venv/bin:/home/tristanr/.local/bin:/usr/local/bin:/usr/bin:/bin .venv/bin/lintrunner -a ``` Authored with an AI assistant. Pull Request resolved: https://github.com/pytorch/pytorch/pull/193239 Approved by: https://github.com/kapilsh ghstack dependencies: #194830, #194831, #194832 * [aoti] Prevent exceptions from crossing the AOTI C ABI (#194613) (#194613) Summary: Prevent exceptions from unwinding out of generated AOTI DSOs without changing the public C API function types. Each exported runtime entrypoint uses a function-try-block that records an escaping exception and returns `AOTI_RUNTIME_FAILURE`; secondary failures while recording the error are contained. Consumers retain potentially throwing function pointers for compatibility with older artifacts. Test Plan: `buck2 build fbcode//caffe2/test/inductor/fb:test_aot_inductor_model_runner_pybind fbcode//sigmoid/core/executor:aoti_model_impl_cpu` `buck2 run fbcode//caffe2/test/inductor:aot_inductor_arrayref_cpu -- caffe2.test.inductor.test_aot_inductor_arrayref.AOTInductorTestABICompatibleCpuWithStackAllocationAndMinimalArrayRefInterface.test_runtime_input_count_error_does_not_cross_c_abi_cpu_with_stack_allocation_and_minimal_arrayref_interface` `buck2 run fbcode//caffe2/test/inductor:aot_inductor_arrayref_cpu_local -- caffe2.test.inductor.test_aot_inductor_arrayref.AOTInductorTestABICompatibleCpuWithStackAllocationAndMinimalArrayRefInterface.test_runtime_input_count_error_does_not_cross_c_abi_cpu_with_stack_allocation_and_minimal_arrayref_interface` Reviewed By: desertfire Differential Revision: D116951104 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194613 Approved by: https://github.com/desertfire * Revert "Enable torch.compile for RNN modules by default (#185902)" This reverts commit a2ed3ae5b69bb2ec3427f3ca18a27ae1c100b229. Reverted https://github.com/pytorch/pytorch/pull/185902 on behalf of https://github.com/jansel due to Causes perf regressions in dashboard ([comment](https://github.com/pytorch/pytorch/pull/185902#issuecomment-5443506876)) * Revert "[CUDA][ROCm] Support pruned lookups in quantized embedding_bag_byte_rowwise_offsets (#192573) (#192573)" This reverts commit a0edfca08d1316d2ce990fc6b892366bb03e1040. Reverted https://github.com/pytorch/pytorch/pull/192573 on behalf of https://github.com/georgehong due to Reverting due to OSS and internal sync issues (ERROR: does NOT match the GitHub patch), will be succeeded by https://github.com/pytorch/pytorch/pull/194934 ([comment](https://github.com/pytorch/pytorch/pull/192573#issuecomment-5443720585)) * [Test] Enable XPU for `test_hop.py` and add hardware classification (#192772) This PR enables XPU coverage for `test_hop.py` by opting `TestHOP` into XPU instantiation with `allow_xpu=True` and classifying it as `ACCELERATOR`. Previously, the file collected zero XPU tests under the XPU CI device filter because `XPUTestBase` was never added. This change ensures the file is exercised on XPU instead of silently being dropped. The tests are already device-generic, so no logic changes were needed beyond the XPU registration/classification. On a single-XPU machine: Before: 44 tests (35 passed, 9 xfailed), all `TestHOPCPU` After: 88 tests (70 passed, 18 xfailed) Test plan: ```bash python -m pytest test/export/test_hop.py -q PYTORCH_TESTING_DEVICE_ONLY_FOR=xpu python -m pytest test/export/test_hop.py -q ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/192772 Approved by: https://github.com/pbielak, https://github.com/etaf, https://github.com/jansel * Reland #188184: derive ROCM_VERSION from torch.version.rocm (#194985) Summary: Relands #188184 with the import-time failure fixed. #188184 changed `ROCM_VERSION` in `torch/utils/cpp_extension.py` to come from `torch.version.rocm` rather than `torch.version.hip`, so extensions can compare against the ROCm release version rather than the HIP runtime version. It was reverted out of Meta's internal monorepo because it raised at module import. `torch.version.rocm` was added by #168097, which wired it through `tools/generate_torch_version.py` and CMake only. A build that produces `torch/version.py` some other way -- Meta's Buck build substitutes a hand-written module -- leaves `rocm` absent or None while `torch.version.hip` is set. #188184 raised `AssertionError` in that state; where the attribute was absent entirely it raised `AttributeError` instead, because `if torch.version.rocm is None:` dereferences the very attribute whose absence it reports, so the assertion message was unreachable. Either way, `torch/testing/_internal/common_utils.py` imports `cpp_extension` at module scope, so this took down test collection rather than a single test. Read `torch/utils/cpp_extension.py` first. `_derive_rocm_version()` prefers `torch.version.rocm` and falls back to the HIP version with a warning instead of raising, which keeps `ROCM_VERSION` set whenever `torch.version.hip` is -- `_get_hipcc_path` compares it with no None guard. That function also moves to the newly added `HIP_VERSION`: it selects between `hipcc.exe` and `hipcc.bat`, which is a HIP SDK layout property, and #188184 left it reading `ROCM_VERSION` after repointing that name at a different quantity. The remaining files put the other readers of the ROCm release version on the same footing. `common_cuda.ROCM_VERSION`, `common_cuda._get_torch_rocm_version()` and the `skip_if_rocm_ver_*` decorators in `common_distributed.py` all derived it from HIP even though their thresholds mirror the C++ `ROCM_VERSION` macro. `blaslt_supported_device()` now uses the shared `common_cuda.ROCM_VERSION` instead of recomputing the same tuple, and keeps its `torch.version.hip` guard; switching that guard, as the original did, makes a build with a falsy `rocm` fall through to `return True` and report hipBLASLt support on every AMD GPU without checking `gcnArchName`. `torch/_inductor/heuristics/template/triton.py` reads the attribute through `getattr`, since its `torch.version.rocm is not None` test is itself an attribute access and raised at import on any build lacking the attribute, CUDA builds included. On shipped ROCm the HIP and ROCm major and minor components agree, so this is behaviour-preserving there. It changes results only on builds where the two diverge, which is what the original change was for. Test Plan: Internal context: D116640565 was the import of #188184 and was backed out by D117079854 after S699381. The internal producer that defines `rocm` in `caffe2/fb/version.py` is already on trunk via D117070955, so this diff needs no `caffe2/fb/**` change -- the consumer-side fallback makes a None `rocm` harmless, which is also why the previously proposed `rocm = hip` producer fallback is not needed. Done locally: - `arc f` and `arc lint` clean. - fbcode and xplat copies verified byte-identical for all six files. - The derivation was exercised over every `(hip, rocm)` state: both set and equal, both set and different, `rocm` None, `rocm` attribute absent, `rocm` empty string, `hip` None, `hip` absent, and a `-rc1` third component. No state raises; a differing `rocm` is preserved rather than overwritten by `hip`. To run before landing (needs AMD hardware): - `buck2 test fbcode//caffe2/test/inductor:max_autotune_amd -- --list-only` - `buck2 test fbcode//deeplearning/aot_inductor/inductor_passes/post_grad/tests:test_move_cat_before_batch_expand_in_bmm_amd -- --list-only` - `buck2 test fbcode//mode/opt-amd-gpu fbcode//dper_lib/silvertorch/core/experimental/mrs_publish/utils/tests:li_utils_test` - the APS checkpoint-compatibility target, testx 844425290139603 - the owning targets for `caffe2/test/test_utils.py` and `caffe2/test/test_linalg.py` The first two are the land-blocking AMD test-listing signals that were red on D116640565 V1 and were landed through anyway. Differential Revision: D117630130 Pull Request resolved: https://github.com/pytorch/pytorch/pull/194985 Approved by: https://github.com/pytorchgreenlight, https://github.com/jeffdaily * Fix nn.functional.threshold OpInfo reference for strict NumPy uint8 casting (#189273) ## Summary Fixes #189267 by updating the `nn.functional.threshold` OpInfo reference inputs for unsigned dtypes. ## Problem NumPy ≥ 2.5 deliberately stopped truncating out-of-range Python ints in `np.where` (the reference path). The threshold OpInfo used `value=-9` against `uint8`, which now raises `OverflowError`. ## Fix Use a non-negative fill (`9`) when `dtype == torch.uint8`; keep `value=-9` for signed and float dtypes to preserve negative-value coverage. This is the code-side fix; #189566 proposes capping numpy `<2.5` in CI requirements as an alternative. Pull Request resolved: https://github.com/pytorch/pytorch/pull/189273 Approved by: https://github.com/jeffdaily * [MPS] enable masked_fill fast path for scalar boolean indexing (#195020) Keeps MPS indexing scalars on CPU, allowing `x[bool_mask] = scalar` to use the existing `masked_fill_` fast path instead of generic `index_put_`. Perf: | dtype | Shape | Layout | Before (us) | After (us) | Speedup | |---|---|---|---:|---:|---:| | bf16 | 256 x 1024 | contiguous | 205.0 | 18.4 | 11.14x | | bf16 | 8 x 512 x 1024 | contiguous | 847.6 | 51.3 | 16.52x | | bf16 | 4096 x 4096 | contiguous | 2291.5 | 535.4 | 4.28x | | bf16 | 2048 x 4096 | transposed dense | 1656.8 | 508.2 | 3.26x | | bf16 | 16 x 128 x 64 x 64 | channels-last | 2602.0 | 962.4 | 2.70x | | bf16 | 4 x 32 x 32 x 32 x 32 | channels-last-3d | 1656.0 | 626.1 | 2.64x | | bf16 | 2048 x 2048 | sliced non-dense | 751.7 | 292.9 | 2.57x | | fp16 | 4096 x 4096 | contiguous | 2375.4 | 544.7 | 4.36x | | fp32 | 4096 x 4096 | contiguous | 2841.0 | 1068.5 | 2.66x | | fp32 | 16 x 128 x 64 x 64 | channels-last | 2639.9 | 989.3 | 2.67x | Smaller shapes | dtype | Shape | Layout | Before (us) | After (us) | Speedup | |---|---|---|---:|---:|---:| | bf16 | 128 | contiguous | 168.7 | 7.0 | 24.24x | | bf16 | 256 | contiguous | 168.4 | 7.5 | 22.61x | | bf16 | 512 | contiguous | 168.3 | 6.7 | 25.19x | | bf16 | 1024 | contiguous | 169.2 | 6.8 | 25.00x | | bf16 | 32 x 32 | contiguous | 168.1 | 5.8 | 28.94x | | bf16 | 64 x 64 | contiguous | 169.7 | 4.0 | 42.95x | | bf16 | 128 x 128 | contiguous | 169.0 | 5.2 | 32.63x | | bf16 | 256 x 256 | contiguous | 173.1 | 6.1 | 28.19x | | fp16 | 128 | contiguous | 168.9 | 2.9 | 57.86x | | fp16 | 256 | contiguous | 169.3 | 3.8 | 45.02x | | fp16 | 512 | contiguous | 168.6 | 2.9 | 57.93x | | fp16 | 1024 | contiguous | 169.3 | 3.0 | 56.05x | | fp16 | 32 x 32 | contiguous | 169.3 | 3.0 | 57.00x | | fp16 | 64 x 64 | contiguous | 169.0 | 2.6 | 64.02x | | fp16 | 128 x 128 | contiguous | 169.3 | 2.7 | 63.42x | | fp16 | 256 x 256 | contiguous | 173.5 | 3.1 | 56.14x | | fp32 | 128 | contiguous | 169.8 | 2.3 | 74.82x | | fp32 | 256 | contiguous | 168.3 | 2.3 | 73.51x | | fp32 | 512 | contiguous | 169.3 | 2.4 | 71.72x | | fp32 | 1024 | contiguous | 169.0 | 2.5 | 68.97x | | fp32 | 32 x 32 | contiguous | 169.0 | 2.5 | 68.15x | | fp32 | 64 x 64 | contiguous | 169.9 | 2.5 | 68.24x | | fp32 | 128 x 128 | contiguous | 168.7 | 2.5 | 66.69x | | fp32 | 256 x 256 | contiguous | 175.1 | 3.2 | 54.71x | Pull Request resolved: https://github.com/pytorch/pytorch/pull/195020 Approved by: https://github.com/malfet * Drop unused includes in c10 (#194959) ## Summary Split out of #194958. Removes unused `#include`s in c10, flagged by clang-tidy's misc-include-cleaner (not enabled by this repo's `.clang-tidy`/lintrunner config). Pull Request resolved: https://github.com/pytorch/pytorch/pull/194959 Approved by: https://github.com/albanD * [Inductor] Reuse DeviceInterface for out-of-tree C++ device options (#190614) Reuse the existing DeviceInterface registration path so out-of-tree backends (e.g. NPU) can provide device-specific C++ compile options without monkey-patching `get_cpp_torch_device_options()`. Changes: - Add an optional `get_cpp_device_options()` method to `DeviceInterface` - `get_cpp_torch_device_options()` queries the registered device interface for device-specific C++ options - Preserve existing behavior when the device interface does not provide C++ options @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo @azahed98 Pull Request resolved: https://github.com/pytorch/pytorch/pull/190614 Approved by: https://github.com/JiasenTian, https://github.com/jansel * Sync requires_grad after _set_tensor_requires_grad under torch.func (#193724) ## Issue Fixes #193279 ## Summary `_create_differentiable` flips `requires_grad` via `_set_tensor_requires_grad`, an in-graph function whose mutation never reached its `TensorVariable`. `call_apply` gates autograd customization on that stale `False` (`misc.py:1118`), so it inlined `forward` and dropped the custom `backward`. The fix re-syncs metadata after fake propagation; the graph node is unchanged. Any `x.requires_grad` read in a transform was stale too - second test. Behaviour change: a custom `vjp` or `jvp` now hits its pre-existing graph break under `func.grad` (`misc.py:1143`, `:1158`); both were previously inlined, so a `fullgraph=True` compile that happened to be correct now errors, as `.backward()` already does. It also fails #193277's new `test_custom_jvp_under_func_grad_does_not_graph_break`; whichever lands second updates it. Happy to rebase behind it (same reporter, forward-mode half). ## Checklist - [x] Passes lint (`spin fixlint`) - [x] Added/updated tests - [ ] Updated documentation (if applicable) - [ ] Included benchmark results (for PRs impacting perf) ## BC-breaking? No --- AI assistance was used (Opus 5) Pull Request resolved: https://github.com/pytorch/pytorch/pull/193724 Approved by: https://github.com/Skylion007 * Move c10/util math helper headers into torch/headeronly (#194673) ## Summary - Move `MathConstants.h`, `copysign.h`, `Load.h`, `BFloat16-math.h`, and `irange.h` from `c10/util` into `torch/headeronly/util` under `torch::headeronly`. - Leave the old `c10/util` paths as one-line forwarding includes, matching `Half.h` / `BFloat16.h`. - Register the new symbols in `torch/header_only_apis.txt` and add header-only aoti_abi_check tests (`copysign`, `irange`, `load`, `MathConstants`, plus `is_reduced_floating_point` coverage in `test_math.cpp`). This is the first leaf-utilities PR toward moving `aten/src/ATen/native/Math.h` and its dependency chain into headeronly. see https://github.com/pytorch/pytorch/pull/194673/changes/1bf583b0c09259f98d88a6c6d9e3a1b216122070 for the updated headeronly diffs. ## Test plan - `TestCopysign.*` - `TestIrange.*` - `TestLoad.*` - `TestMathConstants.*` - `TestMath.TestIsReducedFloatingPoint` ## Test Results Passed Pull Request resolved: https://github.com/pytorch/pytorch/pull/194673 Approved by: https://github.com/janeyx99 * [1/N][Test][Refactor] Make `test_control_flow.py` tests generic. (#190552) This PR is part of refactoring the `test_control_flow.py` test file to be generic. ## Important review notes For easier review, I split changes into several commits, where: | commit | changes | | - | - | | `Only change functions order.` | Only changes the order of test functions - NO FUNCTIONAL CHANGES AT ALL | | `Refactor.` | Applies the real refactor | | ... | Other commits address review comments. | …

## Problem - PyTorch trunk ROCm CI still builds against ROCm 7.14. - ROCm 10.0 packages are published from the new `whl-next` index. - rocSHMEM must build for the requested gfx target instead of substituting gfx90a XNACK variants. ## Changes - Upgrade the shared `rocm-n` trunk image from ROCm 7.14 to 10.0. - Install ROCm 10.0 from `https://stable.repo.amd.com/rocm/whl-next/`. - Cherry-pick the rocSHMEM target fix from ROCm#3583. - Keep wheel, builder, and nightly changes independent in pytorch#194919. ## Validation - Bash syntax and stubbed trunk Docker argument routing pass. - The trunk image selects ROCm 10.0 and the `whl-next` index; the preview image retains its nightly index. - `git diff --check` and targeted IDE lint checks pass for the Docker and CMake changes. - Full image and rocSHMEM builds were not run locally; CI provides end-to-end validation. - [trunk workflow](https://github.com/pytorch/pytorch/actions/runs/33000092249/job/98385742860) ran successfully with ROCm10.0 - [rocm-mi200](https://github.com/pytorch/pytorch/actions/runs/33002840829/job/98340667528) ran successfully with ROCm10.0 Pull Request resolved: pytorch#194794 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily

Build rocSHMEM for the requested gfx target, not gfx90a xnack variants

2e39ed8

pragupta merged commit afe84db into ROCm:release/2.14 Aug 19, 2026
3 checks passed

lucbruni-amd deleted the users/lucbruni-amd/rocshmem-gfx90a-drop-release-2.14 branch August 19, 2026 17:30

ethanwee1 mentioned this pull request Aug 26, 2026

[ROCm][CI] Upgrade trunk to ROCm 10.0 pytorch/pytorch#194794

Closed

Navigation Menu

[release/2.14] Build rocSHMEM for the requested gfx target, not gfx90a xnack variants - #3583

[release/2.14] Build rocSHMEM for the requested gfx target, not gfx90a xnack variants#3583
pragupta merged 1 commit into
ROCm:release/2.14from
lucbruni-amd:users/lucbruni-amd/rocshmem-gfx90a-drop-release-2.14

lucbruni-amd commented Aug 19, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

lucbruni-amd commented Aug 19, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants