[release/2.11] Fix MultiProcContinuousTest subclass hanging on its base class's dead workers - #3526
Open
albmalamd wants to merge 3039 commits into
Open
[release/2.11] Fix MultiProcContinuousTest subclass hanging on its base class's dead workers#3526albmalamd wants to merge 3039 commits into
albmalamd wants to merge 3039 commits into
Conversation
…9165) Pull Request resolved: pytorch#179165 Approved by: https://github.com/aorenste, https://github.com/anijain2305 ghstack dependencies: pytorch#179113, pytorch#179164
…ait_event (pytorch#179166) Pull Request resolved: pytorch#179166 Approved by: https://github.com/anijain2305 ghstack dependencies: pytorch#179113, pytorch#179164, pytorch#179165
Pull Request resolved: pytorch#179172 Approved by: https://github.com/desertfire, https://github.com/shunting314 ghstack dependencies: pytorch#179113, pytorch#179164, pytorch#179165, pytorch#179166
Pull Request resolved: pytorch#179177 Approved by: https://github.com/karthickai ghstack dependencies: pytorch#179113, pytorch#179164, pytorch#179165, pytorch#179166, pytorch#179172
…nd properly assigns stream to fused kernels (pytorch#179438) Pull Request resolved: pytorch#179438 Approved by: https://github.com/karthickai ghstack dependencies: pytorch#179113, pytorch#179164, pytorch#179165, pytorch#179166, pytorch#179172, pytorch#179177
Ensure DynamicInt instances preserve their type through power operations by implementing __pow__ and __rpow__ methods. This ensures that expressions like DynamicInt(2) ** 3 return DynamicInt(8) rather than plain int(8), maintaining dynamism for torch.compile. Added test coverage for: - DynamicInt ** int - int ** DynamicInt - DynamicInt ** DynamicInt - pow(DynamicInt, int) - pow(DynamicInt, int, mod) Fixes issue where power operations would strip DynamicInt wrapper, potentially causing missed recompilations or incorrect guards. Pull Request resolved: pytorch#179868 Approved by: https://github.com/bobrenjc93
…179808) (pytorch#179808) Summary: Pull Request resolved: pytorch#179808 Previously `reinterpret_tensor` was always added to `addmm` in max-autotune mode, even if aten was chosen. This can lead to performance regressions e2e. Instead, we preserve the 1D bias if it is 1D and feed different kernel inputs to NVIDIA and AMD paths. pytorch#177130 did this previously just for amd after discovering the same regression. The test added by that PR was removed as we no longer use `bias_addmm` for 1D biases. Test Plan: `test_addmm_1d_bias_no_reinterpret_tensor` Reviewed By: njriasan Differential Revision: D100157470 Pull Request resolved: pytorch#179808 Approved by: https://github.com/njriasan
…ytorch#179808) (pytorch#179808)" This reverts commit 6c10ca4. Reverted pytorch#179808 on behalf of https://github.com/PaulZhang12 due to creating new pr internally ([comment](pytorch#179808 (comment)))
Add OSDC (ARC) runner support to _docs.yml following the same pattern as _linux-build.yml: a build-docs-osdc job with a matrix for cpp/python doc types, running directly in the container image instead of docker-in-docker. The EC2 path is preserved and gated on !inputs.use-arc. Authored with Claude. Pull Request resolved: pytorch#179994 Approved by: https://github.com/atalman, https://github.com/malfet
…ch#179867) Extend the torch.cuda.graph context manager with an enable_annotations keyword that enables kernel annotation recording on entry and automatically calls resolve_pending_annotations() before the capture ends. Annotations are not cleared on exit so multiple graphs in the same workload can accumulate annotations. Authored with Claude. Pull Request resolved: pytorch#179867 Approved by: https://github.com/ngimel ghstack dependencies: pytorch#179768
pytorch#179718) …ic_trace, traceback Remove mypy suppressions. Add return type, parameter type, and generic type annotations. Backward-compatible APIs use pyrefly ignore comments. Authored with Claude. Pull Request resolved: pytorch#179718 Approved by: https://github.com/aorenste
inner size_t threadgroup_size redaclaration shadowed the outer one, so we computed the optimal inner_size and discarded and still launched max thread groups with many threads being idle which killed perf. See speedups below: <img width="1800" height="480" alt="image" src="https://github.com/user-attachments/assets/09a360c2-c3d6-4688-b06b-9f1f860cec2b" /> Pull Request resolved: pytorch#180173 Approved by: https://github.com/malfet
…ch#180218) `aten.quantize_per_tensor` is a base operator in the aten namespace, not an MKLDNN-specific one. The meta-tensor support should not be conditional on whether the MKLDNN backend is enabled in the build. Pull Request resolved: pytorch#180218 Approved by: https://github.com/Lucaskabela
This reverts commit d90db5c. Reverted pytorch#179994 on behalf of https://github.com/huydhn due to C++ doc push is failing due to a missing awscli https://github.com/pytorch/pytorch/actions/runs/24363515340/job/71150965479 ([comment](pytorch#179994 (comment)))
…LIST (pytorch#180235) Pull Request resolved: pytorch#180235 Approved by: https://github.com/jeffdaily
…ock ordering to fix NCCL symmetric memory mismatch (pytorch#178362) Previously, blocks were ordered by their memory address. However, this caused issues because different ranks might allocate memory at non-uniform addresses, leading to inconsistent block ordering across ranks. This inconsistency could result in misaligned tensor reuse during communication, causing incorrect or corrupted results. To fix this, we replace address-based sorting with an allocation-time counter, which guarantees a globally consistent order of blocks across all ranks. This ensures that tensor block reuse is aligned and deterministic, eliminating communication errors due to block misalignment. This pr is based on pytorch#167662 and comments in issue pytorch#178138. Pull Request resolved: pytorch#178362 Approved by: https://github.com/ngimel
The apt valgrind (3.18.1 on jammy) is sufficient for PyTorch's usage (primarily callgrind-based instruction counting in benchmarks). Building 3.20.0 from source also pulled in unnecessary build deps (asciidoc, docbook-xml, docbook-xsl, xsltproc). This saves ~130 MB in the install_base Docker layer and ~70s of build time. Added suppression for dlopen false-positive, that were fixed in https://sourceware.org/git/?p=valgrind.git;a=commit;h=947388eb043ea1c44b37df94046e1eee790ad776 Authored with Claude. Pull Request resolved: pytorch#180130 Approved by: https://github.com/Skylion007 ghstack dependencies: pytorch#180133
pytorch#179864) … symbolic_shapes, sym_node Remove mypy suppressions. Add return type, parameter type annotations, and parameterize Callable/nullcontext/ValueRanges types. Authored with Claude. Pull Request resolved: pytorch#179864 Approved by: https://github.com/aorenste
## Summary Root cause: The backward cache wrappers duplicated the same `post_compile` implementation in both `CompiledBackward` and `BundledCompiledBackward`, even though the behavior is shared by all generated backwards. Proposed fix: Move the shared `post_compile` implementation into `GenericCompiledBackward`, keep delegating through `super().post_compile(...)` so each subclass still lands on its cache-specific loader via MRO, and fix the `CompiledBackward` docstring to say `backward function`. Why this is the right long term fix: The `torch._dynamo.disable(...)` wrapper is common backward-specific behavior, not something that should be maintained separately per cache transport. Centralizing it in the shared base removes duplication and keeps future changes to backward post-compile handling in one place. ## Testing - `python3 -m compileall torch/_functorch/_aot_autograd/aot_autograd_result.py` - `git diff --check` - Runtime `python3 test/dynamo/test_aot_autograd_cache.py -v -k 'test_vmap or test_regional_inductor_with_backward'` was attempted, but this container does not have a built `torch` or basic Python deps such as `typing_extensions`. - Standalone Python MRO proof script passed for the relevant inheritance order. Drafted via Codex, published after manual review by @bobrenjc93 Pull Request resolved: pytorch#180096 Approved by: https://github.com/aorenste
…ytorch#179185) Pull Request resolved: pytorch#179185 Approved by: https://github.com/anshul-si
## Summary
Refactor `GenericAOTAutogradResult.wrap_post_compile` into four private helpers so the cached post-compile flow is easier to follow.
## Root cause problem
`wrap_post_compile` had grown into a single large method that mixed four distinct phases: cached-graph logging, fw/bw load plus `post_compile`, runtime wrapper installation, and guard evaluation.
## Proposed fix
Split the existing logic into `_log_cached_graphs`, `_load_and_post_compile`, `_apply_runtime_wrappers`, and `_install_guards`, and have `wrap_post_compile` call them in sequence while keeping the `dynamo_timed("AOTAutogradCache.inductor_load")` scope around the load/post-compile step.
## Why the proposed fix is the right long term fix
This keeps behavior unchanged while making each phase explicit and independently maintainable, which reduces the chance of accidental regressions when the cached AOTAutograd path changes again.
## Testing
- `python3 -m compileall torch/_functorch/_aot_autograd/aot_autograd_result.py`
- Ran a focused Python smoke test under local stubs that covered helper orchestration, inference and autograd load/post-compile paths, runtime wrapper application, and guard installation. The native PyTorch test suite could not run in this environment because there is no built/importable local `torch` tree.
Drafted via Codex, published after manual review by @bobrenjc93
Pull Request resolved: pytorch#180097
Approved by: https://github.com/aorenste
…OM (pytorch#180231) The linux-jammy-py3.14t-clang18 / test-osdc (dynamo_wrapped, 1, 3) job has been flakily OOMing since ~April 11. Investigation across 8 failing jobs shows the same pattern: - test_nn.py (531 tests) runs under dynamo wrapping with compiled autograd - After ~11 minutes with zero output, the OSDC Kubernetes pod is OOM-killed (exit code 137 / SIGKILL) - The pod has 64GB RAM (linux.2xlarge → l-x86iavx512-8-64 on OSDC) Free-threaded Python 3.14t has significantly higher per-object memory overhead (per-object locks, biased reference counting) compared to standard CPython. Dynamo compilation creates millions of objects per test, and with 531 tests plus compiled autograd, the cumulative memory from allocator fragmentation exceeds the 64GB pod limit. Pull Request resolved: pytorch#180231 Approved by: https://github.com/malfet
…chmark (pytorch#179926) Summary: What: Adds `alg_id` search to select the best performing hipSPARSELt kernel for each shape at runtime. Why: * Default `alg_id=0` performs poorly for many shapes on MI350X — with 13/18 Fixed K shapes regressing vs. dense * alg_id search recovers significant performance: Fixed K improves from 5/18 to 17/18 shapes beating dense, and the kernel bug shapes (M=N=11264–14336) recover from 0.11x to ~1.30x Test Plan: ``` buck run mode/amd-gpu -c cxx.extra_cxxflags=-Wno-unused-value scripts/gylls/torchao:benchmark_semi_structured_sparsity -- --mode nvidia-fixed-mn --dtype fp8 --backend cusparselt -save ``` | Mode | Avg Speedup | Min | Max | Notes | | -- | | BERT Shapes | 0.95x | 0.72x | 1.11x | 2/4 shapes beat dense; small shapes still regress | | Fixed K=10240 | 1.22x | 0.85x | 1.47x | 17/18 shapes beat dense after alg_id search; kernel bug shapes recover to ~1.30x | | Fixed M=N=10240 | 1.42x | 1.12x | 1.57x | Sparse outperforms dense across all K values | | M4 Target Shapes | 1.05x | 0.84x | 1.29x | Large N (133120) shapes gain 1.17–1.29x; small N (20544) still regresses | * FP8 2:4 sparsity on MI350X shows strong results with alg_id search — Fixed K improves from mostly regressing to 17/18 shapes beating dense (avg 1.22x), and the kernel bug shapes (M=N=11264–14336) recover to ~1.30x using alg_id. * Small M/N shapes (BERT, M4 small-N) remain below 1x for most shapes regardless of alg_id — sparse overhead dominates at these sizes and sparsity should not be enabled by default. Benchmark results with alg_id: https://fburl.com/gdoc/dg622pvx Differential Revision: D99334033 Pull Request resolved: pytorch#179926 Approved by: https://github.com/jerryzh168
…rch#179095) (pytorch#179095) Summary: Add non-TMA persistent matmul template support for addmm on AMD GPUs D96971956 added a non-TMA persistent MM template (persistent_mm_template) as a fallback for AMD GPUs that lack TMA support, but only wired it up for torch.mm. The torch.addmm code path was left unconditionally using persistent_tma_mm_template, which doesn't work on AMD. This applies the same torch.version.hip check to the addmm path in mm.py and adds a corresponding test_max_autotune_regular_addmm_persistent test, mirroring the existing mm test. Changes applied to both fbcode/ and xplat/ mirrors. Test Plan: TORCHINDUCTOR_FORCE_DISABLE_CACHES=1 buck2 run fbcode//mode/opt-amd-gpu fbcode//scripts/robeck/inductor:test_max_autotune_amd -- -r test_max_autotune_regular_addmm_persistent Differential Revision: D99194762 Pull Request resolved: pytorch#179095 Approved by: https://github.com/nmacchioni, https://github.com/drisspg
Remove mypy suppressions. Add return type, parameter type, and generic type annotations. Backward-compatible APIs use pyrefly ignore comments. Authored with Claude. Pull Request resolved: pytorch#179731 Approved by: https://github.com/aorenste
…rch#180180) Run git clone, git checkout, and python install.py as the jenkins user directly (via as_jenkins/conda_run), matching the pattern used by install_huggingface() and install_timm() in the same file. This removes two expensive recursive chown calls on the torchbench dir and /opt/conda. Authored with Claude. Pull Request resolved: pytorch#180180 Approved by: https://github.com/Skylion007 ghstack dependencies: pytorch#180133, pytorch#180130
…wlist (pytorch#180268) (pytorch#180268) Summary: D99020245 generalized FlightRecorder profiling_name to use the actual backend name (e.g. "ncclx:", "gloo:") instead of hardcoded "nccl:". The FR trace analyzer's Op class needs to accept these additional backend prefixes to avoid AssertionError when parsing traces from non-nccl backends. This is the types.py-only portion of D100117074, split out so it can land independently without the C++ TorchComm changes that caused OSS CI failures. Test Plan: Python-only change — adds two strings to an allowlist. No behavioral change for existing "nccl:" and "xccl:" prefixes. Reviewed By: dolpm Differential Revision: D100657879 Pull Request resolved: pytorch#180268 Approved by: https://github.com/atalman
Fix pytorch#166173 ## Summary 1) What is the root cause problem Dynamo already knows which Python side effects it needs to replay after graph execution, but bytecode hooks only receive code objects. That forces consumers like vLLM to guess from transformed bytecode details such as whether `"update"` appears in `co_names`, which is brittle and not tied to Dynamo's actual side-effect classification. 2) What is the proposed fix Attach Dynamo's filtered side-effect source refs to the generated code object before bytecode hooks run, expose `torch._dynamo.convert_frame.get_compiled_code_side_effects()` and `torch._dynamo.convert_frame.compiled_code_has_side_effects()` for hook consumers, preserve that metadata when a hook returns replacement bytecode, and add regression tests that exercise mutating and side-effect-free cases through chained bytecode hooks. 3) Why is the proposed fix the right long term fix This reuses the same side-effect analysis Dynamo already trusts for warning/error handling, removes the need for bytecode-name heuristics in downstream projects, and keeps the bytecode-hook contract extensible without changing its call signature. ## Testing - Added tests for list mutation, dict mutation, tensor-in-container (cudagraphs-relevant), multiple simultaneous side effects, and pure (no side effect) cases - `python3 -m py_compile torch/_dynamo/convert_frame.py torch/_dynamo/output_graph.py test/dynamo/test_reconstruct.py` - `git diff --check` Drafted via Codex, published after manual review by @bobrenjc93 Pull Request resolved: pytorch#180079 Approved by: https://github.com/Lucaskabela
While running tests, I noticed the error, `torch._dynamo.exc.InternalTorchDynamoError: IndexError: pop index out of range` if attempted pop() with an index equal to the length. For example, if `self.assertRaises(IndexError, p.pop, 10)` in tests is changed to `self.assertRaises(IndexError, p.pop, 2)`, it will run into IndexError: pop index out of range. The original assertion with 10 passes because it's higher than the index length. Pull Request resolved: pytorch#179617 Approved by: https://github.com/Skylion007
… builds (pytorch#180293) Work around fo: pytorch#167658 ## Summary Extends the existing CUDA dependency injection to Windows x64 CPU and macOS ARM64 CPU wheel builds. Currently, `pytorch_extra_install_requirements` is populated only for Linux x86/aarch64 CUDA builds. This PR adds the stable CUDA version's dependencies (`PYTORCH_EXTRA_INSTALL_REQUIREMENTS[CUDA_STABLE]`) to Windows CPU and macOS CPU builds as well. The dependencies use `platform_system == 'Linux'` markers so they are only installed on Linux systems. Triton is handled automatically by `binary_populate_env.sh` — since `PYTORCH_EXTRA_INSTALL_REQUIREMENTS` is now non-empty for these builds, the existing triton append logic kicks in without any changes needed. ### Changes - **`.github/scripts/generate_binary_build_matrix.py`**: For CPU builds on `windows` and `macos-arm64`, set `pytorch_extra_install_requirements` to `PYTORCH_EXTRA_INSTALL_REQUIREMENTS[CUDA_STABLE]` - **`generated-windows-arm64-binary-wheel-nightly.yml`**: Remove CUDA metadata injection from windows-arm64 builds (not applicable) ## Test plan - [x] Verified Windows CPU nightly matrix entry includes stable CUDA deps - [x] Verified macOS ARM64 CPU nightly matrix entry includes stable CUDA deps - [x] Verified Windows CUDA/XPU entries are unaffected - [x] Verified Linux CPU entries remain empty - [x] Verified windows-arm64 does not get injection - [x] Confirmed `binary_populate_env.sh` runs on both Windows and macOS builds and will append triton automatically Pull Request resolved: pytorch#180293 Approved by: https://github.com/malfet, https://github.com/seemethere
…'HIPCUB_CCCL_VERSION' (pytorch#3441) Cherry-pick of pytorch#188072 (commit c8e2473) onto `release/2.12`. ## Summary - Introduce dependency on libhipcxx - Guard `FpLimits` for `c10::BFloat16` backport to pre CCCL 3.x - Extend `CUB_VERSION` derivation for hipCUB using `HIPCUB_CCCL_VERSION` ## Notes This PR has merge conflicts that need to be resolved manually (conflict markers committed as requested). ## Test plan - [x] Build with ROCm - We were able to run the core PyTorch tests, and verify that they are passing here: https://github.com/ROCm/TheRock/actions/runs/29543027958 - [x] Verify hipCUB/CCCL compatibility glue compiles Authored with assistance from Cursor Made with [Cursor](https://cursor.com) --------- Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: Stanley Tsang <stanley.tsang@amd.com>
Solves ROCm/TheRock#6723 In the rock core UT workflow, fbscribelogger's transitive dependency thriftpy2 doesn't ship a python 3.14 prebuilt wheel and causing installation to fail because there isn't any compiler available in the environment. This wasn't observed in previous dependency bump PR runs because those were run on the `test_pytorch_wheels_full` workflow, where a compiler is available, but some headers and libraries are missing. Presumably this dependency doesn't need any of those missing components and therefore able to be built. Example core UT workflow run: https://github.com/ROCm/TheRock/actions/runs/29853620514 (installation succeeded)
## Summary - Advance the ROCm Triton pin to the tip of `ROCm/triton` `release/internal/3.8.x`. - Update `.ci/docker/triton_version.txt` from `3.7.1` to `3.8.0`. Triton pin: `285c85383660b2e4aac38d4947d23aa61dab3ce2` Triton branch: `ROCm/triton@release/internal/3.8.x` ## Test plan - [x] `git diff --check -- .ci/docker/ci_commit_pins/triton.txt .ci/docker/triton_version.txt .ci/docker/common/install_triton.sh` - [x] Verified Triton checkout reports `__version__ = '3.8.0'` - [ ] Run PyTorch CI Follows the structure of ROCm#3360. Made with [Cursor](https://cursor.com) --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…#190348) (pytorch#3474) Fix a memory leak when CUDA graph capture fails: `CUDAGraph::reset()` now releases the graph's private mempool even if capture never completed successfully. Previously, a failed capture could leave the pool's use count elevated for the rest of the process, so `empty_cache()` could not reclaim its reserved segments. `capture_begin()` acquires a private mempool and routes allocations to it via `beginAllocateToPool()`. On a successful capture, `capture_end()` ends pool routing, sets `capture_ended_ = true`, and `reset()` releases the pool. If capture fails, `capture_end()` may still call `endAllocateToPool()` (that happens before the capture error is propagated), but `capture_ended_` is only set on success. `reset()` gated pool release on `capture_ended_`, so after a failed capture it skipped releasing the pool entirely. The pool's use count never returned to zero, and its reserved segments leaked until process exit - even with no live allocations (`active.all.current == 0`). Because the leaked segments persist for the rest of the process, this also causes spurious failures in later, unrelated tests that assume a clean allocator after `empty_cache()` (e.g. `test_memory_snapshot_*`) that start from empty memory and assert on the exact set of reserved segments. When such a test runs in the same process after a failed-capture test, the leaked pool pollutes `memory_reserved()` / `_snapshot()`. This fix removes that cross-test contamination. `test_graph_rng_after_failed_capture` covered stream/RNG recovery and live allocation leaks, but not this reserved-segment leak. Track two new flags on `CUDAGraph`: - `allocated_pool_` - set in `capture_begin()` once the private pool is acquired; tells `reset()` it must release the pool regardless of capture success. - `capturing_to_pool_` - set in `capture_begin()` and cleared in `capture_end()` after `endAllocateToPool()`; tells `reset()` whether pool routing is still active and must be ended before release (e.g. capture abandoned before `capture_end()` ran). `reset()` now keys pool cleanup off `allocated_pool_` instead of `capture_ended_`. - [x] Added `test_graph_capture_error_releases_reserved_segments`: triggers a failed capture via `torch.cuda.synchronize()` during capture, calls `reset()`, and asserts `memory_reserved()` returns to baseline after `empty_cache()`. - [x] New test comment explains how it complements `test_graph_rng_after_failed_capture` (reserved segments vs. live allocations). ```bash python test/test_cuda.py TestCuda.test_graph_capture_error_releases_reserved_segments ``` output without fix: ``` test_graph_capture_error_releases_reserved_segments (__main__.TestCuda.test_graph_capture_error_releases_reserved_segments) ... FAIL Expected 0 but got 33554432. ``` output with fix: ``` test_graph_capture_error_releases_reserved_segments (__main__.TestCuda.test_graph_capture_error_releases_reserved_segments) ... ok ``` Pull Request resolved: pytorch#190348 Approved by: https://github.com/jeffdaily (cherry picked from commit 240fca9)
…hot, mempool, OOM retry) (pytorch#3473) Addresses three distinct issues that greendog surfaced together as a single "CUDA allocator/mempool" flake cluster in test_cuda.py. They have three different root causes, so review them independently. 1. test_memory_snapshot (format_flamegraph, torch/cuda/_memory_viz.py) Real helper bug. format_flamegraph downloads flamegraph.pl into a NamedTemporaryFile(delete=True) and then os.rename's it into ~/.cache/flamegraph.pl. When the context manager exits it tries to unlink the file it just renamed away, raising FileNotFoundError. It looked like a flake because the rename succeeds, so on retry the cached script exists and the download path (and the crash) is skipped entirely. Fix: download to a mkstemp file created in the target directory (same filesystem, so os.replace is atomic) with delete=False and manual cleanup in finally. This also removes a latent second bug where the old code silently swallowed a cross-device rename OSError and then ran a subprocess on a script that was never installed. 2. test_mempool_ctx_multithread A CUDA tensor leaked by a prior test that is reachable only through a reference cycle stays active (refcounting cannot free it), so empty_cache cannot reclaim its segment and the "Expected empty pool in the beginning" assertion intermittently sees 1 segment. Fix: gc.collect() before empty_cache(), matching the existing idiom used elsewhere in this file. 3. test_out_of_memory_retry Not a flake: a persistent failure on MI350 (gfx950), where allocating ~1.02x of the reported free memory does not raise an OOM RuntimeError, so assertRaisesRegex fails. Skip on MI350 only (skipIfRocmArch(MI350_ARCH)) so the test keeps running on CUDA and other ROCm archs. Test Plan: Reproduced the format_flamegraph crash and the mempool leak locally on 8xA100, confirmed the fixes, and confirmed the OOM test still passes on CUDA. ``` python -m pytest \ test/test_cuda.py::TestMemoryViz::test_format_flamegraph_download_moves_temp_file \ -x -q rm -f ~/.cache/flamegraph.pl python -m pytest test/test_cuda.py::TestCudaAllocator::test_memory_snapshot -x -q python -m pytest test/test_cuda.py::TestMemPool::test_mempool_ctx_multithread -x -q python -m pytest test/test_cuda.py::TestCuda::test_out_of_memory_retry -x -q ``` The new TestMemoryViz regression test reproduces the exact CI FileNotFoundError before the fix and passes after it. Authored with Claude. Pull Request resolved: pytorch#190160 Approved by: https://github.com/eqy ghstack dependencies: pytorch#190151 (cherry picked from commit 215f570) ### fix flaky test_memory_snapshot (pytorch#190545) `test_memory_snapshot` asserts that the last allocator action after cleanup is `segment_free` or `segment_unmap`. cuBLAS workspace allocations on the default stream can keep segments alive and make that assertion flaky. Call `torch._C._cuda_clearCublasWorkspaces()` before `empty_cache()` in `test_memory_snapshot`, so segment teardown is not blocked by lingering cuBLAS workspace allocations. Fixes pytorch#179745 Pull Request resolved: pytorch#190545 Approved by: https://github.com/jeffdaily (cherry picked from commit 6c68acb) Co-authored-by: Aaron Orenstein <aorenste@fb.com>
…rch#3475) ## Summary Make `test_cuda.py::TestCudaAllocator.test_memory_compile_regions` deterministic by resetting Dynamo's process-global frame counter before `torch.compile()`. Region labels in memory snapshots depend on that counter, so running this test after other compile tests could produce non-zero region ids and fail the golden sequence check. ## Problem 1 `test_memory_compile_regions` records CUDA allocations with `compile_context=True` and asserts the observed compile-region labels match a fixed sequence: ```python [ "Torch-Compiled Region: 0/0", "Torch-Compiled Region: 1/0", "Torch-Compiled Region: 0/0", ] ``` Those region ids come from Dynamo's process-global frame counter. When other `torch.compile` tests run first in the same process, the counter is already non-zero, so the snapshot can contain labels like `"Torch-Compiled Region: 2/0"` instead of `"0/0"`, causing order-dependent failures. Example of a falling test sequence for `test_memory_compile_regions`: ``` python test_cuda.py \ TestFXMemoryProfiler.test_fx_memory_profiler_augmentation \ TestCudaAllocator.test_memory_compile_regions ``` ## Problem 2 `test_memory_plots` and `test_memory_plots_free_segment_stack` are failed due to insufficient teardown in test_memory_compile_regions Example of a falling test sequence for `test_memory_plots`: ``` python test_cuda.py \ TestCudaAllocator.test_memory_compile_regions \ TestCudaAllocator.test_memory_plots \ TestCudaAllocator.test_memory_plots_free_segment_stack ``` ## Fix - test_cuda.py::TestCudaAllocatortest_memory_compile_regions - call `torch._dynamo.reset()` after `empty_cache()` and before compiling, so region ids start from zero and the expected sequence is stable regardless of test order. - Improved teardown to fix `test_cuda.py::TestCudaAllocator::test_memory_plots` and `test_cuda.py::TestCudaAllocator::test_memory_plots_free_segment_stack` Fixes pytorch#163202 Fixes pytorch#179744 Fixes pytorch#179798 Pull Request resolved: pytorch#190358 Approved by: https://github.com/Skylion007, https://github.com/jeffdaily (cherry picked from commit 9c751e3)
) (pytorch#3476) Add `torch._C._cudnn_clear_dropout_state()` to release cached RNN dropout-state buffers Multi-layer RNNs with dropout allocate a backend dropout-state buffer that is cached long-term. These buffers stay alive after the model is deleted because the backend cache still holds references. `torch.cuda.empty_cache()` does not free them, so they appear as persistent CUDA memory segments for the rest of the process. This caused flaky test failures when CUDA memory tests assumed a clean allocator state. In particular, this sequence failed before the fix: ```bash python test_cuda.py -v TestCuda.test_graph_cudnn_dropout TestMemPool.test_mempool_ctx_multithread ``` `test_graph_cudnn_dropout` left a persistent segment from the cached dropout buffer. `test_mempool_ctx_multithread` then failed at startup with message `Expected empty pool in the beginning`. - Refactor the cuDNN dropout-state cache into named accessors and add `_cudnn_clear_dropout_state()` to reset cached buffers and events. - Add `_miopen_clear_dropout_state()` to reset the thread-local MIOpen dropout-state buffer. - Expose a single Python API, `torch._C._cudnn_clear_dropout_state()`, which clears the MIOpen cache on ROCm and the cuDNN cache on CUDA. - Call it at the start of `test_mempool_ctx_multithread` so the test is not order-dependent on prior RNN dropout usage. Fixes pytorch#153460 Pull Request resolved: pytorch#190405 Approved by: https://github.com/jeffdaily (cherry picked from commit dc228b8)
…ch#3495) This fixes deadlock with torch.compile while running bigger models in MAD engine. This is a cherry-pick of upstream PR: pytorch#182948 Made with [Cursor](https://cursor.com)
## Summary Backport the Inductor exhaustive-autotuning module-lifetime fix from pytorch#184285 to ROCm PyTorch `release/2.12`. This branch contains three commits: 1. pytorch#183920 (`5fe7534806f`) — prerequisite compile-result pruning helper used by pytorch#184285. 2. pytorch#184285 (`c0c69117734`) — release benchmark-only Triton modules and other autotune artifacts promptly. 3. A release/2.12-specific adaptation of pytorch#188907 (`77b1231acb2`) — clear retained failed-config exceptions without importing the broader pytorch#181827 launcher refactor. ROCm Triton `release/internal/3.7.x` already contains the required triton-lang/triton#9444 explicit module-unload support as commit `110cd8e2d`. pytorch#188607 is intentionally omitted: release/2.12 does not contain the `_build_fast_launcher` path whose dangling function pointer that PR fixes. ## Motivation Long exhaustive autotuning can retain loaded HIP modules until the driver module table is exhausted, producing errors such as: ``` Triton Error [HIP]: Code: 209 no kernel image is available for execution on the device ``` The fix applies to both static-launcher enabled and disabled paths. ## Validation - `git diff --check` passed. - All changed Python files compile. - Targeted pytorch#188907 reference-cycle regression logic passed. - Built and installed PyTorch successfully: - `torch==2.12.0+git3c39ad9` - commit `3c39ad9882961c54fc4891651a2dab05ddd5b0a0` - Triton `3.7.1` - Focused backport tests: `7 passed`. - Full single-process `test_torchinductor_opinfo_properties.py` run with `TORCHINDUCTOR_USE_STATIC_CUDA_LAUNCHER=0`: - Passed the historical HIP 209 point with no HIP 209. - No host-side `MemoryError` occurred. - Reached 969 / 1,134 tests, then the process exited with `SIGSEGV` while entering `test_unary_ufunc_numerical_abs_backend_aot_eager_decomp_partition_cuda_bfloat16`. - That exact node passed when rerun alone in a fresh process (`1 passed in 20.57s`), indicating accumulated process-state failure rather than a deterministic test failure. A second full single-process run with `TORCHINDUCTOR_USE_STATIC_CUDA_LAUNCHER=1` completed all 1,134 tests: 1,060 passed, 18 skipped, 38 xfailed, and 18 failed. No HIP 209, host-side `MemoryError`, or segmentation fault occurred. ## Test artifacts - Build log: `/home/niromero/docker_workspace/pytorch_backport_184285_build/pytorch_2_12_backport_184285_build_20260722_170801.log` - Static-launcher-disabled reproduction: `/home/niromero/docker_workspace/pytorch_backport_184285_build/backport_184285_opinfo_gpu7_20260722_173514.log` Made with [Cursor](https://cursor.com) --------- Co-authored-by: Jason Ansel <jansel@meta.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
…ch#3497) Remove @xfailCUDAIfSM89OrLaterOnWindows from test_graph_capture_error_releases_reserved_segments and drop its import from `test/test_cuda.py`. That decorator was added when cherry-picking the CUDAGraph pool-leak fix (pytorch#190348) onto release/2.12. On this branch, `xfailCUDAIfSM89OrLaterOnWindows` is not available in `torch.testing._internal.common_cuda`, so importing it raises `ImportError` and the test module fails to load. The upstream test used the decorator for a Windows+SM89+xfail that does not apply to our 2.12 backport context. Removing the decorator and import restores a valid import path
pytorch#3505) Skip for `test_cuda.py::TestCuda::test_hip_device_count` due to rocprofiler-sdk issue `AIPROFSDK-840`
…pytorch#3480) ## Motivation Update CK submodule to include cherrypicked ROCm/rocm-libraries#9333 commit. ## Technical Details Resolves below build failure and will unblock torch 2.12 builds for gfx90c. ``` FAILED: [code=1] caffe2/aten/src/ATen/CMakeFiles/ck_gemm.dir/native/hip/bgemm_kernels/bgemm_kernel_bf16bf16bf16_64_16x16x64_16x16_1x1_8x8x1_8x8x1_1x16x1x4_4_Intrawave_v1.hip.o /__w/rockrel/rockrel/external-builds/pytorch/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:49:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD' 49 | wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD; 4 errors generated when compiling for gfx90c. ``` ## Test Plan 1. Compile PyTorch's bundled CK for gfx90c 2. Trigger TheRock Multi-Arch PyTorch build w/fix branch ## Test Result - Before patch -> Results in same `error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'`. - After patch -> Error is no longer seen - Build passes. - PyTorch 2.10 build passes https://github.com/ROCm/TheRock/actions/runs/30103245986 ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
Skip for `test_nn.py::TestNN::test_Embedding_discontiguous_cuda` The test is flaky and disabled on upstream pytorch#186910 ``` python test/test_nn.py -v TestNN.test_Embedding_discontiguous_cuda test_Embedding_discontiguous_cuda (__main__.TestNN.test_Embedding_discontiguous_cuda) skipped 'skipIfRocm: pytorch#186910' ```
Fixes large amount of stream appearing in the FSDP trace when using CUDA graphs with FSDP2. Notes: 1. This fix will only work on CUDA >= 13.2. 2. This patch relieves the issue but doesn't fully resolve it. There's a lot of edge cases that don't fully work currently (tp, ep, micro-batching) but it does reduce the number of streams. In a follow-up, I can try to test and clean-up remaining cases and land fixes incrementally to avoid things from breaking. 3. Not sure what's a good way to test this. Will add a test if the changes look good. 4. It makes the trace much more readable but still has some quirks like the copies, all-gather, and reduce-scatter spread out across different streams even tho the capture has the same stream. Haven't thought about how to resolve that. For a 16-layer linears model with simple FSDP: Before: <img width="1550" height="872" alt="image" src="https://github.com/user-attachments/assets/4dfd0a0d-9b84-466c-9d7a-f9ea06ae7512" /> After: <img width="1328" height="250" alt="image" src="https://github.com/user-attachments/assets/01b67f3f-197f-40e1-8147-f52a0a718881" /> Related issue pytorch#155679 - torch.compile creating a lot of streams Pull Request resolved: pytorch#183983 Approved by: https://github.com/ngimel, https://github.com/weifengpy Co-authored-by: Cursor <cursoragent@cursor.com>
…catter blocking backward compute (pytorch#186000) FSDP keeps **one** reduce-scatter input buffer in flight: the compute stream must wait on the previous reduce-scatter before the next copy-in (`chunk_cat`) can reuse that buffer (`FSDP::post_backward_rs_wait`). When the reduce-scatter is exposed, that recycle wait stalls backward compute every layer — e.g. 37.6 ms/step of compute-stream idle, 100% gated by a reduce-scatter finishing, on an omnifm_v5 trace. There's no compute→reduce-scatter data dependency (backward never reads the reduced gradient), so the coupling is pure buffer-reuse bookkeeping. This exposes the in-flight buffer count — hardwired to 1 — as a per-module tunable: ```python FSDPModule.set_reduce_scatter_max_input_buffers(max_input_buffers: int, *, recurse: bool = True) ``` Raising it lets the next copy-in write a **fresh** buffer instead of waiting, removing the stall. `1` (default) keeps today's behavior; `2` clears the stall once it's `≥` the reduce-scatter pipeline depth while bounding peak memory; larger caps overlap deeper at higher memory. **Design:** generalize the existing recycle to keep at most `max_input_buffers` buffers — reclaim the oldest (`current_stream.wait_event(oldest_rs)`, then drop the keepalive ref; no `record_stream`) before this layer appends one. `1` drains to zero and reuses the freed buffer → **byte-identical to upstream**, and the reclaim is a no-op once `≥` the pipeline depth. The copy-in stays on the compute stream (`foreach_reduce` unchanged) and grads free promptly, so only the retained input buffers cost memory. The cap is a positive int — no unbounded mode (a large cap already means "retain all"). Co-authored-by: Lei Tian <2119521+leitian@users.noreply.github.com> Pull Request resolved: pytorch#186000 Approved by: https://github.com/anshul-si Co-authored-by: Lei Tian <2119521+leitian@users.noreply.github.com>
…ytorch#186335) By default FSDP2 runs all-gather and reduce-scatter on separate CUDA streams but through the same process group -- one NCCL communicator, which processes one collective at a time and so serializes them on the wire. This adds an opt-in FSDPModule API to give reduce-scatter its own communicator: FSDPModule.set_separate_reduce_scatter_group(enable=True, *, recurse=True) When enabled, FSDP creates a dedicated process group over the shard ranks (dist.new_group), one per distinct set of shard ranks (typically one communicator), so reduce-scatter and all-gather can progress concurrently when the network can sustain it; enable=False resets to the shared group. The default is unchanged -- reduce-scatter shares the shard process group -- so this is purely opt-in and creates no extra communicators unless requested. This redesigns the approach explored in (closed) PR pytorch#177015, which made the separate communicator the unconditional default (and created one even for post-forward meshes); here it is an opt-in toggle. Test Plan: ``` python test/distributed/_composable/fsdp/test_fully_shard_overlap.py \ TestFullyShardOverlap.test_set_separate_reduce_scatter_group \ TestFullyShardOverlap.test_fully_shard_backward_comm_overlap ``` Both pass on 4xH100: - test_set_separate_reduce_scatter_group: default shares the shard PG; enabling creates one dedicated PG shared across same-rank-set meshes; disabling resets to the shared PG. - test_fully_shard_backward_comm_overlap: real backward AG/RS overlap (large matmuls + collectives) is no slower than a serialized single-communicator reference. Authored with Claude. Co-authored-by: Lei Tian <2119521+leitian@users.noreply.github.com> Pull Request resolved: pytorch#186335 Approved by: https://github.com/anshul-si ghstack dependencies: pytorch#186000 Co-authored-by: Lei Tian <2119521+leitian@users.noreply.github.com>
IS_LINUX is used by @unittest.skipIf at line 320 but was not imported, causing a NameError at module load time when running the overlap tests. The import was introduced by upstream commit 9661ae6 which is not cherry-picked onto release/2.12 due to its broad scope (125 test files). Co-authored-by: Cursor <cursoragent@cursor.com>
## Overview * Mini tuning database tested on actual gfx1250 hardware + seqlen limited to 64, 256, 2048 + hdim limited to 16, 32, 64, 128 ,256 * Non-gfx1250 GPU images are copied directly from 0.13b
…#3521) PyTorch does not today promise transactional backward. If backward() throws, .grad is generally undefined. The test asserts a strong rule: no leaf grads after a failed combined backward, including nested reentrant work. That is stricter than what engine behavior gives. Track upstream pytorch#86735
…ailable (pytorch#3524) [release/2.12] [ROCm] Skip DTensor linalg.eig test when MAGMA is unavailable torch.linalg.eig has no hipSOLVER implementation, so on ROCm it is routed unconditionally to MAGMA. TheRock's ROCm wheels are built without MAGMA (torch.cuda.has_magma is False), so the unguarded eig sub-case in test_linalg_ops fails with "requires compiling PyTorch with MAGMA". CUDA builds are unaffected because cuSOLVER provides eig natively. Guarding only the eig block keeps cholesky/inv/lu_factor/solve coverage and matches the TEST_WITH_ROCM and not torch.cuda.has_magma idiom already used in test_linalg.py.
albmalamd
force-pushed
the
fix_rocm_test_replicate_device_id
branch
from
August 3, 2026 13:01
1ca96fe to
f44a4cb
Compare
) (pytorch#3538) Cherry-pick of upstream commit 2314113 (pytorch#187052) onto release/2.12. The LocalTensorMode coordinate cache was keyed by the mesh's Python object id, which is not unique over the lifetime of the mode: when a temporary submesh is destroyed, Python may reuse its id for the next submesh, silently returning stale coordinates for a different mesh dimension. This fixes the cache key to use the actual inputs to the coordinate calculation (ndim, flattened rank map, layout) instead. Fixes ROCM-28605, pytorch#184526. Verified locally: test/distributed/tensor/test_utils.py::TestStridedShardingWithLocalTensor: {test_2d_mesh_strided_sharding,test_2d_mesh_2d_tensor_strided_sharding} fail before this commit and pass after it. Co-authored-by: Alexander Grund <alexander.grund@tu-dresden.de>
Fix MultiProcContinuousTest subclass hanging on its base class's dead workers A subclass of a concrete MultiProcContinuousTest test class inherits the base class's `_processes_spawned` flag through ordinary attribute lookup, so it never spawns its own workers and instead dispatches tests onto the base class's already torn-down queues, blocking forever in `completion_queue.get()`. Consult the class's own `__dict__` in both the spawn and the teardown guard so that every leaf class owns its worker pool. This makes test_replicate.py order-independent. Under pytest the base class ReplicateTest runs before ReplicateFullyShardInit, which then hangs until the harness timeout; under unittest the subclass happens to sort first, which is why the bug stayed latent in CUDA CI. Backport of the test-infra portion of pytorch#189362.
albmalamd
force-pushed
the
fix_rocm_test_replicate_device_id
branch
from
August 4, 2026 15:48
f44a4cb to
0c157f1
Compare
albmalamd
requested review from
jataylo,
jeffdaily,
jithunnair-amd and
pruthvistony
as code owners
August 4, 2026 15:48
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
20 participants
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.
Motivation
Fix MultiProcContinuousTest subclass hanging on its base class's dead workers
A subclass of a concrete MultiProcContinuousTest test class inherits the base class's
_processes_spawnedflag through ordinary attribute lookup, so it never spawns its own workers and instead dispatches tests onto the base class's already torn-down queues, blocking forever incompletion_queue.get(). Consult the class's own__dict__in both the spawn and the teardown guard so that every leaf class owns its worker pool.This makes test_replicate.py order-independent. Under pytest the base class ReplicateTest runs before ReplicateFullyShardInit, which then hangs until the harness timeout; under unittest the subclass happens to sort first, which is why the bug stayed latent in CUDA CI.
Backport of the test-infra portion of pytorch#189362.
Test Plan
python -m pytest distributed/_composable/test_replicate.py -k test_replicate_device_idTest Result