Add a CUDA/HIP runtime compatibility layer

[digantdesai]

Route the AOTI CUDA runtime surface through a compatibility header that maps to HIP when EXECUTORCH_USE_HIP is defined. Keep the Buck header target separate from caller_stream so existing NVIDIA targets gain no additional linkage.

With assistance from Claude Code and Codex.

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/21889

• 📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

❌ 11 Pending, 1 Unclassified Failure

As of commit ba801c5 with merge base a1eee71 ():

UNCLASSIFIED FAILURE - DrCI could not classify the following job because the workflow did not run on the merge base. The failure may be pre-existing on trunk or introduced by this PR:

• Test CUDA Builds / export-model-cuda-artifact (SocialLocalMobile, Qwen3.5-35B-A3B-HQQ-INT4, quantized-int4-tile-packed) / linux-job (gh) (this job did not run on the merge base, so DrCI cannot tell whether the failure is pre-existing)
  RuntimeError: Command bash /__w/_temp/exec_script failed with exit code 1

This comment was automatically generated by Dr. CI and updates every 15 minutes.

[github-actions]

This PR needs a release notes: label

If your change should be included in the release notes (i.e. would users of this library care about this change?), please use a label starting with release notes:. This helps us keep track and include your important work in the next release notes.

To add a label, you can comment to pytorchbot, for example
@pytorchbot label "release notes: none"

For more information, see
https://github.com/pytorch/pytorch/wiki/PyTorch-AutoLabel-Bot#why-categorize-for-release-notes-and-how-does-it-work.

[Copilot]

Pull request overview

Adds an opt-in ROCm/HIP build mode for the existing CUDA/AOTI backend by introducing a small CUDA-runtime compatibility header (runtime_api.h) and wiring CMake/Buck targets so the backend can compile/link against HIP when EXECUTORCH_BUILD_ROCM is enabled (while remaining mutually exclusive with EXECUTORCH_BUILD_CUDA).

Changes:

• Introduce EXECUTORCH_BUILD_ROCM and propagate HIP linkage/defines through CMake (plus exported CMake package dependency metadata).
• Route CUDA-runtime includes through a new compatibility header (extension/cuda/runtime_api.h) used across the CUDA runtime and AOTI slim CUDA surfaces.
• Add ROCm-facing documentation and an AMD/Triton AOTI export example; adjust Python backend shims advertising + tests to respect the ROCm gate.

Reviewed changes

Copilot reviewed 38 out of 38 changed files in this pull request and generated 1 comment.

Show a summary per file

File	Description
tools/cmake/preset/default.cmake	Adds EXECUTORCH_BUILD_ROCM option and conflicts/requirements wiring.
tools/cmake/executorch-config.cmake	Optionally includes generated backend dependency metadata for consumers.
tools/cmake/executorch-backend-dependencies.cmake.in	Template for generated find_dependency(...) lines for enabled backends.
extension/llm/runner/CMakeLists.txt	Links HIP runtime and defines CUDA_AVAILABLE when building with ROCm.
extension/cuda/targets.bzl	Adds a Buck runtime_api header target and makes caller_stream depend on it.
extension/cuda/runtime_api.h	New CUDA↔HIP runtime API compatibility header.
extension/cuda/CMakeLists.txt	Links extension_cuda against HIP when EXECUTORCH_BUILD_ROCM is set; exports EXECUTORCH_USE_HIP.
extension/cuda/caller_stream.h	Switches CUDA runtime include to the new compatibility header.
extension/cuda/BUCK	Updates commentary to reflect new runtime API indirection.
examples/cuda/scripts/export_amd_pointwise.py	New example script exporting an AMD AOTI/Triton pointwise model.
examples/cuda/README.md	Documents the AMD AOTI pointwise example and build/export steps.
CMakeLists.txt	Adds ROCm build flow, HIP dependency export metadata, and enables relevant subdirs for ROCm.
backends/cuda/tests/test_sort_shim.py	Adds coverage ensuring ROCm gating doesn’t change CUDA shim advertising.
backends/cuda/runtime/utils.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/test/test_cuda_mutable_state.cpp	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/test/test_cuda_allocator.cpp	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/targets.bzl	Adds dependency on //executorch/extension/cuda:runtime_api.
backends/cuda/runtime/shims/sort.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/shims/int8_plain_mm.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/shims/int6_plain_mm.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/shims/int5_plain_mm.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/shims/int4mm.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/shims/int4_plain_mm.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/shims/cuda_guard.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/memory_tracker.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/cuda_mutable_state.cpp	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/cuda_delegate_handle.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/cuda_backend.cpp	Routes CUDA runtime include through runtime_api.h and blocks CUDA graphs under HIP.
backends/cuda/runtime/cuda_allocator.h	Routes CUDA runtime include through runtime_api.h.
backends/cuda/runtime/cuda_allocator.cpp	Routes CUDA runtime include through runtime_api.h.
backends/cuda/rocm.md	New documentation for experimental ROCm/HIP support and limitations.
backends/cuda/cuda_backend.py	Gates CUDA-only fallback shim advertising and custom-op shim map when running on ROCm PyTorch.
backends/cuda/CMakeLists.txt	Uses HIP for runtime libs under ROCm; avoids building CUDA-only operator shims.
backends/aoti/slim/cuda/targets.bzl	Adds dependency on //executorch/extension/cuda:runtime_api.
backends/aoti/slim/cuda/guard.h	Routes CUDA runtime include through runtime_api.h.
backends/aoti/slim/c10/cuda/targets.bzl	Adds dependency on //executorch/extension/cuda:runtime_api.
backends/aoti/slim/c10/cuda/Exception.h	Routes CUDA runtime include through runtime_api.h.
backends/aoti/CMakeLists.txt	Enables SlimTensor GPU path for ROCm via HIP linkage/defines.

Suppressed comments (2)

extension/cuda/runtime_api.h:32

• cudaStreamCaptureModeRelaxed is declared as auto under HIP. For a compatibility header, it’s safer to give this constant an explicit CUDA-named type so it behaves like the CUDA runtime API and avoids type surprises in overload resolution/templates.
  extension/cuda/runtime_api.h:136
• The HIP wrapper for cudaStreamBeginCapture takes a HIP-specific enum type in its signature. To keep the CUDA runtime surface consistent across CUDA and HIP builds, the parameter should use the CUDA-named alias type (e.g., cudaStreamCaptureMode).

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[digantdesai]

@claude review this PR.

[claude]

Claude finished @digantdesai's task in 4m 31s —— View job

---

PR Review: CUDA/HIP runtime compatibility layer

Reviewed the full diff against origin/main. This is a clean, well-scoped change: it routes the CUDA runtime surface through a single runtime_api.h shim so the AOTI CUDA backend can be built against HIP behind the opt-in EXECUTORCH_BUILD_ROCM gate, without perturbing existing NVIDIA targets. The include-rewrite is mechanical and consistent, the CMake/Buck gating keeps CUDA and ROCm mutually exclusive, and the tests explicitly assert the ROCm gate doesn't change what CUDA advertises. Nice touch decoupling runtime_api from caller_stream so NVIDIA targets gain no new linkage.

I verified the compatibility header is complete for the HIP path: every cuda* symbol used by files that are actually compiled under ROCm is mapped. The only unmapped symbols (cudaEvent*, cudaGetDeviceProperties, cudaFuncGetAttributes, cudaErrorCudartUnloading, etc.) live exclusively in the .cu/.cuh operator shims and CUDA-only tests, which the CMake correctly excludes when EXECUTORCH_BUILD_ROCM is set. Good.

Below are findings, none of them blocking.

---

Correctness / robustness

1. hipPointerAttribute_t::type is ROCm-version-dependent — extension/cuda/runtime_api.h:123

attributes->type = hip_attributes.type;

The field on hipPointerAttribute_t was named memoryType in older ROCm releases and only became type more recently. Since this is the one place the shim reaches into a HIP struct member (rather than a straight function forward), it's the most likely spot to break on an untested ROCm version. Given support is documented as validated only on MI300X/gfx942, this is acceptable, but a one-line comment pinning the minimum ROCm version that exposes .type would save a future debugging session. Fix this →

2. Graph functions are still exposed on HIP but silently degrade. The header forwards cudaGraphInstantiate → hipGraphInstantiateWithFlags even though (per the rocm.md note and the set_option guard in cuda_backend.cpp:280) HIP ignores the auto-free-on-launch flag. The runtime guard that returns Error::NotSupported for enable_cuda_graph_for_method under EXECUTORCH_USE_HIP is the right defense and the only reachable entry point, so this is fine — just flagging that the header itself provides no protection, which is a reasonable separation of concerns.

Consistency (matches the Copilot reviewer's notes)

3. HIP-branch capture types don't use CUDA-named aliases — extension/cuda/runtime_api.h:31 and :133

inline constexpr auto cudaStreamCaptureModeRelaxed = hipStreamCaptureModeRelaxed;
...
inline cudaError_t cudaStreamBeginCapture(cudaStream_t stream, hipStreamCaptureMode mode);

For a header whose whole purpose is to present a CUDA-named surface, auto for the capture-mode constant and a raw hipStreamCaptureMode parameter are the two spots that leak the HIP type. Adding using cudaStreamCaptureMode = hipStreamCaptureMode; and using it in both places would make the surface uniform. Low priority — the capture path is unreachable on HIP anyway (see #2) — but cheap to do.

Minor / nits

4. Dropped explanatory comment — backends/cuda/cuda_backend.py:410. The refactor into _get_custom_ops_to_c_shim_options kept the try/except AttributeError but dropped the original comment (quantize_op_dispatch not imported — op not registered, skip C shim mapping). Worth restoring so the bare except AttributeError: return {} doesn't read as swallowing an unexpected error.

5. torch re-import removed correctly. The old inline import torch inside the try block is gone and the module-level import torch (line 17) covers the new torch.version.hip / torch.ops.executorch_cuda references. Confirmed — no functional gap.

CI

The one red job — export-model-cuda-artifact (SocialLocalMobile, Qwen3.5-35B-A3B-HQQ-INT4, quantized-int4-tile-packed) — looks unrelated to this change: every other CUDA export (including the same quantized-int4-tile-packed recipe on Gemma/Qwen/Voxtral/Whisper/parakeet), all CUDA unit tests, and all build jobs are green. This PR touches only the ROCm-gated paths and a mechanical include swap that CUDA CI otherwise exercises fully. The failure on a single 35B HQQ model is consistent with a resource/model-specific flake rather than a regression, but worth a re-run to confirm.

---

Overall this is a solid, low-risk foundation for ROCm. The findings above are all optional polish; nothing blocks merge from a correctness standpoint. The main caveat to keep visible for users is the one already captured in rocm.md: no ROCm CI, so the HIP path is only as validated as manual MI300X runs.

[Copilot]

Pull request overview

Copilot reviewed 38 out of 38 changed files in this pull request and generated 1 comment.

[Copilot]

Pull request overview

Copilot reviewed 38 out of 38 changed files in this pull request and generated no new comments.

Suppressed comments (3)

extension/llm/runner/CMakeLists.txt:76

• In the ROCm branch, extension_llm_runner links against extension_cuda, but the runner sources under extension/llm/runner don’t reference caller_stream/extension_cuda. This adds an extra shared-library dependency for ROCm builds (and is inconsistent with the CUDA branch, which doesn’t link extension_cuda). Consider dropping extension_cuda here and explicitly defining EXECUTORCH_USE_HIP for runner sources that need the compatibility header.
  backends/cuda/tests/test_sort_shim.py:25
• This module-level import pulls in quantize_op_dispatch during test collection, which can introduce extra optional dependencies (e.g., torchao) for a test file that otherwise only exercises the sort shim. Guard the import so missing optional deps don’t fail collection; tests that require the custom op registration can then skip when unavailable.

# Register custom ops so the shim-map assertions exercise the ROCm gate.
import executorch.backends.cuda.quantize_op_dispatch  # noqa: F401
import torch

backends/cuda/tests/test_sort_shim.py:157

• This test assumes executorch_cuda custom ops have been registered (so aot_inductor.custom_ops_to_c_shims is present). If the optional registration import fails (e.g., missing torchao), this currently fails with KeyError. Skip the assertion when custom ops aren’t available, so the rest of the CUDA shim tests can still run.

    def test_cuda_shim_map_unchanged_by_rocm_gate(self):
        """Same, for the C shim signatures."""
        with patch.object(torch.version, "hip", None):
            options = CudaBackend.get_aoti_compile_options([])