[AUTOGENERATED] release/2.14_IFU_20260821 - #3591
Open
rocm-repo-management-api-6[bot] wants to merge 87 commits into
Open
rocm-repo-management-api-6[bot] wants to merge 87 commits into
rocm-repo-management-api-6[bot] wants to merge 87 commits into
Conversation
) * [release 2.13] Apply Release only changes to 2.13 branch Release-only changes for the release/2.13 branch cut, produced by running scripts/release/apply-release-changes.sh. The script repoints reusable workflows and composite actions from @main to @release/2.13, rewrites templates to release/2.13 (with checkout_pr_head=False so PRs build the merge base rather than the PR head), pins the XLA checkout to the r2.13 branch, pins the disabled/unstable jobs and disabled-tests S3 JSON blobs to fixed versionIds, sets RELEASE_VERSION_TAG=2.13 for the workflow-regeneration lint check, and drops the pull_request-specific checkout ref from the linux binary build/test workflows. Only files that are tracked in release/2.13 are included. The 2.12 PR additionally touched .ci/manywheel/build_cuda.sh and a few workflow files (torchbench/nitpicker/quantization-periodic) that are not tracked in this checkout, so they are intentionally omitted here. Test Plan: Ran the release script and linter: ``` DRY_RUN=disabled ./scripts/release/apply-release-changes.sh lintrunner -a ``` lintrunner reported only pre-existing ACTIONLINT shellcheck warnings on generated workflow lines unrelated to the release-version edits, and made no changes to the staged files. Verified every staged hunk is a release-only edit (@main -> @release/2.13, main -> release/2.13, XLA r2.13 pin, S3 versionId pins, RELEASE_VERSION_TAG=2.13, and the pull_request ref removal) with no submodule or unrelated changes. This PR was authored with the assistance of Claude Code. * [release 2.13] Pin Linux manywheel builder docker images Release builds should use a fixed, reproducible build toolchain instead of the floating builder image tags that main tracks (e.g. pytorch/manylinux2_28-builder:cuda12.6). For 2.12 the binary build workflows resolved the image dynamically via calculate-docker-image; for 2.13 we freeze that resolved image as a literal pin. The pin is applied in the generator (generate_binary_build_matrix.py) rather than in the generated YAML directly, so re-running .github/regenerate.sh (and the lint job that asserts the generated files are up to date) reproduces the pinned tags. wheel_container_image_tag_prefix() appends the pin only for the linux manywheel OSes (linux, linux-aarch64, linux-s390x), since only those builds run inside these containers; windows and macos keep the plain tag prefix. The pin suffix is the .ci/docker tree hash, f38ba0b (git rev-parse HEAD:.ci/docker), which is exactly the tag that .github/actions/binary-docker-build (and the s390x equivalent) publish as ${prefix}-${CI_FOLDER_SHA}. It matches the image already published on Docker Hub, e.g. pytorch/manylinux2_28_aarch64-builder:cpu-aarch64-f38ba0b10220982e39441d29d203d803a2b56c92 Test Plan: Regenerated the workflows in release mode and verified the result: ``` RELEASE_VERSION_TAG=2.13 python3 .github/scripts/generate_ci_workflows.py ``` - Every `image:` and matrix `docker_image_tag_prefix` in the three linux manywheel generated workflows now carries the `-f38ba0b...` suffix (cpu, cpu-aarch64, cuda12.6/13.0/13.2, rocm7.1/7.2, xpu, cpu-s390x); no linux builder image remains floating. - Windows and macos generated workflows are unchanged (plain `cpu`, `cuda12.6`, ...), confirming the pin is linux-only. - Re-running the generator a second time produced byte-identical output (md5sum unchanged), so the regeneration/lint up-to-date check is stable. Note: the local linters that shell out to `uv` could not run in this environment; the change was format-checked manually. This PR was authored with the assistance of Claude Code.
pytorch#186992) The release-only DOCKER_IMAGE_PIN appends the .ci/docker tree hash to the manywheel builder image tags so the release uses a reproducible toolchain. This was applied to all manywheel OSes, including s390x. Root cause: unlike the x86_64 and aarch64 builder images, the s390x builder image is built locally on the self-hosted linux.s390x runners and is never published to docker.io under the pinned tag. The pinned pull therefore fails: Error: initializing source docker://pytorch/manylinuxs390x-builder:cpu-s390x-<hash>: reading manifest ... requested access to the resource is denied Fix: drop linux-s390x from MANYWHEEL_OSES so s390x falls back to the floating "cpu-s390x" tag, while x86_64 and aarch64 keep their pinned tags. Regenerated the s390x nightly workflow accordingly. Test Plan: - Regenerated workflows and confirmed the pin is gone only for s390x: ``` python3 .github/scripts/generate_ci_workflows.py grep -c "cpu-s390x-f38ba0b" .github/workflows/generated-linux-s390x-binary-manywheel-nightly.yml # 0 grep -m2 docker_image_tag_prefix .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml # still pinned ``` Authored-with: Claude (AI assistant)
Update torch-xpu-ops commit pin (pytorch#186768) Update the torch-xpu-ops commit to [intel/torch-xpu-ops@bc2942](intel/torch-xpu-ops@bc29424), includes: Fixes pytorch#186661 for build warnings: - [-Werror=template-id-cdtor] caused by SYCLTLA with GCC15 - [-Werror=int-in-bool-context] caused by ‘*’ in boolean context - [-Wdeprecated-attributes] caused by deprecated attribute intel::reqd_sub_group_size Migrate std::exp2, std::expm1, std::frexp and std::copysign to sycl:: namespace for SYCL compatibility. Add compiler version check for SymmetricMemory. Pull Request resolved: pytorch#186768 Approved by: https://github.com/guangyey, https://github.com/EikanWang, https://github.com/atalman (cherry picked from commit ddc3c36) Co-authored-by: Cui, Yifeng <yifeng.cui@intel.com> Co-authored-by: Wang, Chuanqi <chuanqi.wang@intel.com>
pytorch#187001) Release candidate builds triggered by a tag (e.g. v2.13.0-rc1) were producing wheels that pinned triton to a dev-style version such as triton==3.7.1+git5d6048aa instead of the release triton==3.7.1, which then fails to install in the test job because that version does not exist on the test index. Root cause: the unified inline manywheel build job (cpu/cpu-aarch64/ cuda/cuda-aarch64) checks out the raw commit via `ref: github.sha` with `fetch-depth: 2`. Checking out a bare SHA does not fetch the tag ref, so `tagged_version()` in .ci/pytorch/binary_populate_env.sh (`git describe --tags --exact`) fails and the build falls through to the `<version>.dev<DATE>` default. binary_populate_env.sh then takes its `.*dev.*` branch and appends `+git<triton-shorthash>` to the triton requirement that gets baked into the wheel metadata. Fix: add `fetch-tags: true` to the inline build job checkout in the linux binary build template so the tag pointing at the checked-out commit is fetched, letting `git describe --tags --exact` succeed and the build resolve to the release version. The ROCm/XPU/s390x builds use the reusable _binary-build-linux.yml whose checkout has no pinned ref (it uses the default tag ref for a tag push) and are unaffected. This is the same class of bug fixed for the older checkout-pytorch path in pytorch#180508; this covers the inlined unified build job introduced since. Test Plan: Regenerated the workflows from the template and confirmed only the two linux manywheel workflows change, each gaining `fetch-tags: true` on its four unified build jobs: ``` python3 .github/scripts/generate_ci_workflows.py git diff --stat .github/workflows grep -c 'fetch-tags: true' \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml ``` Validated the regenerated YAML parses and passes actionlint (the only reported shellcheck warnings are pre-existing on the base files): ``` python3 -c "import yaml; yaml.safe_load(open(f)) for f in (...)" lintrunner --take ACTIONLINT \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml ``` This PR was authored with the assistance of an AI coding assistant.
…ag (pytorch#187055) Full git fetch on tag pushes so release manywheel builds detect the tag (pytorch#187042) The unified inline manywheel build jobs (cpu/cpu-aarch64/cuda/cuda-aarch64) check out a pinned commit SHA with a shallow `fetch-depth: 2`. On a release tag push (e.g. v2.13.0-rc1) the tag ref is never written into that shallow checkout, so `tagged_version()` in .ci/pytorch/binary_populate_env.sh (`git describe --tags --exact`) fails and the build falls through to the `<version>.dev<DATE>` default, producing wheels versioned `2.13.0.devYYYYMMDD` and pinning triton to a `+git<hash>` version that does not exist on the test index. Root cause: pytorch#187001 added `fetch-tags: true` but kept `fetch-depth: 2`. `fetch-tags` is ineffective when `ref` is a bare SHA and the fetch is shallow; git does not write the tag ref locally. The proven path (the checkout-pytorch action) works because it uses `fetch-depth: 0`, where a full fetch brings all tags regardless. Fix: in the linux binary build template, set `fetch-depth` to 0 on tag pushes (`github.ref_type == 'tag'`) and keep 2 for nightlies so they stay fast. With a full fetch the tag is present, `git describe --tags --exact` succeeds, and the build resolves to the release version with a clean triton requirement. Needs to be cherry-picked onto release/2.13 for the RC builds to pick it up. Test Plan: PR: pytorch#187047 Regenerated the workflows from the template and confirmed only the two linux manywheel workflows change, each unified build job getting the conditional fetch-depth and fetch-tags: ``` python3 .github/scripts/generate_ci_workflows.py git diff --stat .github/workflows grep -c "github.ref_type == 'tag' && 0 || 2" \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml ``` Validated the regenerated YAML parses and passes actionlint (the only reported shellcheck warnings are pre-existing on the base files): ``` python3 -c "import yaml, sys; [yaml.safe_load(open(f)) for f in sys.argv[1:]]" \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml lintrunner --take ACTIONLINT \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml ``` This PR was authored with the assistance of an AI coding assistant. Pull Request resolved: pytorch#187042 Approved by: https://github.com/NicolasHug, https://github.com/albanD, https://github.com/jeanschmidt (cherry picked from commit 80bd062) Co-authored-by: atalman <atalman@meta.com>
…etect the tag (pytorch#187058) Followup - Full git fetch on tag pushes so release manywheel builds detect the tag (pytorch#187047) ## Problem The unified inline manywheel build jobs (`cpu` / `cpu-aarch64` / `cuda` / `cuda-aarch64`) check out a pinned commit SHA with a shallow `fetch-depth: 2`. On a release tag push (e.g. `v2.13.0-rc1`) the tag ref is never written into that shallow checkout, so `tagged_version()` in `.ci/pytorch/binary_populate_env.sh` (`git describe --tags --exact`) fails. The build then falls through to the `<version>.dev<DATE>` default, producing wheels versioned `2.13.0.devYYYYMMDD` and pinning triton to a `+git<hash>` version that does not exist on the test index. ## Root cause pytorch#187001 added `fetch-tags: true` but kept `fetch-depth: 2`. `fetch-tags` is ineffective when `ref` is a bare SHA and the fetch is shallow -- git does not write the tag ref locally. The proven path (the `checkout-pytorch` action) works because it uses `fetch-depth: 0`, where a full fetch brings all tags regardless. ## Fix Do a full fetch (`fetch-depth: 0`) on tag pushes and keep the shallow `fetch-depth: 2` for nightlies: ```yaml fetch-depth: ${{ github.ref_type != 'tag' && 2 || 0 }} fetch-tags: ${{ github.ref_type == 'tag' }} ``` ### Why the expression is written this way GitHub Actions has no ternary, so the `A && B || C` idiom is used -- but it only works when `B` is **truthy**. The naive `github.ref_type == 'tag' && 0 || 2` is broken because `0` is falsy: on a real tag push it evaluates `true && 0` -> `0`, then `0 || 2` -> `2`, so it **always** yields `2` and the full-fetch branch is unreachable. An earlier revision shipped this broken form, which is why tag builds kept getting `fetch-depth: 2`. Negating the condition keeps the truthy literal (`2`) as the `&&` operand so the value is never swallowed: | trigger | expression | result | |---|---|---| | tag | `false && 2` -> false -> `false \|\| 0` | **0** (full fetch) | | branch / nightly | `true && 2` -> 2 | **2** (shallow) | With a full fetch the tag is present, `git describe --tags --exact` succeeds, and the build resolves to the release version with a clean triton requirement. > Needs to be cherry-picked onto `release/2.13` (superseding the broken expression) for the RC builds to pick it up. ## Test Plan Regenerated the workflows from the template and confirmed only the two linux manywheel workflows change, each of the 4 unified build jobs getting the conditional `fetch-depth` and `fetch-tags`: ``` python3 .github/scripts/generate_ci_workflows.py git diff --stat .github/workflows grep -c "github.ref_type != 'tag' && 2 || 0" \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml ``` Validated the regenerated YAML parses and passes actionlint (the only reported shellcheck warnings are pre-existing `SC1090` on `source "${BINARY_ENV_FILE}"`): ``` python3 -c "import yaml,sys;[yaml.safe_load(open(f)) for f in sys.argv[1:]]" \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml lintrunner --take ACTIONLINT \ .github/workflows/generated-linux-binary-manywheel-nightly.yml \ .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml ``` This PR was authored with the assistance of an AI coding assistant. Pull Request resolved: pytorch#187047 Approved by: https://github.com/jeanschmidt, https://github.com/albanD (cherry picked from commit 1e97809) Co-authored-by: atalman <atalman@meta.com>
… tag (pytorch#187172) Full git fetch on tag pushes so macOS release wheel builds detect the tag (pytorch#187139) ## Problem Same issue as pytorch#187058, but for the **macOS arm64 binary wheel** build. The `wheel-build` job checks out a pinned commit SHA with the default shallow `actions/checkout` (depth 1, no tags). On a release tag push (e.g. `v2.13.0-rc1`) the tag ref is never written into that checkout, so `tagged_version()` in `.ci/pytorch/binary_populate_env.sh` (`git describe --tags --exact`, invoked per-iteration by `.ci/wheel/build_all_macos_wheels.sh`) fails. The build then falls through to the `<version>.dev<DATE>` default, producing wheels versioned `2.13.0.devYYYYMMDD` instead of the release version. ## Fix Do a full fetch (`fetch-depth: 0`) on tag pushes and keep the shallow checkout for nightlies: ```yaml fetch-depth: ${{ github.ref_type != 'tag' && 1 || 0 }} fetch-tags: ${{ github.ref_type == 'tag' }} ``` The condition is negated (matching pytorch#187058) because GitHub Actions has no ternary and the `A && B || C` idiom only works when `B` is truthy; the naive `github.ref_type == 'tag' && 0 || 2` always yields the shallow value because `0` is falsy. Only the `wheel-build` checkout is changed; the libtorch `*-extract` job uses a sparse checkout of the already-built wheel artifact and does not need the tag. > Needs to be cherry-picked onto `release/2.13` for the RC builds to pick it up. ## Test Plan Regenerated the workflows from the template and confirmed only the macOS wheel workflow changed, with the conditional fetch on the build job's checkout: ``` python3 .github/scripts/generate_ci_workflows.py git diff --stat .github/workflows ``` Validated the regenerated YAML parses and passes actionlint (the only reported shellcheck warning, `SC2086` on `uv python install $DESIRED_PYTHONS`, is pre-existing and unrelated): ``` python3 -c "import yaml; yaml.safe_load(open('.github/workflows/generated-macos-arm64-binary-wheel-nightly.yml'))" lintrunner --take ACTIONLINT .github/workflows/generated-macos-arm64-binary-wheel-nightly.yml ``` *This PR was authored with the assistance of an AI coding agent.* Pull Request resolved: pytorch#187139 Approved by: https://github.com/malfet (cherry picked from commit fb17c34) Co-authored-by: atalman <atalman@meta.com>
pytorch#185760)" (pytorch#187347) Revert "Skip cudagraphs for kernel-free inductor graphs (pytorch#185760)" This reverts commit b1e3edd. Reverted pytorch#185760 on behalf of https://github.com/wdvr due to reverting due to failing signal on D108074417 - discussed with author ([comment](pytorch#185760 (comment))) (cherry picked from commit 82a2694) Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
…gorithms (pytorch#186653)" (pytorch#187349) Revert "Make topk deterministic under deterministic algorithms (pytorch#186653)" This reverts commit 1100802. Reverted pytorch#186653 on behalf of https://github.com/wdvr due to failing some internal ads tests - see D108014200 for signal ([comment](pytorch#186653 (comment))) (cherry picked from commit ca8348e) Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
[BE] Make spmd_type a CI rather than CD dependency (pytorch#187067) Alas, it could not be added to `requirements-ci.txt`, as it depends on torch and we don't want any pre-installed torch wheels inside CI docker builds Pull Request resolved: pytorch#187067 Approved by: https://github.com/pianpwk, https://github.com/atalman, https://github.com/fegin (cherry picked from commit ba71580) Co-authored-by: Nikita Shulga <nshulga@meta.com>
…d symbols in fake traces (pytorch#183837)" (pytorch#187348) Revert "[HOP][Dynamic Shapes] Ignore discarded unbacked symbols in fake traces (pytorch#183837)" This reverts commit 5ccc7ff. Reverted pytorch#183837 on behalf of https://github.com/wdvr due to failing some executorch / other tests on the HOP / control-flow / FlexAttention surface I believe ([comment](pytorch#183837 (comment))) (cherry picked from commit 3d10f15) Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
[CD] Bump 13.0 builds to 13.0.3 (pytorch#179758) CUDA 13.0.3 is released which includes the 1. patch release fix from cublas that fixes the below issue: Quote from [CUBLAS Patch Release notes](https://docs.nvidia.com/cuda/cublas-patch-release-notes/) Fixed an issue in cublasLtMatmul() that could lead to incorrect results when it ran concurrently with another kernel that used Tensor Memory. This issue affected only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with Compute Capability 10.x and 11.x, and had existed since cuBLAS 12.8. [5807900, 5943783] 2. nvjitlink is set to a lower bound, rather than a pin, for better compatibility for RAPIDS wheels nvcc/nvvm/crt compiler stack are pinned to the same CL for better compatibility. It will pull the latest nvjitlink available rather than the one that is pinned to the released CTK. Updating the CTK version to 13.0.3 in place of the drop in fix in pytorch#178333 Component | 13.0.2 | 13.0.3 -- | -- | -- CUDA cuBLAS | 13.1.0.3 | 13.1.1.3 CUDA nvsdm | 580.95.05 | 580.126.20 Pull Request resolved: pytorch#179758 Approved by: https://github.com/malfet, https://github.com/atalman (cherry picked from commit 99fd1c8) Co-authored-by: Ting Lu <tingl@nvidia.com>
…ytorch#187416) Skip TestMultiprocessingDeviceType on XPU due to lack of IPC support (pytorch#187314) # Motivation When pytorch#184737 generalize `test_multiprocess.py`, it helps enable `TestMultiprocessingDeviceType` test on XPU. However, XPU doesn't support this test suite because lack of IPC support currently. So we would like to skip this test case and restore XPU CI to green. Fixes pytorch#186585 Code originates from pytorch#186793 Co-authored-by: @xuhancn Pull Request resolved: pytorch#187314 Approved by: https://github.com/EikanWang, https://github.com/atalman ghstack dependencies: pytorch#181233, pytorch#187137, pytorch#187232, pytorch#187306 (cherry picked from commit 4f26a57) Co-authored-by: Yu, Guangye <guangye.yu@intel.com>
…rch#187417) [xpu][fix] Include kernel_compile_result.h in aoti xpu.h header (pytorch#187137) # Motivation pytorch#184735 introduces AOTI lazy autotuning with the dual-wrapper mechanism, which uncovered some new XPU test failures in CI. # Additonal Context Fix pytorch#187124 Fix pytorch#187123 Fix pytorch#187035 Fix pytorch#187035 Fix pytorch#187036 Fix pytorch#187125 Fix pytorch#187121 This needs to be cherry-picked to the release branch. Pull Request resolved: pytorch#187137 Approved by: https://github.com/EikanWang, https://github.com/atalman ghstack dependencies: pytorch#181233 (cherry picked from commit 6ec77b0) Co-authored-by: Yu, Guangye <guangye.yu@intel.com>
…o driver (pytorch#187422) Handle pyzes import failures gracefully on machines without level zero driver (pytorch#187232) # Motivation Previously, we didn't handle pyzes import failures gracefully on machines without level zero driver. Without this PR: <img width="510" height="436" alt="image" src="https://github.com/user-attachments/assets/5c6db7ff-ee8d-4ebe-9177-db2ff80916c1" /> With this PR: <img width="530" height="179" alt="image" src="https://github.com/user-attachments/assets/9239ac61-3722-42d7-8411-68f61abd2e0a" /> # Additional Context <img width="1004" height="36" alt="image" src="https://github.com/user-attachments/assets/0035d2c6-2e49-4570-97de-3a4b48ae0575" /> This needs to be cherry-picked to the release branch. Pull Request resolved: pytorch#187232 Approved by: https://github.com/EikanWang, https://github.com/atalman ghstack dependencies: pytorch#181233, pytorch#187137 (cherry picked from commit fd79a5f) Co-authored-by: Yu, Guangye <guangye.yu@intel.com>
Revert XPU device-wide synchronization (pytorch#187306) # Motivation This PR reverts pytorch#182630 to fix pytorch#187277 Root cause is `ext_oneapi_wait_and_throw` will introduce SYCL Graph to sync an invalid queue. # Additional Context This needs to be cherry-picked to the release branch. Why isn't it captured on CI? This issue is only found on BMG (Xe2), and the current CI is on Data Center GPU (Xe). Pull Request resolved: pytorch#187306 Approved by: https://github.com/EikanWang, https://github.com/atalman ghstack dependencies: pytorch#181233, pytorch#187137, pytorch#187232 (cherry picked from commit 421abac) Co-authored-by: Yu, Guangye <guangye.yu@intel.com>
…with attn_bias (pytorch#187684) [SDPA] Fix 32-bit offset overflow in mem-efficient attention forward with attn_bias (pytorch#187225) ## Summary `F.scaled_dot_product_attention` with the **EFFICIENT_ATTENTION** backend and an explicit `attn_mask` **silently returns garbage** for every query row past `2^32 / seq_len` once `num_queries * num_keys > 2^32` (i.e. S > 65536 for square self-attention). No error, no warning, no NaN — the corrupted values look plausible, which makes this extremely hard to attribute downstream (we found it as silently diverging RL training metrics). ## Root cause In `kernel_forward.h`, `query_start` is `uint32_t`: https://github.com/pytorch/pytorch/blob/545c55c1724bafed6772a0b85fecfd09f9aa0dbb/aten/src/ATen/native/transformers/cuda/mem_eff_attention/kernel_forward.h#L635 and the per-row offset into the bias tensor is computed as `query_start * p.bias_strideM` where `bias_strideM` is `int32_t`, so the multiply happens in 32-bit unsigned arithmetic and wraps modulo 2^32: https://github.com/pytorch/pytorch/blob/545c55c1724bafed6772a0b85fecfd09f9aa0dbb/aten/src/ATen/native/transformers/cuda/mem_eff_attention/kernel_forward.h#L811 Past `row = 2^32 / seq_len` the kernel reads the mask from wrapped addresses (typically an earlier row of the same tensor), producing wrong attention for those rows. Corruption onset at exactly `2^32 / S` (not `2^31 / S`) matches the unsigned arithmetic. The dropout RNG `skipahead` offset in the same kernel has the same 32-bit multiply (`(query_start + thread_i) * p.num_keys_absolute`). The backward kernel computes the equivalent offset with an `int64_t` `query_start` (`kernel_backward.h`), so with dropout + large seq_len the forward and backward dropout masks diverge past the same boundary, corrupting gradients. The batch/head-level strides (`bias_strideB`/`bias_strideH`) were already widened to `int64_t` previously; the per-row offset was missed. Note these kernels were vendored from xformers, which has since removed its in-tree CUTLASS FMHA sources, so the fix needs to land here. ## Fix Compute both offsets in 64-bit, matching `kernel_backward.h`. Two one-line casts; in-bounds behavior is bit-identical. ## Reproduction (any sm80+ GPU, bf16) ```python import torch import torch.nn.functional as F from torch.nn.attention import SDPBackend, sdpa_kernel S, HD = 73728, 64 # S^2 = 1.27 * 2^32 torch.manual_seed(0) q = torch.randn(1, 1, S, HD, device="cuda", dtype=torch.bfloat16) * 0.1 k = torch.randn(1, 1, S, HD, device="cuda", dtype=torch.bfloat16) * 0.1 v = torch.randn(1, 1, S, HD, device="cuda", dtype=torch.bfloat16) idx = torch.arange(S, device="cuda") causal = (idx[:, None] >= idx[None, :])[None, None] ref = F.scaled_dot_product_attention(q, k, v, is_causal=True) # ground truth with sdpa_kernel(SDPBackend.EFFICIENT_ATTENTION): out = F.scaled_dot_product_attention(q, k, v, attn_mask=causal) B = 2**32 // S # predicted corruption onset = row 58254 print((out - ref)[0, 0, :B].abs().max()) # ~6e-5 (bf16 noise) print((out - ref)[0, 0, B:].abs().max()) # ~0.39 (garbage) ``` Observed on `2.14.0.dev20260612+cu130` (H200, sm90) — corruption onset is exactly `2^32 // S` for every S tested (73728→58254, 81920→52428, 98304→43690, 90000→47721). The MATH backend with the same mask tensor is correct everywhere, confirming the mask contents are fine and the bug is in the kernel's indexing. `is_causal=True` (no mask tensor) is unaffected. ## Test plan - Added `test_mem_eff_attention_large_seq_len_attn_mask_index_overflow` (gated by `@largeTensorTest("12GB", "cuda")`): explicit additive causal mask vs `is_causal=True` on the same backend at seq_len=70000. Fails before this fix (max abs diff 0.54 past row 61356, exact 0 before it), passes after. - Verified the kernel change directly on H200: compiled `kernel_forward.h` (pristine vs patched, same CUTLASS pin) into a standalone extension driving `AttentionKernel<bf16, Sm80, true, 64, 64, 64>` with identical S=73728 inputs. Pristine reproduces the corruption bit-for-bit with the stock binary (max diff 0.39 past row 58254); patched returns bf16-noise-level diff (6.1e-5) everywhere, and is bit-identical to pristine for all in-bounds rows. AI assistance (Claude) was used to locate the overflow and prepare this PR; the change and verification were reviewed end-to-end by the submitter. Pull Request resolved: pytorch#187225 Approved by: https://github.com/Skylion007 (cherry picked from commit 77abe15) Co-authored-by: Bin Lei <leib2765@gmail.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…ytorch#187666) Revert "Preserve FX graph cache guard provenance (pytorch#184193)" (pytorch#187581) This reverts commit 060670a. Pull Request resolved: pytorch#187581 Approved by: https://github.com/zou3519
…pytorch#187973) Fix Windows libtorch x86_64 and arm64 packages overwriting each other (pytorch#187837) Fixes pytorch#187812 ## Problem The Windows libtorch CD packages for x86_64 and arm64 are generated with the **same filename** and uploaded to the **same S3 location**, so the arm64 build clobbers the x86_64 one. Users downloading the Windows x64 libtorch release actually receive Aarch64 binaries (`c10.dll: PE32+ executable (DLL) (GUI) Aarch64`). ## Root cause `compute_zip_prefix()` in `.ci/libtorch/extract_libtorch_from_wheel.py` returned `libtorch-win-shared-with-deps` for every Windows build regardless of architecture (the docstring even named `libtorch-win-arm64-shared-with-deps`, but the code never produced it). Both x86_64 and arm64 are CPU builds, so the upload subfolder (`DESIRED_CUDA=cpu`) and the zip filename were identical, and the S3 object key is the zip filename. The arch-blind naming was introduced in pytorch#174753 (which added the wheel-extraction approach); the collision became live once the windows-arm64 wheel workflow started running the same extraction in pytorch#181586. ## Fix Thread architecture through the extraction so arm64 produces `libtorch-win-arm64-shared-with-deps-{version}.zip` while x86_64 keeps its existing `libtorch-win-shared-with-deps-{version}.zip` name (preserving the public download URL). `generate_libtorch_extraction_configs` also gives the arm64 job a distinct `build_name` so the GitHub artifact and job names no longer alias. ## Test Plan Verified prefix logic: ``` x86_64 -> libtorch-win-shared-with-deps arm64 -> libtorch-win-arm64-shared-with-deps ``` Regenerated the binary build workflows (`python .github/scripts/generate_ci_workflows.py`); the windows-arm64 wheel workflow now passes `--arch arm64` and uploads a distinct `libtorch-arm64-cpu-shared-with-deps-release` artifact, while the x86_64 workflow is unchanged apart from an explicit `--arch x86_64`. This PR was authored with the assistance of an AI coding assistant (Claude Code). Pull Request resolved: pytorch#187837 Approved by: https://github.com/Skylion007, https://github.com/huydhn, https://github.com/malfet (cherry picked from commit 599fb5c) Co-authored-by: atalman <atalman@meta.com>
…ch#187983) Fix bmm outer product Triton launch on non-current CUDA device (pytorch#187733) The Python native bmm outer-product override launches a Triton kernel for CUDA inputs shaped like `(B, M, 1) x (B, 1, N)`. Unlike the generated C++ ATen CUDA path, this Python dispatch path does not get an automatic CUDA device guard before launching the kernel. In vLLM multimodal tests, the model hooks can run the relevant rotary embedding buffers and inputs on `cuda:1` while the process current device is still `cuda:0`; Triton then launches on the wrong device/stream and rejects the `cuda:1` pointers with `ValueError: Pointer argument cannot be accessed from Triton (cpu tensor?)`. Guard the Triton call when the input device is not already current, while leaving the common already-current path direct to avoid unnecessary Python context manager overhead. Also require both bmm inputs to be on the same CUDA device before selecting the override, so mismatched CUDA inputs fall back to native bmm and report the normal same-device error instead of reaching the Triton kernel. I considered guarding unconditionally, but a small same-device benchmark showed that the Python context manager adds measurable overhead for small outer-product bmm calls. The conditional guard keeps the fix scoped to the non-current-device case that needs it. Fixes pytorch#187729 Generated by my agent Benchmark Results: - Command: `CUDA_VISIBLE_DEVICES=0 python - <<'PY' ...` benchmarking 7 trials of 1000 `torch.bmm` calls on `(32, 128, 1) x (32, 1, 512)` CUDA tensors with current device matching the inputs. - Before: median 24.210 us/call; raw `[24.34, 24.365, 24.195, 24.206, 24.21, 24.272, 24.142]`. - After: median 24.946 us/call; raw `[25.045, 24.909, 25.008, 25.042, 24.906, 24.946, 24.926]`. Test Plan: - `CUDA_VISIBLE_DEVICES=0,1 python - <<'PY' ...` minimal repro: before the fix failed with the Triton pointer `ValueError`; after the fix produced a `cuda:1` result and preserved current device `0`. - `CUDA_VISIBLE_DEVICES=0,1 python -m pytest test/test_bmm_outer_product.py -q -rs` passed: `11 passed, 10 subtests passed`. - `lintrunner torch/_native/ops/bmm_outer_product/triton_impl.py test/test_bmm_outer_product.py` passed. - `git diff --cached --check` passed. - `lintrunner -a` was run and failed on unrelated pre-existing clang-tidy findings in untouched C++ files. Pull Request resolved: pytorch#187733 Approved by: https://github.com/slayton58 (cherry picked from commit 5a5e50f) Co-authored-by: Jason Ansel <jansel@meta.com>
[xpu] Fix LSTM oneDNN Integration (pytorch#187334) 1. onednn LSTM on XPU does not support bf16 on bias and cell state 2. add deterministic attribute Pull Request resolved: pytorch#187334 Approved by: https://github.com/EikanWang, https://github.com/atalman (cherry picked from commit 691e8b6) Co-authored-by: Zhang, Jianyi <jianyi.zhang@intel.com>
* Revive CUDA 12.9 nightly binary builds (pytorch#186015) Restores the Linux-only (x86_64 + aarch64) CUDA 12.9 build matrix and `cuda12.9` docker images that were removed in pytorch#179155 and pytorch#179158. Generated workflows regenerated via `.github/regenerate.sh`. Windows continues to skip 12.9; `install_129` and the `build_cuda.sh` 12.9 arch list were already present on this branch. Fixes pytorch#185269. Authored with Claude. * Remove unused noqa: B950 directive in generate_binary_build_matrix.py --------- Co-authored-by: Nikita Shulga <2453524+malfet@users.noreply.github.com>
* [MPS] Refactor inlined Metal reduction logic into shared header (pytorch#187541) Pulls the open-coded identity / replace / combine / simd_reduce / threadgroup_reduce helpers used by the Metal min/max value-reduction kernels (and by the inductor MPS codegen's max/min and argmin/argmax paths) into shared `c10::metal::MaxOp<T>` / `MinOp<T>` structs in `c10/metal/reduction_utils.h`. The follow-up change migrating argmin/argmax will reuse those primitives as well. Fix latent bug in MPSInducotr when min/max would not identify -INFINITY, as identity used to be `-FLT_MIN` in inductor inlined reduction code, that allows one to remove xfails from bunch of dynamic shapes reduction tests Authored with Claude. Pull Request resolved: pytorch#187541 Approved by: https://github.com/Skylion007 * [MPS] Migrate argmin/argmax from MPSGraph to Metal kernels (pytorch#187304) Replaces the MPSGraph-based argmax/argmin implementation with three native Metal kernels mirroring the value_reduction layout introduced for amax/amin/all/any in pytorch#180752: a generic NormParams kernel for arbitrary reduction patterns, an inner-dim kernel (one simdgroup per row, simd_arg_reduce collapse) for contiguous reductions over the last dim, and an outer-dim kernel (TG_X x TG_Y tile, shared-mem pair tree reduce) for contiguous reductions over dim 0. The host-side dispatcher in ReduceOps.mm picks inner/outer when the input and output are contiguous and exactly one of those dims is reduced, and falls back to the generic kernel otherwise; full reduction (dim=None) is handled via a contiguous flatten so the returned linear index matches the standard as-if-contiguous convention. The MPSGraph path used to cast bfloat16 to float32 before the reduction, so bf16 is where the migration wins the most (e.g. 4096x4096 dim=0 argmax goes 1518us -> 294us = 5.17x on M4 Max); fp16/fp32 dim=0 hit 2.4-3.5x and dim=1 hit 1.1-1.3x. Correctness: the natural simd_argmax(val, idx) helper ties on lowest LANE, but when a single lane scans multiple positions its stored idx is not necessarily the lowest one carrying the winning value, so the kernels do a two-step simd_max + simd_min-on-eff_idx pattern instead (NaN lanes count as winners so first-NaN-in-source-order wins). The shared-memory tree reduction in the outer kernel uses the same strict-better-or-equal-with-lower-idx predicate. The NaN-propagating "should-replace" predicates argmax_replace / argmin_replace are factored into c10/metal/reduction_utils.h so they sit next to simd_argmax/simd_argmin, which already encode the same NaN-as-winner logic via simd_ballot. TODO: argmin/argmax and amin/amax now share the (general, inner, outer) kernel shape, the per-thread-scan + simd-collapse skeleton, and the address-math helpers. Unifying them under a single Reducer concept (per-thread accumulator + finalize hook, with NCHAINS as a property of the reducer) would let the three kernel grids be written once and instantiated for both ValueReducer<Op, Load, NCHAINS=8> and ArgReducer<Op, NCHAINS=1>. Worth doing in a follow-up. Benchmark (median of 500 iters on M4 Max, agent_space/bench_argminmax_pr.py): | op | shape | dtype | mode | baseline (us) | new (us) | speedup | |--------|-----------|----------|------|--------------:|---------:|--------:| | argmax | 4096x4096 | bfloat16 | dim0 | 1518 | 294 | 5.17x | | argmax | 4096x4096 | bfloat16 | dim1 | 991 | 315 | 3.15x | | argmax | 4096x4096 | float16 | dim0 | 1024 | 290 | 3.53x | | argmin | 4096x4096 | float16 | dim0 | 1014 | 326 | 3.11x | | argmax | 4096x4096 | float32 | dim0 | 1009 | 424 | 2.38x | | argmin | 4096x4096 | int64 | dim0 | 1075 | 740 | 1.45x | | argmax | 1024x1024 | bfloat16 | dim0 | 192 | 100 | 1.92x | | argmax | 1024x1024 | float32 | dim0 | 194 | 120 | 1.61x | | argmax | 4096x4096 | float32 | full | 5044 | 4926 | 1.02x | Authored with Claude. Fixes pytorch#130295 Pull Request resolved: pytorch#187304 Approved by: https://github.com/kurtamohler, https://github.com/seemethere, https://github.com/izaitsevfb ghstack dependencies: pytorch#187541
Remove setuptools upper bound (pytorch#187262) Upper bound was originally introduced by pytorch#174631, but `pkg_resources` is not used by `torch.utils.cpp_extensions` Test Plan: Install latest setuptools and all `cpp_extensions` test still work Fixes pytorch#187188 Pull Request resolved: pytorch#187262 Approved by: https://github.com/malfet, https://github.com/atalman (cherry picked from commit c776884) Co-authored-by: Yufeng He <40085740+he-yufeng@users.noreply.github.com> Co-authored-by: Nikita Shulga <2453524+malfet@users.noreply.github.com>
…ch#188192) Fix build_with_debinfo.py broken by CONFIGURE_DEPENDS globbing (pytorch#186780) ## Summary `tools/build_with_debinfo.py` recompiles named source files with `-g` so you can get debug info for a targeted file without a full (slow, ABI-incompatible) debug build. It is referenced in `CONTRIBUTING.md`, and it broke on current `main`: running it just prints `Nothing to do` and rebuilds nothing. ## Root cause The tool derived its compile/link commands from `ninja -j1 -v -n torch_python` (a dry run) and rewrote `-O2`/`-O3` to `-g`. The build now uses `file(GLOB ... CONFIGURE_DEPENDS)`, which wires a glob-verification step into `build.ninja`'s own regeneration (`CMakeFiles/VerifyGlobs.cmake` + `regenerate-during-build`). In dry-run (`-n`) mode ninja cannot run that check or reload the regenerated graph, so `ninja -n <target>` only ever reports the regeneration step and never the real compile/link commands. The tool parsed an empty plan and no-op'd. ## Fix Source the commands from places the glob-check does not gate: - the per-file compile command from `build/compile_commands.json` (rewrite `-O2`/`-O3` -> `-g`); - the `libtorch_python` link command from `ninja -t commands torch_python` (a graph walk, not a dry run; strip ninja's `: && ... && :` wrapper). Then symlink `build/lib` into `torch/lib` as before. The tool now rebuilds exactly the files you name plus the relink, rather than touching sources and relying on ninja's staleness detection. The old `>100 stale items` guard was a dry-run artifact and is removed; the tool now requires `compile_commands.json` and an already-built tree. ## Tests - **Unit tests** (`tools/test/test_build_with_debinfo.py`) for the pure plan-derivation helpers: the `-O` -> `-g` rewrite, `compile_commands.json` indexing, the command/arguments entry forms, and the `ninja -t commands` link extraction. - **Build-tree smoke check**: a new `--dry-run` mode plus a step in `.ci/pytorch/build.sh` that, after the Linux build, asserts the emitted plan contains a `-g` compile and the `libtorch_python` relink. This exercises the real `compile_commands.json` and `ninja -t commands` and fails loudly if a future build-system change breaks those data sources -- the class of bug here, which unit tests alone would not catch. `--dry-run` only reads the tree, so it does not dirty the checkout. ## Test plan Validated on a real editable CUDA build of `viable/strict`: ``` python tools/build_with_debinfo.py --verbose torch/csrc/Module.cpp ``` recompiles `Module.cpp.o` with `.debug_info`, relinks `libtorch_python.so` (which then carries `.debug_info`), symlinks it into `torch/lib`, and `import torch` still works. ``` python tools/build_with_debinfo.py --dry-run torch/csrc/Module.cpp ``` exits 0, prints the plan, and leaves the git tree clean; the build.sh `-g` and `libtorch_python` plan assertions both pass against that output. ``` python -m pytest tools/test/test_build_with_debinfo.py ``` -> 9 passed. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Pull Request resolved: pytorch#186780 Approved by: https://github.com/benjaminglass1, https://github.com/atalman (cherry picked from commit 8f9a6c8) Co-authored-by: Klaus Zimmermann <klaus.zimmermann@quansight.com> Co-authored-by: Benjamin Glass <3344239+benjaminglass1@users.noreply.github.com>
…dd fast path (pytorch#184824)" (pytorch#188196) Revert "[Reland] Port D104346887/PR 182675 for index_add fast path (pytorch#184824) (pytorch#184824)" This reverts commit 9e051b3. Reverted pytorch#184824 on behalf of https://github.com/meta-codesync due to Diff reverted internally ([comment](pytorch#184824 (comment))) (cherry picked from commit eaa0ca8) Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
…cross graph breaks (pytorch#184487)" (pytorch#188197) Revert "dynamo: round-trip torch.cuda.stream ctx mgr across graph breaks (pytorch#184487)" This reverts commit b0a67c7. Reverted pytorch#184487 on behalf of https://github.com/williamwen42 due to Breaks internal builds ([comment](pytorch#184487 (comment))) (cherry picked from commit 697514a) Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
…rch#184075)" (pytorch#188277) This reverts commit b35bc89 on the release/2.13 branch. It is the release-branch counterpart of the main-branch revert (pytorch#188219). PR pytorch#184075 reworked Inductor's temporary generalized scatter form to use private OpOverload targets (a torch.library FRAGMENT registering generalized_scatter / inplace_generalized_scatter_ with flattened, encoded view-op arguments) and added matching Piecewise support in the z3 validator. This revert restores the prior representation that passes view ops as a list of ViewOp dataclasses directly as call_function args. The reinplace.py hunks conflicted against the release branch; all were resolved to the pre-PR form with bare asserts (torch/_inductor/** is still S101-exempt on this branch, so no noqa is added), fully removing the encode/decode machinery, the torch.library registration, and the z3 validator additions. Test Plan: ``` lintrunner torch/_inductor/fx_passes/reinplace.py ``` This PR was authored with the assistance of an AI agent.
…188409) Bump DOCKER_IMAGE_PIN (and the generated linux manywheel binary build workflows) from f38ba0b to 78e737a, the latest manylinux2_28 / manylinuxaarch64 builder images built for v2.13.0-rc10. Verified the new images are published on Docker Hub (e.g. pytorch/manylinux2_28-builder:cuda12.6-78e737ad..., cpu-78e737ad...). Equivalent to regenerating the workflows after the pin bump (only the docker hash changes).
…188443) Add CUDA 12.9 to TORCH_CUDA_ARCH_LIST tables (fix cu129 build) CUDA 12.9 is in the release/2.13 build matrix (CUDA_ARCHES), but the arch-list tables introduced in the 2.13 refactor omit it, so manywheel-cuda-cu129-build fails at .ci/manywheel/build_env_setup.py with 'unknown cuda version 12.9'. In 2.12 the arch list was a bash case in build_cuda.sh that handled 12.9 as base + sm_12.0 (Blackwell), i.e. the 13.x set, not the legacy 12.6 set. Restore that by adding a 12.9 entry to both TORCH_CUDA_ARCH_LIST_TABLE (build) and PYTORCH_RELEASES_CODE_CC (runtime), kept in sync per validate_runtime_release_table_consistency(): x86_64: {75, 80, 86, 90, 100, 120} (== 13.0/13.2) aarch64: {80, 90, 100, 120} (13.x aarch64 minus sm_110, which 12.9 never added) validate_runtime_release_table_consistency() passes.
#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 <prachi.gupta@amd.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: Jeff Daily <jeff.daily@amd.com> (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 <noreply@anthropic.com> (cherry picked from commit 10d8e1d)
…h#183014) Upstream PR pytorch#181112 (`[inductor] Unify threads-per-wave (warp_size) extraction`) has now landed and did much of the original broad warp-size plumbing for this PR: it centralized `DeviceProperties.warp_size`, threaded `warp_size` through Inductor's Triton config helpers, and made HIP wave32 devices follow the non-wave64 path. After rebasing on top of that work, this PR is now focused on the remaining stricter ROCm-safe behavior: - Preserve missing `DeviceProperties.warp_size` as `None` instead of synthesizing `32` for CUDA/HIP devices. - Make `DeviceProperties.warp_size_or_default` raise for missing CUDA/HIP warp-size metadata, while preserving the explicit non-GPU fallback. - Skip optional autotune/rblock/fusion/split-reduction heuristics when required warp-size metadata is unavailable instead of guessing. - Keep CUDA launcher generation on the fast `32` path, but use the device warp size for HIP launcher block dimensions. - Use queried device warp size in template register-spill pruning, with graceful fallback on non-CUDA/XPU builds. - Update the relevant Inductor tests for the stricter missing CUDA/HIP metadata behavior. - Rebased cleaned stack onto `upstream/main` at `19791183fec`. - Rebuilt PyTorch from scratch with `/home/niromero/docker_workspace/framework_scripts/pytorch/build.sh`. - Ran impacted UTs: - `python test/inductor/test_triton_heuristics.py TestWarpSizeUnification TestTritonHeuristics.test_autotune_hints_to_configs TestFastLauncherDeviceSupport` - `python test/inductor/test_max_autotune.py TestMaxAutotune.test_max_autotune_includes_max_autotune_pointwise_configs` - `python test/inductor/test_torchinductor_dynamic_shapes.py -k test_embedding_backward_dynamic_shapes_large_grid` - `python test/inductor/test_torchinductor.py -k test_reduction_config_limit` Made with [Cursor](https://cursor.com) Pull Request resolved: pytorch#183014 Approved by: https://github.com/jansel, https://github.com/glen-amd, https://github.com/jeffdaily, https://github.com/eellison (cherry picked from commit 9183299) Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com>
## Summary Align release/2.13 with the ROCm 7.14 compatibility handling required by TheRock CI. This change adds reusable ROCm version guards for standard and multiprocess tests and handles known ROCm 7.14 runtime and packaging incompatibilities found during full-suite validation. ## Test plan - Distributed suite: https://github.com/ROCm/TheRock/actions/runs/30552305532 - Default + Inductor suites: https://github.com/ROCm/TheRock/actions/runs/30585206557 All validation completed successfully. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
Cherry-pick of #3521 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com>
) (#3539) Cherry-pick of upstream commit 2314113 (pytorch#187052) onto release/2.13. 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>
Porting only missing code (gfx1250) from https://github.com/pytorch/pytorch/pull/190703/changes Signed-off-by: Jagadish Krishnamoorthy <jagadish.krishnamoorthy@amd.com>
…3545) Fixes empty HIP_ARCHITECTURE build failure when gfx1250 is the only provided value to PYTORCH_ROCM_ARCH
…span drives (#3550) cherry pick of pytorch#191743 which solves ROCm/TheRock#6963
…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>
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
…ytorch#188002) (#3551) Cherry-pick the upstream commit to fix MI450 test crash caused by following: test_fake_tensor.py::PropagateRealTensorsFakeTensorTest::test_embedding_meta_indices_propagate_real_tensors Verified test is passed after this fix. Co-authored-by: Aaron Orenstein <aorenste@fb.com>
## Summary First step of the `release/2.14` bringup: drop the pre-release `a0` suffix from `version.txt`. - `version.txt`: `2.14.0a0` -> `2.14.0` This mirrors the `release/2.13` bringup (where `version.txt` was set to `2.13.0`). The branch was cut from upstream `pytorch/pytorch` `release/2.14` HEAD (`8e86e0a2`). Follow-up bringup steps (separate PRs): `related_commits` for dependencies, dependency/requirements pinning, and cherry-pick management. Tracking: AIPYTORCH-1003 Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: pablo-garay <pablo-garay@users.noreply.github.com>
…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>
…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 e61531f)
- [#3525](#3525) (`[release/2.13] Stabilize ROCm 7.14 testing on release/2.13`) added a module-level `sys.exit(0)` in `test/distributed/test_shmem_triton.py` for `ROCM_VERSION >= (7, 14)`. - This PR does two things: 1. **Unskip** that 7.14 `sys.exit(0)`. Keep the existing capability skip when the SHMEM backend is unavailable. 2. **Take the finder hunk** from upstream commit [`4a31e29`](pytorch@4a31e29) (`[ROCm][CI] Fix trunk rocSHMEM and CK failures`, from [pytorch#188429](pytorch#188429)): search `$ROCM_PATH/lib` when `ROCM_PATH` is set, else `/opt/rocm/lib`. - We did **not** take the rest of `4a31e29`. That commit also adds `@skipIfRocmVersionAtLeast([7, 14])` on `test/inductor/test_ck_backend.py`. That CK skip is already on `release/2.13` and is unrelated to this UT. Full `cherry-pick -x` of `4a31e29` would just restage it. We also did not backport the rest of pytorch#188429 (TheRock CI/docker switch). ## Test plan - [x] TheRock `test_pytorch_wheels_full.yml`, `test_configs=distributed`, `tests_to_include=distributed/test_shmem_triton` - gfx94X / `linux-gfx942-1gpu-ccs-csp-ossci-rocm`: https://github.com/ROCm/TheRock/actions/runs/32195967020/job/95904033020 ``` distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_barrier PASSED [13.5964s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_fence PASSED [2.5392s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_get_nbi_False PASSED [1.3613s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_get_nbi_True PASSED [1.4139s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_get_ring PASSED [1.7522s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_put PASSED [1.6263s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_put_signal_add PASSED [1.5214s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_put_signal_set PASSED [1.5903s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_quiet PASSED [1.5085s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_signal_wait_until PASSED [0.0122s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_sync PASSED [1.5976s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_wait_until PASSED [1.2603s] ======================= 12 passed, 25 skipped in 32.21s ======================== ``` - gfx950: https://github.com/ROCm/TheRock/actions/runs/32201312766/job/95915621757 ``` distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_barrier PASSED [8.2912s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_fence PASSED [1.2513s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_get_nbi_False PASSED [1.0780s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_get_nbi_True PASSED [1.0367s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_get_ring PASSED [1.1043s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_put PASSED [0.9847s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_put_signal_add PASSED [1.1261s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_put_signal_set PASSED [1.1993s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_quiet PASSED [1.1537s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_signal_wait_until PASSED [0.0080s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_sync PASSED [1.2694s] distributed/test_shmem_triton.py::SHMEMTritonTest::test_triton_wait_until PASSED [1.0119s] ======================= 12 passed, 25 skipped in 22.50s ======================== ``` - [x] Confirm the 12 Triton cases actually run (not a finder crash, not a module-level skip). Multiprocess alltoall/broadcast/reduces stay `@skip_if_rocm_multiprocess`.
…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.
…a xnack variants (#3583) Drops the gfx90a→xnack± special case; builds `torch_rocshmem` for the requested target like every other arch. rocSHMEM ships xnack-agnostic device bitcode (`librocshmem_device_gfx90a.bc`, `target-cpu=gfx90a`), and the mapping matched per-xnack archives that ROCm/rocm-systems#4822 removed in April. Fixes the split `torch_gfx90a:xnack±.kpack` stubs that break torch on MI250 (ROCm/TheRock#7081, ROCM-29326). Scope: gfx90a only, minimal. The `list(FIND ...)` membership test is untouched; it needs bare-arch matching once ASan enables xnack uniformly (ROCm/TheRock#6624), tracked separately.
Cherry-pick of #3584 Co-authored-by: albmalamd <albert.malewski@amd.com>
## Summary Bringup step 2 for `release/2.14`: add `related_commits` pinning the ROCm dependency builds. Mirrors the `release/2.13` `related_commits` layout, advanced to the 2.14 dependency refs. Depends on / part of the same bringup as #3571 (version.txt). Tracking: AIPYTORCH-1003 --------- Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: Jithun Nair <37884920+jithunnair-amd@users.noreply.github.com>
## Summary Pin Python dependencies for the `release/2.14` bringup so CI and builds are reproducible, and make everything resolve cleanly on Python **3.10 / 3.11 / 3.12 / 3.13 / 3.14**. Mirrors the 2.13 bringup (#3445), plus the extra version-conditional bumps needed for 3.13/3.14. Three commits, one per file: - `requirements.txt` — pin all direct dev/install deps (with a `python_version` split for `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.13/3.14. ## Per-version splits (dev/build) - **numpy** (build): `2.2.6` (<3.11), `2.4.6` (==3.11), `2.5.2` (>=3.12). - **networkx** (dev): `3.4.2` (<3.11), `3.6.1` (>=3.11). Every other dev/build dep resolved identically across 3.10–3.14, so those are flat-pinned. ## 3.13 / 3.14 conditional bumps in requirements-ci.txt (new vs 2.13) - **scikit-image**: `0.22.0` for `< 3.13`, `0.26.0` for `>= 3.13` (0.22.0 has no cp313/cp314 wheel). - **networkx**: `2.8.8` for `< 3.13`, `3.6.1` for `>= 3.13` (scikit-image 0.26 requires networkx>=3.0; keeping 2.8.8 forced pip to backtrack to a wheelless scikit-image). - **pywavelets**: adds `1.9.0` for `>= 3.14` (1.7.0 has no cp314 wheel). - **fbscribelogger**: guarded to `< 3.14` (its `thriftpy2` dep has no cp314 wheel and would source-build in the compiler-less test image) — mirrors 2.13. ## How versions were resolved Per the documented process (AIPYTORCH-748): installed each requirements file in a fresh venv inside an `ubuntu:24.04` container (deadsnakes for the non-native Python versions) and pinned the `pip freeze` output. Verified all files install on **3.10–3.14** with no new compiler-dependent source builds (only pre-existing `coremltools` for <3.12 and pure-Python `pwlf`). ## Test plan TheRock CI `multi_arch_build_portable_linux_pytorch_wheels.yml` (gfx94X-dcgpu, ROCm 7.15.0a20260712, `pytorch_git_ref` = the pinned integration branch). The **Build job passed for every Python version** — torch + torchaudio + torchvision + triton + apex all built against the pinned requirements: - [x] Python 3.10 — build ✅ [run 32165499080](https://github.com/ROCm/TheRock/actions/runs/32165499080) - [x] Python 3.11 — build ✅ [run 32165508608](https://github.com/ROCm/TheRock/actions/runs/32165508608) - [x] Python 3.12 — build ✅ [run 32165519879](https://github.com/ROCm/TheRock/actions/runs/32165519879) - [x] Python 3.13 — build ✅ [run 32165528430](https://github.com/ROCm/TheRock/actions/runs/32165528430) - [x] Python 3.14 — build ✅ [run 32165536568](https://github.com/ROCm/TheRock/actions/runs/32165536568) > Note: the downstream GPU **Test** job hit a transient 30-min git-clone timeout in its "Checkout PyTorch Source Repos" step (infra/network, unrelated to these pins). The build — which is what exercises the requirements files — is green across all five versions. --------- Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: pablo-garay <pablo-garay@users.noreply.github.com>
…4_IFU_20260821 # Conflicts: # .ci/docker/ci_commit_pins/triton.txt # .ci/docker/requirements-ci.txt # .ci/pytorch/build.sh # .ci/pytorch/common_utils.sh # .github/ci_commit_pins/xla.txt # .github/scripts/filter_test_configs.py # .github/scripts/generate_binary_build_matrix.py # .github/templates/common.yml.j2 # .github/templates/linux_binary_build_workflow.yml.j2 # .github/templates/macos_binary_build_workflow.yml.j2 # .github/templates/windows_binary_build_workflow.yml.j2 # .github/workflows/_binary-build-flash-attention-wheel-linux.yml # .github/workflows/_binary-build-flash-attention-wheel-windows.yml # .github/workflows/_binary-upload.yml # .github/workflows/_docs.yml # .github/workflows/_lint.yml # .github/workflows/_linux-build.yml # .github/workflows/_linux-test-stable-fa3.yml # .github/workflows/_linux-test.yml # .github/workflows/_mac-build.yml # .github/workflows/_mac-test.yml # .github/workflows/_rocm-test.yml # .github/workflows/_runner-determinator.yml # .github/workflows/_vllm-benchmark.yml # .github/workflows/_win-build.yml # .github/workflows/_win-test.yml # .github/workflows/_xpu-test.yml # .github/workflows/attention_op_microbenchmark.yml # .github/workflows/b200-distributed.yml # .github/workflows/b200-symm-mem.yml # .github/workflows/build-almalinux-images.yml # .github/workflows/build-manywheel-images-s390x.yml # .github/workflows/build-manywheel-images.yml # .github/workflows/build-triton-wheel.yml # .github/workflows/build-vllm-wheel.yml # .github/workflows/claude-autorevert-advisor.yml # .github/workflows/claude-code.yml # .github/workflows/claude-distributed-triage.yml # .github/workflows/claude-issue-triage-run.yml # .github/workflows/close-nonexistent-disable-issues.yml # .github/workflows/create_release.yml # .github/workflows/docker-builds.yml # .github/workflows/docker-cache-rocm.yml # .github/workflows/docker-release.yml # .github/workflows/docs-build.yml # .github/workflows/dtensor.yml # .github/workflows/dynamo-unittest.yml # .github/workflows/generated-linux-aarch64-binary-manywheel-nightly.yml # .github/workflows/generated-linux-binary-manywheel-nightly.yml # .github/workflows/generated-linux-s390x-binary-manywheel-nightly.yml # .github/workflows/generated-macos-arm64-binary-wheel-nightly.yml # .github/workflows/generated-windows-arm64-binary-libtorch-debug-nightly.yml # .github/workflows/generated-windows-arm64-binary-wheel-nightly.yml # .github/workflows/generated-windows-binary-libtorch-debug-nightly.yml # .github/workflows/generated-windows-binary-wheel-nightly.yml # .github/workflows/h100-cutlass-backend.yml # .github/workflows/h100-distributed.yml # .github/workflows/h100-symm-mem.yml # .github/workflows/inductor-micro-benchmark-x86.yml # .github/workflows/inductor-micro-benchmark.yml # .github/workflows/inductor-nightly.yml # .github/workflows/inductor-pallas.yml # .github/workflows/inductor-perf-compare.yml # .github/workflows/inductor-perf-test-b200.yml # .github/workflows/inductor-perf-test-nightly-aarch64.yml # .github/workflows/inductor-perf-test-nightly-h100.yml # .github/workflows/inductor-perf-test-nightly-rocm-mi300.yml # .github/workflows/inductor-perf-test-nightly-rocm-mi350.yml # .github/workflows/inductor-perf-test-nightly-x86-zen.yml # .github/workflows/inductor-perf-test-nightly-x86.yml # .github/workflows/inductor-perf-test-nightly-xpu.yml # .github/workflows/inductor-perf-test-nightly.yml # .github/workflows/inductor-periodic.yml # .github/workflows/inductor-rocm-mi200.yml # .github/workflows/inductor-rocm-mi300.yml # .github/workflows/inductor-rocm-mi350.yml # .github/workflows/inductor-unittest.yml # .github/workflows/inductor.yml # .github/workflows/lint-bc.yml # .github/workflows/lint.yml # .github/workflows/llm_td_retrieval.yml # .github/workflows/nightly-s3-uploads.yml # .github/workflows/nightly.yml # .github/workflows/operator_benchmark.yml # .github/workflows/operator_microbenchmark.yml # .github/workflows/operator_microbenchmark_compare.yml # .github/workflows/periodic-rocm-mi200.yml # .github/workflows/periodic-rocm-mi300.yml # .github/workflows/periodic-rocm-mi350.yml # .github/workflows/periodic.yml # .github/workflows/pull.yml # .github/workflows/riscv64.yml # .github/workflows/rocm-mi200.yml # .github/workflows/rocm-mi300.yml # .github/workflows/rocm-mi350.yml # .github/workflows/rocm-navi31.yml # .github/workflows/rocm-preview.yml # .github/workflows/slow-rocm-mi200.yml # .github/workflows/slow.yml # .github/workflows/target-determination-indexer.yml # .github/workflows/target_determination.yml # .github/workflows/test-b200.yml # .github/workflows/test-check-binary.yml # .github/workflows/test-h100.yml # .github/workflows/tools-unit-tests.yml # .github/workflows/torchtitan.yml # .github/workflows/trunk-rocm-sandbox.yml # .github/workflows/trunk.yml # .github/workflows/tsan.yml # .github/workflows/unstable.yml # .github/workflows/update-viablestrict.yml # .github/workflows/update_pytorch_labels.yml # .github/workflows/upload-test-stats-while-running.yml # .github/workflows/upload-test-stats.yml # .github/workflows/upload-torch-dynamo-perf-stats.yml # .github/workflows/upload_test_stats_intermediate.yml # .github/workflows/vllm-benchmark.yml # .github/workflows/weekly.yml # .github/workflows/xpu.yml # aten/src/ATen/native/cuda/GroupedBlas.cpp # aten/src/ATen/native/cuda/ScaledBlas.cpp # aten/src/ATen/native/cuda/UpSampleNearest2d.cu # aten/src/ATen/native/mps/kernels/ReduceOps.metal # aten/src/ATen/native/mps/operations/ReduceOps.mm # aten/src/ATen/native/sparse/cuda/cuSPARSELtOps.cpp # aten/src/ATen/native/transformers/cuda/mem_eff_attention/kernel_forward.h # c10/xpu/XPUFunctions.cpp # c10/xpu/XPUStream.cpp # cmake/Dependencies.cmake # cmake/External/aotriton.cmake # cmake/public/LoadHIP.cmake # related_commits # requirements-build.txt # requirements.txt # setup.py # test/distributed/_composable/fsdp/test_fully_shard_training.py # test/distributed/launcher/test_run.py # test/distributed/tensor/test_attention.py # test/inductor/test_aot_inductor.py # test/inductor/test_compile_subprocess.py # test/inductor/test_cudagraph_trees.py # test/inductor/test_gpu_cpp_wrapper.py # test/inductor/test_torchinductor.py # test/inductor/test_torchinductor_dynamic_shapes.py # test/inductor/test_torchinductor_opinfo_properties.py # test/test_bmm_outer_product.py # test/test_cuda.py # test/test_linalg.py # test/test_multiprocessing.py # test/test_nn.py # test/test_ops.py # test/test_sort_and_select.py # test/test_utils.py # test/test_xpu.py # third_party/composable_kernel # third_party/xpu.txt # tools/build_with_debinfo.py # tools/stats/import_test_stats.py # tools/test/test_build_with_debinfo.py # torch/_dynamo/variables/streams.py # torch/_higher_order_ops/utils.py # torch/_inductor/codecache.py # torch/_inductor/compile_fx.py # torch/_inductor/fx_passes/reinplace.py # torch/_inductor/runtime/triton_heuristics.py # torch/_native/ops/bmm_outer_product/triton_impl.py # torch/csrc/distributed/c10d/symm_mem/nccl_dev_cap.hpp # torch/cuda/graphs.py # torch/testing/_internal/common_cuda.py # torch/testing/_internal/common_methods_invocations.py # torch/testing/_internal/opinfo/definitions/linalg.py # torch/xpu/__init__.py # version.txt
rocm-repo-management-api-6
Bot
requested review from
jataylo,
jeffdaily,
jithunnair-amd and
pruthvistony
as code owners
August 21, 2026 17:53
This was referenced Aug 21, 2026
jithunnair-amd
pushed a commit
that referenced
this pull request
Aug 25, 2026
…l) (#3594) ## 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](#3503)). Analyzes the full IFU delta from [PR #3591](#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](https://amd.atlassian.net/wiki/spaces/MLSE/pages/1819593560): - **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 (pytorch#180927) ? not needed on 2.14. Upstream [pytorch#180927](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 (#3434) ? 2.14 already has zero `rocm_smi` refs. - Enable linalg via hipSOLVER, remove Magma (#3455) ? all 6 upstream PRs already in 2.14. - Stabilize ROCm 7.14 testing (#3525) ? 2.14 already carries the guards + skip messages. - Add gfx1250 support in GroupedBlas (#3532) ? already in 2.14 via upstream pytorch#190703 (`ScaledBlasDeviceUtils`). - Fix uninitialized embedding indices `test_embedding_meta_indices` (pytorch#188002) ? upstream already in 2.14. - 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 - [x] 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: - py3.10 https://github.com/ROCm/TheRock/actions/runs/32863299013 - py3.11 https://github.com/ROCm/TheRock/actions/runs/32863305363 - py3.12 https://github.com/ROCm/TheRock/actions/runs/32863311391 - py3.13 https://github.com/ROCm/TheRock/actions/runs/32863318109 - py3.14 https://github.com/ROCm/TheRock/actions/runs/32863325370 - [x] 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): - py3.10 https://github.com/ROCm/TheRock/actions/runs/32818564344 - py3.11 https://github.com/ROCm/TheRock/actions/runs/32818568751 - py3.12 https://github.com/ROCm/TheRock/actions/runs/32818572249 - py3.13 https://github.com/ROCm/TheRock/actions/runs/32818576288 - py3.14 https://github.com/ROCm/TheRock/actions/runs/32818580229 ## Processed Commits List Spreasheet with each commit evaluated for cherry-pick & resolution (drop/include) with reasoning: [pr3594-commit-outcomes.xlsx](https://github.com/user-attachments/files/31390971/pr3594-commit-outcomes.xlsx) --------- Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: Iurii Paikov <Iurii.Paikov@amd.com> Co-authored-by: Chinmay Dattanand Kuchinad <40351312+chinmaydk99@users.noreply.github.com> Co-authored-by: rocm-repo-management-api-6[bot] <212817015+rocm-repo-management-api-6[bot]@users.noreply.github.com> Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com> Co-authored-by: srinivamd <52507740+srinivamd@users.noreply.github.com> Co-authored-by: Subramanian <ssubrama1@amd.com> Co-authored-by: zjliu-amd <zhuojin.liu@amd.com> Co-authored-by: Andres Lugo <108368282+alugorey@users.noreply.github.com> Co-authored-by: pablo-garay <pablo-garay@users.noreply.github.com> Co-authored-by: Bo Li <110066325+BLOrange-AMD@users.noreply.github.com> Co-authored-by: Anatoliy Litvinenko <alitvine@amd.com> Co-authored-by: Jeff Daily <jeff.daily@amd.com> Co-authored-by: Luca Bruni <luca.bruni@amd.com> Co-authored-by: albmalamd <albert.malewski@amd.com>
jithunnair-amd
force-pushed
the
release/2.14
branch
from
August 27, 2026 18:18
40ed5e2 to
d42c3b5
Compare
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.
rocm_base: aafe5bd