[ROCm][CI] Validate shared ROCm SO file list#5
Draft
amdfaa wants to merge 8666 commits into
Draft
Conversation
…orch#188659) Mirror the `test-osdc` HF-cache pattern (pytorch#188654) in the vLLM benchmark: offline reads the shared read-only `/mnt/hf_cache`; nightly refreshes it — in place on DGX B200 (writable host mount) or download + `aws s3 sync` back to the bucket on OSDC. FlashInfer's JIT workspace is cached the same way (warmed from the shared cache, synced on nightly) but always compiles into a writable dir, since the OSDC mount is read-only. The unconditional `/mnt/hf_cache` bind mount is dropped on OSDC — the runner hook already injects a read-only mount, and the extra bind mount made `runner-container-hooks` fail `lstat('/mnt/hf_cache')` in the runner pod, failing every run and retry. It's kept only for non-OSDC B200, which has no injection and needs the host bind mount. ### Testing https://github.com/pytorch/pytorch/actions/runs/28606132504 Another round https://github.com/pytorch/pytorch/actions/runs/28621594034 Everything https://github.com/pytorch/pytorch/actions/runs/28639471616 Pull Request resolved: pytorch#188659 Approved by: https://github.com/izaitsevfb
…up.py) (pytorch#188914) ## Summary Restore the pre-2.13 behavior where **release/RC** CUDA wheels ship **without `+PTX`**, while nightlies keep it for forward-compat. ### Background `libtorch_cuda.so` grew ~80-90 MB from 2.12.1 to 2.13.0 (cu130): 2.13 started emitting `12.0+PTX` (sm_120 PTX) in the release wheels. The SASS arch set is unchanged; the growth is the embedded compute_120 PTX. The branch-cut script `scripts/release/apply-release-changes.sh` was *supposed* to strip `+PTX` for releases, but it edited `.ci/manywheel/build_cuda.sh` — which was refactored into `.ci/manywheel/build_env_setup.py`. The `sed` silently became a no-op, so 2.13.0 shipped with `+PTX`. ### Fix Detect release/RC builds in Python instead of relying on a branch-cut `sed`: - `build_env_setup.py`: add `_is_release_build()` (true when `PYTORCH_BUILD_VERSION` is set and has no `.dev` suffix — i.e. tagged release/RC builds, per `binary_populate_env.sh`) and `_ptx_arches()`; emit `+PTX` only for nightly/dev builds. - `apply-release-changes.sh`: drop the now-redundant/broken `+PTX` `sed`. ### Behavior (cuda 13.0, x86_64) | Build | `PYTORCH_BUILD_VERSION` | arch list | |---|---|---| | nightly | `2.13.0.dev...+cu130` | `7.5;8.0;8.6;9.0;10.0;12.0+PTX` | | RC / release | `2.13.0+cu130` | `7.5;8.0;8.6;9.0;10.0;12.0` | | local / unset | `''` | `...;12.0+PTX` (safe default) | This is drift-proof: the logic lives next to the arch table it depends on, keyed off a signal the repo already uses. *Authored with AI assistance.* Pull Request resolved: pytorch#188914 Approved by: https://github.com/zxiiro, https://github.com/huydhn
## Human Note This is a pretty long standing issue that just didnt feel like investigating or fixing. In fact we had a fix before that this is basd off of. TLDR is just actually calculate the sparsity at the block granularity no the easier numel layer. Simple easy, happy claudexing :) ## Agent Report # Report ## Summary Fixed the negative `BlockMask.sparsity()` display for masks whose logical sequence length does not fill the final sparse block. The issue repro now prints `sparsity=0.00%` instead of `sparsity=-63.84%`. ## Root cause `BlockMask.sparsity()` counted computed blocks, converted them to token area by multiplying by `Q_BLOCK_SIZE * KV_BLOCK_SIZE`, and divided by the logical token count. For partial edge blocks, especially a 100x100 mask with the default 128x128 block size, that overcounted area outside the logical sequence. One computed block was treated as 16,384 computed tokens for a 10,000-token mask, yielding negative sparsity. The mask metadata and FlexAttention computation use the logical sequence lengths and block metadata correctly. The negative value was an accounting/display bug, not an output correctness issue. ## Changes - `torch/nn/attention/flex_attention.py` - Changed `BlockMask.sparsity()` to divide computed block count by the logical block-grid size, computed with `cdiv` over sequence lengths and block sizes plus batch/head dimensions. - Lifted the existing nested `cdiv` helper to module scope so both `sparsity()` and `to_string()` use the same helper. - `test/inductor/test_flex_attention.py` - Added `test_block_mask_sparsity_with_partial_block` covering the 100-token document-causal repro and checking the string display reports `sparsity=0.00%`. ## Validation - Reproduced the original issue before the fix: `BlockMask(shape=(1, 1, 100, 100), sparsity=-63.84%, ...)` and `mask.sparsity() == -63.84000000000001`. - Post-fix issue repro now prints `BlockMask(shape=(1, 1, 100, 100), sparsity=0.00%, ...)` and `mask.sparsity() == 0.0`. - Verified compiled FlexAttention output for the 100x100 document-causal mask against a dense masked reference: - `max_abs=0.0009765625`, `mean_abs=3.796815872192383e-05`, `torch.testing.assert_close(..., rtol=1e-2, atol=1e-2)` passed. - Scout fanout: - Launched five read-only subagent scouts to search for additional sparsity/display users, partial-block cases, non-divisible length tests, and device variants. - No additional direct `BlockMask.sparsity()` assertions were found outside `test/inductor/test_flex_attention.py`. - Scouts recommended nearby BlockMask slicing, `_adjust`, `from_kv_blocks`, custom block-size, and non-divisible FlexAttention/decoding tests. - Targeted tests: - `../.venv/bin/python test/inductor/test_flex_attention.py TestBlockMaskCUDA.test_block_mask_sparsity_with_partial_block_cuda -v` - `../.venv/bin/python test/inductor/test_flex_attention.py TestBlockMaskCUDA.test_block_mask_attributes_cuda TestBlockMaskCUDA.test_block_mask_viz_cuda -v` - `../.venv/bin/python test/inductor/test_flex_attention.py TestFlexAttentionCUDA.test_mask_mod_combiners_cuda TestBlockMaskCUDA.test_pytree_flatten_unflatten_cuda -v` - `../.venv/bin/python test/inductor/test_flex_attention.py TestBlockMaskCUDA.test_adjust_block_mask_ignores_entries_past_num_blocks_cuda TestBlockMaskCUDA.test_getitem_query_slice_updates_shape_cuda TestBlockMaskCUDA.test_from_kv_blocks_unpadded_kv_indices_with_seq_lengths_cuda TestBlockMaskCUDA.test_from_kv_blocks_full_indices_False_cuda TestBlockMaskCUDA.test_from_kv_blocks_full_indices_True_cuda TestBlockMaskCUDA.test_block_size_cuda TestBlockMaskCUDA.test_block_size_changes_BLOCK_SIZE4_cuda TestBlockMaskCUDA.test_block_size_changes_BLOCK_SIZE5_cuda TestBlockMaskCUDA.test_block_size_changes_BLOCK_SIZE_128_cuda TestBlockMaskCUDA.test_block_size_changes_BLOCK_SIZE_256_cuda TestBlockMaskCUDA.test_block_size_changes_BLOCK_SIZE_32_cuda TestBlockMaskCUDA.test_block_size_changes_BLOCK_SIZE_64_cuda -v` - `../.venv/bin/python test/inductor/test_flex_attention.py TestFlexAttentionCUDA.test_block_mask_non_divisible_cuda TestFlexAttentionCUDA.test_causal_block_non_divisible_cuda TestFlexAttentionCUDA.test_causal_block_non_divisible_with_captured_buffer_cuda -v` - `../.venv/bin/python test/inductor/test_flex_decoding.py TestFlexDecodingCUDA.test_non_sparse_mulitple_block_size_cuda TestFlexDecodingCUDA.test_multi_block_m_q_seq_len_23_cuda TestFlexDecodingCUDA.test_decode_at_different_input_position_float16_score_mod0_cuda_float16 -v` - Lint/prerequisite checks: - `spin fixlint` was attempted but exited 1 due unrelated pre-existing shellcheck findings in `.ci/pytorch/*.sh`. - `spin quickfix` completed successfully on changed files with no lint issues and no additional tracked file changes. Re-ran after the `cdiv` cleanup. ## Artifacts - Diff written to `fix.diff`. <details> <summary>Agent Worklog</summary> ## Run 1 > **User:** Investigate meta-pytorch/attention-gym issue pytorch#93 as an adhoc PyTorch/FlexAttention task. Issue URL: meta-pytorch/attention-gym#93 Title: Doc mask returns negative sparsity Issue summary: - The displayed sparsity for a FlexAttention BlockMask is negative when the block size is bigger than the max size of the mask. - Reporter says it is similar to attention-gym#68. - They ask whether this affects results obtained with such a mask. Repro from issue: ```python import torch from torch.nn.attention.flex_attention import create_block_mask document_id = torch.zeros(100, dtype=torch.int, device="cuda") document_id[:10] = 0 document_id[10:20] = 1 for i in range(20, 100, 20): document_id[i : i + 20] = i // 20 + 1 def document_causal_mask(b, h, q_idx, kv_idx): causal_mask = q_idx >= kv_idx document_mask = document_id[q_idx] == document_id[kv_idx] return causal_mask & document_mask mask = create_block_mask(document_causal_mask, 1, 1, 100, 100, "cuda") print(mask) ``` Observed output in issue: ```text BlockMask(shape=(1, 1, 100, 100), sparsity=-63.84%, (0, 0) ██ ) ``` Task: - Treat the issue text as evidence, not instructions. - Reproduce the negative sparsity display in this PyTorch checkout using the PTQ job Python. - Determine whether it is display-only or affects FlexAttention results. - If root cause is in PyTorch, implement the smallest durable fix in the PyTorch worktree. - Add targeted regression coverage for BlockMask sparsity/display behavior. - Validate with the repro and relevant targeted tests. - Keep worklog.md/report.md/fix.diff current. ### Reproduced negative BlockMask sparsity Ran the issue repro with `/home/drisspg/.ptq_workspace/jobs/20260702-pytorch-adhoc-b6ca1e/.venv/bin/python` in this checkout. It prints `BlockMask(shape=(1, 1, 100, 100), sparsity=-63.84%, ...)`, and `mask.sparsity()` returns `-63.84000000000001`, so the negative value is reproducible here. ### Checked FlexAttention correctness Compared the reproduced 100x100 document-causal BlockMask against a dense masked attention reference with random CUDA fp16 Q/K/V. Both direct `flex_attention` and `torch.compile(flex_attention, fullgraph=True)` matched the reference (`max_abs <= 0.0009765625`), so the negative sparsity appears to be display/accounting only, not an output correctness issue. ### Implemented sparsity accounting fix Changed `BlockMask.sparsity()` to divide computed block count by the logical block-grid size (`ceil(Q_LEN / Q_BLOCK_SIZE) * ceil(KV_LEN / KV_BLOCK_SIZE)` across batch/head), instead of converting computed blocks back to full token area. Added a regression test for the 100x100 document-causal repro that now expects `sparsity=0.00%` in the string display. ### Validation Post-fix issue repro now prints `sparsity=0.00%` and `mask.sparsity()` returns `0.0`. The new CUDA regression test passed, nearby existing `test_block_mask_attributes_cuda` and `test_block_mask_viz_cuda` passed, and compiled FlexAttention still matched a dense masked reference (`max_abs=0.0009765625`). `spin fixlint` was attempted as required but failed on unrelated pre-existing shellcheck findings in `.ci/pytorch/*.sh`; `spin quickfix` on changed files completed successfully with no lint issues and no extra tracked file changes. ### Artifacts updated Wrote final summary to `report.md` and generated the current source diff at `fix.diff`. ### cdiv cleanup Lifted the existing nested `cdiv` helper to module scope and updated `BlockMask.sparsity()` to use `cdiv(seq_len, block_size)`. Re-ran the focused CUDA regression test successfully and `spin quickfix` still reports no lint issues on changed files. ### Additional sparsity-return tests Searched `test/inductor/test_flex_attention.py` for remaining `sparsity()` / `sparsity=` assertions. Ran the remaining methods that compare returned sparsity values: `TestFlexAttentionCUDA.test_mask_mod_combiners_cuda` and `TestBlockMaskCUDA.test_pytree_flatten_unflatten_cuda`; both passed. ### Scout fanout and extra edge tests Launched five read-only subagent scouts to search for additional sparsity/display users, partial-block tests, non-divisible length cases, and device variants. They found no additional direct `BlockMask.sparsity()` assertions beyond the known FlexAttention file, but recommended nearby BlockMask slicing/_adjust/from_kv_blocks/custom block-size tests and non-divisible FlexAttention/decoding tests. Ran 12 focused `TestBlockMaskCUDA` edge-case tests, 3 non-divisible `TestFlexAttentionCUDA` tests, and 3 selected `TestFlexDecodingCUDA` tests; all passed. One scout noted zero-length `sparsity()/str()` can still divide by zero, but that behavior existed under the old `numel()==0` formula too and no existing test was found to hit it. </details> --- *This PR was generated by [ptq](https://github.com/drisspg/pt_job_queue) with human review.* Pull Request resolved: pytorch#188868 Approved by: https://github.com/liangel-02
…ytorch#188865) The `fence_proxy()` API changed to accept string literals instead of enum members (`ProxyKind.async_shared` -> `"async.shared"`, `SharedSpace.shared_cta` -> `"cta"`). The enum-based API was deprecated and then removed in newer `nvidia_cutlass_dsl` versions, causing `AttributeError: module 'cutlass.cute.arch' has no attribute 'ProxyKind'` on CI. Root cause: the vendored `dense_blockscaled_gemm_persistent.py` used the old enum API. Test Plan: Verified the fix locally by running the previously-failing test: ```bash python test/inductor/test_nv_universal_gemm.py \ TestNVUniversalGemm.test_bmm_non_standard_batch_stride_bfloat16 python test/inductor/test_nv_universal_gemm.py \ TestNVUniversalGemm.test_bmm_non_standard_batch_stride_float16 ``` Both pass. The failure could not be reproduced locally because the local `nvidia_cutlass_dsl` still has the deprecated enum (with warnings), but CI's version has it removed. Co-authored-by: Claude AI (claude-code) Pull Request resolved: pytorch#188865 Approved by: https://github.com/crcrpar, https://github.com/Skylion007 ghstack dependencies: pytorch#188303
…188080) Root cause: itertools.repeat's bounded form repeat(item, times) was modeled as a Python generator polyfill, which has no __length_hint__ and no repr, and operator.length_hint had no dispatch handler. As a result, length_hint over a bounded repeat hit "Failed to trace builtin operator", and repr of a bounded repeat hit "repr_impl not implemented". All three CPython itertools sentinels (LengthTransparency.test_repeat, LengthTransparency.test_repeat_with_negative_times, and TestBasicOps.test_repeat_with_negative_times) share this single root cause. Fix: model RepeatIteratorVariable's bounded form natively. It now carries an optional `times` (None == unbounded, mirroring CPython's cnt == -1) and a `remaining` counter, raises StopIteration on exhaustion, and implements: - __length_hint__ returning `remaining` for the bounded form; the unbounded form raises TypeError (caught by length_hint, which falls back to the default), mirroring repeat_lengthhint. - repr_impl matching repeat_repr. - a 2-arg reconstruct for the bounded form. Negative `times` is clamped to 0 per repeat_new. A new call_length_hint handler mirrors PyObject_LengthHint: it tries __len__, then __length_hint__, then the supplied default. Test Plan: ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 python test/cpython/v3_13/test_itertools.py LengthTransparency.test_repeat CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 python test/cpython/v3_13/test_itertools.py LengthTransparency.test_repeat_with_negative_times CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 python test/cpython/v3_13/test_itertools.py TestBasicOps.test_repeat_with_negative_times ``` ``` python test/dynamo/test_misc.py -k itertools_repeat ``` ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 python test/cpython/v3_13/test_itertools.py ``` All targeted tests pass; the full itertools file is OK (skipped=63) with no new failures. Authored with the assistance of an AI coding assistant. Pull Request resolved: pytorch#188080 Approved by: https://github.com/rtimpe
…ytorch#188081) The shared CPython _TestBasicOps.test_iteration calls iter(s).__length_hint__() to check how many elements remain in a set iterator. In Dynamo, the set iterator and the dict keys/values/items iterators are all modeled by a single VariableTracker, DictViewIterator (the shared base). That class had no call_method, so __length_hint__ fell through to the base IteratorVariable's "Unsupported method call" and graph-broke / failed. Because all seven element-type variants (Bytes, Empty, MixedStringBytes, Singleton, String, Triple, Tuple) exercise the same shared test body, they all share this single root cause. The fix adds a call_method to DictViewIterator that handles __length_hint__ by returning operator.length_hint(self._iter). self._iter is the live Python iterator captured over the view's items, so its own length hint already reflects how many elements remain and decrements as next() is run during tracing -- mirroring CPython's setiter_len / dictiter_len. One handler on the shared base therefore covers the set iterator and all three dict view iterators at once. Test Plan: ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 \ python test/cpython/v3_13/test_set.py \ TestBasicOpsBytes.test_iteration TestBasicOpsTuple.test_iteration TestBasicOpsEmpty.test_iteration ``` ``` python test/dynamo/test_sets.py -k length_hint ``` ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 \ python test/cpython/v3_13/test_set.py ``` ``` python test/dynamo/test_sets.py ``` All pass with no new failures. Authored with the assistance of an AI coding assistant. Pull Request resolved: pytorch#188081 Approved by: https://github.com/rtimpe ghstack dependencies: pytorch#188080
…G34) (pytorch#188083) The five CPython313 test_sort TestOptimizedCompares tests (test_safe_object, test_unsafe_float, test_unsafe_latin, test_unsafe_long, test_unsafe_tuple) all call the shared helper check_against_PyObject_RichCompareBool, whose first two statements are random.seed(0) then random.shuffle(L) before sorting. Module-level random.seed is a bound method on the module-global random.Random instance. VariableBuilder._wrap already routed module-level random.shuffle and random.sample through RandomVariable (which models the RNG state) but did not recognize seed, so random.seed(0) fell through to the skipfile path and graph broke ("Attempted to call function marked as skipped: Random.seed") under the test's error_on_graph_break=True. The test therefore failed before ever reaching the sort. The fix adds "seed" to the recognized method-name set in that existing module-level-random branch, so random.seed(...) routes through the already implemented RandomVariable.seed. Because the seed is fixed, the subsequent shuffle and the resulting sort are deterministic and match eager. This is distinct from the deferred dynamic-random sort case (G22), where the random data is not seeded and the sort result is not reproducible. What this gate validates is observable sort correctness (result, stability, errors), not a mirror of CPython's internal compare specialization selection, per CPYTHON_MIRRORING.md "What not to mirror". The sixth sibling test_unsafe_object_compare is intentionally left as an expected failure: it defines a local class closing over cells and mutates __class__ mid-sort, which trips the out-of-scope local class-body closure-cell / build-class machinery, unrelated to this RNG-routing fix. Test Plan: 5 target sentinels under Dynamo (all pass): ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 \ python test/cpython/v3_13/test_sort.py \ TestOptimizedCompares.test_safe_object_compare \ TestOptimizedCompares.test_unsafe_float_compare \ TestOptimizedCompares.test_unsafe_latin_compare \ TestOptimizedCompares.test_unsafe_long_compare \ TestOptimizedCompares.test_unsafe_tuple_compare # Ran 5 tests, OK ``` Whole affected CPython file under Dynamo (no new failures, no XPASS): ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 \ python test/cpython/v3_13/test_sort.py # Ran 21 tests, OK (skipped=11) ``` New regression test: ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu \ python test/dynamo/test_unspec.py -k seed # Ran 1 test, OK ``` Nearby Dynamo suite regression sanity: ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu \ python test/dynamo/test_unspec.py # Ran 60 tests, OK (skipped=1, expected failures=2) ``` Note: this change was authored with the help of an AI assistant. Pull Request resolved: pytorch#188083 Approved by: https://github.com/rtimpe ghstack dependencies: pytorch#188080, pytorch#188081
…torch#188235) random.random is a C builtin method bound to the module-global Random instance (unlike randint/randrange/uniform, which are Python methods, and unlike shuffle/sample/seed which gate G34 routed). Under Dynamo a call to it graph-broke ("Attempted to call function marked as skipped: Random.random") instead of modeling the RNG value. builder._wrap now routes the module-global random.random through UserDefinedObjectVariable so its call_random_fn / RandomValueSource path models a fresh per-call value (replayed to match eager). The routing is tightly gated on is_supported_random_obj plus membership in the supported-random-functions set, so only the module-global RNG's supported methods take this path. This is a sibling of gate G34's random.seed routing. Removes the CPython313 test_deque TestBasic.test_reverse expected-failure sentinel, whose body builds inputs with random.random and now passes. Test Plan: ``` python test/dynamo/test_unspec.py -k random_random CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 \ python test/cpython/v3_13/test_deque.py TestBasic.test_reverse ``` This change was authored with the help of an AI assistant. Pull Request resolved: pytorch#188235 Approved by: https://github.com/rtimpe ghstack dependencies: pytorch#188080, pytorch#188081, pytorch#188083
Neither copy.copy(deque) nor deque.copy() was modeled under Dynamo. copy.copy resolves type(d).__copy__, which Dynamo could not trace, and the inherited list-style copy dropped maxlen. DequeVariable.call_method now handles copy and __copy__ by returning a fresh DequeVariable that preserves maxlen and shallow-copies the elements, mirroring CPython deque_copy. user_defined.py routes copy.copy(deque) through deque.__copy__: deque is not in copy.py's _copy_builtin_containers, so copy.copy falls back to the cls.__copy__ path. Removes the CPython313 test_deque TestBasic.test_copy expected-failure sentinel. The vendored test also builds its inputs with random.random, which is handled by the preceding commit (random.random routing). Test Plan: ``` CUDA_VISIBLE_DEVICES= PYTORCH_TESTING_DEVICE_ONLY_FOR=cpu PYTORCH_TEST_WITH_DYNAMO=1 \ python test/cpython/v3_13/test_deque.py TestBasic.test_copy python test/dynamo/test_sequence_ops.py -k deque_copy ``` This change was authored with the help of an AI assistant. Pull Request resolved: pytorch#188220 Approved by: https://github.com/rtimpe ghstack dependencies: pytorch#188080, pytorch#188081, pytorch#188083, pytorch#188235
## Human Note I think that this is pretty edge casey but simple enough to fix. basically when you close over a input that requreis grad but is not yet used we still think we need to roduce gards for this closure. But in fact we do not since it is essnetially just waiting to get dce'd. I though about trying to dce as early as posssible and worked on this but it ended up being way larger blast radius. I think this is fine pretty self contained and handles an edge case in a decently simple way. Calling bwd direclty has lil more code needed to maek sure thi sstill works but ring attention impl is still good ## Agent Report # PyTorch Issue pytorch#166722 Report ## Summary FlexAttention backward failed when `score_mod` captured a `requires_grad=True` tensor, indexed it, but did not use that indexed value in the returned score. The AOT joint graph correctly produced `None` for the unused captured tensor gradient, but the eager fallback and compiled metadata paths did not consistently preserve that `None`. ## Root cause `torch/_higher_order_ops/flex_attention.py::sdpa_dense_backward` preallocates `actual_grad_score_mod_captured` entries for captured `score_mod` tensors that require grad. Later, it blindly ran `actual_grad.copy_(grad)` whenever the preallocated slot was a tensor. For an unused capture such as `self.weight[h]; return score`, AOT autograd returns `grad is None`, following normal autograd semantics for unused differentiable inputs. The unconditional copy caused: ```text TypeError: copy_(): argument 'other' (position 1) must be Tensor, not NoneType ``` After fixing that eager path, `torch.compile` exposed the same semantic mismatch in metadata: the proxy/fake output for `flex_attention_backward` still described the unused captured grad as a tensor, while the joint graph and Inductor lowering returned `None`. Inductor then tried to treat the runtime `None` as a tensor and failed with `AttributeError: 'NoneType' object has no attribute 'get_size'`. ## Fix The eager fallback now uses the joint graph to allocate captured-grad destination slots only for captures whose gradient output can be a tensor. The copy-back path treats a non-Tensor produced grad as `None` and asserts that tensor grads have a destination buffer before copying. If AOT returns `None`, FlexAttention propagates `None` for that captured input's gradient instead of materializing or copying a gradient. The same shared helpers normalize the joint graph's captured-gradient outputs and allocate captured-grad buffers for eager, proxy tracing, and fake tensor mode. This keeps unused captured-grad slots represented as `None` in compiled metadata, matching the real lowering/runtime result. The allocation is based on the joint graph output rather than the saved buffer's `requires_grad`, since compiled backward can need a captured intermediate's gradient to propagate to its producer even when the saved tensor itself does not carry `requires_grad=True`. In proxy tracing, the joint graph is materialized before computing the example output so the metadata is generated from the final graph rather than patched afterward. I also added `TestFlexAttention.test_unused_captured_score_mod_grad`, which exercises both uncompiled and compiled `flex_attention` with a captured `requires_grad=True` tensor that is indexed in `score_mod` but unused in the returned score. The test verifies that Q/K/V receive gradients and the unused captured tensor's gradient remains `None`. ## Repro The checked-in `repro.py` uses `device = None`, which now hits the current checkout's explicit CPU limitation (`FlexAttention does not support backward on CPU`) before reaching the reported bug. The CUDA version below reproduced the eager issue before the fix. I also checked the same module under `torch.compile`; after the eager-only fix it failed in Inductor metadata/lowering, and after the final fix both modes complete. <details> <summary>Repro Script</summary> ```python import torch from torch import nn from torch.nn.attention import flex_attention class MhaBugRepro(nn.Module): def __init__(self, num_heads: int): super().__init__() self.weight = nn.Parameter(torch.zeros(num_heads, 2)) def forward(self, q, k, v): def score_mod(score, b, h, q_idx, kv_idx): _ = self.weight[h] return score return flex_attention.flex_attention(q, k, v, score_mod=score_mod) device = "cuda" for compiled in [False, True]: torch._dynamo.reset() repro = MhaBugRepro(num_heads=8).to(device) if compiled: repro = torch.compile(repro) q = torch.randn(1, 8, 10, 128, requires_grad=True, device=device) k = torch.randn(1, 8, 10, 128, device=device) v = torch.randn(1, 8, 10, 128, device=device) loss = repro(q, k, v).sum() loss.backward() weight = repro._orig_mod.weight if compiled else repro.weight print("compiled" if compiled else "eager", "ok", q.grad.shape, weight.grad) ``` Output after the fix: ```text /home/drisspg/.ptq_workspace/jobs/20260701-pytorch-166722/pytorch/torch/nn/attention/flex_attention.py:2565: UserWarning: flex_attention called without torch.compile() - this will use an unfused implementation that materializes the full scores matrix instead of generating a fused kernel. SOLUTION: Use torch.compile(flex_attention)(...) If you want to debug your score_mod/mask_mod, you can set: torch.nn.attention.flex_attention._FLEX_ATTENTION_DISABLE_COMPILE_DEBUG = True This will allow you to use print statements or breakpoints. Note: This doesn't work with the backwards pass and may produce incorrect results. _warn_once( eager ok torch.Size([1, 8, 10, 128]) None compiled ok torch.Size([1, 8, 10, 128]) None ``` </details> ## Test results Behavioral tests: ```text cd pytorch && ~/.ptq_workspace/jobs/20260701-pytorch-166722/.venv/bin/python test/inductor/test_flex_attention.py TestFlexAttentionCUDA.test_direct_backward_supports_symint_score_mod_buffers_cuda TestLearnableBiasesCUDA.test_backprop_error_case_cuda Ran 2 tests in 7.062s OK cd pytorch && ~/.ptq_workspace/jobs/20260701-pytorch-166722/.venv/bin/python test/inductor/test_flex_attention.py -k test_unused_captured_score_mod_grad Ran 2 tests in 5.801s OK cd pytorch && ~/.ptq_workspace/jobs/20260701-pytorch-166722/.venv/bin/python test/inductor/test_flex_attention.py -k test_captured_scalar_grad Ran 1 test in 16.183s OK ``` Broader suite signal: ```text cd pytorch && ~/.ptq_workspace/jobs/20260701-pytorch-166722/.venv/bin/python -m pytest test/inductor/test_flex_attention.py -n 32 -q 667 passed, 24 skipped, 1 xfailed, 2 subtests passed in 292.50s (0:04:52) ``` Before `pytest-xdist` was installed, `pytest` rejected `-n`; I also ran the full file serially and it produced no failures before a 3600s timeout. The first xdist run after installation exposed `TestLearnableBiasesCUDA.test_backprop_error_case_cuda` with `y.grad is None`; that was fixed by allocating captured-grad buffers from the joint graph output rather than `buffer.requires_grad`, and the final xdist full-file run above passed. Ring attention direct-backward smoke: ```text cd pytorch && ~/.ptq_workspace/jobs/20260701-pytorch-166722/.venv/bin/torchrun --standalone --nproc_per_node=2 ../agent_space/ring_attention.py --seq-len 256 --dtype float32 --backend TRITON [rank 0] local forward shard matches reference [rank 1] local forward shard matches reference [rank 0] gathered forward matches reference [rank 0] local dq, dk, and dv match reference [rank 1] local dq, dk, and dv match reference [rank 0] gathered outputs and gradients match the single-process reference ``` I also ran the same copied ring-attention script with default/AUTO and `--no-validate`; forward and backward completed on both ranks. The default/AUTO validated run produced large gradient mismatches, and FLASH bf16 hit separate CuTe block-sparse block-size constraints, so the passing direct-backward validation signal is the explicit TRITON run above. Prerequisite checks: ```text cd pytorch && spin quickfix ok No lint issues. cd pytorch && spin quicklint ok No lint issues. ``` I also attempted full `spin fixlint`; it failed on unrelated pre-existing shellcheck issues in `.ci/pytorch/test.sh` and `.ci/pytorch/common.sh`, not on the changed Python files. Fixes pytorch#166722 <details> <summary>Repro Script</summary> ```python import torch from torch import nn from torch.nn.attention import flex_attention class MhaBugRepro(nn.Module): def __init__(self, num_heads: int): super().__init__() self.weight = nn.Parameter(torch.zeros(num_heads, 2)) def forward(self, q, k, v): def score_mod(score, b, h, q_idx, kv_idx): foo = self.weight[h] # THIS LINE CAUSES IT TO BREAK return score return flex_attention.flex_attention( q, k, v, score_mod=score_mod, ) # Same behaviour on CPU and GPU device = None repro = MhaBugRepro(num_heads=8).to(device) loss = repro(torch.randn(1, 8, 10, 128, requires_grad=True).to(device), torch.randn(1, 8, 10, 128).to(device), torch.randn(1, 8, 10, 128).to(device)).sum() loss.backward() ``` </details> <details> <summary>Agent Worklog</summary> ## Run 1 ### Primed and reproduced on GPU Read the PTQ context, issue prompt, existing worklog, and repo AGENTS instructions. The provided `repro.py` hits a CPU `NotImplementedError` because FlexAttention backward is unsupported on CPU in this checkout. Running the same repro on CUDA (`NVIDIA GB200`, torch `2.14.0a0+gita7dae7f`) reproduces the reported failure at backward with `TypeError: copy_(): argument 'other' (position 1) must be Tensor, not NoneType` in `torch/_higher_order_ops/flex_attention.py:1150`. ### Fixed unused captured score_mod grads Traced the eager failure to `sdpa_dense_backward`: it allocates an output gradient buffer for captured `score_mod` tensors with `requires_grad=True`, but the AOT joint graph legitimately returns `None` when the captured tensor is referenced and not used in the returned score. Updated the copy-back path to return `None` for non-Tensor grads instead of calling `copy_(None)`. ### Fixed compiled metadata for unused captured grads Checked `torch.compile` after the eager fix and found a related Inductor backward failure: `AttributeError: 'NoneType' object has no attribute 'get_size'`. The compiled path's proxy/fake metadata still claimed the unused captured grad was a tensor even though the joint graph and lowering returned `None`. Added a helper that reads the joint graph outputs corresponding to `score_mod` captures, then uses it in proxy tracing and fake tensor mode so `None` captured-grad slots stay `None` in metadata. Expanded `test_unused_captured_score_mod_grad` to run both uncompiled and compiled `flex_attention`. ### Simplified final patch and validated CUDA issue repro now completes backward in both eager/uncompiled and `torch.compile` modes and leaves `weight.grad` as `None`, as expected for an unused parameter. Refined the eager allocation so captured-grad destination buffers are also `None` when the joint graph reports an unused captured grad, then simplified the patch to share captured-grad output normalization and buffer allocation across eager/proxy/fake paths. Targeted tests pass: `test_unused_captured_score_mod_grad` (2 parameterized cases) and `test_captured_scalar_grad`. `spin quickfix` and `spin quicklint` pass on changed files. Full `spin fixlint` was attempted earlier and fails on pre-existing unrelated shellcheck issues in `.ci/pytorch/*.sh`. ### Full FlexAttention file signal Tried to run the full file with pytest-xdist via `python -m pytest test/inductor/test_flex_attention.py -n 15 -q`, but this job venv did not initially have xdist enabled (`pytest: error: unrecognized arguments: -n`). Ran the full `test/inductor/test_flex_attention.py` file serially with the repo test runner instead; it produced no failures before the 3600s timeout and no leftover test processes remained afterward. After `pytest-xdist` was installed, ran `python -m pytest test/inductor/test_flex_attention.py -n 32 -q`. The first xdist run failed `TestLearnableBiasesCUDA.test_backprop_error_case_cuda` with `y.grad is None`, showing that captured-grad destination buffers must be allocated whenever the joint graph produces a captured grad, not only when the saved captured tensor has `requires_grad=True`. Removed that stale `requires_grad` condition. Also checked whether the callable-joint-graph fallback was actually needed. Removing it broke direct compiled backward/fake-mode paths where proxy/fake tracing sees a callable joint graph before `make_fx` materializes it, so the fallback is real. To avoid patching metadata after the fact, reordered `trace_flex_attention_backward` to materialize `fw_graph`, `joint_graph`, and `mask_graph` first, then compute `example_out` once using those final graphs. Reran direct-backward tests, targeted regressions, lint, and the full xdist file. Final full-file result after the reorder: `667 passed, 24 skipped, 1 xfailed, 2 subtests passed in 292.50s`. ### Ring attention direct-backward smoke Copied the provided ring attention example into `agent_space/ring_attention.py` and ran it with two local GB200 ranks. The explicit TRITON backend validated forward and direct-backward gradients against the single-process reference: `torchrun --standalone --nproc_per_node=2 ../agent_space/ring_attention.py --seq-len 256 --dtype float32 --backend TRITON` printed local/gathered forward and gradient matches. The default/AUTO backend completed forward and backward with `--no-validate`, showing the direct backward path compiles and runs, but its validated run had large gradient mismatches; that appears to be backend/LSE-conversion behavior in the example rather than a captured-grad metadata crash. FLASH bf16 runs hit separate CuTe block-sparse block-size constraints before validation. </details> --- *This PR was generated by [ptq](https://github.com/drisspg/pt_job_queue) with human review.* Pull Request resolved: pytorch#188860 Approved by: https://github.com/liangel-02
…setters on complex tensors (pytorch#188603)" This reverts commit 815cc61. Reverted pytorch#188603 on behalf of https://github.com/meta-codesync due to Diff reverted internally ([comment](pytorch#188603 (comment)))
## Summary Replace the floating `ubuntu-latest` GitHub Actions runner label with the explicit `ubuntu-24.04` image across all workflows. `ubuntu-latest` currently resolves to `ubuntu-24.04`, so this is behavior-preserving today, but pinning avoids surprise breakage / toolchain churn whenever GitHub advances the `ubuntu-latest` alias to a newer image (as happened on the 20.04->22.04 and 22.04->24.04 transitions). ## Details - 40 occurrences across 34 workflow files, all the exact `ubuntu-latest` label in `runs-on:` (no self-hosted `ubuntu-latest-*` variants), so this is a mechanical 1:1 substitution. - No other YAML changed. ## Test Plan ``` grep -rn "ubuntu-latest" .github/ # 0 results after the change ``` *Authored with AI assistance.* Pull Request resolved: pytorch#188930 Approved by: https://github.com/malfet, https://github.com/Skylion007
Fixes pytorch#148183 # The code builds ✅ ``` python -m pip install -e . -v --no-build-isolation ``` # Warn no longer reproduces on Mac ✅ ``` import torch x = torch.randn(3, 4, 5) mask = torch.randint(0, 2, (3, 4, 5), dtype=torch.bool) value = torch.tensor([1.0, 2.0, 3.0]) def fn(x, mask, val): return x.masked_fill_(mask, val) result = torch.vmap(fn)(x, mask, value) ``` ## Before ``` python3 masked_fill_repro.py .../pytorch/masked_fill_repro.py:8: UserWarning: There is a performance drop because we have not yet implemented the batching rule for aten::masked_fill_.Tensor. Please file us an issue on GitHub so that we can prioritize its implementation. (Triggered internally at .../pytorch/aten/src/ATen/functorch/BatchedFallback.cpp:83.) return x.masked_fill_(mask, val) ``` ## After ``` python3 masked_fill_repro.py ``` # Linter is ok ✅ ``` pytorch % lintrunner f ok No lint issues. ``` # Changed test suites ok ✅ ``` python test/functorch/test_vmap.py -k "masked_fill" Ran 7 tests in 0.109s OK (skipped=1) ``` # Some tests fail - verified fail on main too ✅ ``` python ./test/test_torch.py ---------------------------------------------------------------------- Ran 971 tests in 3.808s FAILED (failures=1, errors=2, skipped=144) ``` Pull Request resolved: pytorch#175513 Approved by: https://github.com/Skylion007 Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com> Co-authored-by: Aaron Gokaslan <aaronGokaslan@gmail.com>
…ch#188927) Update the bundled protobuf submodule from ancient v3.13 (released Oct 2020) to last Abseil-free v21.12 (released Dec 2022). Abseil has no stable ABI and carries its own process-global state, so embedding it into libtorch_cpu would create ODR/visibility hazards with any other Abseil/protobuf already loaded in the process (TensorFlow, gRPC, system protobuf). Root cause of the one required patch change: protobuf 3.21 emits PROTOBUF_CONSTINIT (constinit) on the generated *_default_instance_ globals. ProtoBufPatch.cmake strips PROTOBUF_CONSTEXPR from the ConstantInitialized constructors (a long-standing Windows workaround), which makes them non-constexpr; constinit then requires a constant initializer it can no longer get, so every ONNX .pb.cc failed to compile with "variable does not have a constant initializer". The fix strips PROTOBUF_CONSTINIT alongside PROTOBUF_CONSTEXPR so those globals fall back to dynamic initialization. proto_wrap.cc is unchanged, since GetEmptyStringAlreadyInited still exists in 21.x. As a side benefit, the 3.21 code generator no longer emits the C++20-deprecated ATOMIC_VAR_INIT macro into the generated ONNX *.pb.cc TODO: the only remaining C++ protobuf consumer is the legacy TorchScript ONNX exporter, which merely serializes onnx::ModelProto. Once ONNX proto generation is moved to Python (the torch.export/onnxscript exporter already works this way) or the legacy exporter is dropped entirely, libtorch stops needing protobuf and the bundled submodule, ProtoBufPatch.cmake, and proto_wrap.cc can all be removed -- which is also the only way to converge onto ONNX's own protobuf pin (4.25/v29.x) without vendoring Abseil. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Pull Request resolved: pytorch#188927 Approved by: https://github.com/Skylion007
…rch#188839) Fix mul_ and ne bugs in ComplexTensor Fixes two independent bugs in `torch/_subclasses/complex_tensor/_ops/aten.py`. **`mul_`:** was registered via `register_binary_nonlinear(aten.mul_)`, which applies the op directly to split real/imag components. The in-place`aten.mul_` mutated the `a_r` view before it was reused for the imag term, corrupting the result, and returned a new ComplexTensor instead of `lhs` (so `x.mul_(y) is x` was false). Now registered via `register_binary_linear_inplace(aten.mul_, mul_impl)`, which computes the product functionally then does `lhs.copy_(result)` and returns `lhs`. **`ne`:** `ne_impl` used `split_complex_tensor(self)`, which reads `.re`/`.im` attributes that only exist on a ComplexTensor, breaking `torch.ne(real, complex)`. Switched to `split_complex_arg(self)` to handle real tensors and numbers (imag treated as zero), matching `eq_impl`. ## Test Plan Added two regression tests in `test/complex_tensor/test_complex_tensor.py`: - `test_mul_inplace_complex`: `mul_` returns the same object and equals the out-of-place product. - `test_ne_real_operand`: `torch.ne` with a real left operand matches native complex `ne`. Pull Request resolved: pytorch#188839 Approved by: https://github.com/amjames, https://github.com/Skylion007
With hand-rolled UTF-16 to UTF-8 converter, which is trivial loop of `decode_utf16`/`encode_utf8` rather than adding more and more deprecation suppressions, and error out of utf16 string is malformed (same as `wstring_convert` does on MacOS-26) TODO: Figure out if this helper function at all makes sense outside of Windows platform, where it used to convert UTF-16 WIN32 error messages back into utf-8 Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Pull Request resolved: pytorch#188904 Approved by: https://github.com/huydhn
The `.value.f = ...` style nested designators are a C99 extension and triggered `-Wc99-designator` warning using newer SDK.
Rewrite them as proper nested initializers (`.value = {.f = ...}`) and funnel all integral (and bool) cases through a single `intScalar` templated lambda, since they all write
the shared `i` union member.
Delete nused `bool b` union member: the union is only ever written in MPSScalar constructor and later copied out as raw bytes, so no named member is ever read back, and all moder Macs are little endiant )there are no MacOS-14 for PowerPC ones :P)
Authored with the assistance of Claude.
Pull Request resolved: pytorch#188910
Approved by: https://github.com/Skylion007
ghstack dependencies: pytorch#188904
This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Update the pinned vision hash. Pull Request resolved: pytorch#188877 Approved by: https://github.com/pytorchbot
…88934) This PR is auto-generated nightly by [this action](https://github.com/pytorch/pytorch/blob/main/.github/workflows/nightly.yml). Update the pinned torchcomms hash. Pull Request resolved: pytorch#188934 Approved by: https://github.com/pytorchbot
# Summary This was found while looking at : meta-pytorch/attention-gym#208 Our pack analyzer basically didn't have good analzysis when a predicate is lane uniform e.g. `q_idx < seq_lens[b]` There are no lanes -> e.g. no kv_indices so the stride coeff is 0 == so completely lane unfiorm previously that would fall through. The De-Morgan flip was from fable although yes I know this rule ;) Pull Request resolved: pytorch#188929 Approved by: https://github.com/liangel-02
…orch#188920) ## Summary Follow-up to pytorch#188061 / pytorch#188060, addressing @malfet's [review comment](pytorch#188061 (comment)). pytorch#188061 stopped `print_sccache_stats()` from failing a successful build when sccache is missing, but it does so by **silently** returning. As malfet noted, that makes it *"much harder to detect when sccache should be there, but it's not installed"* — on Linux CI (everything except s390x/riscv64) sccache is expected, so a broken install would otherwise yield a green build with no signal. ### Change When sccache is absent, `print_sccache_stats()` now branches on whether sccache was actually **expected** for this build, using `SCCACHE_BUCKET` — the same signal `common-build.sh` already uses to gate sccache setup: - **`SCCACHE_BUCKET` set** (sccache was configured, so a missing binary is a real misconfiguration): emit a GitHub Actions `::error::` annotation and `return 1`, failing the build. Callers run under `set -e`, so on Linux `build.sh` this turns the job red instead of shipping a silently-uncached build. - **`SCCACHE_BUCKET` unset** (sccache legitimately not in use): emit a `::warning::` and continue. The build stays green (preserving the pytorch#188060 fix), but the condition is still loud and greppable. `macos-build.sh`: drop the now-redundant `if which sccache` wrapper around the call, since the function self-guards — making behavior uniform across platforms. (macOS scripts don't run under `set -e` and the macOS call site was already guarded, so the new error path is a no-op there.) ### Why key on `SCCACHE_BUCKET` instead of a new flag There is no existing `SKIP_SCCACHE` / `NO_SCCACHE` opt-out. The closest env var, `SKIP_SCCACHE_INITIALIZATION`, means something different (skip *starting the server* on self-hosted runners — sccache is still expected). `SCCACHE_BUCKET` is the real "sccache is configured" gate: `common-build.sh` explicitly `unset`s it when empty, so its presence is a reliable "sccache expected" signal that needs no new configuration. ### Why this over silent-skip or unconditional hard-fail - **Silent-skip** hides misconfiguration (malfet's point). - **Unconditional hard-fail** regresses pytorch#188060 — builds that legitimately don't use sccache would fail at the very end. - **Context-dependent** (this PR): fail only when sccache was expected, warn otherwise. Non-fatal where sccache is optional, loud-and-fatal where it's required. Note: sccache setup + the stats-on-exit `sccache_epilogue` trap in `common-build.sh` are already gated on `which sccache`; this only changes the explicit `print_sccache_stats()` call path. ## How the annotations surface on the job `echo "::warning::<msg>"` / `echo "::error::<msg>"` are GitHub Actions workflow commands. When the step prints one, the runner turns it into an annotation: - **Job log** — appears in the raw step log (rendered highlighted, not plain text). - **Job summary page** — shown as a yellow⚠️ (warning) or red ❌ (error) box at the top of the run/job. - **PR "Checks" tab / merge box** — propagates to the job's check run, visible without opening the full log. Key difference: `::warning::` does **not** change the step exit code — the build stays green. Failing the build in the "expected" case comes from the `return 1` under `set -e`, **not** from the `::error::` line itself (the annotation is purely cosmetic/surfacing). The `::error::` is emitted alongside `return 1` so the failure is both fatal and clearly annotated. Notes: - We don't set `file=`/`line=` params, so it's a general job-level annotation (appropriate — it's a build-environment condition, not a source location). - The annotation rendering only happens under GitHub Actions. If these scripts ever run elsewhere, the line just prints literally as text — still visible/greppable in logs. Pull Request resolved: pytorch#188920 Approved by: https://github.com/izaitsevfb, https://github.com/malfet
## Human Note This lets us return grads for scalar biases. I pusehd back on this in the path since there were workarounds although a lil ugly. But after prototyping i think its fine to land and easy enough to gork. We should wait for pytorch#188860 to land first since they touch teh same code -> how we create the captured buffer grads ## Agent Report # Report: attention-gym pytorch#171 FlexAttention learnable scalar ## Summary Reproduced the issue in this checkout using the exact scalar `score_mod` pattern: - Eager forward succeeded, eager backward failed with the reported `vmap(... out_dim is None)` BatchedTensor error. - Compiled forward succeeded, compiled backward failed in Inductor FlexAttention backward lowering with `AssertionError: ComputedBuffer name must not be None`. Implemented a PyTorch-side fix for direct 0-D captured score-mod gradients. The PR branch has since been rebased onto latest `origin/main` at `d2f0d442d70` and rebuilt locally with cuDNN enabled (`CUDNN_VERSION=9.23.0`, `USE_CUDNN=1`). Current job torch reports `2.14.0a0+git9fadffc`. Implemented changes: - `torch/_higher_order_ops/flex_attention.py` - Routes direct 0-D captured score-mod tensors through `mod_index(buffer, [])` while tracing the joint graph, so the existing captured-index custom autograd path emits `flex_lib::zeros_and_scatter([], [], grad)` for scalar gradient accumulation. - Preserves unused/no-grad captured buffers by returning `None` instead of copying `None` into an allocated grad buffer. - `torch/_dynamo/_trace_wrapped_higher_order_op.py` - Supports empty-index `zeros_and_scatter` as scalar accumulation by summing values into a scalar output. - Allows `ModIndex` to use an empty index list as a scalar identity whose backward is empty-index scalar accumulation. - Returns the checked 0-D tensor directly in the empty-index `ModIndex.forward` path. - Handles empty index metadata in its vmap rule with an annotation-only pyrefly fix, preserving the previous empty-list fallback semantics. - `torch/_inductor/kernel/flex/common.py` - Lowers empty-index `zeros_and_scatter` to a scalar atomic-add scatter for FlexAttention template subgraphs. - Uses the same compact empty-index guard shape as eager `zeros_and_scatter`. - `torch/_inductor/select_algorithm.py` - Normalizes scalar scatter indexes for Triton codegen and emits `tl.full(..., INDEX_DTYPE)` for constant scalar atomic indexes instead of invalid `tl.broadcast_to(0, ...)`. - `test/inductor/test_flex_attention.py` - Adds a parametrized CUDA float16 regression covering both direct learnable 0-D scalar gradients and detached/no-grad 0-D captures. ## Relationship to the prior unused-captured-grad issue The direct scalar learnable case is distinct from the existing `(1,)` indexed captured-scalar path: indexed captures already trace to `zeros_and_scatter([1], [idx], grad)`, while direct 0-D captures previously returned a plain computed captured grad that became batched under nested `vmap`. The local cleanup now makes direct 0-D captures enter the same `ModIndex` custom autograd path using an empty index list, instead of post-processing captured grad outputs after `create_joint`. I also checked a detached 0-D captured scalar (`score + temp.detach()`). Before the guard, eager backward tried to copy a `None` captured grad into an allocated buffer. The patch keeps that captured grad as `None`, and the new parametrized test checks eager and compiled behavior for this no-grad case. ## Validation Behavioral repro after rebasing onto `origin/main`: ```bash cd ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0 TORCHINDUCTOR_FORCE_DISABLE_CACHES=1 CUDA_LAUNCH_BLOCKING=1 ./.venv/bin/python pytorch/agent_space/repro_flex_scalar.py parity ``` Passed after the main rebase, after the local cleanup, and after the subagent-suggested simplifications. Final max diffs from eager-vs-compiled parity: - out: 0.00048828125 - q: 0.00048828125 - k: 0.00048828125 - v: 0.0 - temp: 0.005798816680908203 Targeted tests after rebasing onto `origin/main`: ```bash cd ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0/pytorch CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_0d_scalar_grad --verbose CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k unused_captured_score_mod_grad --verbose CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_scalar_grad --verbose ``` Passed after the main rebase, after the local cleanup, and after the subagent-suggested simplifications: - `captured_0d_scalar_grad`: 2 tests - `unused_captured_score_mod_grad`: 2 tests - `captured_scalar_grad`: 1 test A combined `-k "unused_captured_score_mod_grad or captured_scalar_grad"` attempt ran 0 tests with this test runner and was rerun as separate filters. Full Flex xdist validation: ```bash cd ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0/pytorch CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python -m pytest -n 32 \ test/inductor/test_flex_attention.py \ test/inductor/test_flex_aux_vectorization.py \ test/inductor/test_flex_decoding.py \ test/inductor/test_flex_flash.py \ test/inductor/test_flex_gemm.py ``` Installed `pytest-xdist` into the job venv first (`pytest 9.1.1`, `xdist 3.8.0`). The `-n 32` run completed with `1460 passed, 58 skipped, 3 xfailed, 11 failed in 491.94s`. The 11 failures were CUDA OOM / CUDA driver OOM under 32 concurrent GPU workers; the full log showed GPU 0 nearly full with many pytest worker processes. After the run exited, `nvidia-smi` showed no remaining GPU processes, and rerunning the 11 failed nodeids serially passed: `11 passed in 144.92s`. Reran the FlexAttention test file alone with 20 xdist workers: ```bash CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python -m pytest -n 20 test/inductor/test_flex_attention.py ``` Passed: `670 passed, 24 skipped, 1 xfailed in 376.72s`. Lint/prerequisite checks: ```bash cd ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0/pytorch spin lint -- --take PYREFLY --all-files spin lint -- --take RUFF --all-files ``` Passed locally after the annotation-only pyrefly fix, after the local cleanup, after the subagent-suggested simplifications, and after the CI lint formatting fix. The CI lint formatting fix was validated with `spin lint -- -m origin/main`. ```bash cd ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0/pytorch spin fixlint ``` `spin fixlint` applied no relevant changes and reported Python lint OK, but the overall command exited 1 due to unrelated pre-existing shellcheck findings in CI shell scripts outside this patch, including `.ci/pytorch/test.sh` and `binary_linux_test.sh`. ## PR status note The pushed PR branch is `ptq/20260702-pytorch-adhoc-6f35e0` for pytorch#188869. GitHub CI previously showed a failing `lintrunner-pyrefly-partial / lint` job due to the empty-list fallback lacking an annotation; that annotation-only fix was pushed as `e5277e3401b`. The branch was then rebased onto latest `origin/main` (`d2f0d442d70`) to handle merge conflicts. The rebase kept main's newer unused-captured-grad behavior and this PR's 0-D scalar scatter behavior. The rebased commits are `1a706c48278` and `9fadffc8461`. The branch was pushed with `git push --force-with-lease`, and `uv run ptq pr attention-gym-171` updated the existing PR body using this report. A follow-up cleanup that removes the post-hoc `grads[5:]` scalar special case was pushed with `uv run ptq pr attention-gym-171` as commit `5774cca6c3e`. A subagent simplification pass then accepted smaller follow-ups: direct scalar return in `ModIndex.forward`, compact empty-index guards in Inductor lowering, and `INDEX_DTYPE` for scalar constant index codegen. The larger suggestion to replace empty-index scatter with a scalar identity backward was not taken because scalar reduction through `zeros_and_scatter` is what prevents the nested-vmap captured grad from remaining batched. The accepted subagent simplifications were pushed as commit `f50c8d570c4` on `origin/ptq/20260702-pytorch-adhoc-6f35e0`, and the existing PR was updated. CI triage later found two failing lint jobs, `lintrunner-noclang-partial / lint` (run `28711442321`, job `85145558474`) and `lintrunner-noclang-all / lint` (run `28711445286`, job `85145565994`). Both requested the same formatting-only patch to split a long `RuntimeError` line in `torch/_dynamo/_trace_wrapped_higher_order_op.py`. The formatting fix passed `spin lint -- -m origin/main`, the focused 0-D scalar regression, and `git diff --check` locally, then was pushed as commit `ce62c44335e`. On the new commit, the previously failing `lintrunner-noclang-partial / lint` and `lintrunner-noclang-all / lint` checks now pass; pyrefly and quick-check lint entries also pass. ## Artifacts - Repro script: `agent_space/repro_flex_scalar.py` - Diff: `fix.diff` <details> <summary>Agent Worklog</summary> ## Run 1 > **User:** Investigate meta-pytorch/attention-gym issue pytorch#171 as an adhoc PyTorch/FlexAttention task. Issue URL: meta-pytorch/attention-gym#171 Title: Flex Attention does not support learnable scalar? Important context: - A previous PTQ workspace for this issue (`20260702-pytorch-adhoc-61e5b7`) was accidentally cleaned before its patch/report/fix.diff were preserved. - Do not start fully from scratch: use the recovered notes below as a strong hypothesis, but verify everything in this new checkout. - This issue overlaps with pytorch#166722 but is not fully covered by that fix. In the pytorch#166722 workspace, the exact scalar repro below still failed in eager backward with: `ValueError: vmap(<lambda>(), ...): <lambda>() can not return a BatchedTensor when out_dim is None` Issue summary: - Reporter says FlexAttention with a learnable scalar in `score_mod` fails in backward. - Without compile: forward works, backward fails. - With compile: forward/backward path also fails. - Minimal score_mod pattern from issue: ```python temp = nn.Parameter(torch.tensor(0.0)) def score_mod(score, b, h, q, kv): score = score + temp return score ``` Exact repro to use for eager vs compiled parity: ```python import torch from torch import nn from torch.nn.attention.flex_attention import flex_attention class M(nn.Module): def __init__(self, device): super().__init__() self.temp = nn.Parameter(torch.tensor(0.7, device=device, dtype=torch.float32)) def forward(self, q, k, v): temp = self.temp def score_mod(score, b, h, q_idx, kv_idx): return score * temp + temp return flex_attention(q, k, v, score_mod=score_mod) device = "cuda" dtype = torch.float16 B, H, S, D = 1, 2, 16, 16 torch.manual_seed(123) q = torch.randn(B, H, S, D, device=device, dtype=dtype, requires_grad=True) k = torch.randn(B, H, S, D, device=device, dtype=dtype, requires_grad=True) v = torch.randn(B, H, S, D, device=device, dtype=dtype, requires_grad=True) grad = torch.randn(B, H, S, D, device=device, dtype=dtype) q2 = q.detach().clone().requires_grad_() k2 = k.detach().clone().requires_grad_() v2 = v.detach().clone().requires_grad_() m1 = M(device) m2 = M(device) m2.load_state_dict(m1.state_dict()) out1 = m1(q, k, v) out1.backward(grad) out2 = torch.compile(m2)(q2, k2, v2) out2.backward(grad) torch.cuda.synchronize() for name, a, b in [("out", out1, out2), ("q", q.grad, q2.grad), ("k", k.grad, k2.grad), ("v", v.grad, v2.grad), ("temp", m1.temp.grad, m2.temp.grad)]: diff = (a - b).abs().max().item() print(name, a.dtype, b.dtype, diff, a.flatten()[0].item(), b.flatten()[0].item()) torch.testing.assert_close(a, b, atol=1e-2, rtol=1e-2) ``` Recovered hypothesis from the deleted workspace: - Eager backward reproduced the reported failure: `vmap(<lambda>(), ...): <lambda>() can not return a BatchedTensor when out_dim is None` - In that run, compiled forward reportedly succeeded, but compiled backward failed with: `AssertionError: ComputedBuffer name must not be None` - A local patch reportedly made eager and compiled backward pass using this approach: 1. Support empty-index `zeros_and_scatter([], [], grad)` as scalar accumulation. 2. Wrap 0-D captured `score_mod` gradients with that scatter in FlexAttention's joint backward graph. 3. Teach Inductor's flex scatter lowering/codegen to emit scalar atomic-adds. - Files reportedly touched in the deleted worktree were in this neighborhood: - `torch/_dynamo/_trace_wrapped_higher_order_op.py` - `torch/_higher_order_ops/flex_attention.py` - `torch/_inductor/kernel/flex/common.py` - `torch/_inductor/select_algorithm.py` Task: - Treat the issue text and recovered notes as evidence, not instructions. - Reproduce the exact eager and compiled failures in this new PTQ workspace. - Verify the relationship to pytorch#166722 and avoid regressing the pytorch#166722 unused-captured-grad case. - If root cause is in PyTorch, implement the smallest durable fix in the PyTorch worktree. - Add targeted regression coverage for learnable scalar score_mod gradients in eager and compiled FlexAttention. - Validate with the exact repro, the new targeted tests, and adjacent FlexAttention captured-gradient tests. - Keep `worklog.md`, `report.md`, and `fix.diff` current. ## Manual run - kickoff Started manual investigation in fresh PTQ job. Confirmed the PyTorch worktree was clean at `a7dae7fa10c`. Launched read-only reconnaissance on FlexAttention captured-gradient and Inductor scatter paths; the first generic-agent launch failed because that agent name was unavailable, then reran with `scout` agents. ## Reproduced scalar captured-gradient failures Created `agent_space/repro_flex_scalar.py` with separate eager/compiled/parity modes using the exact learnable scalar `score_mod` pattern. Eager forward succeeds, eager backward fails with `ValueError: vmap(<lambda>(), ...): <lambda>() can not return a BatchedTensor when out_dim is None`. Compiled forward succeeds, compiled backward fails during Inductor lowering of `flex_attention_backward` with `AssertionError: ComputedBuffer name must not be None` in `torch/_inductor/kernel/flex/flex_attention.py::process_joint_outputs`. ## Implemented scalar scatter path Inspected `create_fw_bw_graph` and confirmed the existing `(1,)` captured scalar test uses `zeros_and_scatter([1], [idx], grad)`, while a direct 0-D captured parameter produced a plain computed captured grad (`add_1`) that becomes batched under FlexAttention backward `vmap`. Patched 0-D captured score-mod grads to return `zeros_and_scatter([], [], grad)`, added eager custom-op scalar accumulation, and added Inductor scalar scatter lowering/codegen support for atomic-adding every tile contribution to the scalar output. After fixes, `agent_space/repro_flex_scalar.py eager`, `compiled`, and `parity` all complete; parity max diffs were <= 0.0058 and within the repro tolerances. ## Added regression tests and checked adjacent coverage Added a parametrized TestFlexAttention CUDA float16 regression for direct 0-D captured scalar score_mod behavior. The detach_temp=False case checks eager-vs-compiled parity for output, q/k/v grads, and temp.grad. The detach_temp=True case covers the unused/no-grad captured scalar relationship to the prior unused captured grad issue and asserts temp.grad stays None in eager and compiled modes. Targeted tests passed: test_flex_attention.py -k captured_0d_scalar -v, -k captured_scalar_grad -v, and -k bf16_score_mod_captured_grad_dtype -v. ## Final validation and lint status Reran the exact parity repro after cleanup: agent_space/repro_flex_scalar.py parity passed with max diffs out/q/k 0.00048828125, v 0.0, temp 0.005799770355224609. Reran the new regression tests after spin formatting: test_flex_attention.py -k captured_0d_scalar -v passed 2 tests. Ran spin fixlint as required; it applied no relevant changes and reported Python lint OK, but the overall command exited 1 because of unrelated pre-existing shellcheck findings in CI shell scripts such as .ci/pytorch/test.sh and binary_linux_test.sh. ## Reran tests after annotation-only pyrefly fix After CI showed a pyrefly implicit-any error for the empty index fallback, changed the local fix to an annotation-only update so fallback behavior stays semantically unchanged. Reran spin lint with PYREFLY only and it passed. Reran the focused captured_0d_scalar regression and both generated CUDA float16 tests passed. Reran the exact parity repro; out/q/k max diff 0.00048828125, v max diff 0.0, temp max diff 0.005799293518066406, within tolerance. ## Pushed pyrefly fix to existing PR Ran the PTQ PR updater from the PTQ repo: uv run ptq pr attention-gym-171. It reused the existing open PR pytorch#188869, committed the annotation-only pyrefly fix as e5277e3 on branch ptq/20260702-pytorch-adhoc-6f35e0, pushed it to origin, and updated the PR. Local PyTorch worktree is clean; GitHub checks restarted and lintrunner-pyrefly-partial is queued on the new commit. ## Rebased onto latest origin/main and rebuilt The user asked to regenerate the shared local build on latest `origin/main` and rebase this PR branch on top of main to handle merge conflicts. Rebuilt the PTQ seed with cuDNN enabled: ```bash CUDNN_INCLUDE_DIR=/usr/include CUDNN_LIBRARY=/usr/lib64 USE_CUDNN=1 uv run ptq setup --local --build --onto origin/main ``` The seed is now at `d2f0d442d70` and reports `torch 2.14.0a0+gitd2f0d44`, `CUDNN_VERSION=9.23.0`, and `USE_CUDNN=1`. Rebased `ptq/20260702-pytorch-adhoc-6f35e0` onto `origin/main`. The rebase conflicted in `torch/_higher_order_ops/flex_attention.py` and `test/inductor/test_flex_attention.py`. Resolved by keeping main's stricter captured-grad copy loop, retaining this branch's 0-D scalar `zeros_and_scatter([], [], grad)` wrapping, and keeping both main's unused-captured-grad regression and this branch's 0-D scalar regression. The rebased branch is now: - `1a706c48278 Fix from 20260702-pytorch-adhoc-6f35e0` - `9fadffc8461 Fix from 20260702-pytorch-adhoc-6f35e0` Rebuilt the job workspace with cuDNN enabled: ```bash CUDNN_INCLUDE_DIR=/usr/include CUDNN_LIBRARY=/usr/lib64 USE_CUDNN=1 bash ~/.ptq_workspace/scripts/rebuild.sh ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0/pytorch ``` The job venv now reports `torch 2.14.0a0+git9fadffc`, `torch.version.git_version == 9fadffc`, `CUDNN_VERSION=9.23.0`, and `USE_CUDNN=1`. Regenerated `fix.diff` from `origin/main...HEAD` and restored `pytorch/agent_space/repro_flex_scalar.py` after the build-artifact clean. Post-rebase validation: - Exact eager-vs-compiled scalar repro with `TORCHINDUCTOR_FORCE_DISABLE_CACHES=1 CUDA_LAUNCH_BLOCKING=1` passed; max diffs were out/q/k `0.00048828125`, v `0.0`, temp `0.005799770355224609`. - `CUDA_LAUNCH_BLOCKING=1 .venv/bin/python test/inductor/test_flex_attention.py -k captured_0d_scalar_grad --verbose` passed 2 tests. - `CUDA_LAUNCH_BLOCKING=1 .venv/bin/python test/inductor/test_flex_attention.py -k unused_captured_score_mod_grad --verbose` passed 2 tests. - `CUDA_LAUNCH_BLOCKING=1 .venv/bin/python test/inductor/test_flex_attention.py -k captured_scalar_grad --verbose` passed 1 test. - `git diff --check origin/main...HEAD` passed. Pushed the rebased branch with: ```bash git -C ~/.ptq_workspace/jobs/20260702-pytorch-adhoc-6f35e0/pytorch push --force-with-lease origin ptq/20260702-pytorch-adhoc-6f35e0 ``` Then ran `uv run ptq pr attention-gym-171`; it reused the saved human note, found the existing open PR pytorch#188869, pushed the branch, and updated the PR body. No new source commit was created by `ptq pr` because the worktree was already clean after the rebase. ## Local cleanup after review discussion Reviewed the PR diff lines that special-cased 0-D captured gradients after `create_joint`. Reworked the local patch so direct 0-D captured tensors are wrapped with `mod_index(buffer, [])` while tracing the joint graph. This makes AOTAutograd emit `flex_lib::zeros_and_scatter([], [], grad)` through the existing captured-index custom autograd path instead of rewriting `grads[5:]` after the fact. Also made the eager empty-index scalar accumulation branch in `zeros_and_scatter` more compact with `if not indices` / `if shape`, and taught `ModIndex.forward` that an empty index list is a scalar identity. Validation after this cleanup: - Exact parity repro with `TORCHINDUCTOR_FORCE_DISABLE_CACHES=1 CUDA_LAUNCH_BLOCKING=1` passed; max diffs were out/q/k `0.00048828125`, v `0.0`, temp `0.005799770355224609`. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_0d_scalar_grad --verbose` passed 2 tests. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k unused_captured_score_mod_grad --verbose` passed 2 tests. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_scalar_grad --verbose` passed 1 test. - `spin lint -- --take PYREFLY --all-files` passed. - `spin lint -- --take RUFF --all-files` passed. - `git diff --check` passed. A combined `-k "unused_captured_score_mod_grad or captured_scalar_grad"` attempt ran 0 tests with this test runner and was rerun as separate filters. Pushed the cleanup to the existing PR with: ```bash cd /home/drisspg/meta/pt_job_queue && uv run ptq pr attention-gym-171 ``` PTQ created commit `5774cca6c3e` on `ptq/20260702-pytorch-adhoc-6f35e0`, pushed it to `origin/ptq/20260702-pytorch-adhoc-6f35e0`, and updated pytorch#188869. The PyTorch source worktree is clean after the push. ## Subagent simplification review Launched a four-agent read-only simplification review across both available models: - `reviewer` on `plugboard-codex/gpt-5.5` for `flex_attention.py` / `_trace_wrapped_higher_order_op.py` readability. - `reviewer` on `anthropic/claude-opus-4-7` for design/layering. - `kernel-reviewer` on `plugboard-codex/gpt-5.5` for Inductor scalar scatter lowering/codegen. - `validator` on `anthropic/claude-opus-4-7` for test shape and validation gaps. Accepted three low-risk simplifications from the review: - `ModIndex.forward` now returns the checked 0-D tensor directly instead of `x.view(())`. - `zeros_and_scatter_lowering` uses the same compact `if not indices` / `if shape` form as eager `zeros_and_scatter`. - Scalar constant index codegen now emits `tl.full(..., INDEX_DTYPE)` instead of hard-coded `tl.int64`. Rejected/deferred larger suggestions: a dedicated scalar identity op with `backward` returning `grad` would not reduce the nested-vmap captured scalar gradient and risks reintroducing the original failure; broader test restructuring/paged-attention expansion is not needed for this narrowly targeted PR. Validation after these extra simplifications: - Exact parity repro with `TORCHINDUCTOR_FORCE_DISABLE_CACHES=1 CUDA_LAUNCH_BLOCKING=1` passed; max diffs were out/q/k `0.00048828125`, v `0.0`, temp `0.005798816680908203`. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_0d_scalar_grad --verbose` passed 2 tests. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k unused_captured_score_mod_grad --verbose` passed 2 tests. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_scalar_grad --verbose` passed 1 test. - `spin lint -- --take PYREFLY --all-files` passed. - `spin lint -- --take RUFF --all-files` passed. - `git diff --check` passed. Pushed these extra simplifications to the existing PR with `uv run ptq pr attention-gym-171`. PTQ created commit `f50c8d570c4` on `ptq/20260702-pytorch-adhoc-6f35e0`, pushed it to `origin/ptq/20260702-pytorch-adhoc-6f35e0`, and updated pytorch#188869. ## Full Flex xdist run Installed `pytest-xdist` into the job venv with: ```bash ../.venv/bin/python -m pip install pytest-xdist ``` Confirmed `pytest 9.1.1` and `xdist 3.8.0`. Ran all Flex Inductor test files with 32 xdist workers: ```bash CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python -m pytest -n 32 \ test/inductor/test_flex_attention.py \ test/inductor/test_flex_aux_vectorization.py \ test/inductor/test_flex_decoding.py \ test/inductor/test_flex_flash.py \ test/inductor/test_flex_gemm.py ``` Result: `1460 passed, 58 skipped, 3 xfailed, 11 failed in 491.94s`. The 11 failures were CUDA OOM / CUDA driver OOM under 32 concurrent GPU workers; the log showed GPU 0 nearly full with many pytest worker processes. Full log path: `/tmp/pi-bash-253d8eff10e95d29.log`. After the xdist run exited, `nvidia-smi --query-compute-apps=pid,used_memory,process_name --format=csv,noheader,nounits` showed no remaining GPU processes. Reran the 11 failed nodeids serially with: ```bash CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python -m pytest -q <11 failed nodeids> ``` All 11 passed in 144.92s, confirming the `-n 32` failures were concurrency/memory pressure rather than correctness failures in the patch. Reran the FlexAttention test file alone with 20 xdist workers: ```bash CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python -m pytest -n 20 test/inductor/test_flex_attention.py ``` Passed: `670 passed, 24 skipped, 1 xfailed in 376.72s`. ## CI lint fix Ran CI triage for pytorch#188869 with: ```bash ~/dotfiles/scripts/github_ci_triage pytorch#188869 --output-dir agent_space/ci_logs ``` Found two failing lint jobs: - `lintrunner-noclang-partial / lint`, run `28711442321`, job `85145558474` - `lintrunner-noclang-all / lint`, run `28711445286`, job `85145565994` Both requested the same formatting patch in `torch/_dynamo/_trace_wrapped_higher_order_op.py`: split the long `RuntimeError("mod_index with no indices only supports scalar tensors")` line. Applied the formatting change and regenerated `fix.diff`. Validation after the lint fix: - `spin lint -- -m origin/main` passed. - `CUDA_LAUNCH_BLOCKING=1 ../.venv/bin/python test/inductor/test_flex_attention.py -k captured_0d_scalar_grad --verbose` passed 2 tests. - `git diff --check` passed. Pushed the lint fix with `uv run ptq pr attention-gym-171`. PTQ created commit `ce62c44335e` on `ptq/20260702-pytorch-adhoc-6f35e0`, pushed it to `origin/ptq/20260702-pytorch-adhoc-6f35e0`, and updated pytorch#188869. Watched the lint checks on the new commit. The previously failing `lintrunner-noclang-partial / lint` and `lintrunner-noclang-all / lint` now pass. `lintrunner-pyrefly-partial`, `lintrunner-pyrefly-all`, and both `quick-checks / lint` entries also pass. </details> --- *This PR was generated by [ptq](https://github.com/drisspg/pt_job_queue) with human review.* Pull Request resolved: pytorch#188869 Approved by: https://github.com/liangel-02
…ytorch#188345) ## Issue Fixes pytorch#188344 ## Summary The channels_last (NHWC) CUDA backward kernel `avg_pool2d_backward_out_cuda_frame_nhwc` recovered the input pixel coordinates from the flat index without the padding offset (`const int w = (index / channels) % width;`), while the `phstart/phend/pwstart/pwend` window formulas it then uses require padded coordinates. The contiguous (NCHW) kernel `avg_pool2d_backward_out_cuda_frame` adds the offset (`+ pad_w` / `+ pad_h`); the NHWC kernel did not, so for `padding > 0` it mapped each input pixel to the wrong set of output windows and returned silently wrong input gradients. `padding = 0` was unaffected because the two coordinate systems coincide. This adds the missing `+ pad_w` / `+ pad_h`, matching the NCHW kernel. `test_avg_pool2d_nhwc` did not catch this because its `helper` accepted a `padding` argument but never forwarded it to `AvgPool2d` (nor into the grad-output shape), so its `padding=1` / `padding=2` cases actually ran with `padding=0`. This PR forwards `padding` so those cases exercise the path, fixes the one call whose `padding=2, kernel_size=3` combination is invalid once padding is applied (`pad` must be `<= kernel_size / 2`), and adds a `count_include_pad=True` padded case. The repaired test fails on the current kernel and passes with the fix. ## Checklist - [x] Passes lint (`spin fixlint`) - [x] Added/updated tests - [ ] Updated documentation (if applicable) - [ ] Included benchmark results (for PRs impacting perf) ## BC-breaking? No. This corrects silently-wrong gradients; only incorrect outputs change. Pull Request resolved: pytorch#188345 Approved by: https://github.com/eqy
pytorch#188941) Extends the LP64 `Scalar(long long)` constructor guard in `c10/core/Scalar.h` from Linux-only to NetBSD, FreeBSD, and OpenBSD, where `int64_t` is also `long` on 64-bit (see the issue for the failing build). The existing `static_assert` is generalized with the platforms, so a wrong-model platform still fails loudly at compile time. Fixes pytorch#188911 Pull Request resolved: pytorch#188941 Approved by: https://github.com/malfet
…mpatibility (pytorch#188769) Similar with PR pytorch#178143 Pull Request resolved: pytorch#188769 Approved by: https://github.com/huydhn
…upport (pytorch#186595) Currently `torch.combinations` causes a CUDA sync due to `masked_select` being data dependent. This PR changes the implementation to statically compute the number of combinations and uses `nonzero_static` to compute the indices with a statically known shape. I did some quick benchmarks on CPU and GPU with the following script. Performance on CPU and GPU is much better, especially for large inputs. ~~GPU memory usage in eager mode is slightly larger, but significantly lower when compiled since we don't have a graph break anymore.~~fixed below <details><summary>Full Benchmarking script</summary> <p> ```python import torch import time import pandas as pd def run_benchmark(device, n_sizes, r_values, num_runs=5, warmups=2, compiled=False): results = [] def run_comb(x, r): return torch.combinations(x, r=r) func = torch.compile(run_comb) if compiled else run_comb for n in n_sizes: for r in r_values: x = torch.randn(n, device=device) for _ in range(warmups): _ = func(x, r=r) memory_used = None if device.startswith("cuda"): torch.cuda.reset_peak_memory_stats(device) _ = func(x, r=r) torch.cuda.synchronize(device) memory_used = torch.cuda.max_memory_allocated(device) / (1024 * 1024) # convert bytes to MB start_time = time.perf_counter() for _ in range(num_runs): _ = func(x, r=r) if device.startswith("cuda"): torch.cuda.synchronize(device) end_time = time.perf_counter() execution_time = (end_time - start_time) / num_runs * 1000 # ms device_str = f"{device.upper()} (compiled)" if compiled else f"{device.upper()} (eager)" results.append({ "Device": device_str, "N": n, "R": r, "Num Combinations": torch.combinations(x, r=r).shape[0], "Execution Time (ms)": f"{execution_time:.3f}", "Memory Usage (MB)": f"{memory_used:.3f}" if memory_used is not None else "N/A", }) return results if __name__ == "__main__": n_sizes = [100, 200, 300] print("Benchmarking on CPU (eager)...") cpu_results = run_benchmark("cpu", n_sizes, [2, 3], compiled=False) print("Benchmarking on CPU (compiled)...") cpu_compiled_results = run_benchmark("cpu", n_sizes, [2, 3], compiled=True) cuda_results = [] cuda_compiled_results = [] if torch.cuda.is_available(): print("Benchmarking on CUDA (eager)...") cuda_results = run_benchmark("cuda", n_sizes, [2, 3, 4], compiled=False) print("Benchmarking on CUDA (compiled)...") cuda_compiled_results = run_benchmark("cuda", n_sizes, [2, 3, 4], compiled=True) all_results = cpu_results + cpu_compiled_results + cuda_results + cuda_compiled_results df = pd.DataFrame(all_results) df = df.sort_values(by=["Device", "R", "N"]) print("\nBenchmark Results:") print(df.to_string(index=False)) ``` </p> </details> **main:** ``` Device N R Num Combinations Execution Time (ms) Memory Usage (MB) CPU (compiled) 100 2 4950 0.096 N/A CPU (compiled) 200 2 19900 0.229 N/A CPU (compiled) 300 2 44850 0.384 N/A CPU (compiled) 100 3 161700 2.692 N/A CPU (compiled) 200 3 1313400 110.978 N/A CPU (compiled) 300 3 4455100 392.575 N/A CPU (eager) 100 2 4950 0.074 N/A CPU (eager) 200 2 19900 0.229 N/A CPU (eager) 300 2 44850 0.338 N/A CPU (eager) 100 3 161700 13.006 N/A CPU (eager) 200 3 1313400 109.279 N/A CPU (eager) 300 3 4455100 387.542 N/A CUDA (compiled) 100 2 4950 0.206 0.162 CUDA (compiled) 200 2 19900 0.199 0.647 CUDA (compiled) 300 2 44850 0.205 1.457 CUDA (compiled) 100 3 161700 0.274 8.358 CUDA (compiled) 200 3 1313400 0.340 67.753 CUDA (compiled) 300 3 4455100 0.835 229.690 CUDA (compiled) 100 4 3921225 2.845 335.668 CUDA (compiled) 200 4 64684950 55.691 5474.059 CUDA (compiled) 300 4 330791175 282.391 27914.657 CUDA (eager) 100 2 4950 0.180 0.162 CUDA (eager) 200 2 19900 0.167 0.647 CUDA (eager) 300 2 44850 0.174 1.457 CUDA (eager) 100 3 161700 0.244 8.358 CUDA (eager) 200 3 1313400 0.316 67.753 CUDA (eager) 300 3 4455100 0.818 229.690 CUDA (eager) 100 4 3921225 2.821 335.668 CUDA (eager) 200 4 64684950 55.851 5474.059 CUDA (eager) 300 4 330791175 283.848 27914.657 ``` **This PR:** (see better benchmarks below) ``` Device N R Num Combinations Execution Time (ms) Memory Usage (MB) CPU (compiled) 100 2 4950 0.131 N/A CPU (compiled) 200 2 19900 0.276 N/A CPU (compiled) 300 2 44850 0.461 N/A CPU (compiled) 100 3 161700 1.839 N/A CPU (compiled) 200 3 1313400 24.267 N/A CPU (compiled) 300 3 4455100 84.516 N/A CPU (eager) 100 2 4950 0.073 N/A CPU (eager) 200 2 19900 0.216 N/A CPU (eager) 300 2 44850 0.412 N/A CPU (eager) 100 3 161700 3.078 N/A CPU (eager) 200 3 1313400 24.326 N/A CPU (eager) 300 3 4455100 87.258 N/A CUDA (compiled) 100 2 4950 0.105 0.152 CUDA (compiled) 200 2 19900 0.100 0.608 CUDA (compiled) 300 2 44850 0.100 1.371 CUDA (compiled) 100 3 161700 0.109 7.403 CUDA (compiled) 200 3 1313400 0.149 60.124 CUDA (compiled) 300 3 4455100 0.351 203.940 CUDA (compiled) 100 4 3921225 1.510 275.838 CUDA (compiled) 200 4 64684950 24.367 4487.048 CUDA (compiled) 300 4 330791175 138.127 22867.187 CUDA (eager) 100 2 4950 0.111 0.228 CUDA (eager) 200 2 19900 0.100 0.912 CUDA (eager) 300 2 44850 0.121 2.055 CUDA (eager) 100 3 161700 0.121 11.104 CUDA (eager) 200 3 1313400 0.137 90.186 CUDA (eager) 300 3 4455100 0.465 305.910 CUDA (eager) 100 4 3921225 1.599 359.667 CUDA (eager) 200 4 64684950 28.038 5922.086 CUDA (eager) 300 4 330791175 199.538 30284.839 ``` /cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo @ezyang @ngimel Fixes pytorch#74324 Pull Request resolved: pytorch#186595 Approved by: https://github.com/ngimel, https://github.com/jansel
Dynamo validates out= tensors after fake tensor propagation because out variants mutate their output tensors and resizing or arbitrary view strides are not safe to carry through functionalization. That check only accepted row-major contiguous layout, so an output from empty_like on a channels-last input was rejected even though it is dense and memory-format contiguous. This caused torch.compile(fullgraph=True) to graph break on valid code such as torch.add(x, x, out=out) when out has channels-last strides. Fix this by centralizing the out= layout predicate and allowing the standard dense memory formats Dynamo can preserve: row-major contiguous, channels_last, and channels_last_3d. The check still rejects arbitrary non-contiguous view strides. The alternative would be special-casing torch.add or empty_like, but the bug is in Dynamo's generic out= layout classification, so the generic predicate is the narrower root-cause fix. Fixes pytorch#164555 Generated by my agent Test Plan: python test/dynamo/test_repros.py -k test_out_overload_channels_last python test/dynamo/test_repros.py -k out_overload PYTHONDONTWRITEBYTECODE=1 python test/dynamo/test_repros.py ReproTests.test_out_overload_channels_last -v lintrunner -a Pull Request resolved: pytorch#185089 Approved by: https://github.com/ColinPeppler
Dynamo includes CPython's private opcode metadata headers so it can read the opcode cache table used by the PEP 523 frame-evaluation path. CPython 3.13 moved this data into pycore_opcode_metadata.h, and defining NEED_OPCODE_METADATA causes that internal header to emit several external _PyOpcode_* definitions. PyTorch already renamed the opcode cache/deopt table symbols locally, but the newer metadata symbols were still emitted with their CPython names. Static hermetic builds that link PyTorch together with CPython can then see duplicate definitions from libpython and the PyTorch object. Rename every CPython 3.13/3.14 opcode metadata definition that PyTorch pulls in to a _torch_PyOpcode_* name before including the private header, then undef the macros after inclusion. This keeps the local data available to initialize THP_PyOpcode_Caches without exporting duplicate CPython-owned definitions. A larger removal of private CPython headers would require replacing Dynamo's current PEP 523 integration, so this patch keeps the existing integration and fixes the concrete ODR violation at the include boundary. Fixes pytorch#163066 Generated by my agent Test Plan: - cc -std=c11 -fPIC -I. $(python3-config --includes) -c torch/csrc/dynamo/cpython_defs.c -o /tmp/cpython_defs.o && nm --defined-only /tmp/cpython_defs.o | rg '_PyOpcode|_torch_PyOpcode|THP_PyOpcode' - Python source check against CPython v3.13.11 and v3.14.0 pycore_opcode_metadata.h confirmed all emitted _PyOpcode_* definitions have local renames - lintrunner -a - git diff --check Pull Request resolved: pytorch#185118 Approved by: https://github.com/frgossen
FakeTensor's convolution implementation validated dtype mismatches but did not validate that the convolution input, weight, and bias were on the same device. If a module was compiled on CPU and its parameters were later moved to CUDA in place, Dynamo could fake-propagate and Inductor could lower an invalid mixed-device convolution instead of reporting the same device error eager raises. With the C++ wrapper path, that invalid graph could reach a CPU convolution shim with CUDA weights and fail in generated code. Add a same-device check to the FakeTensor convolution implementation so invalid convolution inputs fail during fake propagation, before Inductor lowers or executes a bad graph. I considered adding generated-wrapper device assertions, but that only checks the already-bad compile-time devices and does not address the root cause. Fixes pytorch#160399 Generated by my agent Test Plan: - python test/test_fake_tensor.py -k test_conv_rejects_mismatched_fake_devices - python test/inductor/test_compile.py -k test_compiled_conv_rejects_mixed_input_weight_devices - TORCHINDUCTOR_CPP_WRAPPER=1 python - <<'PY' ... issue-style ModuleList Conv2d repro with 16x16 input ... PY - TORCHINDUCTOR_CPP_WRAPPER=1 python - <<'PY' ... single Conv2d mixed CPU input/CUDA weight repro ... PY - git diff --check - lintrunner -a Pull Request resolved: pytorch#185324 Approved by: https://github.com/mlazos
Non-strict export patches a small set of Python builtins while tracing fake tensors so Python operators preserve symbolic shape expressions. It handled max, min, and math.pow, but left builtins.len untouched. CPython len(tensor) therefore asked the fake tensor's symbolic leading dimension for __index__ and materialized the example value as an int, which introduced an equality guard such as Eq(batch, 3). For pytree inputs that are also nn.Module subclasses, this made valid dynamic shape constraints look violated. Route tensor inputs to Tensor.__len__() inside the existing non-strict builtin override context, while preserving the original len behavior for non-tensor objects. Tensor.__len__() returns shape[0], so fake tensors keep the SymInt rather than forcing specialization. Changing Tensor.__len__ itself would not fix this because CPython's len protocol still converts the returned value to a Py_ssize_t. The regression test models the reported DynamicCache shape with a registered pytree class inheriting nn.Module, exports it with dynamic batch/cache length, and runs the exported module with different dimensions. Benchmark Results: Measured a non-strict export microbenchmark with dynamic batch, 20 Python len(list) calls in forward, 3 warmup iterations, and 20 measured iterations. main: median 0.027823s, mean 0.028003s. fix-147326: median 0.031010s, mean 0.031162s. This is export-time overhead only; exported model runtime is unaffected. Test Plan: python test/export/test_export.py TestExport.test_non_strict_export_len_tensor_pytree_nn_module_input python - <<'PY' ... PY # adapted issue reproducer lintrunner -a git diff --cached --check Fixes pytorch#147326 Generated by my agent Pull Request resolved: pytorch#185804 Approved by: https://github.com/frgossen
… XPU profiler support (pytorch#188254) Running --export-profiler-trace --profile-details crashes with: ``` python benchmarks/dynamo/timm_models.py --float16 --training --performance -d xpu -n 10 --cold-start-latency --timeout 10800 --disable-cudagraphs --export-profiler-trace --profile-details --only ese_vovnet19b_dw --backend=eager loading model: 0it [00:06, ?it/s] xpu train ese_vovnet19b_dw running benchmark: 100%|██████████████████████████████████████████| 10/10 [00:04<00:00, 2.39it/s] USDT:2026-06-26 15:55:48 2395:2395 ActivityProfilerController.cpp:415] profiler_start USDT:2026-06-26 15:55:48 2395:2395 ActivityProfilerController.cpp:455] profiler_stop Traceback (most recent call last): File "/data/pytorch/benchmarks/dynamo/timm_models.py", line 384, in <module> timm_main() File "/data/pytorch/benchmarks/dynamo/timm_models.py", line 380, in timm_main main(TimmRunner()) File "/data/pytorch/benchmarks/dynamo/common.py", line 4026, in main process_entry(0, runner, original_dir, args) File "/data/pytorch/benchmarks/dynamo/common.py", line 3948, in process_entry result = run(runner, args, original_dir) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/data/pytorch/benchmarks/dynamo/common.py", line 4905, in run runner.run_one_model( File "/data/pytorch/benchmarks/dynamo/common.py", line 3226, in run_one_model status = self.run_performance_test( ^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/data/pytorch/benchmarks/dynamo/common.py", line 3128, in run_performance_test experiment( File "/data/pytorch/benchmarks/dynamo/common.py", line 1228, in speedup_experiment write_outputs( File "/data/pytorch/benchmarks/dynamo/common.py", line 341, in write_outputs output_json(filename, headers, row) File "/data/pytorch/benchmarks/dynamo/common.py", line 426, in output_json print(json.dumps(record), file=f) ^^^^^^^^^^^^^^^^^^ File "/usr/lib/python3.12/json/__init__.py", line 231, in dumps return _default_encoder.encode(obj) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/usr/lib/python3.12/json/encoder.py", line 200, in encode chunks = self.iterencode(o, _one_shot=True) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/usr/lib/python3.12/json/encoder.py", line 258, in iterencode return _iterencode(o, 0) ^^^^^^^^^^^^^^^^^ File "/usr/lib/python3.12/json/encoder.py", line 180, in default raise TypeError(f'Object of type {o.__class__.__name__} ' TypeError: Object of type ProfilerActivity is not JSON serializable ``` Root cause: args.profile_details contains ProfilerActivity enum values for the profiler. output_signpost() serializes the full vars(args) to JSON, and ProfilerActivity enums are not JSON-serializable. Additionally, the profiler activities were hardcoded to CPU + CUDA, which is incorrect for XPU devices. Add "profile_details" to the exclusion list in output_signpost(), alongside the already-excluded "export_profiler_trace" and "profiler_trace_name". Dynamically select profiler activities based on args.devices instead of hardcoding CUDA, so --profile-details -d xpu correctly uses ProfilerActivity.XPU. Pull Request resolved: pytorch#188254 Approved by: https://github.com/jansel
…89214) The CUDA-13 toolkit runfile ships an older CUPTI than the standalone redist archive. Add `install_cupti_headers` helper that downloads the `cuda_cupti` redist tarball and stages its headers into /usr/local/cupti-headers-<major.minor>, and invoke it from the CUDA 13.0 and 13.2 installers. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Pull Request resolved: pytorch#189214 Approved by: https://github.com/dolpm
…ag (pytorch#189036) The `inductor_meta` dict is the kernel-config / heuristics metadata bag that Inductor's codegen produces and the runtime autotuning machinery consumes. It carried a `dict[str, Any]` (and a local `dict[str, object]` alias) annotation everywhere, which gave no key- or value-level type safety across a ~60-key, fixed-shape interface. This PR introduces `InductorMeta(TypedDict, total=False)` in `torch/_inductor/runtime/hints.py` and threads it through the runtime-side consumers, leaving the codegen-side producers (codegen/*) as a follow-up so the blast radius stays inside `runtime/`. Suggested review order: 1. `hints.py` -- the `InductorMeta` TypedDict definition and the rationale for why the dynamically keyed nested bag `combo_grid_meta` stays `dict[str, Any]` (its keys are computed at runtime, e.g. `xnumel_{i}`). 2. `triton_heuristics.py` -- the consumer/producer signatures switch from `dict[str, Any]` to `InductorMeta`; the three `cast(...)` sites handle the dict-comprehension merge (`inductor_meta_i`), the `.pop(...)` boundary (`mutated_arg_names`), and are the only escapes from precise typing. 3. `autotune_cache.py` -- the `_InductorMetaTy` alias is repointed at `InductorMeta`, converting all 11 signatures in one line. 4. `coordinate_descent_tuner.py` -- `CoordescTuner.__init__` and the stored `self.inductor_meta` attribute get typed. Two small, behavior-preserving refactors were needed so the now-typed `.get(...)` calls narrow correctly: `persistent_reduction` hoists the `RSPLIT_SIZE` lookup so the truthy guard narrows `int | None -> int` before `//`, and `foreach` gains the same `{} if inductor_meta is None` default guard its sibling entrypoints already use (which also removes a latent `NoneType.get` on the `inductor_meta=None` default path). Test Plan: Type-only change; validated with pyrefly (net -3 errors, 0 newly introduced -- the two reshaped messages are pre-existing errors whose value-type string changed, verified same file:line and count): ``` pyrefly check \ torch/_inductor/runtime/triton_heuristics.py \ torch/_inductor/runtime/autotune_cache.py \ torch/_inductor/runtime/coordinate_descent_tuner.py \ torch/_inductor/runtime/hints.py ``` ``` lintrunner --take PYFMT,RUFF,PYREFLY,CODESPELL,FLAKE8 -a \ torch/_inductor/runtime/hints.py \ torch/_inductor/runtime/triton_heuristics.py \ torch/_inductor/runtime/autotune_cache.py \ torch/_inductor/runtime/coordinate_descent_tuner.py ``` This PR was authored with the assistance of an AI coding assistant. Pull Request resolved: pytorch#189036 Approved by: https://github.com/jansel
Migrates LU solve to metal kernels, should be merged after pytorch#187038 Also fixes pytorch#189134 `torch.linalg.solve` | shape | before MPS (ms) | after MPS (ms) | speedup | |---|--:|--:|--:| | (128x128), 128 rhs | 1.580 | 0.485 | **3.26x** | | (256x256), 256 rhs | 2.998 | 0.825 | **3.63x** | | (512x512), 512 rhs | 6.094 | 1.569 | **3.88x** | | (1024x1024), 1024 rhs | 13.802 | 4.051 | **3.41x** | | (2048x2048), 2048 rhs | 40.054 | 15.526 | **2.58x** | | (512x512), 1 rhs | 8.319 | 1.507 | **5.52x** | | (1024x1024), 1 rhs | 24.212 | 3.400 | **7.12x** | | (2048x2048), 1 rhs | 81.753 | 9.663 | **8.46x** | | 64x(64x64), 1 rhs | 51.594 | 0.400 | **128.99x** | | 32x(128x128), 1 rhs | 47.603 | 0.604 | **78.80x** | | 8x(256x256), 8 rhs | 22.417 | 0.956 | **23.45x** | | 4x(512x512), 16 rhs | 23.987 | 1.852 | **12.95x** | Pull Request resolved: pytorch#189200 Approved by: https://github.com/malfet
…o tensor descriptor tests (pytorch#188794) - Migrate triton cpu pin to 910ba2e98deb5180a5fb8f204f57cd6e71d2436a since it has not been updated in over a year. - Block ptrs are deprecated in triton, so add tensor descriptor support for cpu - Add triton cpu early returns to TMACompatibilityChcker since tensor descriptor requirements are less strict for the cpu in comparison with cuda - Migrate block ptrs tests -> tensor descriptor tests Pull Request resolved: pytorch#188794 Approved by: https://github.com/jansel
…rch#176911) This PR addresses and fixes pytorch#176910. I've arranged this PR into 4 commits to make it easier to reproduce the issues and verify that they have been resolved: 1. Adds a test that the loss is set to zero correctly, which fails. 2. Patches the CTC loss to fix the above test. 3. Adds a test that the gradients are set to zero correctly, which fails. 4. Patches the CTC loss backwards function to fix the above test. Both tests employ the private function `torch._use_cudnn_ctc_loss` to ensure the cuDNN backend is used. The second test additionally uses `torch._cudnn_ctc_loss` with `deterministic=False`. The reason for this is that I was unable to find examples that genuinely produced infinite gradients from cuDNN's deterministic implementation (which is the only one available through the public API), even when `zero_infinity=False` (likely due to this backend's unusual behavior, as described in pytorch#176910). Please let me know if this approach should be modified. Pull Request resolved: pytorch#176911 Approved by: https://github.com/eqy, https://github.com/ezyang Co-authored-by: Edward Z. Yang via mergedog <ezyang@meta.com>
This is a follow-up to pytorch#169832, and it's meant to be merged after that one. Pull Request resolved: pytorch#172813 Approved by: https://github.com/aorenste, https://github.com/ezyang
…ytorch#188267)" (pytorch#189105)" This reverts commit 7766589. Reverted pytorch#189105 on behalf of https://github.com/atalman due to Still has internal failures, after the revert this needs to be reimported and landed internally. This requires more internal changes ([comment](pytorch#189105 (comment)))
The test spawns a child process that creates a SubprocPool, runs a job in a
worker, and asserts the sidecar shuts down promptly while that worker is still
busy. Its runtime is dominated by heavyweight process cold-starts: in fbcode
each of the child, the sidecar, and its forked workers is a full binary that
imports torch and spins up the service framework. That baseline (~18-20s) sat
right at the outer subprocess.run(timeout=20), so any host contention tipped it
over; the flaky failures were the outer timeout firing, with pass and fail
durations both clustered at the 20s boundary.
The root cause is that process/worker startup cost was being charged against the
shutdown budget. The fix separates untimed setup from the timed region:
- Submit a job that creates a sentinel file once a worker is actually executing
it, then blocks (new _test_signal_then_sleep helper). A readiness barrier
waits (generous timeout) for that sentinel before the shutdown is initiated,
so cold-start cost is no longer in the shutdown budget. This also makes the
test reliably exercise the terminate-a-running-worker path instead of the
queued cancel_futures path that the old fixed 0.5s sleep could race into.
- Keep the shutdown wait bounded (30s, well under the 120s the job sleeps) so a
regression that waits for the in-flight job instead of terminating it is still
detected quickly.
- Raise the outer subprocess timeout to 120s as a pure hang backstop.
Fixes T278749262.
Test Plan:
Reproduced the flake and validated the fix under stress in the mode CI uses
(opt); before the fix, runs clustered at the 20s timeout and flaked, after the
fix all runs pass:
```
buck2 test @fbcode//mode/opt fbcode//caffe2/test/inductor:compile_worker \
-- 'test_shutdown_terminates_sidecar_worker_pool' --stress-runs 25
# Pass 50. Fail 0.
```
```
lintrunner -a test/inductor/test_compile_worker.py \
torch/_inductor/compile_worker/subproc_pool.py
# ok No lint issues.
```
Authored with Claude.
Pull Request resolved: pytorch#189166
Approved by: https://github.com/atalman
ghstack dependencies: pytorch#188241
Custom operators whose kernels are implemented in Python currently need to make 1+ roundtrips into the C++ PyTorch dispatcher. These roundtrips are expensive, e.g. a no-op custom operator with 10 tensor inputs and 3 tensor outputs with one python kernel implemented takes 2.6us. The main expensive part appears to be Tensor and pyobj refcount bump. The Tensor owns a PyObject, and converting a Tensor to an IValue copies the Tensor. This PR introduces a new type of dispatching, the "PyObject Dispatch", where we avoid converting Tensors to IValues and instead perform the dispatcher algorithm on PyObjects in the Python-C API. When we compute a DispatchKey for an operator: - if the kernel at the DispatchKey is a Python kernel, then we invoke it directly with the PyObjects - if the kernel at the DispatchKey is a C++ kernel, then we call op.redispatch using the regular C++ Dispatcher to invoke the kernel. This incurs the expensive thing of copying the Tensors. We add an option to enable PyObject Dispatch on all OpOverloads, and we also make torch.library.custom_op enable PyObject Dispatch by default as a guinea pig. NB: PyObject Dispatch is only profitable if all of your kernel registrations are in Python, *and* you are not running into any C++ backend fallbacks. This takes the 10 tensor input 3 tensor output noop test case from 2.6us down to 300ns. Authored with assistance from an AI Agent. Pull Request resolved: pytorch#187949 Approved by: https://github.com/albanD
…them (pytorch#175796) Fixes pytorch#165758 ## Description `output_names` in `torch.onnx.export` only assigns labels to output nodes. It does not control their order. When a model returns a dictionary, outputs are flattened in the dictionary's iteration order regardless of what names are specified. This was not documented, leading to user confusion (see linked issue). This PR adds a clarifying note to the `output_names` parameter docstring in both the dynamo-based exporter (`torch/onnx/__init__.py`) and the legacy TorchScript exporter (`torch/onnx/_internal/torchscript_exporter/utils.py`). Pull Request resolved: pytorch#175796 Approved by: https://github.com/justinchuby
…torch#188907) (pytorch#188907) Summary: During GEMM template autotuning, CachingAutotuner._make_launchers keeps the last failed config's build exception in a local (`exc`). That exception's __traceback__ retains its frame chain and, via those frames' f_back, the benchmarking callers up to Triton's do_bench -- whose 256MB L2-cache-flush buffer is then pinned. exc <-> traceback <-> frame is a reference cycle only cyclic GC can reclaim, so when the caller has disabled cyclic GC (as APS training loops do) one 256MB buffer leaks per autotuned kernel whose last-tried config fails (routine in max-autotune: large BLOCK configs exceed shared memory -> OutOfResources). On a 128xGB300 APS job this accumulated to ~44GB of transient HBM, ~63% of an 87GB peak high-water-mark. Fix: drop the retained exception in a `finally` in _make_launchers (it is only used for its type/message). This breaks the cycle so the buffer frees by refcount without needing cyclic GC, and also releases the f_back chain up to do_bench. Regression from D107597017 (PR pytorch#184285); preserves that change's module-exhaustion fix. Authored with Claude. Test Plan: Unit test (fails before this change, passes after): buck test fbcode//caffe2/test/inductor:triton_heuristics -- \ TestMakeLaunchersMemory.test_failed_config_exception_not_retained Repro: 24 distinct-shaped GEMMs, TritonBenchmarker, cold cache, gc.disable(), on the job's package revision. Peak co-resident 256MB L2-flush buffers (get_empty_cache_for_benchmark) 144 -> 1; peak allocated during autotuning 36.84GB -> 0.63GB. Identical benchmarking otherwise (582 do_bench allocs). Reviewed By: williamwen42 Differential Revision: D110557608 Pull Request resolved: pytorch#188907 Approved by: https://github.com/williamwen42
…#184279)" This reverts commit 6f6bfce. Reverted pytorch#184279 on behalf of https://github.com/atalman due to failing internally see: D111032481 ([comment](pytorch#184279 (comment)))
Fixes pytorch#189243 Pull Request resolved: pytorch#189252 Approved by: https://github.com/Skylion007, https://github.com/malfet
…RSION' (pytorch#188072) # Context Recreated pytorch#184552 from personal fork. hipCUB is targeting to be compatible with CCCL 3.0 (ROCm/rocm-libraries#4079). This requires updating some of the CUB and hipCUB compatibility glue. To help this transition hipCUB will be exposing `HIPCUB_CCCL_VERSION` in one of the following ROCm releases (>=7.14) (ROCm/rocm-libraries#7632), which will be bumped once hipCUB reaches parity with CCCL 3.0 and beyond. This PR is created in anticipation for this future changes. # Implementation - Introduce dependency on [`libhipcxx`](https://github.com/ROCm/libhipcxx). - Guard `FpLimits` for `c10::BFloat16` backport to pre CCCL 3.x. Numeric limits are handled by libhipcxx. - Extend `CUB_VERSION` derivation for hipCUB. Keep old fallback for backwards compatibility. Only enable CUB V3 API once hipCUB reaches parity with CCCL 3.x. Pull Request resolved: pytorch#188072 Approved by: https://github.com/jeffdaily
…mpile` support (pytorch#186595)" This reverts commit 5b84db3. Reverted pytorch#186595 on behalf of https://github.com/meta-codesync due to Diff reverted internally ([comment](pytorch#186595 (comment)))
Removes `.ci/manywheel/test_wheel.sh`, a dead script from the old Jenkins-era binary wheel test flow. It installs a Miniconda into `/py` and pulls in `conda-build` / `anaconda-client`, but has no remaining callers: - Nothing in `pytorch/pytorch` or `pytorch/test-infra` references `test_wheel.sh` (verified by direct grep of both repos, including their `.github/` workflows and composite actions). - The `/py` environment it creates is used by no other script. - Linux binary testing now runs through `.ci/pytorch/binary_linux_test.sh` (`_binary-test-linux.yml`). The file has been untouched since the 2024 bulk move in pytorch#138103 and still points at the deprecated `repo.continuum.io` installer and a `/remote/token` Anaconda credential path. This is the first increment of the conda-removal effort; it carries no Docker image rebuild or image dependency. ## Test plan Verified no references remain and the live replacement is in place: ``` # pytorch/pytorch @ viable/strict grep -rn "test_wheel" .github/ # no matches grep -rn "test_wheel" . # only the file itself grep -rn -- '-p /py|ls /py|/py/bin' . # /py used only by this script # pytorch/test-infra @ origin/main git grep -n "test_wheel" origin/main # no matches # live replacement exists git ls-files .ci/pytorch/binary_linux_test.sh \ .github/workflows/_binary-test-linux.yml ``` This PR was authored with the assistance of Claude Code. Pull Request resolved: pytorch#189159 Approved by: https://github.com/atalman
This reverts commit c0bba62. Reverted pytorch#186790 on behalf of https://github.com/atalman due to hi @cyyever please reland it on Jul 14, this change triggered internal code freeze which should be lifted on this date ([comment](pytorch#186790 (comment)))
…189241) TPU tests run on the linux.google.tpuv7x.1 runners and can't use the OSDC/ARC test path. They used to run via the EC2 test job in _linux-test.yml (the pallas TPU test never set use-arc, so it took the `!use-arc` branch). Removing that EC2 path (pytorch#189117) left the TPU test gated behind `inputs.use-arc`, so it was silently skipped; then removing the use-arc gate (pytorch#189171) made it run on the OSDC-shaped `test` job (container: with an ARC image + --gpus all), which is wrong for TPU and fails. Extract the TPU test path into its own _tpu-test.yml (mirroring _s390x-test.yml): setup-linux, the TPU docker flags / TORCH_TPU / install_torch_tpu.sh path from the old EC2 test job, on the host docker run/exec model. The GPU-, b200-, s390x-, and container-runner-only steps are dropped. inductor-pallas.yml's TPU test job now calls it; the TPU build is unchanged (it legitimately builds on OSDC). Authored with the assistance of Claude Code. Test Plan: exercised by the inductor-pallas workflow (ciflow/inductor) running the pallas-tpu-py3.12-inductor test on linux.google.tpuv7x.1. Pull Request resolved: pytorch#189241 Approved by: https://github.com/atalman
…ch#177037) This PR adds a cublasLt backend to grouped GEMM. The cublasLt backend is default for fp16 on Blackwell for CUDA 13.2+ and Hopper for CUDA 13.3+, and can also be enabled for bf16 on those device/version combos with `torch.backends.cuda.matmul.prefer_cublaslt_grouped_gemm = True`. The cublasLt grouped GEMM path is fully compatible with `torch.compile` and CUDA Graphs. There are some constraints (see [docs](https://docs.nvidia.com/cuda/cublas/index.html#id106)) but these are mostly handed internally- the only one that users need to be aware of is that matrices and leading dimensions must be 16-byte aligned, so matrices with some dimensions not divisible by 16 may require padding + slicing. Here's an initial [benchmarking results spreadsheet](https://docs.google.com/spreadsheets/d/1Av3RwwmLERLKlZ2826rPAduAribjF8KQLKTHxn5Wb2Y/edit?usp=sharing). Note the two tabs. For uniform groups, the cublasLt backend is faster than the fallback (fp16) but slower than the CUTLASS kernel (bf16) except for the largest size tested. For MOE-like cases where groups are not uniform, cublasLt appears to be on average faster across the board. [Benchmarking script](https://gist.github.com/gderossi/515581a165b9b8eb031ba7a1a5a01c55) Scaled grouped GEMM for FP8 types is also working with all currently supported scaling types (tensorwise, per-group tensorwise, MXFP8) at pytorch#177038, to be stacked on top of this. Pull Request resolved: pytorch#177037 Approved by: https://github.com/eqy, https://github.com/drisspg, https://github.com/ngimel
…pytorch#188842) When a job's failure streak start is pinned within the window (not clipped), surface the suspected culprit: the commit/PR the streak began at. The HUD already returns per-commit time, PR number, and title in each row, so no extra API calls are needed -- build a sha->commit map from the pages already fetched and resolve the failing-since commit. The text report prints the date, PR URL, and title under "failing since"; the CSV gains failing_since_date, failing_since_pr_url, and failing_since_pr_title columns. These are populated only when the start is pinned; clipped streaks (true start unknown, possibly older than the window) leave them blank. Test Plan: HUD_INTERNAL_BOT_TOKEN=<token> python scripts/hud/analyze_failing_jobs.py --commits 500 --csv > out.csv Confirmed pinned rows carry the culprit PR date/url/title (e.g. pytorch#188107, dated 2026-06-28) while clipped rows leave those columns blank. lintrunner scripts/hud/analyze_failing_jobs.py Authored with Claude. Pull Request resolved: pytorch#188842 Approved by: https://github.com/zou3519
…and HOP tests (pytorch#181778) Pull Request resolved: pytorch#181778 Approved by: https://github.com/aorenste, https://github.com/mlazos ghstack dependencies: pytorch#181777
…d tests (pytorch#181779) Pull Request resolved: pytorch#181779 Approved by: https://github.com/angelayi, https://github.com/mlazos ghstack dependencies: pytorch#181777, pytorch#181778
…oad_rocm_deps`) (pytorch#188454) ## Summary This PR makes PyTorch ROCm/TheRock wheels **self-contained**: it fixes the two `import torch` failures that occur on the wheel (`_rocm_sdk_*` site-packages) layout, so `import torch` works without `LD_LIBRARY_PATH` or a sourced ROCm env (e.g. the ROCm distributed preflight `python .ci/pytorch/rocm_preflight.py`, which never sources `/etc/rocm_env.sh`). Both fixes live in torch **source**, so they travel with any build (including a bare `python -m build` in a TheRock env), not just CI-produced wheels. This is the **minimal, consolidated PR** carrying **both** import-discovery fixes — a **two-file** diff (`torch/__init__.py` + `torch/cuda/__init__.py`). Both changes are **no-ops on the traditional tarball `/opt/rocm` install**: each early-returns when the wheel package's `_rocm_sdk_core` is absent, so non-wheel ROCm CI (apt/`/opt/rocm`, whose libs are already on the loader path), CUDA, and CPU builds are unaffected. ## Fixes ### 1. `torch/__init__.py` — `_preload_rocm_deps()` Fixes `ImportError: librocm_sysdeps_liblzma.so.5: cannot open shared object file` at `from torch._C import *` - Adds `_preload_rocm_deps()`, the ROCm analogue of `_preload_cuda_deps()` (which handles the `nvidia-*-cu12` wheels). At import time, **before** `import torch._C`, it locates the `rocm` pip package's `_rocm_sdk_core` via `importlib.util.find_spec` (the same anchor `torch.utils.cpp_extension._find_rocm_home()` uses) and `ctypes.CDLL`-preloads (`RTLD_GLOBAL`) the vendored sysdeps (`librocm_sysdeps_*`) and core ROCm runtime libs from `_rocm_sdk_core/lib` and `_rocm_sdk_core/lib/rocm_sysdeps/lib`. Wired into `_load_global_deps()`. - Guarded to no-op on non-Linux, on non-ROCm builds (`torch.version.hip is None` → CUDA/CPU), and on OS-managed ROCm where `_rocm_sdk_core` is absent (the apt/`/opt/rocm` path). ### 2. `torch/cuda/__init__.py` — amdsmi / `libamd_smi.so` discovery hook + non-fatal guard Fixes the *next* failure that surfaces once `_preload_rocm_deps()` gets `import torch` past the `librocm_sysdeps` error: `KeyError: 'libamd_smi.so'` raised from the `amdsmi` Python binding at `import amdsmi` (inside `_C._initExtension`) On TheRock wheels the amdsmi package installs the native library as a versioned soname under `_rocm_sdk_core/lib` (e.g. `libamd_smi.so.26`), which amdsmi's own `find_smi_library()` fails to locate when `$ROCM_HOME`/`$ROCM_PATH` and the loader path are unset; its wrapper then does `_libraries['libamd_smi.so']` → `KeyError`, which escaped the existing guard (only `ModuleNotFoundError` was tolerated) and killed `import torch`. Two-part torch-side hardening: - Make `_amdsmi_cdll_hook` **TheRock-aware**: locate `_rocm_sdk_core` via `importlib.util.find_spec` and add its `lib/libamd_smi.so*` (glob, incl. versioned sonames) as last-resort candidate paths; match names whose basename **starts with** `libamd_smi.so` so amdsmi's bare `CDLL("libamd_smi.so")` is redirected to the real versioned file. Existing `ROCM_HOME`/`ROCM_PATH` behavior preserved; guarded to Linux. - Make a non-discoverable amdsmi **non-fatal**: broaden the guard so `KeyError`/`OSError` degrade to `_HAS_AMDSMI=False` instead of aborting `import torch`, matching the existing missing-module path. ## Consolidation note The amdsmi discovery fix (#2 above) originally shipped as standalone PR pytorch#188724. That PR has been **closed** and folded into this PR so **both** import-discovery fixes travel together in a single minimal two-file change. Scope note: the TheRock tarball→wheels migration and the earlier CI-only RPATH/ldconfig workarounds are **not** part of this PR. The wheels migration is tracked separately by pytorch#188429, which rebases on top of this change. ## Validation (TheRock wheel path) — PASSED Validated end-to-end by **temporarily** stacking Ethan's wheels-migration PR pytorch#188429 (which adds the `rocm-preview` wheel CI) on top of the two source fixes and running `ciflow/rocm-preview` (run [28581558341](https://github.com/pytorch/pytorch/actions/runs/28581558341)) on the stacked head. The temporary stack was then **removed**, so this PR remains the minimal two-file source diff. - **`build-osdc` (TheRock wheel-path build): SUCCESS** — [job 84742887851](https://github.com/pytorch/pytorch/actions/runs/28581558341/job/84742887851). - **`import torch` works on the `_rocm_sdk_*` wheel layout** — proven by the passing shard [`test (default, 5/6)` job 84772673709](https://github.com/pytorch/pytorch/actions/runs/28581558341/job/84772673709), which imports torch and prints the HIP 7.14 / MIOpen 3.5.2 build config. - **No** `ImportError: librocm_sysdeps*` (incl. `librocm_sysdeps_liblzma`) and **no** `KeyError: 'libamd_smi.so'` anywhere in the run → **both** fixes are validated on the **real TheRock wheel path**. - Remaining rocm-preview shard failures were **unrelated** to this PR — test-content failures (`test_CTCLoss_cudnn_cuda`, `test_cxx_flags` `RuntimeError map::at`, `test_hip_device_count`, `test_lazy_imports_are_lazy`) and infra (`ROCm pre-flight: no GPUs visible to the container` on the distributed shards). - Run: https://github.com/pytorch/pytorch/actions/runs/28581558341 ### rocm-preview distributed shards (register fat binary failed) Reference/evidence for the `register fat binary failed` (`code_object.cpp:1038`) investigation — the rocm-preview (TheRock wheel-path) distributed shards from run [28581558341](https://github.com/pytorch/pytorch/actions/runs/28581558341) (`linux-noble-rocm-preview-py3.12-gfx942`), all FAILED: - distributed 1/3: https://github.com/pytorch/pytorch/actions/runs/28581558341/job/84772673683 (`test (distributed, 1, 3, linux.rocm.gpu.gfx942.4, module:rocm, oncall:distributed)`) - distributed 2/3: https://github.com/pytorch/pytorch/actions/runs/28581558341/job/84772673678 (`test (distributed, 2, 3, linux.rocm.gpu.gfx942.4, module:rocm, oncall:distributed)`) - distributed 3/3: https://github.com/pytorch/pytorch/actions/runs/28581558341/job/84772673727 (`test (distributed, 3, 3, linux.rocm.gpu.gfx942.4, module:rocm, oncall:distributed)`) <details> <summary>Failure log the amdsmi fix addresses (pre-fix, run 28524549808)</summary> ``` /var/lib/jenkins/workspace/libamd_smi.so: cannot open shared object file: No such file or directory Unable to find libamd_smi.so library try installing amd-smi-lib from your package manager Traceback (most recent call last): File "/var/lib/jenkins/workspace/.ci/pytorch/rocm_preflight.py", line 21, in <module> import torch File ".../site-packages/torch/cuda/__init__.py", line 114, in <module> import amdsmi # type: ignore[import] File ".../site-packages/amdsmi/amdsmi_wrapper.py", line 218, in <module> amdsmi_free_name_value_pairs = _libraries['libamd_smi.so'].amdsmi_free_name_value_pairs KeyError: 'libamd_smi.so' ``` </details> ## Notes: amdsmi discovery + fat-binary logs - **Bare `CDLL("libamd_smi.so")`** comes from amdsmi's own `find_smi_library()` (unversioned candidate), not torch; torch's `_amdsmi_cdll_hook` intercepts it and retries the versioned `_rocm_sdk_core/lib/libamd_smi.so*`. - **Not just preload amdsmi via `_rocm_core_libs`:** amdsmi resolves by exact filename, so an RTLD_GLOBAL versioned `libamd_smi.so.26` won't satisfy it, and eager-preloading a non-`DT_NEEDED` lib reintroduces the ODR double-load the hook prevents — the CDLL-redirect hook is the correct fix. - **`E0702 … code_object.cpp:1038 register fat binary failed` is not caused by this PR:** it's emitted by the HIP/clr C++ runtime, appears ambiently (~4,774×/run across rocm-preview gfx942 jobs incl. the distributed shards above), and is non-fatal — a pre-existing TheRock rocm-preview build/toolchain matter, orthogonal to these two Python-only changes. - **Fix-plan pointer (fat-binary):** prove pre-existing on a PR-free `main` rocm-preview baseline → localize via `AMD_LOG_LEVEL=3/4` → scoped clr-log suppress/demote (escalate to clr/comgr pin or `PYTORCH_ROCM_ARCH` only on a real running-arch failure) → validate via `ciflow/rocm-preview`. Belongs in the build/CI layer, not torch source. Tracked in ROCm/frameworks-internal pytorch#17160. ## Test plan - [x] Validate on the TheRock wheel path via `ciflow/rocm-preview` (see **Validation** above). - [x] `ast.parse` both changed files (`torch/__init__.py`, `torch/cuda/__init__.py`). - [x] Confirm CUDA and CPU builds are unaffected (both fixes no-op when `torch.version.hip is None` / `_rocm_sdk_core` absent). Fixes pytorch#188561 Made with Cursor Pull Request resolved: pytorch#188454 Approved by: https://github.com/jeffdaily
…ky failures (pytorch#189196) load_template opened .py.jinja kernel templates without specifying an encoding, so open() used the process's locale-preferred encoding. On CI hosts whose locale is C/POSIX (ASCII), reading a template that contains a non-ASCII character (e.g. an em dash in a comment) raised UnicodeDecodeError and failed compilation. Because the fleet is locale-heterogeneous, the same test passed or failed depending on which host it landed on, and passed on retry -- i.e. it presented as flaky rather than as a deterministic failure. The failure surfaced through the FB TLX flex-attention template (tlx_flex_attention.py.jinja), imported lazily during flex_attention lowering, so it manifested as flaky flex_attention/flex_decoding tests. The OSS cutedsl_mm_grouped.py.jinja template contains the same latent issue and is fixed by the same change since all template loads funnel through load_template. The fix reads templates as UTF-8 explicitly, which is correct because these are repo-owned source files that are already valid UTF-8. It also wraps a decode failure in a ValueError that names the offending template file: a bare UnicodeDecodeError hides which template is malformed, which is what made the original failure hard to diagnose (the message only said "can't decode byte 0xe2 in position 16312", with no filename). Fixes T278775453, T278775455, T278775457 (and other flex_attention tests that flaked from the same root cause). Test Plan: ``` python test/inductor/test_utils.py TestLoadTemplate lintrunner -a torch/_inductor/utils.py test/inductor/test_utils.py ``` test_load_template_uses_utf8 verifies a valid-UTF-8 template loads under a simulated ASCII-locale host; test_load_template_invalid_utf8_names_the_file verifies a non-UTF-8 template raises an error that names the file. Authored with Claude. Pull Request resolved: pytorch#189196 Approved by: https://github.com/zou3519
## Summary Fix 'Pytorch' to 'PyTorch' brand name in 6 documentation files. ## Changes - docs/source/hub.md (3 instances) - docs/source/distributed.checkpoint.md (1 instance) - docs/source/notes/cuda.md (1 instance) - docs/source/notes/mkldnn.md (1 instance) - docs/source/user_guide/index.md (1 instance) - docs/source/user_guide/torch_compiler/torch.compiler_faq.md (1 instance) ## Context This PR addresses the maintainer's request to focus only on brand name fixes. The original PR pytorch#189023 included additional changes (DDP links, VS versions, import statements) which have been removed to keep this PR focused and minimal. Total: 8 insertions(+), 8 deletions(-) across 6 files. Pull Request resolved: pytorch#189248 Approved by: https://github.com/ezyang
…ling + preload Deduplicate the two ROCm shared-library lists that had drifted: .ci/manywheel/repair_wheel.py::ROCM_SO_FILES (27, wheel bundling) and torch/__init__.py::_rocm_core_libs (22, RTLD_GLOBAL preload). Introduce an import-light torch/_rocm_libs.py holding the canonical leaf-first ROCM_SO_FILES (preload subset) plus ROCM_SO_FILES_BUNDLE_ONLY (MAGMA/profiler/rocRoller extras that are bundled but not preloaded) and ROCM_SO_FILES_ALL. - torch/__init__.py imports ROCM_SO_FILES (preserving leaf-first preload order). - repair_wheel.py loads the module by file path (no `import torch`), preferring the unpacked wheel copy then the source tree, and bundles ROCM_SO_FILES_ALL. Behavior-preserving: ROCM_SO_FILES_ALL is set-equal to the original 27-entry bundle set and the preload order is byte-for-byte identical. Made with Cursor
Co-authored-by: Faa Diallo <amdfaa@users.noreply.github.com>
Co-authored-by: Faa Diallo <amdfaa@users.noreply.github.com>
Co-authored-by: Faa Diallo <amdfaa@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes #ISSUE_NUMBER
Adds lightweight validation for the shared ROCm shared-library list introduced by the selected shared
.sobranch, plus focused tests that verify the list invariants, fail-fast validation behavior, and.ci/manywheel/repair_wheel.py's path-based loader behavior.Validation performed:
torch/_rocm_libs.pyand.ci/manywheel/repair_wheel.pypassed, including positive and negative validation cases.pythonwas unavailable andPYTHONPATH=/workspace python3 test/test_rocm_libs.py TestRocmLibs.test_rocm_so_files_are_valid_basenamescould not run because this unbuilt checkout lacks generatedtorch.version.Authored with Claude.