Add HIP backend for AMD ROCm support - #663
Open
paralin wants to merge 2 commits into
Open
Conversation
paralin
force-pushed
the
rocm-hipify-fused-trainer
branch
from
August 26, 2026 02:24
14e0ca1 to
16d4d73
Compare
paralin
marked this pull request as draft
August 26, 2026 02:25
Translate the CUDA training sources to HIP with HIPIFY and build them with hipcc, so the fused trainer runs on AMD GPUs. The CUDA/cuDNN build path is unchanged; everything new is additive. - src-hip/ holds the hipified bindings, kernels, models, muon, ocean, and pufferlib sources, plus a cuda->HIP forwarding shim used by the C env libraries and a stub NVML header (utilization reads as zeros). - The conv encoder path has dual implementations: NVIDIA builds keep src/cudnn_conv2d.cu (cuDNN) untouched via ocean.cu, while AMD builds use src-hip/cudnn_conv2d.hip.cpp (MIOpen) with an identical interface via ocean.hip.cpp. Forward is miopenConvolutionForward with bias + in-place ReLU; backward uses the MIOpen weights/data kernels. The nethack encoder is not ported. - src/vecenv.h maps its small CUDA surface onto HIP equivalents under __HIP_PLATFORM_AMD__; other platforms keep the existing declarations. - build-hip.sh mirrors build.sh: it compiles the static env library, then builds pufferlib/_C with hipcc --offload-arch=<arch> linking amdhip64, hipblas, hiprand, rocrand, rccl, and MIOpen. Verified on an RX 7700 XT (gfx1101) with ROCm 7.2.2: cartpole trains through the fused rollout path at ~6M SPS.
paralin
force-pushed
the
rocm-hipify-fused-trainer
branch
from
August 26, 2026 02:27
16d4d73 to
8ecbb7b
Compare
paralin
changed the title
paralin
marked this pull request as ready for review
August 26, 2026 02:27
Remove the duplicate convolution include guard and link the MIOpen library explicitly. This lets HIP environments compile when the fused trainer includes the convolution encoder source. Signed-off-by: Christian Stewart <christian@aperture.us>
paralin
force-pushed
the
rocm-hipify-fused-trainer
branch
from
August 26, 2026 14:59
437889e to
2a52210
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
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.
Translate the CUDA training sources to HIP with HIPIFY and build them with hipcc, so the fused trainer runs on AMD GPUs. The CUDA build path is unchanged; everything new is additive.
Verified on an RX 7700 XT (gfx1101) with ROCm 7.2.2: cartpole trains through the fused rollout path at ~6M SPS.