[release/2.13] IFU dedup: backport refined release/2.12 ROCm commits#3503
Open
pablo-garay wants to merge 13 commits into
Open
[release/2.13] IFU dedup: backport refined release/2.12 ROCm commits#3503pablo-garay wants to merge 13 commits into
pablo-garay wants to merge 13 commits into
Conversation
…d via issue) (pytorch#180927) [ROCm] skip test_autoheuristic in-code (already disabled via issue) (cherry picked from commit 366eac9)
…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 <[email protected]> (cherry picked from commit f8efdb3)
…collision (#3303) This file is excluded from hipify (see tools/amd_build/build_amd.py), so it ships its own self-contained __nv_bfloat16 type plus the __float2bfloat16/__bfloat162float conversions for the runtime-compiled fuser kernel. Starting with ROCm 7.x, hiprtc auto-injects a runtime header (hiprtc_runtime.h) that also defines __float2bfloat16 returning __hip_bfloat16. Since C++ cannot overload on return type alone, the two definitions collide. Rather than gate on ROCM_VERSION (a build-time value that does not reliably track whether the *runtime* hiprtc actually provides these symbols), we keep PyTorch's own round-to-nearest-even implementation and rename our private definitions via the preprocessor so they cannot clash. hiprtc's definitions remain (parsed earlier in the translation unit) but go unused. These #defines must stay in effect for the rest of the kernel source -- do NOT #undef them, since the emitted kernel body calls these conversions after this header is spliced in. Pull Request resolved: pytorch#185656 Approved by: https://github.com/jeffdaily (cherry picked from commit 5ca8c13) https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Andy Lugo <[email protected]> (cherry picked from commit 6aad4ef)
Before ROCm 7.14, HIP does not report about errors during stream capture, need to use check hipStreamCaptureStatus to reject capture-unsafe work early. From ROCm 7.14+, HIP reports capture error through, so that pre-check is unnecessary Fixes pytorch#180232 Pull Request resolved: pytorch#187110 Approved by: https://github.com/jeffdaily, https://github.com/naromero77amd (cherry picked from commit 1047ba5) (cherry picked from commit caf158f)
…7613) Cherry-pick of #3336 Co-authored-by: zjliu-amd <[email protected]> Co-authored-by: Xinya Zhang <[email protected]> (cherry picked from commit daea54b)
…ed_callables (#3354) ### Summary This PR updates `torch.cuda.make_graphed_callables` so warmup and graph capture use the same stream, avoiding ROCm graph capture failures tied to hipBLASLt. ### Why On ROCm, hipBLASLt handle caching is per (device, stream). If `stream=` is omitted, `torch.cuda.graph` uses its own default capture stream, which is not the stream used for warmup. The first real use on that capture stream will trigger a lazy hipblasLtCreate there. Handle creation runs capture‑unsafe internal allocation/setup, which fail with stream-capture errors (and sometimes hang) during capture. Running warmup on the same stream passed into `torch.cuda.graph(..., stream=...)` creates and caches the hipBLASLt handle before capture, so hipblasLtCreate (and its associated allocations) are not hit mid‑capture. Fixes pytorch#179961 Pull Request resolved: pytorch#187745 Approved by: https://github.com/jeffdaily (cherry picked from commit d9d756c) (cherry picked from commit eaf6098)
Backporting pytorch#180310 to release 2.12 (cherry picked from commit 9ffbd03)
…ts (pytorch#188580) Cherry-pick of #3384 Co-authored-by: zjliu-amd <[email protected]> (cherry picked from commit 6e2a13b)
…behavior change Cherry-pick of #3394 Co-authored-by: zjliu-amd <[email protected]> (cherry picked from commit 618bbf8)
…y repeated call… (#3403) This is a cherry-pick of pytorch#188673 Removing a repeated call of hipMemMap which is causing the errors in SymmetricMemory. (cherry picked from commit 7a68a17)
#3421) Validation theRock Run: https://github.com/ROCm/TheRock/actions/runs/29045360944 Cherry-pick of upstream [pytorch#188597](pytorch#188597) onto `release/2.12` to add initial ROCm support for **gfx1250 (CDNA5)**. Co-authored-by: @glen-amd - Gates gfx1250-specific behavior on ROCm 7.14+ (the support floor), applied consistently across CUDABlas.cpp, ScaledBlas.cpp, CUDAHooks.cpp, and the test helpers. - Adds gfx1250 to the hipBLASLt preferred/supported arch lists and the hipSparseLt support check (ROCm 7.14+). - Extends the scaled GEMM / MX-format paths (block-wise Float8_e8m0fnu, mxfp8, mxfp4) to gfx1250. - Sets the gfx1250 shared-memory limit to 320 KB (vs 160 KB on gfx950), matching on the suffix-stripped arch name. - Adds CDNA2/CDNA3/CDNA5 "or later" arch helpers and updates FP8 / MX GEMM test gating in common_cuda.py. - Builds Composable Kernel (CK) GEMM as a separate ck_gemm library with gfx1250 filtered out of HIP_ARCHITECTURES, so a multi-arch build that includes gfx1250 keeps CK GEMM for the other archs instead of failing to compile for gfx1250 (composable_kernel has no gfx1250 support yet). The filter is removed once CK supports gfx1250. - Extends the gfx942 nontemporal vectorized-load path in MemoryAccess.cuh to gfx1250. - Refactors duplicated arch-check logic in CUDABlas.cpp, ScaledBlas.cpp, and cuSPARSELtOps.cpp per review feedback. - Writes the CublasHandlePool workspace archs in full (gfx942/gfx950/gfx1250). Enabled via the main merge: - Flash attention and memory-efficient attention on gfx1250 via AOTriton 0.12.1b (from pytorch#188242), which this branch merges in. What is not enabled yet: - CK SDPA on gfx1250: composable_kernel has no gfx1250 support; the CK SDPA target auto-filters to gfx942/gfx950. - CK GEMM on gfx1250: filtered out (see above) until composable_kernel supports it. - FP8 grouped GEMM on gfx1250: MSLK builds only gfx942/gfx950. - int4 mm on gfx1250: the tinygemm MFMA kernel needs a WMMA port; the ops return a clear "not supported yet" error and the int4 unit tests skip gfx1250. CI note: - gfx1250 is not added to the ROCm 7.2 (rocm-n) docker image, whose HIP compiler cannot target gfx1250. gfx1250 CI belongs on the nightly ROCm image and is added in a follow-up. Co-authored-by: @glen-amd - Upstream PR: pytorch#188597 - Related AOTriton bump: pytorch#188242 --------- Co-authored-by: Prachi Gupta <[email protected]> Co-authored-by: Claude Opus 4.8 (1M context) <[email protected]> Co-authored-by: Jeff Daily <[email protected]> (cherry picked from commit 13eab2b)
Validation theRock Run: https://github.com/ROCm/TheRock/actions/runs/29045354548 [ROCm] Bump AOTriton to 0.13b Notable new features compared with 0.12.1b: * Debug printing can be enabled with env var `AOTRITON_DEBUG_LEVEL`, without re-compiling + We use the same logging level as of `AMD_LOG_LEVEL` - CAVEAT: AOTriton itself does not read `AMD_LOG_LEVEL` + Set `AOTRITON_DEBUG_LEVEL=5` to print everything * **BREAKING** `.aks2` files are packed into **uncompressed** zip files + `chcp 65001` is not required to build on Win32 any more. * TheRock 7.15 support Bug Fixes: Add `check_requires_grad_and_nested` back to `can_use_mem_efficient_attention`, which was removed accidentally. *Other changes are intentionally left out: they consist mainly of internal refactoring that is not perceivable to pytorch users.* [ROCm] Enable gfx1250 SDPA with AOTriton 0.12.1b (pytorch#188612) Notable New Features: * Tech Preview support of gfx1250 hardware + For broader testing, `TORCH_ROCM_AOTRITON_ENABLE_EXPERIMENTAL=1` is NOT NEEDED for gfx1250 at the moment. + This support tier is tracked by new function `aotriton::isArchTechPreview()` [ROCm] Bump AOTriton to 0.12b (pytorch#184288) Notable new features: * **BREAKING** Varlen LSE tensor shape changes to (H, Total_seqlen) * Support head_dim != head_dim_v * Support `use_deterministic_algorithims` * Support seqused_k in test/test_varlen_attention.py * gfx1100 and gfx1151 promoted out of experimental * Partial FAv3 support on gfx950 Bug Fixes: * GQA kernel failed to read bias tensor with the right offset. Known Issues * gfx950's Triton kernel has problem handling hdim=16's fwd, in addition to hdim=48/80's bwd. * Disables gfx90a's CK SDPA support due to GPU Segfault. --------- Co-authored-by: Claude Opus 4.8 <[email protected]> (cherry picked from commit 10d8e1d)
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
9 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.
Summary
IFU (integrate-from-upstream) dedup pick: backports the refined set of
release/2.12ROCm commits that are still missing fromrelease/2.13, on top of the currentrelease/2.13tip (6ae7dbe2aab, #3499).Started from a ~79-commit candidate set, deduplicated by patch-id and semantic overlap against what
release/2.13already carries, landing 13 commits. Conflicts were resolved keepingrelease/2.13's newer logic where it already superseded 2.12 (test decorators, SDPA constraints, dependency pins, CK/mslk build config).Commits (top = newest)
575b373f95cBump AOTriton from 0.11.2b to 0.13b ([release/2.12] Bump AOTriton from 0.11.2b to 0.13b #3419)af2e80a240a[ROCm] Add initial support for gfx1250 ([ROCm] Add initial support for gfx1250 pytorch/pytorch#188597) ([release/2.12] [ROCm] Add initial support for gfx1250 (#188597) #3421)3a7e84c7cdbFixed memory errors in SymmetricMemory caused by repeated calls ([release/2.12] [ROCm] Fixed memory errors in SymmetricMemory caused by repeated call… #3403)1d99b803c53Skip test_is_pinned_no_context on python 3.14+ (behavior change)8f26b5065cd[ROCm] Use portable compiler names in cpp_extension compatibility tests ([ROCm] Use portable compiler names in cpp_extension compatibility tests pytorch/pytorch#188580)7da221a5b39Backport upstream PR 180310 ([release/2.12] Backport upstream PR 180310 #3352)62df3cdfb2freuse one stream for warmup and captures in make_graphed_callables ([release/2.12] reuse one stream for warmup and captures in make_graphed_callables #3354)9de5161e4aa[ROCm] Fix test_cross_entropy_loss_2d_out_of_bounds_class ([ROCM] Fix test_cross_entropy_loss_2d_out_of_bounds_class pytorch/pytorch#187613)ca4a752528cFix graph capture error handling for ROCm 7.14+ ([release/2.12] Fix graph capture error handling for ROCm 7.14+ #3323)640f125845eRename duplicate hiprtc names for jit to avoid symbol collision ([release/2.12] Rename duplicate hiprtc names for jit to avoid symbol collision #3303)64d9e9dd4abGuard NCCL one-sided API behind device support ([release/2.12] Guard NCCL one-sided API behind device support #3295)f01e0e507e0[ROCm] Skip profiler check for foreach tests on ROCm (ROCM-21749) ([ROCm][release/2.12] Skip profiler check for foreach tests on ROCm (ROCM-21749) #3262)8eb83e68471[ROCm][RELEASE_ONLY] skip test_autoheuristic in-code (already disabled via issue) ([ROCm][RELEASE_ONLY] skip test_autoheuristic in-code (already disabled via issue) pytorch/pytorch#180927)Notes
.ci/docker/requirements-ci.txtverified unchanged from therelease/2.13base — dependency pins from release/2.13 bringup: pin requirements.txt #3445 preserved, no unintended pin drift.mix_order_reduction_allow_multi_stagesROCm evaluation.Test plan
release/2.13