[release/2.14] Carry-forward ROCm cherry-picks from release/2.13 (full) by pablo-garay · Pull Request #3594 · ROCm/pytorch

pablo-garay · 2026-08-21T19:33:17Z

Summary

Carry-forward of ROCm-specific commits from release/2.13 onto release/2.14 ? the 2.14 analog of the 2.13 "IFU dedup" (PR #3503).

Analyzes the full IFU delta from PR #3591 at release/2.13's tip (6596a0d): 81 commits in release/2.13 not in release/2.14, each either dropped (already in 2.14, or 2.13-only) or cherry-picked here.

Method

Same rules as the release/2.13 cherry-pick process:

Drop ? already in 2.14: patch-id equivalence (git cherry), upstream PR already present, or content already relocated/absorbed into 2.14 (verified during cherry-pick).
Drop ? 2.13-only: version bump, related_commits, dependency/requirements pins, release-line CD/build changes.
Carry-forward: ROCm-specific fixes not present in 2.14 in any form.

Commits carried forward (13)

Source	Notes
Fix reentrant deadlock in `torch.cuda._lazy_call` (#3496)	clean
Skip ROCm-failing tests in test_cuda/test_nn (#3504)	conflict resolved ? kept 2.14's `test_graph_capture_error...` cleanup + took `@skipIfRocm`; kept 2.14's `test_allocator_backend` (already fixes LD_LIBRARY_PATH via `subprocess_env()`)
skip test_hip_device_count (rocprofiler-sdk)	clean
skip flaky test_Embedding_discontiguous_cuda	clean
Skip profiler check for foreach tests on ROCm (#3262)	clean
Skip test_is_pinned_no_context on python 3.14+	clean
Skip flaky test_reentrant_parent_error_on_cpu	clean
Filter out CK gemms for gfx1250-only arch (#3545)	clean
Fix hipify skipping sources under Windows subst drives (ROCM-29365) (#3561)	clean
Fixed gfx1250_preferred_backend test (#3546)	clean
[ROCm] Optimize AMD normalization backward kernel (pytorch#189405)	clean
[release/2.13] Build rocSHMEM for the requested gfx target (#3582)	clean
fix PicklingError handling for Python 3.14	clean

Removed from this PR (per team)

[RELEASE_ONLY] skip test_autoheuristic in-code ([ROCm][RELEASE_ONLY] skip test_autoheuristic in-code (already disabled via issue) pytorch/pytorch#180927) ? not needed on 2.14. Upstream pytorch/pytorch#180927 is RELEASE_ONLY, merged to upstream release/2.12 only (never on main or upstream release/2.14).

Reviewed and excluded (team decision: not needed for 2.14)

Commit	Reason
Update composable_kernel submodule with gfx1033 support (#3499)	2.14's CK pin already includes gfx1033; carrying would move the submodule pointer backward.
Revert "Skip cudagraphs for kernel-free inductor graphs" (pytorch#187347)	2.14 already has the fixed upstream re-land on top of the revert; carrying = regressing 2.14.
Revert "[Reland] Port index_add fast path" (pytorch#188196)	2.14 already has the newer re-land (pytorch#189506) on top of the revert; carrying = regressing 2.14.

Dropped as already in 2.14 (verified during cherry-pick)

rocm-smi deprecation ([release/2.12] rocm-smi deprecation #3434) ? 2.14 already has zero rocm_smi refs.
Enable linalg via hipSOLVER, remove Magma ([ROCm][release/2.13] Enable linalg tests via hipSOLVER and remove Magma dependency #3455) ? all 6 upstream PRs already in 2.14.
Stabilize ROCm 7.14 testing ([release/2.13] Stabilize ROCm 7.14 testing on release/2.13 #3525) ? 2.14 already carries the guards + skip messages.
Add gfx1250 support in GroupedBlas (Add missing gfx1250 support in GroupedBlas #3532) ? already in 2.14 via upstream Deduplicate scaled_mm arch gate into ScaledBlasDeviceUtils pytorch/pytorch#190703 (ScaledBlasDeviceUtils).
Fix uninitialized embedding indices test_embedding_meta_indices (Fix uninitialized embedding indices in test_embedding_meta_indices pytorch/pytorch#188002) ? upstream already in 2.14.
Unskip ROCSHMEM triton UTs on rocm 7.14 ([release/2.13] Unskip ROCSHMEM triton UTs on rocm 7.14 #3578) ? targets a 2.13-only 7.14 module skip that 2.14 never had; tests already run on 2.14.

Remaining commits drop as 2.13-only (version/pins/related_commits/CD) or as already-in-2.14 by patch-id/upstream.

Test plan

TheRock full-stack dev build (torch + torchaudio + torchvision + triton) on gfx94X-dcgpu for Python 3.10-3.14, all Build jobs green (ROCm 10.1.0a20260822, on the aws-linux-scale-rocm-prod pool), after bringing the branch up to its base so related_commits resolves torchaudio/torchvision:
TheRock CI wheel build on gfx94X-dcgpu for Python 3.10-3.14 all green (ROCm 10.1.0a20260822, torch-only, on the aws-linux-scale-rocm-prod pool, ~1h40m each):

Processed Commits List

Spreasheet with each commit evaluated for cherry-pick & resolution (drop/include) with reasoning:

pr3594-commit-outcomes.xlsx

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) (cherry picked from commit 24cf333)

Adds @skipIfRocm to 4 tests that consistently fail in the TheRock wheel test lane (test_pytorch_wheels.yml) on gfx94X/MI300 with torch==2.13.0a0+rocm7.14.0. - test_allocator_backend - test_graph_capture_error_releases_reserved_segments - test_lazy_call_reentrant_set_rng_state_does_not_deadlock - test_linear_cross_entropy_loss_default Also removes the now-dead @skipIfRocmVersionLessThan((7, 14)) from test_graph_capture_error_releases_reserved_segments — with @skipIfRocm on top it is unreachable on ROCm. Test plan - Validated via ROCm/TheRock test_pytorch_wheels.yml run #30218037322 on gfx94X-dcgpu with torch==2.13.0a0+rocm7.14.0 — all 4 tests now SKIPPED, run passes - Core UTs passing on 2.13: https://github.com/ROCm/TheRock/actions/runs/30218037322 (cherry picked from commit 523576b)

Cherry-pick of #3505 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> (cherry picked from commit ae1ef9e)

Cherry-pick of #3513 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> (cherry picked from commit 2396d45)

…OCM-21749) (#3262) ## Summary Fixes ROCM-21749: `test_big_num_tensors__foreach_norm` fails on all ROCm GPUs. ## Root Cause `ForeachFuncWrapper.__call__` profiles a `_foreach_norm` call and checks for `multi_tensor_apply_kernel` in profiler output to verify the multi-tensor fast-path was taken. On CUDA, Kineto captures demangled kernel names like `multi_tensor_apply_kernel[...]`. On ROCm, ROCTracer/rocprofiler captures different symbols (`hipLaunchKernel`, `vectorized_elementwise_kernel`, etc.), so `mta_called` is always `False` and the assertion fails. ## Fix Add `or TEST_WITH_ROCM` to `skip_profiler_check` so the profiler-based assertion is bypassed on ROCm, consistent with the existing skip for CUDA 12.6/12.8 flakiness. ## Test Plan - `test_big_num_tensors__foreach_norm_*` should now pass on ROCm GPUs (MI210, MI300X, MI350X) - No change in behavior on CUDA Jira: https://amd-hub.atlassian.net/browse/ROCM-21749 Co-authored-by: Subramanian <ssubrama1@amd.com> (cherry picked from commit f8efdb3) (cherry picked from commit e143818)

…behavior change Cherry-pick of #3394 Co-authored-by: zjliu-amd <zhuojin.liu@amd.com> (cherry picked from commit 618bbf8) (cherry picked from commit 16fa274)

Cherry-pick of #3521 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> (cherry picked from commit d0fcdf9)

…3545) Fixes empty HIP_ARCHITECTURE build failure when gfx1250 is the only provided value to PYTORCH_ROCM_ARCH (cherry picked from commit bdbef9c)

…3561) ## Summary On Windows ROCm CI, the `release/2.13` torchaudio wheel build fails compiling GPU sources with `fatal error: 'cuda_runtime_api.h' file not found` (ROCM-29365 / ROCm/TheRock#7266). Root cause is in `CUDAExtension`'s hipify step, not in torchaudio. `CUDAExtension` scopes hipify to `os.getcwd()` and filters with `includes=[build_dir/*]`. The Windows CI runs the build under a `subst` drive (`CHECKOUT_ROOT: B:/src`), but torchaudio computes its source directory with `Path(...).resolve()`, which follows the `subst` alias back to the real drive (`C:\...`). So `build_dir` (`B:\...`) and the source paths (`C:\...`) use different drive spellings, the `includes` filter matches nothing, and hipify silently skips the GPU sources — leaving the raw `#include <cuda_runtime_api.h>` (via `cuda_utils.h`) for hipcc. This is the same `B:` vs `C:` mount split as ROCM-28908; that fix removed the earlier hard crash, which let the build progress far enough to expose this. ## Fix Normalize both sides with `os.path.realpath` so the build directory and the source paths share one canonical form (collapsing symlinks/junctions and `subst` drive aliases). With that, `includes=[build_dir/*]` matches again and the header-recursion that rewrites `cuda_runtime_api.h` → `hip/hip_runtime_api.h` runs. ```python build_dir = os.path.realpath(os.getcwd()) ... extra_files=[os.path.realpath(s) for s in sources], ... s_abs = os.path.realpath(source) ``` ## Evidence / validation - Failing run 31459420855: hipcc compiles the original `compute.cu` (not `.hip`); torchaudio's `CUDAExtension` hipify passes report `replaced kernel launches: 0`; `hip_utils.h`/`compute.hip` never appear in the log. - Local repro of the exact `CUDAExtension` hipify call on the real torchaudio GPU sources: - cwd == source tree → `cuda_utils.h → hip_utils.h`, `compute.cu → compute.hip`, `replaced kernel launches: 1`. - cwd != source tree (mimics `B:` vs `C:`) → 0 conversions, no `.hip` — matches the failing CI. - `realpath` collapses the `subst` alias so both paths agree, restoring the working (cwd == source) behavior. No effect on Linux, where the paths already agree. ## Test plan - [x] Windows ROCm `release/2.13` wheel pipeline builds torchaudio successfully across supported Python versions. - [x] Linux ROCm extension builds unaffected (regression check). - [x] ROCM-28908 relpath fix remains intact. Refs: ROCM-29365, ROCm/TheRock#7266, ROCM-28908. ## Successful Passing Test Validated end-to-end with a full Windows ROCm multi-arch PyTorch wheel build on this branch (same TheRock workflow that fails on `release/2.13`, only `pytorch_git_ref` repointed to this branch). **Build job succeeded:** [Build Multi-Arch Windows PyTorch Wheels (dev, 3.12, garayp/rocm-29365-hipify-realpath)](https://github.com/ROCm/TheRock/actions/runs/31667526330) — torch, torchaudio, and torchvision all built on Windows ROCm with the fix. This is definitive confirmation. The exact ROCM-29365 sources now hipify and compile (build log, lines 714709-714718): ``` cuda_utils.h -> hip_utils.h [ok] rnnt/gpu/compute.cu -> rnnt/gpu/compute.hip [ok] forced_align/gpu/compute.cu -> forced_align/gpu/compute.hip [ok] Total number of replaced kernel launches: 11 Total number of replaced kernel launches: 8 ``` Before the fix (on `release/2.13`), these passes reported `Total number of replaced kernel launches: 0` and the CUDA header was left unhipified, so `hipcc` failed with `cuda_runtime_api.h` not found. With the fix, `hipcc` compiles `compute.hip` and resolves `hip_utils.h` cleanly. _Note: the downstream GPU test jobs in that run failed at dependency install (`No matching distribution found for rocm==10.1.0a20260810` from the dev-releases index) — an unrelated TheRock test-harness/package-index gap, not this change. The tests never ran any torchaudio code._ Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: pablo-garay <pablo-garay@users.noreply.github.com> (cherry picked from commit 3a4c3a1)

Fixes https://amd-hub.atlassian.net/browse/ROCM-29096. Once merged, will cherry-pick to other branches. test_cuda.py::TestCuda::test_preferred_blas_library_settings (cherry picked from commit 4fba733)

…two pass implementations where they are most efficient. (pytorch#189405) Hugging Face model dropped **5-10%** after switching to tiled kernel (10+ model and tests). Implemented mixed approach, using combination of Tiled and Two Pass: Hugging Face (huggingface_bart) performance is back: Legacy Two pass performance=1652, 1644 Tiled performance =**1574, 1568** Mixed performance =1649, 1652 Synthetic reproducer where performance from chess board like became uniform, keeping benefits from both implementation: # Layer Norm Backward Benchmark: Tiled only vs Tiled+Two pass **Device:** AMD Instinct MI350X | **Warmup:** 20 | **Iters:** 100 | **Runs averaged:** 3 ## Summary (avg µs over 3 runs) | Benchmark | Op | Shape | dtype | Branch | Tiled only | Tiled+Two pass | Δ (µs) | Speedup | Winner | |-----------|-----|-------|-------|--------|------------|----------------|--------|---------|--------| | tile8_small | layer_norm | (32, 512) | float16 | Tile-8 | 7.94 | 8.16 | -0.22 | 0.97× | ~tie | | tile64_medium | layer_norm | (96, 768) | float16 | Tile-64 | 20.71 | 20.78 | -0.07 | 1.00× | ~tie | | tile128_medium | layer_norm | (192, 1024) | bfloat16 | Tile-128 | 8.20 | 8.51 | -0.31 | 0.96× | ~tie | | tile256_large | layer_norm | (4096, 1024) | float16 | Tile-256 | 31.07 | 13.12 | +17.95 | 2.37× | Two pass | | tile256_bert | layer_norm | (1024, 768) | float16 | Tile-256 BERT | 11.15 | 11.05 | +0.10 | 1.01× | ~tie | | tile256_large_bf16 | layer_norm | (4096, 1024) | bfloat16 | Tile-256 | 31.28 | 13.10 | +18.18 | 2.39× | Two pass | | tile256_large_fp32 | layer_norm | (4096, 1024) | float32 | Tile-256 | 30.87 | 15.90 | +14.97 | 1.94× | Two pass | | two_pass_huge_M | layer_norm | (131072, 64) | float16 | Two-pass M-parallel | 26.48 | 27.70 | -1.22 | 0.96× | ~tie | | llm_hidden_4096 | layer_norm | (32, 4096, 4096) | bfloat16 | Tile-256 LLM | 509.65 | 513.38 | -3.73 | 0.99× | ~tie | | gpt2_style | layer_norm | (8, 1024, 768) | float16 | Tile-256 GPT2 | 58.22 | 16.56 | +41.66 | 3.52× | Two pass | | rms_tile256 | rms_norm | (4096, 1024) | float16 | Tile-256 rms | 22.34 | 11.18 | +11.16 | 2.00× | Two pass | | rms_llm | rms_norm | (32, 4096, 4096) | bfloat16 | Tile-256 rms LLM | 476.18 | 476.03 | +0.15 | 1.00× | ~tie | | rms_two_pass | rms_norm | (131072, 64) | float16 | Two-pass rms | 16.73 | 17.29 | -0.56 | 0.97× | ~tie | Δ = Tiled only − Tiled+Two pass (negative = Two pass faster). Speedup = Tiled only / Tiled+Two pass. Also result of generated test for correcteness: tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile128_medium] PASSED [ 4%] tests/test_correctness.py::test_rms_norm_autograd_matches_cpu_reference PASSED [ 9%] tests/test_correctness.py::test_edge_case_shape_M_eq_1 PASSED [ 13%] tests/test_correctness.py::test_rms_norm_gamma_backward_matches_cpu_reference[rms_two_pass] PASSED [ 18%] tests/test_correctness.py::test_output_mask_selects_expected_grads[dgamma_and_dbeta] PASSED [ 22%] tests/test_correctness.py::test_rms_norm_gamma_backward_matches_cpu_reference[rms_tile256] PASSED [ 27%] tests/test_correctness.py::test_noncontiguous_input_matches_reference PASSED [ 31%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[two_pass_huge_M] PASSED [ 36%] tests/test_correctness.py::test_output_mask_selects_expected_grads[dbeta_only] PASSED [ 40%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile256_large_fp32] PASSED [ 45%] tests/test_correctness.py::test_rms_norm_gamma_backward_matches_cpu_reference[rms_llm] PASSED [ 50%] tests/test_correctness.py::test_configs_cover_all_tile_branches PASSED [ 54%] tests/test_correctness.py::test_layer_norm_autograd_matches_cpu_reference PASSED [ 59%] tests/test_correctness.py::test_output_mask_selects_expected_grads[dgamma_only] PASSED [ 63%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[gpt2_style] PASSED [ 68%] tests/test_correctness.py::test_edge_case_shape_N_eq_1 PASSED [ 72%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile256_bert] PASSED [ 77%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile256_large] PASSED [ 81%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile256_large_bf16] PASSED [ 86%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[llm_hidden_4096] PASSED [ 90%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile8_small] PASSED [ 95%] tests/test_correctness.py::test_layer_norm_gamma_beta_backward_matches_cpu_reference[tile64_medium] PASSED [100%] ========== 22 passed in 9.26s =================================== Used AI assistance from Cursor. Short performance reproducer: [layer_norm_gamma_beta_backward_reproducer.py](https://github.com/user-attachments/files/31031711/layer_norm_gamma_beta_backward_reproducer.py) [run.sh](https://github.com/user-attachments/files/31031712/run.sh) Pull Request resolved: pytorch#189405 Approved by: https://github.com/jeffdaily Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit 29e524b)

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

Cherry-pick of #3584 Co-authored-by: albmalamd <albert.malewski@amd.com> (cherry picked from commit 6596a0d)

…d_commits + requirement pins) Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com>

jithunnair-amd · 2026-08-25T18:23:24Z

NOTE: This PR was accidentally merged using "Squash and merge", so it was manually rebased on top of release/2.14 to retain the individual commits.

iupaikov-amd and others added 13 commits August 21, 2026 12:30

[release/2.13] Fix reentrant deadlock in torch.cuda._lazy_call (#3496)

9380746

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) (cherry picked from commit 24cf333)

skip test_hip_device_count due to rocprofiler-sdk issue

51b85c5

Cherry-pick of #3505 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> (cherry picked from commit ae1ef9e)

skip flaky test_Embedding_discontiguous_cuda

5375774

Cherry-pick of #3513 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> (cherry picked from commit 2396d45)

Skip test_is_pinned_no_context on python 3.14 and above due to known …

2c57184

…behavior change Cherry-pick of #3394 Co-authored-by: zjliu-amd <zhuojin.liu@amd.com> (cherry picked from commit 618bbf8) (cherry picked from commit 16fa274)

Skip flaky test_reentrant_parent_error_on_cpu

afa20e9

Cherry-pick of #3521 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> (cherry picked from commit d0fcdf9)

[ROCm] - Filter out CK gemms for when ONLY gfx1250 arch is provided (#…

646541e

…3545) Fixes empty HIP_ARCHITECTURE build failure when gfx1250 is the only provided value to PYTORCH_ROCM_ARCH (cherry picked from commit bdbef9c)

Fixed gfx1250_preferred_backend test (#3546)

a128789

Fixes https://amd-hub.atlassian.net/browse/ROCM-29096. Once merged, will cherry-pick to other branches. test_cuda.py::TestCuda::test_preferred_blas_library_settings (cherry picked from commit 4fba733)

fix PicklingError handling for Python 3.14

3bb3f9d

Cherry-pick of #3584 Co-authored-by: albmalamd <albert.malewski@amd.com> (cherry picked from commit 6596a0d)

pablo-garay force-pushed the pablo-garay/release-2.14-carryforward branch from 1410f57 to 3bb3f9d Compare August 24, 2026 19:04

Merge release/2.14 into carryforward (bring branch up to base: relate…

e200bc4

…d_commits + requirement pins) Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com>

jithunnair-amd merged commit 6fa6aef into release/2.14 Aug 25, 2026
3 checks passed

jithunnair-amd deleted the pablo-garay/release-2.14-carryforward branch August 25, 2026 18:06

jithunnair-amd restored the pablo-garay/release-2.14-carryforward branch August 25, 2026 18:08

jithunnair-amd mentioned this pull request Aug 25, 2026

[release/2.14] Cherry-pick Windows HIPIFY realpath fix for torchaudio (ROCM-29365) #3595

Closed

4 tasks

Navigation Menu

[release/2.14] Carry-forward ROCm cherry-picks from release/2.13 (full) - #3594

[release/2.14] Carry-forward ROCm cherry-picks from release/2.13 (full)#3594
jithunnair-amd merged 14 commits into
release/2.14from
pablo-garay/release-2.14-carryforward

pablo-garay commented Aug 21, 2026 •

edited

Loading

Uh oh!

Uh oh!

jithunnair-amd commented Aug 25, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

Conversation

pablo-garay commented Aug 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Method

Commits carried forward (13)

Removed from this PR (per team)

Reviewed and excluded (team decision: not needed for 2.14)

Dropped as already in 2.14 (verified during cherry-pick)

Test plan

Processed Commits List

Uh oh!

Uh oh!

jithunnair-amd commented Aug 25, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

pablo-garay commented Aug 21, 2026 •

edited

Loading