Conversation
Adds a ROCm/HIP build of the GPU compute backend (rmagine::cuda) behind a new top-level USE_HIP CMake option, so the correspondence, statistics, reduction, and noise compute runs on AMD GPUs. USE_HIP defaults to OFF, and with it off the CUDA/NVIDIA build is unchanged. The .cu sources keep their CUDA spelling and are compiled as HIP via LANGUAGE HIP. A small force-included compatibility header, src/rmagine_cuda/include/rmagine/util/cuda/cuda_to_hip.h, maps the runtime and driver-context symbols used onto their HIP spellings. The cuRAND-to-hipRAND mapping lives in a sibling curand_to_hiprand.h included only by the two headers that name curandState in declarations (random.cuh, NoiseCuda.hpp), so every public header keeps its upstream include footprint and stays compilable by a plain host C++ compiler. Build system changes, all inside the USE_HIP branch: - enable_language(HIP), with -fgpu-rdc and --hip-link because svd/umeyama_transform are cross-translation-unit __device__ functions (the CUDA build uses CUDA_SEPARABLE_COMPILATION for the same reason). - Links hip::host and hip::hiprand publicly; hip::device stays private so the offload flags do not leak onto plain C++ consumers. The cublas/cusolver link references, which have no call sites, are dropped in this configuration. - CMAKE_HIP_ARCHITECTURES defaults to gfx90a only when unset; a caller passing another architecture is left alone. - The installed package config (rmagine-cuda-config.cmake.in) now records which backend the component was built against and, in a HIP build, resolves hip and hiprand instead of a CUDA toolkit, so find_package(rmagine) works from an installed prefix on a machine with no CUDA. Previously the generated config always fell through to the legacy FindCUDA module, which aborts the downstream configure when no CUDA is installed. Verified end to end with a standalone consumer project on a CUDA-free host. - CudaContext maps the driver context API onto hipCtx*; cuCtxSetSharedMemConfig/cuCtxGetSharedMemConfig are guarded out under HIP (no configurable shared-memory bank width on AMD). Wave-size hardening, applied only on the HIP path (the CUDA path is byte-identical): the block reductions behind rm::sum, rm::mean, rm::cov, and sumBatched ran their __syncthreads tree only down to 32 lanes and finished with a volatile warp-synchronous tail that assumes a 32-lane lockstep warp. That assumption is not guaranteed on a 64-lane wavefront (gfx90a/gfx942), so under HIP the full block-wide __syncthreads tree runs to the end, with the same add order, which is correct on any wavefront size. Three bugs surfaced by AMD hardware but undefined behavior on CUDA as well, fixed unconditionally: - The reduction kernels seeded shared memory with `sdata[tid] *= 0.0`, which reads uninitialized shared memory first. On AMD that garbage is routinely NaN, which survives the multiply and poisoned rm::sum, rm::mean, and sumBatched results. Each lane now seeds a true typed zero. - multNx1(Quaternion, Quaternion) dispatched multNxN_kernel, which reads past its size-1 second buffer; the Transform and Matrix3x3 overloads already used multNx1_kernel. CUDA tolerated the out-of-bounds read, AMD faults. Switched to multNx1_kernel. - RMAGINE_FUNCTION in shared_functions.h was keyed on __CUDA_ARCH__ (device pass only), which the stricter HIP/clang front end rejects as a host/device attribute mismatch when device code calls the shared helpers. It is now keyed on __CUDACC__/__HIPCC__, which are defined in both compilation passes on each toolchain; a plain g++ build defines neither, so the CPU path is unchanged. The MemoryView::raw definitions get the matching annotation. New tests: - cuda_math_reduction_correctness computes rm::sum/rm::mean/rm::cov over a non-power-of-two input, asserts each component against a double-precision CPU reference, and asserts the GPU result is bit-identical across two runs. The pre-existing cuda_math* tests print their reduction outputs without asserting on them; this test failed with NaN before the shared-memory seed fix. - cuda_public_headers includes all 16 installed rmagine_cuda headers at the top of a plain host C++ translation unit, which catches a device-side header leaking into the public include footprint. It compiles on both backends. Authored with the assistance of Claude, an AI coding agent. Test Plan: ``` cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=ON \ -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ -DRMAGINE_EMBREE_DISABLE=ON -DRMAGINE_OPTIX_DISABLE=ON \ -DRMAGINE_VULKAN_DISABLE=ON -DRMAGINE_VULKAN_CUDA_INTEROP_DISABLE=ON \ -DRMAGINE_OUSTER_DISABLE=ON -DRMAGINE_BUILD_TESTS=ON \ -DRMAGINE_BUILD_TOOLS=OFF cmake --build build -j ctest --test-dir build --output-on-failure -R '^cuda_' ctest --test-dir build --output-on-failure -R '^core_' ``` Run on real AMD GPUs at the tip of this series, with the cuda suite run twice to confirm bit-identical results: - AMD Radeon Pro W7800 (gfx1100, Linux, ROCm 7.2.3): cuda 8/8 and 8/8, core 12/12; also install plus find_package(rmagine) from a separate consumer project built with plain g++ and run on the GPU. - AMD Instinct MI250X (gfx90a, Linux): cuda 8/8 and 8/8, core 12/12. - AMD Instinct MI300X (gfx942, Linux): cuda 8/8 and 8/8, core 12/12. - AMD Radeon 8060S (gfx1151, Windows 11, ROCm 7.14): cuda 8/8 and 8/8, core 12/12. - At an earlier revision of this series (7 cuda tests existed then): AMD Radeon RX 9070 XT (gfx1201, Windows) and AMD Radeon PRO V710 (gfx1101, Windows), each cuda 7/7 twice and core 12/12. The CUDA path was verified unchanged with a full USE_HIP=OFF build using nvcc 12.8.93, compiling all 77 targets cleanly, including the two new tests. hipRAND output is not bitwise-identical to cuRAND (expected), so the noise/random paths are validated statistically rather than bitwise.
Adds src/rmagine_hiprt, an experimental GPU ray-casting component for AMD hardware that reimplements the OptiX simulator set on the HIPRT SDK: PinholeSimulatorHiprt, SphericalSimulatorHiprt, O1DnSimulatorHiprt, and OnDnSimulatorHiprt. Acceleration structures are built with hiprtCreateGeometry and hiprtBuildGeometry; the raygen/closest-hit/miss programs become a single trace kernel per sensor type, JIT-compiled with hiprtBuildTraceKernels and launched through the HIP driver API, with hiprtGeomTraversalClosest for traversal. For O1Dn and OnDn the host model holds RAM-side direction/origin arrays that cannot be passed to a kernel, so device-side model structs carry the uploaded pointers instead. This component is a proof of concept: it has no install rules yet, meshes are merged into a single geometry (no multi-mesh instancing), and face normals are not computed. It is built only when USE_HIP=ON and the HIPRT SDK is found (located via the HIPRT_PATH environment variable), and is skipped by CMake with a warning otherwise, so it adds no cost or risk to any other configuration. The only change outside the new directory is the top-level CMakeLists.txt hook that adds the subdirectory under USE_HIP. Authored with the assistance of Claude, an AI coding agent. Test Plan: This component is experimental and has no in-tree tests yet. Building it requires the HIPRT SDK: ``` export HIPRT_PATH=/path/to/hiprt cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DUSE_HIP=ON \ -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ cmake --build build -j ``` It was exercised on an AMD Instinct MI250X (gfx90a, Linux, ROCm 7.2.1) with a standalone harness, not included in the repository, tracing rays against a quad at z=2: Pinhole 25/64 hits with exact center range 2.0, Spherical 25/25 hits at phi=pi/2 with exact center range 2.0, O1Dn 4/4 hits, and OnDn 4/4 hits with exact range 2.0.
Collaborator
Author
|
To approve this port, leave a comment containing this line by itself: To send it back to the porter instead: Both are commands because GitHub greys out the Approve and Request Changes buttons for a pull request's author, and this PR was opened on your credentials. Your latest command is the one that stands, and a command quoted in a code fence -- like the two above -- is ignored. Either box works: a review comment (Review changes -> Comment, or The title and body above are what gets opened upstream, verbatim, so approving here approves all three: the code, the title and the body. Anything pushed afterwards, or any edit to the title or body, voids it and needs a fresh one. |
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.
Adds a HIP build path so the CUDA compute backend runs on AMD GPUs through ROCm, and adds an experimental HIPRT ray-casting backend. Today the README's backend table reaches AMD only through Vulkan, and on the
rmclpath that route is itself wired through CUDA external-memory interop, so there is no standalone AMD path. This adds one.The work was done with the assistance of an AI coding agent, reviewed and validated by a person before submission.
What changed
rmagine_cudagains aUSE_HIPCMake option. When it is off, nothing changes: the NVIDIA path compiles the same sources withnvccexactly as before. When it is on, the existing.cufiles are markedLANGUAGE HIPand compiled byamdclang++. No.cufile is renamed, no host C++ is restructured, and no CUDA source is deleted.The CUDA-to-HIP symbol mapping lives in one small compatibility header covering the runtime API actually used here: allocation, memcpy, memset, streams, device properties, and error handling, all of which are 1:1. cuRAND device calls map to hipRAND.
cuBLASandcuSOLVERappeared only as CMake link targets, never as call sites -- the SVD and Cholesky helpers inlinalg.cuare hand-written kernels -- so those link references are simply not added on the HIP path.CudaContextuses the CUDA driver API, and eachcuCtx*call has ahipCtx*equivalent exceptcuCtxSetSharedMemConfig. Shared-memory bank width is not configurable on AMD hardware, so that single call is guarded out under HIP.The reduction fix
statistics.cuandmath_batched.cuuse the classic warp-synchronous reduction: the__syncthreads()-guarded tree stops at 32 elements, then an unsynchronizedvolatiletail folds the last 32 lanes assuming they advance in lockstep.That assumption is a 32-lane warp. AMD CDNA hardware runs 64-wide wavefronts, where the low 32 lanes are not guaranteed lockstep across those unsynchronized steps, so the reduction produces wrong and run-to-run non-deterministic results -- which is precisely what the existing
cuda_math_statisticstest is there to catch.Under
USE_HIPthe tail is dropped and the barrier-synchronized tree runs all the way to 1. Same summation order, block-wide barrier at every step, correct at any wavefront width. The CUDA path is untouched.cuda_math_reduction_correctnesschecksrm::sum,meanandcovagainst a CPU reference and is run twice per platform to confirm the result is bit-identical between runs.Scope
Only
rmagine_coreandrmagine_cudaare ported. The OptiX and Vulkan ray-tracing backends are unchanged and remain NVIDIA-only and cross-vendor respectively. The second commit addsrmagine_hiprtas an experimental HIPRT backend; CMake skips it with a warning when the HIPRT SDK is absent, which is the configuration every platform below was validated in, so it does not affect the results here.Test Plan
Configure and build on ROCm, disabling the backends that are out of scope for this change:
Run the GPU suite twice to check determinism, then the CPU suite as a no-regression check:
On Windows the same configure runs with
amdclang/amdclang++as the C and C++ compilers and-D_USE_MATH_DEFINES -DNOMINMAX, with assimp, TBB and their dependencies supplied by vcpkg:Results
Every platform below built cleanly and ran the full suite on a real GPU: 8 of 8
cuda_tests and 12 of 12core_tests, with the GPU suite bit-identical across two consecutive runs.cuda_tests:cuda_math,cuda_memory,cuda_memory_slicing,cuda_math_svd,cuda_math_statistics,cuda_math_reduction,cuda_math_reduction_correctness,cuda_public_headers.The NVIDIA build was also configured and built with
USE_HIP=OFFagainst nvcc 12.8 targetingsm_80, 77 of 77 targets, to confirm the CUDA path still compiles unchanged.The only compiler warnings are pre-existing
nodiscardand deprecated-API notices onhipMemset,hipCtxGetCurrent,hipCtxSetCurrent,hipStreamDestroyand the runtime-version queries. None are introduced by this change.This pull request was prepared with the help of an AI assistant acting as a coding agent and was read and approved by a person before it was opened. It comes from an ongoing effort to add AMD GPU support to widely used CUDA projects, one repository at a time: https://github.com/AMD-Ecosystem/moat -- that repository describes how the work is done and what a person checks before anything is submitted.
If you would rather not receive pull requests from this effort, say so here or open an issue at https://github.com/AMD-Ecosystem/moat/issues/new/choose and we will close this and stop. That can cover this repository alone or everything you own, whichever you prefer.