[ROCm] Add HIP/ROCm support for AMD GPUs#10
Open
jeffdaily wants to merge 2 commits into
Open
Conversation
Add AMD GPU support to the CUDA scan-matcher (ICP point-cloud registration) via HIP. A single src/cuda_to_hip.h compatibility header maps the CUDA runtime and OpenGL-interop symbols used by the project to their HIP equivalents, so the same sources compile for NVIDIA (default) and AMD (-DUSE_HIP=ON) GPUs. The GL interop is updated to the modern graphics-resource API (hipGraphicsGLRegisterBuffer and friends), which HIP provides and which deprecates the old cudaGL* entry points. GLM device-function support is enabled through src/glm_device.h, which defines the qualifier macros GLM expects so its math functions get __host__ __device__ under hipcc. rocThrust is a drop-in for CUDA Thrust; it requires C++17, so the host standard is raised accordingly. Include order matters: Thrust must precede GLM so rocThrust selects the HIP backend. To review, start with src/cuda_to_hip.h and the USE_HIP branch of CMakeLists.txt, then the GL interop and device-math call sites in src/main.cpp and the .cu sources. The README documents the AMD build alongside the existing Windows and Linux CUDA instructions. This work was authored with assistance from Claude, an AI assistant by Anthropic. Test Plan: ``` cmake -S . -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_BUILD_TYPE=Release cmake --build build -j HIP_VISIBLE_DEVICES=0 ./build/bin/cis565_ScanMatching # 10 ICP iterations complete, ~440us NN timing per step on gfx90a (MI250X) ``` Built and run on gfx90a (MI250X, CDNA2), gfx1100 (RDNA3), and gfx1201 (RDNA4); the headless ICP run completes 10 iterations on each.
The default-gfx90a block preempted and duplicated CMake's own host-GPU detection. Removing it lets enable_language(HIP) honor an explicit -DCMAKE_HIP_ARCHITECTURES, auto-detect the host GPU, or error on a no-GPU build host. The block was also dead code: it ran after enable_language(HIP), by which point CMAKE_HIP_ARCHITECTURES was already set, so its guard was always false. This work was authored with the Claude AI assistant.
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.
This adds AMD GPU support to the CUDA scan-matcher (ICP point-cloud registration) through HIP, so the project builds and runs on ROCm in addition to CUDA.
A single
src/cuda_to_hip.hcompatibility header maps the CUDA runtime and OpenGL-interop symbols the project uses to their HIP equivalents. The same sources compile for NVIDIA (the default) and AMD (configure with-DUSE_HIP=ON); the CUDA build path is unchanged. The GL interop is updated to the modern graphics-resource API (hipGraphicsGLRegisterBufferand friends), which HIP provides and which supersedes the oldcudaGL*entry points.GLM device-function support is enabled through
src/glm_device.h, which defines the qualifier macros GLM expects so its math functions get__host__ __device__under hipcc. rocThrust is a drop-in for CUDA Thrust; it requires C++17, so the host language standard is raised to match. Include order matters: Thrust headers must precede GLM so rocThrust selects the HIP backend.Build
CMAKE_HIP_ARCHITECTURESselects the target GPU (for examplegfx90afor CDNA2,gfx1100for RDNA3,gfx1201for RDNA4); it defaults togfx90aif unset. The README documents the AMD build alongside the existing Windows and Linux CUDA instructions.Validation
Built and run on gfx90a (MI250X, CDNA2), gfx1100 (RDNA3), and gfx1201 (RDNA4). The headless ICP run (
VISUALIZE 0) completes 10 iterations on each, with stable nearest-neighbor timing per step and no errors.Notes
This work was authored with assistance from Claude, an AI assistant by Anthropic.