Repository navigation
[release/2.15] Use fp32 tolerance for fp32 grads in test_batchnorm_nhwc_cpu (#202210) - #3934
Open
amd-maradosa wants to merge 30 commits into
Open
amd-maradosa wants to merge 30 commits into
amd-maradosa wants to merge 30 commits into
Conversation
## Summary Bringup step for `release/2.15`: add `related_commits` pinning the ROCm dependency builds. - apex: new `ROCm/apex` `release/1.15.0` (cut from master, version 1.15.0) - torchvision: upstream `pytorch/vision` `release/0.29` - torchaudio: new `ROCm/audio` `release/2.11.1`: upstream `pytorch/audio` `release/2.11.1` plus the ROCm build support from `release/2.11.0.3` and the ROCm version check (ROCm/audio#21) Same layout as the `release/2.14` `related_commits` (#3573). Tracking: AIPYTORCH-1326 --------- Signed-off-by: pablo-garay <pgarayfe@amd.com>
## Summary Pin Python dependencies for the `release/2.15` bringup so CI and builds are reproducible, and make everything resolve cleanly on Python **3.10 / 3.11 / 3.12 / 3.13 / 3.14 / 3.15**. Mirrors the 2.14 bringup (#3574), plus the 3.15 handling that is new in this release. Three commits, one per file: - `requirements.txt`: pin all direct dev/install deps (with `python_version` splits for `filelock`, `hypothesis` and `networkx`). - `requirements-build.txt`: pin all build-system deps (with a `python_version` split for `numpy`). - `.ci/docker/requirements-ci.txt`: pin the previously-loose entries (`click`, `psutil`, `pytest-rerunfailures`, `unittest-xml-reporting`, `redis`, `pyzstd` split, `six`, `cuda-bindings`, `cupti-python`, `tqdm`, `uv`, `librosa`) and add `python_version` conditionals for 3.14/3.15. ## Per-version splits (dev/build) - **numpy** (build): `2.2.6` (<3.11), `2.4.6` (==3.11), `2.5.4` (>=3.12). - **networkx** (dev): `3.4.2` (<3.11), `3.6.1` (==3.11), `3.7` (>=3.12). - **filelock** (dev): `4.1.0` (<3.11), `4.2.0` (>=3.11). - **hypothesis** (dev): `6.168.5` (<3.11), `6.169.2` (>=3.11). Every other dev/build dep resolved identically across 3.10-3.15, so those are flat-pinned. ## 3.14 / 3.15 conditionals in requirements-ci.txt Restored from 2.14 (lost in the 2.15 branch): - **pywavelets**: `1.9.0` for 3.14 (1.7.0 has no cp314 wheel). - **fbscribelogger**: guarded to `< 3.14` (its `thriftpy2` dep has no cp314/cp315 wheel and would source-build in the compiler-less test image). New for 3.15 (earliest release with a cp315 wheel): - **numpy** `2.5.2` (the existing `2.4.6` pin for 3.15 has no cp315 wheel), **pandas** `3.0.6`, **optree** `0.20.0`, **pywavelets** `1.10.0`, **aiohttp** `3.14.4`. - **scikit-image**: skipped on 3.15 (no release ships a cp315 wheel yet). ## How versions were resolved Per the documented process: installed each requirements file in a fresh venv inside an `ubuntu:24.04` container with no C compiler (deadsnakes for the non-native Python versions) and pinned the `pip freeze` output. All three files install on **3.10-3.15** with `pip check` clean and no new compiler-dependent source builds (only the pre-existing pure-Python `pwlf` and `coremltools` sdists, plus `pyyaml` on 3.15, which builds without its optional C extension). ## Test plan TheRock CI `multi_arch_build_portable_linux_pytorch_wheels.yml` (gfx94X-dcgpu, ROCm 10.2.0a20261013, `pytorch_git_ref` = `pablo-garay/release-2.15-integration`, which is this PR plus the related_commits from #3816): - [x] Python 3.11: build passed, [run 38705512288](https://github.com/ROCm/TheRock/actions/runs/38705512288) - [x] Python 3.12: build passed, [run 38705525582](https://github.com/ROCm/TheRock/actions/runs/38705525582) - [x] Python 3.13: build passed, [run 38705535781](https://github.com/ROCm/TheRock/actions/runs/38705535781) - [x] Python 3.14: build passed, [run 38705545732](https://github.com/ROCm/TheRock/actions/runs/38705545732) - Python 3.10: n/a. Triton 3.9 requires Python >=3.11, and TheRock dropped 3.10 release builds in ROCm/TheRock#8396 ([run 38705500960](https://github.com/ROCm/TheRock/actions/runs/38705500960) fails installing the triton wheel). - Python 3.15: n/a until ROCm/TheRock#8394 adds 3.15 to the build image ([run 38705556076](https://github.com/ROCm/TheRock/actions/runs/38705556076) fails at setup: `Could not find python: /opt/python/cp315-cp315`). The 3.15 pins were validated by the container install above. --------- Signed-off-by: pablo-garay <pgarayfe@amd.com>
…#3561) (#3832) ## Summary Cherry-pick the hipify `subst`-drive fix into `release/2.15`. It's the same change as `release/2.13` (#3561) and `release/2.14` (f11950f); upstream nightly has it via pytorch#194641. Windows CI builds under a `subst` drive (`B:` aliasing `C:\scratch-b`), so `CUDAExtension` compares `B:\...` with `C:\...` and hipify silently skips torchaudio's GPU sources. Resolving both sides with `os.path.realpath` fixes the match. No effect on Linux. ## Test plan - Failing without this: Windows py3.12 build of `release/2.15` ([run 38716492477](https://github.com/ROCm/TheRock/actions/runs/38716492477)) fails in torchaudio `rnnt/gpu/compute.cu` and `forced_align/gpu/compute.cu` with `'cuda_runtime_api.h' file not found`. - The diff is identical to the 2.14 fix, which builds torchaudio on Windows. - Re-running the Windows build on `pablo-garay/release-2.15-integration` with this commit for ROCm/TheRock#9031. Signed-off-by: pablo-garay <pgarayfe@amd.com>
Cherry-pick of #3808 Co-authored-by: albmalamd <albert.malewski@amd.com>
…dSmiImport (pytorch#200967) (#3846) ## Linked issue or supporting maintainer Supported by @jeffdaily Related: ROCm/TheRock#8888 (the CI lane where the test fails), pytorch#200650 (the first fix for this test) and pytorch#197482, which added the test. ## BC-breaking? No. ## Summary (human written only) Still getting failures after pytorch#200650 landed, this should fix the next failure that was hidden by the first. ## Agent details (if applicable) `test_amdsmi_bare_unversioned_path_uses_loader_first` still fails in the TheRock nightly wheel tests after pytorch#200650, now with exit code 1 during `import torch` in the child: `FileNotFoundError: Could not find rocm library 'amd_comgr' at path '/tmp/.../_rocm_sdk_core/lib'`. `_run_amdsmi_import` writes a stub `_rocm_sdk_core` package into its temp dir, which is the child's cwd and first on `PYTHONPATH`. TheRock wheels preload ROCm libraries from the real `_rocm_sdk_core` during `import torch`, so the stub shadows it and the preload fails. This creates the stub only when no real `_rocm_sdk_core` is importable. When one is installed, its versioned `libamd_smi.so` is the SDK fallback, and the test's assertion (a bare `libamd_smi.so` request uses the loader path first) is unchanged. Upstream CI has no `_rocm_sdk_core`, so nothing changes there. Validation on gfx942, done with Claude, using the TheRock nightly `torch 2.16.0a0+rocm10.2.0a20261014` (git ca883d5) with the test file at that commit. Before: `test_amdsmi_bare_unversioned_path_uses_loader_first` fails with the `amd_comgr` error above and the other four TestAmdSmiImport tests pass. After: all five pass in three runs; the child imports the real `_rocm_sdk_core` (which ships `libamd_smi.so.27`) and the bare request resolves through the loader path. The no-SDK path was not rerun, since the guard is a no-op there. lintrunner clean on main. Pull Request resolved: pytorch#200967 Approved by: https://github.com/jeffdaily Co-authored-by: Jeff Daily <jeff.daily@amd.com>
…ort (pytorch#200650) (#3845) ## Linked issue or supporting maintainer Supported by @jeffdaily Related: ROCm/TheRock#8888 (the CI lane where the test fails) and pytorch#197482, which added the test. ## BC-breaking? No. ## Summary (human written only) `test_amdsmi_bare_unversioned_path_uses_loader_first` (added in pytorch#197482) fails with `Expected 0 but got 127` when the Python running the tests needs `LD_LIBRARY_PATH` to find libpython, as with an interpreter installed by actions/setup-python. `_run_amdsmi_import` replaced `LD_LIBRARY_PATH` with its temp loader directory before spawning `sys.executable`, so the child died in the dynamic loader before running any Python. This prepends the loader directory to the inherited value instead. It is still first on `LD_LIBRARY_PATH`, and the test's assertion (the bare name loads, the SDK fallback is not used) is unchanged. ## Agent details (if applicable) Validation on gfx942 at 6c3af02, done with Claude; the helper and the amdsmi hook are unchanged between that commit and this branch's base. With the stock interpreter all five TestAmdSmiImport tests pass before and after. Under a locally built launcher that reaches libpython only through `LD_LIBRARY_PATH` (it exits 127 without it), `test_amdsmi_bare_unversioned_path_uses_loader_first` fails with `Expected 0 but got 127` before the change and all five pass after. The patched test also passes with a non-loadable `libamd_smi.so` in an inherited `LD_LIBRARY_PATH` directory. This reproduces the failure mode, not the TheRock CI lane itself. lintrunner clean. Pull Request resolved: pytorch#200650 Approved by: https://github.com/jeffdaily Co-authored-by: Jeff Daily <jeff.daily@amd.com>
… libomp140 (#3836) Cherry-picks the two upstream Windows ROCm wheel DLL install fixes onto `release/2.15` (scikit-build-core backend): `aotriton_v2.dll` + `liblzma.dll` (pytorch#200931) and `libomp140.x86_64.dll` (pytorch#191629). Without them, once TheRock's temporary force-includes are removed the DLLs drop out of the wheel and `import torch` fails with WinError 126. The aotriton/liblzma install uses `OPTIONAL` so the existence check runs at install time, because aotriton's ExternalProject only generates those DLLs during the build. ### Validation Branch merged with current `release/2.15` (now includes `related_commits` from #3816). Windows ROCm wheel build (gfx1100, py3.12) on this head (`a119523c5bc`) with ROCm/TheRock#9360, which passes no force-include flags, so every DLL can only ship via CMake install: https://github.com/ROCm/TheRock/actions/runs/39287495719. Build and the `import torch` sanity check passed. Related PR: - pytorch#191629: install `libomp140.x86_64.dll` via CMake - pytorch#200931: install `aotriton_v2.dll` + `liblzma.dll` via CMake - ROCm/TheRock#9055: drop the `_rocm_init.py` wheel force-include - ROCm/TheRock#9058: drop the `aotriton_v2.dll` + `liblzma.dll` wheel force-includes - ROCm/TheRock#9360: drop the `libomp140.x86_64.dll` wheel force-include - #3835: release/2.14 cherry-pick of pytorch#191629 + pytorch#200931 - #3836: release/2.15 cherry-pick of pytorch#191629 + pytorch#200931 (this PR) - Issue: ROCm/TheRock#6965 Tracked in AIPYTORCH-1516.
…ytorch#200682) Cherry-pick of #3796 Co-authored-by: albmalamd <albert.malewski@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
## Summary - Update the apex pin in `related_commits` for the 2.15 release. - Move `ROCm/apex` `release/1.15.0` from `3310832c` to `11867be7`, the `release/1.15.0` backport of [ROCm/apex#341](ROCm/apex#341). It is the only commit between the two. - ROCm/apex#341 removes the hard-coded `-std=c++17` from apex's `op_builder/builder.py`, so PyTorch's `cpp_extension` picks the C++ standard. PyTorch 2.14+ headers require C++20, so with the old pin the apex ops (`amp_C`, `fused_layer_norm_cuda`, ...) fail to build, both when prebuilt and when JIT-compiled on first use. - The torchvision and torchaudio pins in `related_commits` are unchanged. Apex release commit: 11867be79cb16fe53bd8b1b898ea38b7508721ff ## Test plan - [x] Validate the PyTorch build through the Linux and Windows build CI jobs: https://github.com/ROCm/TheRock/actions/runs/40002541100 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
## Linked issue or supporting maintainer Supported by @jeffdaily ## BC-breaking? No ## Summary (human written only) This change fixes communicator teardown hangs and an abort-path resource leak. Since NCCL 2.31, graceful communicator teardown (`ncclCommFinalize`/`ncclCommDestroy`) waits in a barrier for all communicator ranks on the same node. Aborted communicators skip this barrier. - `ProcessGroupNCCL::shutdown()` finalizes communicators one at a time. Ranks may process them in different orders and wait on different communicator barriers forever. This change finalizes them together in one NCCL group. - `_abort_process_group()` did not finalize TorchComms communicators or release NVSHMEM team pools. This change releases these resources after aborting, avoiding collective communicator destruction. If a finalize or release fails, the error is logged and cleanup continues. It also removes the ROCm 10.2 skips that pytorch#199316 added for five teardown tests, because pytorch#200371 and pytorch#200682 already fixed them. ## Agent details (if applicable) > Authored with assistance from an AI coding assistant (Cursor). > > **Current scope after rebasing.** NCCL 2.31 added a barrier among a > communicator's host-local ranks to graceful teardown > (`bootstrapIntraNodeBarrier` in `commDestroySync`). Only an aborted > communicator skips it. RCCL 2.32 contains the same behavior. > > This PR contains three changes: > > 1. `ProcessGroupNCCL::shutdown()` finalizes a group's communicators one at a > time, in `unordered_map` order. Ranks can hold different communicators > (for example, lazily created P2P communicators or one communicator per > device) and enter different teardown barriers. The fix finalizes all > communicators inside one `ncclGroupStart/End`, as > `abortCommsFromMap()` already does for abort. > 2. `_abort_process_group()` did not release the TorchComms comms or NVSHMEM > team pools released by `destroy_process_group()`. It now finalizes > TorchComms and releases NVSHMEM teams after aborting, in both the > single-group and WORLD paths. Finalizing an already-aborted TorchComm > skips its collective communicator destroy and frees only local resources. > Releasing NVSHMEM after the abort matters because the release may > synchronize the device. The TorchComms finalize loop is shared with > `destroy_process_group()` and still deduplicates comms by identity. > Post-abort cleanup is best-effort: a TorchComms finalize failure (for > example a timed-out work item) or an NVSHMEM release failure (for example > a sticky CUDA error from the release's device synchronize) is logged per > communicator or group, and the remaining cleanup still runs. > `destroy_process_group()` stays fail-fast. > 3. pytorch#199316 skipped five teardown tests on ROCm >= 10.2 (ROCM-31513). The > failures they cover were already fixed in the tests by pytorch#200371 (nccl2 > timeout tests) and pytorch#200682 (monitored_barrier allreduce-hang tests), so > the skips are removed. With the skips removed, all five pass on 2 and > 8 MI300X GPUs, with or without the rest of this PR. Reverting the two > test fixes brings the hang and the `ncclRemoteError` failure back. > > **Related upstream work by other contributors.** While this PR was under > review, upstream/main gained the overlapping fixes: > > - [@fduwjj's pytorch#194775](pytorch#194775), > landed as `32575604cf2`, aborts the original group for `SHRINK_ABORT` and > documents that excluded ranks must also abort or exit. > - [@d4l3k's pytorch#200371](pytorch#200371), > landed as `e63e88f27ce`, aborts timed-out nccl2 subgroups on every rank in > the timeout tests. > > The equivalent two commits were dropped from this PR during the rebase. > Consequently, this PR no longer introduces the `SHRINK_ABORT` behavior > change and is not BC-breaking. > > **Validation.** Before the rebase, the same changes were source-built for > gfx942 and tested on 8x MI300X with RCCL 2.32.3. > > - Without grouped finalize, the three `test_send_recv_nccl` variants timed > out after 300 seconds and `test_nccl_backend_4gpu_module` timed out after > 300 seconds. With it, they passed in 75 seconds total and 42 seconds, > respectively. > - The new > `test_abort_finalizes_comms_and_releases_nvshmem_teams` verifies subgroup > and WORLD cleanup order: abort, then TorchComms finalize, then NVSHMEM > release. It fails if cleanup is reverted or moved before the abort. > - `test_abort_cleanup_continues_when_finalize_and_release_raise` verifies > that a `"Work timed out during finalize"` exception and a > `"CUDA error: unspecified launch failure"` from the NVSHMEM release are > logged without preventing subgroup or WORLD state cleanup. > - After rebasing, all seven focused TorchComms destroy/abort unit tests pass. > > ``` > cd test > TEMP_DIR=$T BACKEND=nccl WORLD_SIZE=8 python -m pytest distributed/test_distributed_spawn.py -v -s -k test_send_recv_nccl > python -m pytest distributed/test_c10d_nccl.py -v -s -k test_nccl_backend_4gpu_module > python -m pytest distributed/test_c10d_torchcomms.py -v -k TestC10dTorchCommsDestroyDedup > ``` > > An RCCL-only reproducer confirms the grouped-finalize mechanism: two ranks > finalizing two communicators in opposite orders hang, while the same calls > inside `ncclGroupStart/End` complete. > > **Not covered:** no CUDA run here. `spin fixlint` was unavailable in the > agent environment; IDE diagnostics and `git diff --check` report no errors. Pull Request resolved: pytorch#200673 Approved by: https://github.com/jeffdaily (cherry picked from commit adc4a41) Co-authored-by: Cursor <cursoragent@cursor.com>
…dsmi (pytorch#200937) (#3848) ## Linked issue or supporting maintainer Supported by @jeffdaily ## BC-breaking? No. ## Summary (human written only) On ROCm, the test wrongly treats normal RCCL memory usage as an extra GPU context. This causes failures even when everything works correctly. The change uses AMD SMI to directly check which processes use the GPU. If AMD SMI is unavailable, the unreliable memory check is skipped. ## Agent details (if applicable) The following investigation and implementation details were produced with Cursor and must be reviewed by the author before submission. On ROCm, the fallback used when `pynvml` is unavailable produces both false positives and false negatives: - The first RCCL kernel launch allocates approximately 1.1 GiB of queue scratch on rank 0's own device. This can exceed the fallback's 70% growth threshold even though no other rank created a context there. - A real extra context created by rank 1 on device 0 can consume far less than 1.1 GiB, so the same heuristic can miss the condition under test. The change uses `amdsmi_get_gpu_process_list`, via `torch.cuda._get_amdsmi_handler`, as the ROCm equivalent of the existing NVML process-count check. If `amdsmi` is unavailable, the test skips on ROCm instead of executing the invalid memory heuristic. CUDA behavior is unchanged. The rejected alternatives were increasing the threshold and warming up RCCL before taking the baseline. Both would hide the symptom without directly checking whether another process created a context. Validation was performed on MI300X GPUs with no foreign VRAM users, using the ROCm 2.13 nightly wheel: - upstream stock helper: failed 2/2 because of the memory heuristic; - patched helper with `amdsmi`: passed 3/3; - patched helper with an intentionally injected extra context: failed 2/2 and reported two processes on device 0; - patched helper without `amdsmi`: skipped as intended. The equivalent release/2.13 change was also tested directly: - stock: failed 2/2 with a memory delta byte-identical to the reported CI failure; - patched with `amdsmi`: passed 5/5; - intentionally injected extra context: failed 3/3. Related ROCm CI failure: https://github.com/ROCm/rockrel/actions/runs/37731484533 Pull Request resolved: pytorch#200937 Approved by: https://github.com/jeffdaily (cherry picked from commit 7c67769) Co-authored-by: Cursor <cursoragent@cursor.com>
…imeout tests (pytorch#200371) (#3854) Fixes pytorch#200326 Root cause: in the nccl2 timeout tests rank 0's communicator times out and is aborted, while the other ranks call destroy_process_group(). NCCL 2.31+ runs bootstrapIntraNodeBarrier in commDestroySync unless the comm was aborted, so the survivors wait forever for rank 0 at the teardown barrier and the test hangs until the class timeout. Fix: every rank aborts the timed-out group via _abort_process_group() instead of a graceful destroy, in the five tests that leave a timed-out subgroup behind. Test plan: PYTHONPATH=/tmp/flk/nccl232 python test/distributed/test_c10d_nccl2.py -k WatchdogNoTearDown -k BlockingWait -k DumpOnTimeout -k DumpTimeoutBound # NCCL 2.32.3, 15x pass; hung before python test/distributed/test_c10d_nccl2.py -k WatchdogNoTearDown -k BlockingWait -k DumpOnTimeout -k DumpTimeoutBound # NCCL 2.30.7, 5x pass Authored with an AI assistant (Claude). Pull Request resolved: pytorch#200371 Approved by: https://github.com/fduwjj (cherry picked from commit e63e88f) Co-authored-by: Tristan Rice <rice@fn.lc>
…200054) (#3868) Cherry-picks two upstream commits onto release/2.15 to bring ROCm native AsyncTP support (the CK-tile fused_all_gather_matmul persistent-GEMM path) to this branch: - pytorch#177961 "[ROCm] Enable native AsyncTP" -- adds the CK-tile GEMM kernel (hip/AsyncMM.hip) and ROCm host dispatch for the native AsyncTP path of fused_all_gather_matmul, gated to gfx942/gfx950. - pytorch#200054 "[ROCm] Build AsyncMM.hip only for gfx942/gfx950" -- follow-up fix for multi-arch builds (e.g. TheRock's nightly arch list): the CK-tile headers AsyncMM.hip includes don't compile for every arch, so this passes --no-offload-arch=<arch> for every arch outside gfx942/gfx950. A third commit on top fixes two lint-only issues introduced while cherry-picking pytorch#177961's import-list diff against release/2.15's already-diverged test_symmetric_memory.py: a duplicate instantiate_parametrized_tests import (RUFF F811) and an unused get_cycles_per_ms import pulled in as unrelated diff context (RUFF F401). No functional changes. Both cherry-picks were verified hunk-for-hunk against their upstream commits and are otherwise byte-identical. Same pattern as #3840 (the release/2.14 backport of the same two commits). ## Test plan Built with `pip install -e . -v --no-build-isolation`, `PYTORCH_ROCM_ARCH=gfx950`, on an 8x MI355X (gfx950) box, ROCm 7.17. Confirmed `USE_ROCM_ASYNC_MM` compiled into `build/compile_commands.json`. Ran `AsyncTPTest` in test/distributed/test_symmetric_memory.py at world size 2, 4, and 8 (`HIP_VISIBLE_DEVICES=0,1` / `0,1,2,3` / `0,1,2,3,4,5,6,7`): ``` python test/distributed/test_symmetric_memory.py -k AsyncTPTest ``` 25 passed / 1 skipped (multimem, no multicast support on this box) at each world size, repeated 3x at ws=8. `lintrunner` on the touched files: clean. ## Agent details This PR was prepared with an AI assistant (Claude Code): rebasing the cherry-picks onto the current release/2.15 tip, running the build and the AsyncTPTest suite at ws=2/4/8, and drafting this description. All test runs and lint were verified directly, not taken on faith from the assistant. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Albert Malewski <albert.malewski@amd.com> Co-authored-by: Marko Radosavljevic <maradosa@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
… skip in test_c10d_pypg (pytorch#201852) (#3880) ## Linked issue or supporting maintainer Supported by @jeffdaily ## BC-breaking? No ## Summary (human written only) This PR fixes two failures in test/distributed/test_c10d_pypg.py: 1. The test class `MyWork` overrides `dist._Work.wait` as `wait(self, timeout)` without a default value, while the base method defaults timeout to 0. The manual backward finalization tests (pytorch#196651) call `dist.all_gather`, which calls `work.wait()` with no arguments and fails with `TypeError: MyWork.wait() missing 1 required positional argument: 'timeout'`. The fix gives the override the same default: `timeout=timedelta(0)`. 2. DDP pickling is only supported with the default process group, so pytorch#199109 added an override that skips the inherited `test_manual_backward_finalization_pickle_compatibility` in these classes. pytorch#199110 then renamed the inherited test in `test_c10d_common.py` to `test_ddp_pickle_compatibility`, so the override no longer matches and the test fails with DDP `Pickling/Unpickling are only supported … default process group`. The fix renames the override to match. ## Agent details (if applicable) > The text below was generated by an AI assistant (Claude Code) and reviewed by the author. > AbstractDDPSingleRank inherits the DDP tests from > CommonDistributedDataParallelTest, and two problems surfaced there. > > First, MyWork overrides dist._Work.wait as `wait(self, timeout)`, but the > binding declares `timeout=kNoTimeout`. The manual-backward-finalization tests > (pytorch#196651) call dist.all_gather on the process group, which calls > `work.wait()` without arguments, so > test_manual_backward_finalization_with_comm_hook and > test_mixed_precision_uses_post_backward_finalization fail on > TestDDPWithWorkSubclass with > `TypeError: MyWork.wait() missing 1 required positional argument: 'timeout'`. > Give the override the base class default. > > Second, pytorch#199109 skipped the inherited pickle test for these classes, because > DDP pickling requires the default process group. pytorch#199110 then renamed the > inherited test from test_manual_backward_finalization_pickle_compatibility to > test_ddp_pickle_compatibility, so the override no longer shadows it and both > TestDDPWithWorkSubclass and TestDDPWithWorkWrapper fail with > "DDP Pickling/Unpickling are only supported when using DDP with the default > process group". Rename the override to match. > > Test Plan: > > Validated on ROCm release/2.15 (560e250), where both test files match main > at the changed lines, on 8x MI300X. Before the change the four tests above fail; > after it: > > ``` > python -m pytest test/distributed/test_c10d_pypg.py > # 97 passed, 2 skipped > ``` > > This PR was authored with an AI assistant (Claude Code). > 🤖 This PR was authored with an AI assistant (Claude Code) Pull Request resolved: pytorch#201852 Approved by: https://github.com/jeffdaily (cherry picked from commit 04706f2) Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…201949) (#3881) Go through the following checklist: - [x] Passes lint (`lintrunner -a`) - [x] Added/updated tests ## Summary (human written only) Sort split ranks did not account for preserve_rank_order correctly and sorted subgroup ranks by default. This was happening in eager path ## Agent details (AI-generated) ### Root cause The `new_group` DeviceMesh path sorts subgroup ranks by default and preserves mesh order when `preserve_rank_order=True`. The eager path instead passed mesh order directly to `split_group`, which intentionally preserves the supplied order. Default subgroup order therefore depended on the initialization path for every backend that supports splitting. ### Fix Sort ranks before every `split_group` call when `preserve_rank_order=False`. Keep mesh order when `preserve_rank_order=True`, matching the documented DeviceMesh contract and the existing `new_group` path. ### Behavior change Eager legacy NCCL, Gloo, XCCL, NCCL-lazy, TorchComms, and custom split backends now sort permuted subgroup ranks by default. This aligns them with non-eager initialization. Callers that need mesh order should pass `preserve_rank_order=True`. ### Validation The rank-order tests verify: - Default subgroup ranks are sorted. - `preserve_rank_order=True` follows mesh order. - The local rank matches the corresponding mesh coordinate. It covers the configured eager backend, explicit eager legacy NCCL, bound eager Gloo, and the existing eight-rank non-eager case. All four tests pass with the fix. Removing the sorting line makes both the legacy NCCL and Gloo regressions fail with reversed local ranks. ```bash CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7 agent_space/device-mesh-nightly/bin/python test/distributed/test_device_mesh.py DeviceMeshTest.test_device_mesh_preserve_rank_order_eager DeviceMeshTest.test_device_mesh_preserve_rank_order_eager_nccl_legacy DeviceMeshTestGlooBackend.test_device_mesh_preserve_rank_order_eager DeviceMeshTestNDim.test_device_mesh_preserve_rank_order -v ``` ```bash .venv/bin/spin quicklint ``` Pull Request resolved: pytorch#201949 Approved by: https://github.com/d4l3k (cherry picked from commit 020ab1c) Co-authored-by: Kapil Sharma <kapilsharma.frm@gmail.com>
… (#3882) Partial cherry-pick of c489527 ("[ROCm] Bump trunk ROCm version to 10.1"): only the test/distributed/_composable/fsdp/test_fully_shard_overlap.py hunk. ROCm 10.1 removes high priority queue, which reduces the total number of queues from 8 (4 high + 4 normal) to 4 (4 normal), causing some streams can't overlap. (cherry picked from commit c489527) Co-authored-by: Z Liu <zhuojin.liu@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
…3879) Backported [pytorch#199899](pytorch#199899) for ROCM-31963 Co-authored-by: Cursor <cursoragent@cursor.com>
…nels for the device (pytorch#200850) (#3872) ## Linked issue or supporting maintainer Supported by @jeffdaily ## BC-breaking? No ## Summary (human written only) **The issue** `TestCuda::test_graph_capture_tunableop_cublas_workspace` fails consistently on gfx103X (failed every retry, FAILED CONSISTENTLY). It raises a hipBLASLt runtime error on a plain matrix multiply. Context: ROCm/TheRock#8782 ## Agent details (if applicable) > Drafted with an AI assistant (Claude); reviewed by the author. > > **Problem.** With TunableOp enabled, the first matmul raises on GPUs that hipBLASLt ships no kernels for (e.g. gfx103X): > > ``` > RuntimeError: hipblaslt error: HIPBLAS_STATUS_INVALID_VALUE when calling `hipblaslt_ext::getAllAlgos(...)` > ``` > > TunableOp lists hipBLASLt candidates with `getAllAlgos` and throws on any error, before the rocBLAS and default candidates are registered. The non-tunable path already avoids hipBLASLt on these GPUs via `Context::blasPreferredBackend()`. On gfx1036, 29 of the 30 non-skipped `test_linalg -k tunable` tests fail this way, as does `test_graph_capture_tunableop_cublas_workspace` in CI. > > **Fix:** > - `INVALID_VALUE` from `getAllAlgos` means no hipBLASLt candidates; other statuses still raise. In hipBLASLt this status only comes from a missing kernel library; an uncovered type returns success with zero algorithms. The device is remembered so it isn't queried again. > - `getAllAlgos` uses the pooled handle `HipblasltGemmOp::Call` already uses, instead of a private handle that leaked when the check threw. > - A handle can't be created mid-capture, so during stream capture TunableOp only asks the pool for an existing one: the thread's own handle for that stream, or a spare that `CUDAGraph::capture_begin` stocked in advance. If neither exists, it offers no hipBLASLt candidates and the GEMM is tuned with rocBLAS. For this, `DeviceThreadHandlePool::reserve` gets a `create` flag (default `true`), exposed as `getCurrentCUDABlasLtHandleIfAvailable()`. Outside capture nothing changes. > - A warning is raised only if `PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED=1` was set explicitly, as in `blasPreferredBackend()`. > > Gating on `getHipblasltSupportedArchs()` was the alternative, both for the fix and for the capture check. It was rejected because that list doesn't match what hipBLASLt ships, and it would drop hipBLASLt tuning on a supported GPU whenever an unsupported one is visible, rather than only when the unsupported GPU is the first to use a GEMM type. > > **Testing** (ROCm 10.1; gfx1036, gfx1100 and gfx1201 in one system): > > | | gfx1036 | gfx1100 | gfx1201 | > |---|---|---|---| > | `test_linalg -k tunable`, before | 29 failed | 29 passed | 34 passed | > | `test_linalg -k tunable`, after | 29 passed | 29 passed | 34 passed | > | `test_graph_capture_tunableop_cublas_workspace`, after | passed | passed | skipped (hipBLASLt wins) | > > `test_call_count_tunableop` fails on all three, before and after. Supported GPUs match an unpatched nightly, including the fp8 scaled-GEMM TunableOp tests on gfx1201. > > The failure reproduces on any GPU by pointing hipBLASLt at an empty kernel directory, which could serve as a CI regression test: > > ``` > HIPBLASLT_TENSILE_LIBPATH=/tmp/empty PYTORCH_TUNABLEOP_ENABLED=1 \ > python -c "import torch; a = torch.randn(512, 512, device='cuda'); a @ a" > ``` 🤖 Generated with [Claude Code](https://claude.com/claude-code) Pull Request resolved: pytorch#200850 Approved by: https://github.com/jeffdaily Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> (cherry picked from commit be1b713) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…lized (pytorch#202072) (#3883) ## Linked issue or supporting maintainer Supported by @jeffdaily ## BC-breaking? No ## Summary (human written only) `split_group(new_pg, ...)` rejects a group created by `new_group(..., device_id=dev)` with the error `Parent process group backend is not initialized`, although its communicator is already usable. The reason is that `NCCLComm::split` doesn't set the `initialized_` flag (which `NCCLComm::create` does), so a communicator created by `ncclCommSplit` in blocking mode appears uninitialized until the first collective, and the check introduced in pytorch#195758 rejects it. The fix sets the flag the same way `create()` does, `comm->initialized_ = !comm->nonBlocking_;`, so a blocking communicator is marked ready immediately and a non-blocking one behaves as before. ## Agent details (if applicable) > The text below was generated by an AI assistant (Claude Code) and reviewed by the author. > NCCLComm::create sets `initialized_ = !nonBlocking_`: a blocking ncclCommInit > returns a ready communicator. NCCLComm::split never sets the flag, so a > blocking child created by ncclCommSplit reports isInitialized() == false until > something calls getNcclComm(), which happens on the first collective. > > Since pytorch#195758, ProcessGroup::splitGroup rejects a parent whose backend is not > initialized. So with the legacy backend, a group created eagerly by splitting > the bound default group (new_group(..., device_id=dev)) cannot itself be split > until a collective has run on it. > test_comm_recursive_split_group (NCCL_BACKEND = "nccl-legacy") fails on its last > split_group(new_pg, ...) with "Parent process group backend is not > initialized; pass device_id when creating it or run a collective before > split_group". Its earlier splits pass only because a broadcast runs on each > child first. > > Set initialized_ in split the same way create does. Non-blocking children > still start uninitialized and are marked ready by getNcclComm() after > waitReady(), as before. > > Test Plan: > > Validated on ROCm release/2.15 (560e250), where NCCLUtils.cpp and > ProcessGroup.cpp match main at the relevant lines, on 8x MI300X with RCCL > 2.32.3. Without the change, `new_pg._get_backend(dev)._is_initialized()` is > False right after new_group(..., device_id=dev), and the test fails on all > 8 ranks. With it, the flag is True and: > > ``` > python -m pytest test/distributed/test_c10d_nccl.py -k test_comm_recursive_split_group > # 1 passed > python -m pytest test/distributed/test_c10d_nccl.py -k "split or shrink" > # 20 passed, 1 failed (test_shrink_group_flags, which times out identically without this change) > ``` > > This PR was authored with an AI assistant (Claude Code). > This PR was authored with an AI assistant (Claude Code). Pull Request resolved: pytorch#202072 Approved by: https://github.com/jeffdaily, https://github.com/kapilsh, https://github.com/Skylion007 (cherry picked from commit 1aa1528) Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…lone on ROCm wheels (pytorch#201771) (#3869) ROCm wheels setenv() ROCPROFILER_REGISTER_LIBRARY during GPU init. The child inherits it: test_region_device_time records no kernels, and the test_load_standalone binary aborts before main. Pass an environment rebuilt from os.environ, same as pytorch#197421. Test Plan: ``` python test/profiler/test_profiler.py TestProfiler.test_region_device_time python test/test_utils.py TestStandaloneCPPJIT.test_load_standalone ``` Pull Request resolved: pytorch#201771 Approved by: https://github.com/jeffdaily (cherry picked from commit 95f55ae)
…ce subprocesses (pytorch#201009) (#3886) Supported by @jeffdaily No On TheRock CI, all 9 TestExportOnFakeDevice tests exit 127 because the test replaces the subprocess env with only CUDA_VISIBLE_DEVICES="". That drops the LD_LIBRARY_PATH that Rock's setup-python interpreter needs to find libpython. This change inherits the parent env, sets CUDA_VISIBLE_DEVICES="" and removes HIP_VISIBLE_DEVICES and ROCR_VISIBLE_DEVICES, which TheRock CI sets; ensuring no GPUs are visible and preserving the test's original intent N/A Pull Request resolved: pytorch#201009 Approved by: https://github.com/jeffdaily (cherry picked from commit a60aa50) Co-authored-by: Chinmay Kuchinad <ChinmayDattanand.Kuchinad@amd.com>
…97248) (#3889) ## Issue Fixes pytorch#197246. ## Summary `xfailCUDAIfSM89OrLaterOnWindows` was introduced for NVIDIA Windows failures, but ROCm gfx capability tuples also satisfy `SM89OrLater`. Gate the decorator on `not TEST_WITH_ROCM` so passing Windows ROCm tests are no longer turned into unexpected-success failures. Add a regression test for the Windows + ROCm + capability >= 8.9 case. ## Test Plan - `python -m py_compile test/test_testing.py torch/testing/_internal/common_cuda.py` - `git diff --check` - The focused runtime test requires a built PyTorch tree; this checkout does not contain generated `torch/version.py`, so CI will execute it. ## Checklist - [ ] Passes lint (`spin fixlint`) — lintrunner is not installed locally. - [x] Added/updated tests - [x] Updated documentation (not applicable) - [x] Included benchmark results (not applicable; test-infrastructure-only change) ## BC-breaking? No. ## AI assistance disclosure > Tool: OpenAI Codex > Model: GPT-5.6-sol > Scope: root-cause analysis, implementation, regression test, and drafting the issue and PR text. The contributor reviewed the failure analysis, selected this targeted fix, and is submitting the change under their responsibility. Pull Request resolved: pytorch#197248 Approved by: https://github.com/jeffdaily (cherry picked from commit aad4504) Co-authored-by: Filip Jankovic <filip.jankovic@amd.com>
…freed (pytorch#202143) (#3888) Supported by @jeffdaily No Dynamo checks that the saved tensor hooks haven't changed by comparing their addresses (id()), but a freed hook's address can be reused by a new, different hook, so Dynamo can run a graph compiled for the old hooks. This PR registers both hooks with id_ref, as ID_MATCH guards do, so freeing either hook invalidates the cache entry, and adds a test for it. Used claude to run the pr-review skill internally to verify the diff. The regression test was added as per the review suggestion. Pull Request resolved: pytorch#202143 Approved by: https://github.com/jeffdaily Backport to release/2.15: this branch still checks the hook ids in a Python lambda guard, so id_ref is applied there, and the guards.cpp comment hunk is dropped because this branch has no C++ AUTOGRAD_SAVED_TENSORS_HOOKS guard. (cherry picked from commit a9261ad)
…the caching allocator (pytorch#199196) (#3898) Batched real eigh on ROCm calls rocsolver_*syevd_strided_batched. The solver handle starts with rocBLAS-managed device memory, and growing that workspace uses hipMalloc. When the PyTorch caching allocator already holds the device, hipMalloc fails with rocblas_status_memory_error. rocblas_start_device_memory_size_query resets the query accumulator to 0, so rocblas_status_size_increased only means this call needs more than 0 bytes. It is not a comparison with the workspace installed on the handle. Query the size, allocate that many bytes with getCUDADeviceAllocator() whenever the queried size is nonzero, and pass the pointer to rocblas_set_workspace. rocSOLVER still determines the required size, while PyTorch controls the allocation. Reset the handle to (nullptr, 0) before that block is freed so the pooled solver handle does not keep the pointer. This covers workspace growth after the rocBLAS GEMM code objects are already loaded. A warm-up at a small size followed by a larger eigh can fail for a different reason: the first use of a new GEMM kernel loads a code object while the device is full and returns rocblas_status_internal_error. Binding a workspace does not fix that. ```bash python test_linalg.py -v -k test_eigh_workspace_growth_with_full_caching_allocator_cuda ``` Without the fix that test fails with rocblas_status_memory_error. Example to reproduce the workspace-growth failure on ROCm. Warm up batch-4 eigh at n=1024, fill the device and free those blocks without empty_cache, then run batch-4 eigh at n=2048. On main this raises rocblas_status_memory_error. With this change it succeeds. ```bash python - << 'PY' import torch def spd(n): a = torch.randn(4, n, n, device="cuda") return a @ a.mT + torch.eye(n, device="cuda") warmup = spd(1024) torch.linalg.eigh(warmup) torch.cuda.synchronize() del warmup chunks = [] try: for size in (1 << 30, 64 << 20, 2 << 20, 256 << 10): try: while True: chunks.append(torch.empty(size, dtype=torch.uint8, device="cuda")) except torch.OutOfMemoryError: pass del chunks chunks = [] values, _ = torch.linalg.eigh(spd(2048)) torch.cuda.synchronize() assert not torch.isnan(values.cpu()).any() print("passed") finally: del chunks torch.cuda.empty_cache() PY ``` Pull Request resolved: pytorch#199196 Approved by: https://github.com/jeffdaily (cherry picked from commit b15a13a)
…lit_flex_attention` (pytorch#198273) seem to be reduction differences causing 1/2048 mismatch authored with codex Pull Request resolved: pytorch#198273 Approved by: https://github.com/liangel-02 (cherry picked from commit 1edde00)
…AS library cannot be selected (pytorch#201047) (#3860) Fixes ROCm/TheRock#8781 Supported by @jeffdaily ## Issue `test_cublas_workspace_cached_lazy_reallocation` and `test_cublaslt_workspace_eager_resize_and_zero` (enabled on ROCm in pytorch#196920) fail consistently on Windows ROCm nightly: ``` test_cublas_workspace_cached_lazy_reallocation_backend_cublas (gfx110X, gfx1151, gfx120X) AssertionError: 83887104 not greater than 83887104 test_cublas_workspace_cached_lazy_reallocation_backend_cublaslt (gfx103X) AssertionError: 33555456 not greater than 33555456 test_cublaslt_workspace_eager_resize_and_zero (gfx103X) AssertionError: 32505856 not less than 524288 ``` ## Details The tests assume `torch.backends.cuda.preferred_blas_library(backend)` takes effect. It doesn't in two cases: - On Windows the setter only warns (pytorch#125080), so the arch default runs: hipBLASLt on gfx1100/1101/1151/1200/1201, hipBLAS elsewhere. - On archs without hipBLASLt (e.g. gfx103X), cublaslt falls back to hipBLAS, on Linux too. The test then resizes the workspace of a library that isn't running, so memory never changes. ## Fix Add `skip_unless_blas_library_selectable(backend)` to `test/test_cuda.py`. It requests the backend, reads back the one that will actually run, and skips the test if they differ. Where the backend can be selected, the tests run unchanged. ## Testing Windows 11, torch 2.16.0a0+rocm10.2.0a20261014: | | gfx1100 | gfx1036 | |---|---|---| | `..._lazy_reallocation_backend_cublas` | skipped | passed | | `..._lazy_reallocation_backend_cublaslt` | passed | skipped | | `test_cublaslt_workspace_eager_resize_and_zero` | passed | skipped | | full `test/test_cuda.py` | no new failures | no new failures | Pull Request resolved: pytorch#201047 Approved by: https://github.com/jeffdaily Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
## Human Written Summary For PyTorch 2.15, use the ROCm Triton repository and align the pinned ROCm Triton 3.9 revision with the current tip of `ROCm/triton`'s `release/internal/3.9.x` branch. This advances the pin over the ROCm Triton 3.9 release commits that add gfx1250-strict support and improve the TDM path, while retaining ROCm Triton version `3.9.0`. ## Summary - Change the Triton installer source from `triton-lang/triton` to `ROCm/triton`, matching `release/2.14`. - Advance the ROCm Triton pin from `aad2a60d958f945420c5c40936b37d411fb02e89` to `0a6e13e6e31f18fa75a17ac72dee0dd1bcb507b8`. - Keep `.ci/docker/triton_version.txt` at `3.9.0`. - No block-pointer compatibility backport is needed: `release/2.15` already contains upstream PyTorch commit `7b45c7a09bf79ec4e2826d324714c523147c0e75` from pytorch#191007. ## Notable ROCm Triton changes included The old pin is a direct ancestor of the new pin. The seven intervening commits include: - gfx1250-strict target support. - The AMD TDM shared-order fix that ignores unit dimensions. - ROCm Triton 3.9 compatibility, release, and deprecation-warning fixes. ## Validation - Verified `ROCm/triton` `release/internal/3.9.x` points exactly to `0a6e13e6e31f18fa75a17ac72dee0dd1bcb507b8`. - Verified the old pin is an ancestor of the new pin with seven intervening commits and no divergence. - Verified the ROCm Triton repository URL and target branch resolve successfully. - Verified the installer source matches ROCm PyTorch `release/2.14`. - `bash -n .ci/docker/common/install_triton.sh` passed. - `git diff --check` passed. - No linter errors were reported for the changed files. - No build or runtime tests were run; these changes only update the installer source and commit pin. ## Related Similar to #3831 for `release/2.14`. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…rch lists (#3924) ## Linked issue or supporting maintainer Supported by @jeffdaily ## BC-breaking? No. ## Summary Add gfx1170 and gfx1171 to the supported hipBLASLt architecture lists starting with ROCm 10.3. Enable dense FP8 scaled matrix multiplication for these targets and update the corresponding test capability detection. AOTriton (SDPA) support for gfx117x needs an AOTriton bump and will come in a separate PR. ## Test Plan - Built Triton and PyTorch wheels for gfx1170 and gfx1171. - Verified that the complete wheel build succeeds. - Ran: pytest test/test_cuda.py. - Ran: pytest test/inductor/test_fp8.py. - Ran: pytest test/test_scaled_matmul_cuda.py. - Ran: pytest test/quantization/core/experimental/test_floatx.py. - Verified that basic end-to-end workflows complete on Linux and Windows. - ROCm version: 10.3 - GFX version: gfx1170/1 gfx1170 results: ``` - pytest test/test_cuda.py - using `PYTHONSAFEPATH=1` - 10 failed / 396 passed / 177 skipped / 3 xfailed - pytest test/inductor/test_fp8.py - 428 passed / 231 skipped - pytest test/test_scaled_matmul_cuda.py. - 192 passed / 1842 skipped - pytest test/quantization/core/experimental/test_floatx.py. - 110 passed ``` gfx1171 results: ``` - pytest test/test_cuda.py - using `PYTHONSAFEPATH=1` - 16 failed / 390 passed / 177 skipped / 3 xfailed - pytest test/inductor/test_fp8.py - 428 passed / 231 skipped - pytest test/test_scaled_matmul_cuda.py. - 192 passed / 1842 skipped - pytest test/quantization/core/experimental/test_floatx.py. - 110 passed ``` Tested separately on gfx1170 and gfx1171 with ROCm 10.3: ``` - pytest test/test_cuda.py TestCuda.test_preferred_blas_library_settings - gfx1170: passed; hipBLASLt was selected as the preferred backend. - gfx1171: passed; hipBLASLt was selected as the preferred backend. - pytest test/test_scaled_matmul_cuda.py -k test_float8_basics -v - gfx1170: all 10 CUDA variants passed, including all supported FP8 dtype and output-dtype combinations and layout permutations. - gfx1171: all 10 CUDA variants passed with the same coverage. - pytest test/test_scaled_matmul_cuda.py -k test_scaled_mm_vs_emulated -v - gfx1170: 26 passed and 28 skipped overall; all 13 applicable CUDA numerical variants passed. - gfx1171: 26 passed and 28 skipped overall; all 13 applicable CUDA numerical variants passed. - Skips were expected for Hopper-only blockwise cases, unsupported FP16/FP32 rowwise cases, and CPU instances of accelerator-only tests. ``` Pull Request resolved: pytorch#200757 Approved by: https://github.com/jeffdaily Co-authored-by: Milica Stankovic <mstankov@amd.com> Co-authored-by: Jeff Daily <jeff.daily@amd.com> Co-authored-by: Milica Stankovic <mstankov@amd.com> Co-authored-by: Jeff Daily <jeff.daily@amd.com>
…#202210) Supported by @jeffdaily `test_batchnorm_nhwc_cpu` fails with 1–2 OpenMP threads: ``` BatchNorm3d (4, 8, 2, 10, 10), float16 input, float32 params (mixed_dtype=True) self.assertEqual(bn.weight.grad, ref_bn.weight.grad, atol=precision, rtol=precision) Mismatched elements: 1 / 8 (12.5%) Greatest absolute difference: 2.3365020751953125e-05 at index (0,) (up to 1e-05 allowed) ``` The test compares the channels_last and contiguous CPU batch norm backward kernels. Both reduce `weight`/`bias` grads over N×spatial elements in fp32, but in different orders: [channels_last](https://github.com/pytorch/pytorch/blob/0564ffec41075bbb70b8113eb79e27fc8668decd/aten/src/ATen/native/cpu/batch_norm_kernel.cpp#L1204) accumulates each channel serially per thread chunk, and [contiguous](https://github.com/pytorch/pytorch/blob/0564ffec41075bbb70b8113eb79e27fc8668decd/aten/src/ATen/native/cpu/batch_norm_kernel.cpp#L1055) accumulates across vector lanes. Against a float64 reference, both are well within fp32 summation error; they just disagree with each other by more than the fp32 default `atol=1e-5`. The 3D fp32 cases already pass `precision=1e-4`, and pytorch#189180 applied it to `bias.grad` as well. But the 2D loop passes no precision, and the 3D float16 entry is `None`. So with `mixed_dtype=True`, the fp32 params' grads are held to the fp32 default. This change applies `1e-4` whenever the params are fp32 and no precision was given. It surfaced after pytorch#175988, which changed CPU fp16/bf16 `randn`/`uniform_`, so the test's fixed seed draws different inputs. On the failing inputs, the kernels give bit-identical grads in 2.5.1 and on current main. Upstream CI doesn't hit it because [`.ci/pytorch/test.sh` floors `OMP_NUM_THREADS` at 4](https://github.com/pytorch/pytorch/blob/0564ffec41075bbb70b8113eb79e27fc8668decd/.ci/pytorch/test.sh#L213) on ARC runners, naming this test. Test plan: ``` OMP_NUM_THREADS=1 python test/test_nn.py -k test_batchnorm_nhwc_cpu # fails before, passes after python test/test_nn.py -k test_batchnorm_nhwc_cpu ``` Pull Request resolved: pytorch#202210 Approved by: https://github.com/jeffdaily (cherry picked from commit 64ef93a)
jithunnair-amd
force-pushed
the
release/2.15
branch
from
November 4, 2026 16:58
dde42be to
61dff56
Compare
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Cherry-pick of upstream pytorch#202210 (commit 64ef93a) into release/2.15.
Why:
test_nn.py::TestNN::test_batchnorm_nhwc_cpuexists on this branch and fails on some GPUs (e.g. gfx1103) because channels_last and contiguous kernels reduce fp32 weight/bias grads in different orders, exceeding the default tolerance.Applied cleanly; changed lines are identical to upstream.