Repository navigation
Conversation
Include the resolved CPU runtime thread count in the FX graph cache key when C++ codegen may specialize on it, including no-input CPU factory graphs, so cached serial kernels are not reused under a different thread count. Making the resolved thread count part of the FX graph cache key fixes all reported cases where a CPU graph compiled under one `torch.set_num_threads()` value was incorrectly served from cache after the thread count changed. Fixes pytorch#110611 Fixes pytorch#167463 Fixes pytorch#160812 Generated by my agent Pull Request resolved: pytorch#184444 Approved by: https://github.com/oulgen
I think a debug logging statement was left inadvertently in the code. Pull Request resolved: pytorch#186475 Approved by: https://github.com/Skylion007
Non-strict export already rewrites tensor indexing with SymInt bounds into traceable select and slice operations. Scalar integer tensors used as index, slice start/stop, or slice step fell through to Tensor.__getitem__, where Python indexing tried to extract a concrete integer from an unbacked symbolic scalar. In while_loop body capture this produced GuardOnDataDependentSymNode failures like the u3 scalar slice-bound error discussed in pytorch#173169. Teach the non-strict export torch function handler and the proxy tensor metadata/predispatch modes to recognize 0-d integer tensor indices. The rewrite lowers those values through item() into traceable torch.select or aten.slice calls, including scalar tensor steps. The slice no-op fast path now only checks literal Python int step 1 so scalar tensor steps are not truth-tested during tracing. Bool scalar tensors are deliberately left out of the rewrite because aten.slice expects integer-like bounds, not SymBool. The latest full while_loop repro can still fail later on `if patch_size < z_window_size`; that is dynamic Python control flow inside a traced while_loop body and needs to be expressed with branch-free tensor operations or structured control flow. This change fixes the scalar tensor slicing root cause surfaced before that later intended limitation. Fixes pytorch#173169 Generated by my agent Test Plan: python test/export/test_export.py TestExport.test_non_strict_scalar_tensor_slice_bounds TestExport.test_non_strict_scalar_tensor_slice_step TestExport.test_non_strict_bool_scalar_tensor_slice_not_rewritten TestExport.test_non_strict_while_loop_scalar_tensor_slice_bounds TestExport.test_non_strict_while_loop_scalar_tensor_slice_step TestExport.test_non_strict_while_loop_bool_scalar_tensor_slice_not_rewritten lintrunner -a git diff --cached --check Pull Request resolved: pytorch#184925 Approved by: https://github.com/oulgen
Fake parameters are still FakeTensor instances marked with _is_param. Some downstream module hooks reconstruct a parameter by taking type(fake_param), converting the existing fake tensor to a new elem, and splatting known FakeTensor state from __dict__. FakeTensor.__new__ only supported the internal (fake_mode, elem, device) form, so reconstruction failed on fake_device/_fake_device kwargs or, after downstream remapping, could create a FakeTensor without initialized fake_device state. This updates FakeTensor.__new__ to accept a bounded elem-first reconstruction form that requires fake_mode, accepts device/fake_device/_fake_device aliases when they agree, rejects unknown state, recomputes memo descriptor state, preserves _is_param, and checks FakeTensor elems use the same mode. Keeping this in FakeTensor is preferable to changing Parameter because fake parameters dispatch directly to FakeTensor.__new__. Fixes pytorch#161636 Generated by my agent Test Plan: python test/test_fake_tensor.py -k parameter_state_reconstruction -v python test/test_fake_tensor.py -k fake_device_property_normalization -v lintrunner -a Pull Request resolved: pytorch#185176 Approved by: https://github.com/aorenste
…)" This reverts commit dcf9f6f. Reverted pytorch#184444 on behalf of https://github.com/pytorch-auto-revert due to Reverted automatically by pytorch's autorevert, to avoid this behaviour add the tag autorevert: disable ([comment](pytorch#184444 (comment)))
This reverts commit a827601. Reverted pytorch#185176 on behalf of https://github.com/pytorch-auto-revert due to Reverted automatically by pytorch's autorevert, to avoid this behaviour add the tag autorevert: disable ([comment](pytorch#185176 (comment)))
…torch#178466)" This reverts commit cefb7af. Reverted pytorch#178466 on behalf of https://github.com/atalman due to failing internally ([comment](pytorch#178466 (comment)))
This reverts commit fbecdb3. Reverted pytorch#184327 on behalf of https://github.com/facebook-github-tools due to Diff reverted internally ([comment](pytorch#184327 (comment)))
…de) (pytorch#186247) to prepare moving FakeTensor to C++, first changing all trivial HOPs to use `@register_fake` instead of `.py_impl(FakeTensorMode)` previously these ops were never cached, so adding `skip_cache = True` for all HOPs that are being migrated to preserve original behaviour in pytorch#183776 i change register_fake to also register w/ DispatchKey.Fake Pull Request resolved: pytorch#186247 Approved by: https://github.com/zou3519
…ackends (pytorch#184737) ## Summary Refactors `test/test_multiprocessing.py` to make device-agnostic tests run automatically on all available backends using `instantiate_device_type_tests`. Previously, device coverage was achieved by manually duplicating tests (e.g. `test_variable_sharing` + `test_cuda_variable_sharing`), which meant out-of-tree backends had no coverage. This PR: - Extracts a `_MultiprocessingTestMixin` with 4 shared helpers used by both test classes - Creates `TestMultiprocessingDeviceType` with 8 device-generic test methods that replace 13 previously duplicated CPU/CUDA methods - Fixes `_test_autograd_sharing` to derive device from `var.device` instead of hardcoding `"cuda" if var.is_cuda else "cpu"` - Moves 3 `@staticmethod` child-process targets to module level for better pickle safety with `mp.Process` All 13 CUDA-IPC-specific tests (events, IPC handles, `_share_cuda_()`, fork-poisoning, multi-GPU) remain unchanged in `TestMultiprocessing` with their original names and guards -- these use CUDA-only APIs with no accelerator-generic equivalent. ### Tests moved to `TestMultiprocessingDeviceType` | New method | Replaces | |---|---| | `test_simple_sharing` | `test_cuda_simple` | | `test_empty_tensor_sharing` | `test_empty_tensor_sharing` + `test_empty_tensor_sharing_cuda` | | `test_variable_sharing` | `test_variable_sharing` + `test_cuda_variable_sharing` | | `test_parameter_sharing` | `test_parameter_sharing` + `test_cuda_parameter_sharing` | | `test_integer_parameter_serialization` | `test_integer_parameter_serialization_cpu` + `_cuda` | | `test_leaf_variable_sharing` | `test_leaf_variable_sharing` (had internal device loop) | | `test_non_leaf_variable_sharing` | `test_non_leaf_variable_sharing` (had internal device loop) | | `test_is_shared_accelerator` | `test_is_shared_cuda` | ## Test plan - CI should generate `TestMultiprocessingDeviceTypeCPU` and `TestMultiprocessingDeviceTypeCUDA` classes (visible with `-v` flag) - All existing `TestMultiprocessing` tests pass unchanged - On machines with XPU or OpenReg, additional device variants run automatically - `python test/test_multiprocessing.py -v` to verify test discovery Pull Request resolved: pytorch#184737 Approved by: https://github.com/jansel Co-authored-by: PyTorch MergeBot <pytorchmergebot@users.noreply.github.com>
## Summary `PYTORCH_EXTRA_INSTALL_REQUIREMENTS` becomes the torch wheel's `Requires-Dist`, so `SPMD_TYPES_REQUIREMENT="spmd_types==0.2.0"` (added in pytorch#180880) emitted `Requires-Dist: spmd_types==0.2.0`. `download.pytorch.org` serves the dependency under its PEP 503-normalized name (`spmd-types/`), so when installing with `--index-url https://download.pytorch.org/whl/nightly/<arch>` pip looks up the normalized path. Emit the normalized requirement `spmd-types==0.2.0` so the metadata matches the index path. PyPI treats `spmd_types` and `spmd-types` as equivalent, so default-index installs are unaffected. ## Companion change The matching mirror/index fixes on `download.pytorch.org` are in pytorch/test-infra#8157. ## Note Changing `PYTORCH_EXTRA_INSTALL_REQUIREMENTS` requires triggering `ciflow/binaries` to validate the binary builds. Pull Request resolved: pytorch#186545 Approved by: https://github.com/Skylion007, https://github.com/malfet
Fixes pytorch#163759 When compiled by \�ot_eager\ with \dynamic=True\, \ orch.combinations\ failed with \RuntimeError: Cannot call numel() on tensor with symbolic sizes/strides\ because the C++ implementation used \ umel()\ instead of \sym_numel()\. This commit fixes it by using \sym_numel()\ and passing \c10::SymInt\ to the mask generation. Pull Request resolved: pytorch#186329 Approved by: https://github.com/Skylion007
) This change ensures that oneDNN is used for int8 quantization on AArch64 when using the x86 quantization engine. This fixes the failing test seen in pytorch#180546, as well as other untracked tests under `test_quantization.py`. Pull Request resolved: pytorch#185232 Approved by: https://github.com/aditew01, https://github.com/jansel
…84555) ## Summary This PR enables Armv9-A as a supported target for torch.compile on AArch64 platforms. It extends the existing AArch64 backend configuration to recognize and properly generate code for Armv9-A CPUs. This PR replaces pytorch#175326 in favour of recent support of SVE 128. ## Implementation + Extend architecture detection logic for AArch64 to add Armv9-A to supported `-march` mappings + Ensure Inductor codegen correctly propagates the target triple and feature set + Preserve default behavior for non-Armv9-A environments + No behavior change for x86 or other architectures + 128-bit and 256-bit SVE support ## Testing AArch64 Armv9-A hardware (Neoverse V2) Verified: + `torch.compile` succeeds + Inductor-generated kernels compile without error Pull Request resolved: pytorch#184555 Approved by: https://github.com/aditew01, https://github.com/jansel, https://github.com/jgong5
Keep strided to_dense inside the Python FunctionalTensor wrapper so AOTAutograd metadata collection never sees a raw C++ functional tensor from the no-op path. Fixes pytorch#183268 Generated by my agent Pull Request resolved: pytorch#184586 Approved by: https://github.com/ezyang
Route generated integer floor division through guarded helper semantics so INT_MIN / -1 matches eager instead of trapping in generated C++. Fixes pytorch#184406 Generated by my agent Pull Request resolved: pytorch#184497 Approved by: https://github.com/leslie-fang-intel, https://github.com/jgong5, https://github.com/mlazos
custom_op and triton_op autograd registrations wrap user formulas in a generated autograd.Function. That wrapper passes an internal metadata object through Generated.apply and then appends a None gradient for it after calling the user backward formula. When the user backward formula returned the wrong number of gradients, validation was left to autograd.Function after the wrapper appended the hidden metadata gradient. As a result, the error counted the internal metadata input and reported off-by-one expected/actual values. Validate the user backward formula arity before appending the hidden metadata gradient, using the real positional operator inputs for the expected count. This keeps the internal metadata slot hidden from the diagnostic while preserving the existing metadata gradient path for valid formulas. Fixes pytorch#137219 Generated by my agent Test Plan: - python test/test_custom_ops.py -k incorrect_num_gradients - python test/test_custom_ops.py -k register_autograd - python test/test_custom_ops.py -k supports_tensorlist - lintrunner -a Pull Request resolved: pytorch#186226 Approved by: https://github.com/zou3519
Route debug sync emission through wrapper-specific codegen so cpp_wrapper emits checked CUDA/HIP synchronization instead of Python source. This fixes both `triton.debug_sync_graph` and `triton.debug_sync_kernel` paths and reenables those config fuzzer combinations. Fixes pytorch#140219 Fixes pytorch#140220 Generated by my agent Pull Request resolved: pytorch#184217 Approved by: https://github.com/yushangdi
The online softmax preparation pattern computes softmax max and sum as a multi-output prim. When this prim appears in an AOTAutograd joint graph, the min-cut partitioner previously did not know that it was a recomputable/fusible operation like the decomposed amax/sum pattern it replaces. That skewed the cut and could make the forward save a full softmax activation instead of the much smaller max/sum tuple outputs. Add prims.prepare_softmax_online to the partitioner's recomputable op set. The existing post-grad placement is still the normal path, but this makes the partitioner robust if the online softmax preparation is introduced before partitioning. I considered relying only on the post-grad pattern placement that originally avoided the regression, but that leaves the partitioner's cost model wrong for this prim and makes future joint-graph uses fragile. Benchmark Results: A focused AOTAutograd min-cut reproducer that simulates joint-graph online softmax placement saved 212,992 activation bytes before the fix and 149,504 after the fix, a 29.8% reduction. The old cut saved an extra full [2, 128, 128] fp16 softmax activation; the new cut saves the small [2, 128, 1] fp32 max/sum tuple outputs instead. Test Plan: - python test/inductor/test_online_softmax.py -k test_prepare_softmax_after_partitioning - python test/inductor/test_online_softmax.py -k test_prepare_softmax - python ../state/148141/benchmark_partition_saved_bytes.py - lintrunner -a torch/_functorch/partitioners.py test/inductor/test_online_softmax.py - git diff --check --cached Fixes pytorch#148141 Generated by my agent Pull Request resolved: pytorch#185783 Approved by: https://github.com/shunting314
…186558) Also fixes kernel compilation warning Pull Request resolved: pytorch#186558 Approved by: https://github.com/Skylion007
Add an MPS implementation as device-side assertions work on MPS, matching CUDA's async semantics: a Metal kernel records the assertion message into the per-stream error buffer when the single input element is zero, and the host surfaces it as a c10::AcceleratorError at the next commitAndWait sync point. Authored with the assistance of Claude (claude-opus-4-8). Pull Request resolved: pytorch#186562 Approved by: https://github.com/Skylion007 ghstack dependencies: pytorch#186558
Serialize child modules when their repr is not constructor-valid Python so generated Dynamo minifier repros remain syntactically valid for modules that include nested module display lines. Add a regression test covering FusedMovingAvgObsFakeQuantize. Fixes pytorch#108110 Generated by my agent Pull Request resolved: pytorch#184449 Approved by: https://github.com/ezyang
When replace_random fuses RNG seed and offset helper nodes, copy source metadata from the original per-RNG helper onto the fused node before recomputing tensor metadata. Fixes pytorch#129813 Generated by my agent Pull Request resolved: pytorch#184316 Approved by: https://github.com/choijon5
The max reads fusion test asserts scheduler/codegen metrics that are only populated during fresh compilation. Disable Inductor caches for this test so a warm FX graph cache cannot bypass the metric-producing path; prior abandoned context is chuanqi129#5. Fixes pytorch#181699 Generated by my agent Pull Request resolved: pytorch#183695 Approved by: https://github.com/shunting314
…orch#185232)" This reverts commit c1039f2. Reverted pytorch#185232 on behalf of https://github.com/pytorch-auto-revert due to Reverted automatically by pytorch's autorevert, to avoid this behaviour add the tag autorevert: disable ([comment](pytorch#185232 (comment)))
…185447) C++ symbolic shape guards already fall back to Python guards when C++ guard compilation is unavailable, but the Windows compiler probe raised a plain RuntimeError when cl or CXX was missing. That bypassed the existing InvalidCxxCompiler fallback in SHAPE_ENV and surfaced as an InternalTorchDynamoError for dynamic-shape compilation even with backends that should not require a user-visible C++ compiler. Fix the Windows missing-compiler probe to raise InvalidCxxCompiler, preserving the existing user-facing message when the exception is not handled. This keeps actual C++ compilation failures visible while allowing the intended Python shape-guard fallback for optional C++ guard generation. I considered catching RuntimeError in SHAPE_ENV, but normalizing the toolchain error at the source is narrower and avoids hiding unrelated RuntimeErrors from CppCodeCache.load. Fixes pytorch#157458 Generated by my agent Test Plan: - python test/dynamo/test_guard_manager.py -k test_cpp_shape_guard_missing_windows_compiler_falls_back - python test/dynamo/test_guard_manager.py - lintrunner -a - git diff --cached --check Pull Request resolved: pytorch#185447 Approved by: https://github.com/eellison
The functional aten.zero.default meta kernel allocated its result with self.new_empty(self.shape). new_empty uses a default contiguous layout for the requested shape, so fake/meta execution lost non-contiguous input strides even though eager zero preserves the input layout because it is the functional form of zero_. Use torch.empty_like(self) instead so the meta kernel mirrors eager's layout semantics. Add a regression test covering both direct meta tensors and FakeTensorMode for a transposed input. I considered a broad mechanical replacement of new_empty calls in meta registrations, but many of those kernels either change shape or intentionally return contiguous outputs. This patch fixes the reproduced stride mismatch for the functional zero kernel directly. Fixes pytorch#159660 Generated by my agent Test Plan: python test/test_meta.py -k test_zero_preserves_input_strides python test/test_meta.py TestMetaKernelRegistrations lintrunner -a git diff --check Pull Request resolved: pytorch#185360 Approved by: https://github.com/zou3519
Generated Python tensor-method bindings pack TensorOptions and eagerly call maybe_initialize_device before dispatch. That is correct for ordinary tensors, but it initializes CUDA before Python-dispatched tensor subclasses can handle the method. For a CUDA FakeTensor used outside its creating FakeTensorMode, x.new_empty(device='cuda:0') therefore entered torch.cuda._lazy_init and failed on CUDA-hidden or CPU-only processes even though FakeTensor should not allocate real CUDA memory. Fix this in torchgen by guarding method-level TensorOptions device lazy init with !at::impl::tensor_has_dispatch(self). Ordinary tensors still take the eager lazy-init path before real CUDA allocation; Python-dispatched tensors reach their __torch_dispatch__ first. The Python variable methods template includes TorchDispatchUtils for the generated guard, and the regression test covers new_empty, new_empty_strided, new_full, new_ones, and new_zeros in a CUDA-hidden subprocess while checking CUDA remains uninitialized. I considered special-casing FakeTensor, but the binding layer should only care about dispatch ordering, not a particular Python subclass. The dispatch-key guard fixes the root ordering problem while keeping the generated bindings generic. Fixes pytorch#159566 Generated by my agent Benchmark Results: - main: python -m timeit -n 200000 -r 7 -s 'import torch; x = torch.empty(())' 'x.new_empty((1,))' -> 1.14 usec per loop - fixed: python -m timeit -n 200000 -r 7 -s 'import torch; x = torch.empty(())' 'x.new_empty((1,))' -> 1.13 usec per loop Test Plan: - python -m py_compile torchgen/api/python.py - python -m tools.setup_helpers.generate_code --native-functions-path aten/src/ATen/native/native_functions.yaml --tags-path aten/src/ATen/native/tags.yaml --gen-dir . --subset pybindings - ninja -C build torch_python - python test/test_fake_tensor.py FakeTensorTest.test_cuda_fake_tensor_new_methods_no_device_init - python test/test_fake_tensor.py FakeTensorOperatorInvariants.test_fake_gpu_no_init - lintrunner -a - git diff --cached --check && git diff --check Pull Request resolved: pytorch#185368 Approved by: https://github.com/zou3519
Dynamo resume functions pass restored stack values and locals as normal positional arguments. CPython keeps those argument references alive for the whole resume call, so a `DELETE_FAST` in the resumed bytecode can remove the local name without actually releasing the tensor. That makes compiled graph breaks hold deleted intermediates longer than eager execution. Detect resume functions that may execute `DELETE_FAST` after the resume target and use a boxed frame-values argument for those calls. The resume prologue now loads each saved stack/local value from the list, stores locals where needed, and clears each list slot immediately. Resume functions without a reachable `DELETE_FAST` keep the previous positional calling convention. Nested resume calls need to preserve the child's calling convention, so nested boxed callees receive the frame-values list as one argument while non-boxed callees keep the existing list extension behavior. This fixes the actionable resume-frame/intermediate lifetime issue and the issue-shaped list-cleared input path. Bare direct temporary inputs remain owned by the outer `torch.compile` wrapper's `*args` tuple until the compiled call returns; a pure Python `def wrapper(*args): return fn(*args)` has the same lifetime behavior, so this patch does not try to make `del x` release such bare inputs before return. Fixes pytorch#153701 Generated by my agent Test Plan: - python test/dynamo/test_subgraphs.py SubGraphTests.test_nested_resume_del_releases_tensor SubGraphTests.test_resume_del_releases_tensor SubGraphTests.test_issue_shape_list_clear_and_intermediate_del_release_tensors SubGraphTests.test_del_compiled_only_local_before_graph_break - python test/dynamo/test_subgraphs.py - python test/dynamo/test_nested_graph_breaks.py -k test_step_graph_break_frame_values_not_corrupted - python test/dynamo/test_repros.py ReproTests.test_weakref_reconstruct ReproTests.test_weakref_del ReproTests.test_weakref_callback - lintrunner -a Benchmark Results: Tiny CPU graph-break runtime microbenchmark with `DELETE_FAST`, backend="eager", 7 samples of 3000 cached calls each: - Before: samples_us [26.349, 26.217, 26.076, 26.089, 26.544, 26.578, 26.099], median 26.217 us - After: samples_us [28.897, 29.262, 30.618, 29.686, 31.57, 32.353, 33.83], median 30.618 us (+16.8%) Pull Request resolved: pytorch#185561 Approved by: https://github.com/IvanKobzarev
IndexFillBackward1 uses aten._unique.default on the index tensor when computing the gradient for the fill value. Under compiled autograd, FakeTensor did not have a dedicated implementation for aten._unique.default, and the operator was not tagged as dynamic-output-shape. That allowed tracing/Inductor to see a bogus fixed unique length and compile the value-gradient path against the wrong shape, which could silently produce an incorrect scalar gradient when size assertions were disabled in the inductor-wrapped OpInfo tests. Add a FakeTensor implementation for aten._unique.default that shares the existing dynamic unique shape logic with aten._unique2.default, and tag the schema as dynamic_output_shape for generated operator metadata. The regression test enables dynamic-output-shape capture and verifies compiled autograd compiles through the unique path without graph-breaking, preserving the correct value gradient for noncontiguous index_fill inputs. Fixes pytorch#160664 Generated by my agent Test Plan: python -m pytest -q test/inductor/test_compiled_autograd.py::TestCompiledAutograd::test_index_fill_value_grad_noncontiguous python -m pytest -q test/test_fake_tensor.py::FakeTensorTest::test_unique_output_dtype test/test_fake_tensor.py::FakeTensorTest::test_unique_default_dynamic_output_shape PYTORCH_TEST_WITH_INDUCTOR=1 python - <<'PY' import pytest import torch._dynamo.config torch._dynamo.config.compiled_autograd = True raise SystemExit(pytest.main([ '-q', 'test/test_ops.py::TestCommonCPU::test_noncontiguous_samples_index_fill_cpu_float32', ])) PY lintrunner -a Pull Request resolved: pytorch#185318 Approved by: https://github.com/eellison
Use the same numpy pins as other release branches
## Summary - Advance the ROCm Triton pin to the tip of `ROCm/triton` `release/internal/3.8.x`. - Update `.ci/docker/triton_version.txt` from `3.7.1` to `3.8.0`. Triton pin: `4cff872ced001ea92d9fcf05b3f6517e2b486d19` Triton branch: `ROCm/triton@release/internal/3.8.x` ## Test plan - [x] `git diff --check -- .ci/docker/ci_commit_pins/triton.txt .ci/docker/triton_version.txt` - [x] Verified Triton checkout reports `__version__ = '3.8.0'` - [ ] Run PyTorch CI Follows the structure of ROCm#3451. Made with [Cursor](https://cursor.com) Co-authored-by: Cursor <cursoragent@cursor.com>
Cherry-pick of ROCm#3505 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com>
Cherry picked .ci/docker/requirements-ci.txt from 5379cbe, resolved conflict by accepting changes explicitly made in https://github.com/ROCm/pytorch/pull/3445/changes. Test run: https://github.com/ROCm/TheRock/actions/runs/30402630662/job/90424764289 Successfully installs dependency Co-authored-by: Ethan Wee <Ethan.Wee@amd.com> Co-authored-by: Jithun Nair <37884920+jithunnair-amd@users.noreply.github.com>
…collision (ROCm#3510) This file is excluded from hipify (see tools/amd_build/build_amd.py), so it ships its own self-contained __nv_bfloat16 type plus the __float2bfloat16/__bfloat162float conversions for the runtime-compiled fuser kernel. Starting with ROCm 7.x, hiprtc auto-injects a runtime header (hiprtc_runtime.h) that also defines __float2bfloat16 returning __hip_bfloat16. Since C++ cannot overload on return type alone, the two definitions collide. Rather than gate on ROCM_VERSION (a build-time value that does not reliably track whether the *runtime* hiprtc actually provides these symbols), we keep PyTorch's own round-to-nearest-even implementation and rename our private definitions via the preprocessor so they cannot clash. hiprtc's definitions remain (parsed earlier in the translation unit) but go unused. These #defines must stay in effect for the rest of the kernel source -- do NOT #undef them, since the emitted kernel body calls these conversions after this header is spliced in. Pull Request resolved: pytorch#185656 Approved by: https://github.com/jeffdaily (cherry picked from commit 5ca8c13) Co-authored-by: Andy Lugo <Andy.LugoReyes@amd.com>
…lity tests (ROCm#3511) Use util function to get compiler executable names instead of hardcoded `g++` Pull Request resolved: pytorch#188580 Approved by: https://github.com/jeffdaily (cherry picked from commit e6e48bf) Co-authored-by: Z Liu <zhuojin.liu@amd.com>
…end (ROCm#3512) test_open_device_registration uses a FileCheck to assert the generated CPU C++ wrapper contains a scalar load expression. The expected string was " = in_ptr0[static_cast<long>(i0)];", but inductor CPU codegen now emits " = in_ptr0[static_cast<int64_t>(x0)];": the index cast changed from long to int64_t and the loop variable was renamed from i0 to x0. The golden string was never updated, so the check fails on the nogpu_NO_AVX2 periodic shard, which forces the scalar (non-vectorized) load path via ATEN_CPU_CAPABILITY=default. Rebaseline the expected string to match the current codegen output. Test Plan: Reproduced the shard's codegen path locally by forcing the scalar load branch, which fails before this change and passes after: ``` ATEN_CPU_CAPABILITY=default python -m pytest \ test/inductor/test_extension_backend.py::ExtensionBackendTests::test_open_device_registration -v # -> 1 passed ``` Authored with Claude. Pull Request resolved: pytorch#188509 Approved by: https://github.com/zou3519, https://github.com/Skylion007 (cherry picked from commit 0e8ec23) Co-authored-by: Frederik Gossen <frgossen@meta.com>
…ted call of hipMemMap (ROCm#3517) Repeated call of hipMemMap is causing the errors in SymmetricMemory. Fixing that. Pull Request resolved: pytorch#188673 Approved by: https://github.com/Skylion007, https://github.com/jeffdaily (cherry picked from commit 25825f6) Co-authored-by: Albert Malewski <albert.malewski@amd.com>
Cherry-pick of ROCm#3513 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com>
…d via issue) (pytorch#180927) [ROCm] skip test_autoheuristic in-code (already disabled via issue) (cherry picked from commit 366eac9)
…OCM-21749) (ROCm#3262) ## Summary Fixes ROCM-21749: `test_big_num_tensors__foreach_norm` fails on all ROCm GPUs. ## Root Cause `ForeachFuncWrapper.__call__` profiles a `_foreach_norm` call and checks for `multi_tensor_apply_kernel` in profiler output to verify the multi-tensor fast-path was taken. On CUDA, Kineto captures demangled kernel names like `multi_tensor_apply_kernel[...]`. On ROCm, ROCTracer/rocprofiler captures different symbols (`hipLaunchKernel`, `vectorized_elementwise_kernel`, etc.), so `mta_called` is always `False` and the assertion fails. ## Fix Add `or TEST_WITH_ROCM` to `skip_profiler_check` so the profiler-based assertion is bypassed on ROCm, consistent with the existing skip for CUDA 12.6/12.8 flakiness. ## Test Plan - `test_big_num_tensors__foreach_norm_*` should now pass on ROCm GPUs (MI210, MI300X, MI350X) - No change in behavior on CUDA Jira: https://amd-hub.atlassian.net/browse/ROCM-21749 Co-authored-by: Subramanian <ssubrama1@amd.com> (cherry picked from commit f8efdb3)
…3323) Before ROCm 7.14, HIP does not report about errors during stream capture, need to use check hipStreamCaptureStatus to reject capture-unsafe work early. From ROCm 7.14+, HIP reports capture error through, so that pre-check is unnecessary Fixes pytorch#180232 Pull Request resolved: pytorch#187110 Approved by: https://github.com/jeffdaily, https://github.com/naromero77amd (cherry picked from commit 1047ba5) (cherry picked from commit caf158f)
…ed_callables (ROCm#3354) ### Summary This PR updates `torch.cuda.make_graphed_callables` so warmup and graph capture use the same stream, avoiding ROCm graph capture failures tied to hipBLASLt. ### Why On ROCm, hipBLASLt handle caching is per (device, stream). If `stream=` is omitted, `torch.cuda.graph` uses its own default capture stream, which is not the stream used for warmup. The first real use on that capture stream will trigger a lazy hipblasLtCreate there. Handle creation runs capture‑unsafe internal allocation/setup, which fail with stream-capture errors (and sometimes hang) during capture. Running warmup on the same stream passed into `torch.cuda.graph(..., stream=...)` creates and caches the hipBLASLt handle before capture, so hipblasLtCreate (and its associated allocations) are not hit mid‑capture. Fixes pytorch#179961 Pull Request resolved: pytorch#187745 Approved by: https://github.com/jeffdaily (cherry picked from commit d9d756c) (cherry picked from commit eaf6098)
Backporting pytorch#180310 to release 2.12 (cherry picked from commit 9ffbd03)
ROCm#3421) Validation theRock Run: https://github.com/ROCm/TheRock/actions/runs/29045360944 Cherry-pick of upstream [pytorch#188597](pytorch#188597) onto `release/2.12` to add initial ROCm support for **gfx1250 (CDNA5)**. Co-authored-by: @glen-amd - Gates gfx1250-specific behavior on ROCm 7.14+ (the support floor), applied consistently across CUDABlas.cpp, ScaledBlas.cpp, CUDAHooks.cpp, and the test helpers. - Adds gfx1250 to the hipBLASLt preferred/supported arch lists and the hipSparseLt support check (ROCm 7.14+). - Extends the scaled GEMM / MX-format paths (block-wise Float8_e8m0fnu, mxfp8, mxfp4) to gfx1250. - Sets the gfx1250 shared-memory limit to 320 KB (vs 160 KB on gfx950), matching on the suffix-stripped arch name. - Adds CDNA2/CDNA3/CDNA5 "or later" arch helpers and updates FP8 / MX GEMM test gating in common_cuda.py. - Builds Composable Kernel (CK) GEMM as a separate ck_gemm library with gfx1250 filtered out of HIP_ARCHITECTURES, so a multi-arch build that includes gfx1250 keeps CK GEMM for the other archs instead of failing to compile for gfx1250 (composable_kernel has no gfx1250 support yet). The filter is removed once CK supports gfx1250. - Extends the gfx942 nontemporal vectorized-load path in MemoryAccess.cuh to gfx1250. - Refactors duplicated arch-check logic in CUDABlas.cpp, ScaledBlas.cpp, and cuSPARSELtOps.cpp per review feedback. - Writes the CublasHandlePool workspace archs in full (gfx942/gfx950/gfx1250). Enabled via the main merge: - Flash attention and memory-efficient attention on gfx1250 via AOTriton 0.12.1b (from pytorch#188242), which this branch merges in. What is not enabled yet: - CK SDPA on gfx1250: composable_kernel has no gfx1250 support; the CK SDPA target auto-filters to gfx942/gfx950. - CK GEMM on gfx1250: filtered out (see above) until composable_kernel supports it. - FP8 grouped GEMM on gfx1250: MSLK builds only gfx942/gfx950. - int4 mm on gfx1250: the tinygemm MFMA kernel needs a WMMA port; the ops return a clear "not supported yet" error and the int4 unit tests skip gfx1250. CI note: - gfx1250 is not added to the ROCm 7.2 (rocm-n) docker image, whose HIP compiler cannot target gfx1250. gfx1250 CI belongs on the nightly ROCm image and is added in a follow-up. Co-authored-by: @glen-amd - Upstream PR: pytorch#188597 - Related AOTriton bump: pytorch#188242 --------- Co-authored-by: Prachi Gupta <prachi.gupta@amd.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: Jeff Daily <jeff.daily@amd.com> (cherry picked from commit 13eab2b)
Validation theRock Run: https://github.com/ROCm/TheRock/actions/runs/29045354548 [ROCm] Bump AOTriton to 0.13b Notable new features compared with 0.12.1b: * Debug printing can be enabled with env var `AOTRITON_DEBUG_LEVEL`, without re-compiling + We use the same logging level as of `AMD_LOG_LEVEL` - CAVEAT: AOTriton itself does not read `AMD_LOG_LEVEL` + Set `AOTRITON_DEBUG_LEVEL=5` to print everything * **BREAKING** `.aks2` files are packed into **uncompressed** zip files + `chcp 65001` is not required to build on Win32 any more. * TheRock 7.15 support Bug Fixes: Add `check_requires_grad_and_nested` back to `can_use_mem_efficient_attention`, which was removed accidentally. *Other changes are intentionally left out: they consist mainly of internal refactoring that is not perceivable to pytorch users.* [ROCm] Enable gfx1250 SDPA with AOTriton 0.12.1b (pytorch#188612) Notable New Features: * Tech Preview support of gfx1250 hardware + For broader testing, `TORCH_ROCM_AOTRITON_ENABLE_EXPERIMENTAL=1` is NOT NEEDED for gfx1250 at the moment. + This support tier is tracked by new function `aotriton::isArchTechPreview()` [ROCm] Bump AOTriton to 0.12b (pytorch#184288) Notable new features: * **BREAKING** Varlen LSE tensor shape changes to (H, Total_seqlen) * Support head_dim != head_dim_v * Support `use_deterministic_algorithims` * Support seqused_k in test/test_varlen_attention.py * gfx1100 and gfx1151 promoted out of experimental * Partial FAv3 support on gfx950 Bug Fixes: * GQA kernel failed to read bias tensor with the right offset. Known Issues * gfx950's Triton kernel has problem handling hdim=16's fwd, in addition to hdim=48/80's bwd. * Disables gfx90a's CK SDPA support due to GPU Segfault. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> (cherry picked from commit 10d8e1d)
…h#183014) Upstream PR pytorch#181112 (`[inductor] Unify threads-per-wave (warp_size) extraction`) has now landed and did much of the original broad warp-size plumbing for this PR: it centralized `DeviceProperties.warp_size`, threaded `warp_size` through Inductor's Triton config helpers, and made HIP wave32 devices follow the non-wave64 path. After rebasing on top of that work, this PR is now focused on the remaining stricter ROCm-safe behavior: - Preserve missing `DeviceProperties.warp_size` as `None` instead of synthesizing `32` for CUDA/HIP devices. - Make `DeviceProperties.warp_size_or_default` raise for missing CUDA/HIP warp-size metadata, while preserving the explicit non-GPU fallback. - Skip optional autotune/rblock/fusion/split-reduction heuristics when required warp-size metadata is unavailable instead of guessing. - Keep CUDA launcher generation on the fast `32` path, but use the device warp size for HIP launcher block dimensions. - Use queried device warp size in template register-spill pruning, with graceful fallback on non-CUDA/XPU builds. - Update the relevant Inductor tests for the stricter missing CUDA/HIP metadata behavior. - Rebased cleaned stack onto `upstream/main` at `19791183fec`. - Rebuilt PyTorch from scratch with `/home/niromero/docker_workspace/framework_scripts/pytorch/build.sh`. - Ran impacted UTs: - `python test/inductor/test_triton_heuristics.py TestWarpSizeUnification TestTritonHeuristics.test_autotune_hints_to_configs TestFastLauncherDeviceSupport` - `python test/inductor/test_max_autotune.py TestMaxAutotune.test_max_autotune_includes_max_autotune_pointwise_configs` - `python test/inductor/test_torchinductor_dynamic_shapes.py -k test_embedding_backward_dynamic_shapes_large_grid` - `python test/inductor/test_torchinductor.py -k test_reduction_config_limit` Made with [Cursor](https://cursor.com) Pull Request resolved: pytorch#183014 Approved by: https://github.com/jansel, https://github.com/glen-amd, https://github.com/jeffdaily, https://github.com/eellison (cherry picked from commit 9183299) Signed-off-by: pablo-garay <pablo-garay@users.noreply.github.com>
## Summary Align release/2.13 with the ROCm 7.14 compatibility handling required by TheRock CI. This change adds reusable ROCm version guards for standard and multiprocess tests and handles known ROCm 7.14 runtime and packaging incompatibilities found during full-suite validation. ## Test plan - Distributed suite: https://github.com/ROCm/TheRock/actions/runs/30552305532 - Default + Inductor suites: https://github.com/ROCm/TheRock/actions/runs/30585206557 All validation completed successfully. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
Cherry-pick of ROCm#3521 Co-authored-by: Dmitry Nikolaev <139769634+dnikolaev-amd@users.noreply.github.com>
) (ROCm#3539) Cherry-pick of upstream commit 2314113 (pytorch#187052) onto release/2.13. The LocalTensorMode coordinate cache was keyed by the mesh's Python object id, which is not unique over the lifetime of the mode: when a temporary submesh is destroyed, Python may reuse its id for the next submesh, silently returning stale coordinates for a different mesh dimension. This fixes the cache key to use the actual inputs to the coordinate calculation (ndim, flattened rank map, layout) instead. Fixes ROCM-28605, pytorch#184526. Verified locally: test/distributed/tensor/test_utils.py::TestStridedShardingWithLocalTensor: {test_2d_mesh_strided_sharding,test_2d_mesh_2d_tensor_strided_sharding} fail before this commit and pass after it. Co-authored-by: Alexander Grund <alexander.grund@tu-dresden.de>
Porting only missing code (gfx1250) from https://github.com/pytorch/pytorch/pull/190703/changes Signed-off-by: Jagadish Krishnamoorthy <jagadish.krishnamoorthy@amd.com>
…OCm#3545) Fixes empty HIP_ARCHITECTURE build failure when gfx1250 is the only provided value to PYTORCH_ROCM_ARCH
alugorey
requested review from
jataylo,
jeffdaily,
jithunnair-amd and
pruthvistony
as code owners
August 7, 2026 15:23
This branch has not been deployed
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.
Add gfx1250 to whitelist for hipblasLT