Adds AMD backend to loops - #46
Merged
Merged
Conversation
add include/loops/util/xpu.hxx selecting cuda or hip and route every cudaXxx call (streams, events, memcpy, device attrs, occupancy, cooperative launch, errors) through xpu. group_mapped stays cuda-only since hip lacks the cooperative-groups block_tile_memory path. Co-authored-by: Cursor <cursoragent@cursor.com>
extend sm_flag_t with cdna (gfx906/908/90a/942/950) and rdna (gfx1030/1100/1200) bits, select the target via LOOPS_TARGET_GFX on hip builds, and give spmv wavefront/cache-aware launch params: 256-thread (4x wavefront-64) blocks on cdna, wider tiles on large-llc gfx942/gfx950. Co-authored-by: Cursor <cursoragent@cursor.com>
branch the build on LOOPS_BACKEND (cuda default, hip via hipcc/rocm): enable_language(HIP), link hip::host + roc::rocthrust, compile the shared .cu sources as hip, and map a single gfx target to LOOPS_TARGET_GFX. Co-authored-by: Cursor <cursoragent@cursor.com>
cub/block/block_scan.cuh was included but never used; removing it also unblocks the hip build (no cub on rocm). Co-authored-by: Cursor <cursoragent@cursor.com>
note the per-arch numbers driving the spmv tiles (cdna wavefront 64, 256- thread blocks; gfx90a no last-level cache vs 256 mb infinity cache on gfx942/gfx950) and add the rdna4 gfx1201 slot. Co-authored-by: Cursor <cursoragent@cursor.com>
release-mi210/mi300x/mi350x presets (gfx90a/gfx942/gfx950) and -Wno- unused-result on the hip pass for the fire-and-forget xpu wrappers. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
build the hip backend compile-only in a rocm 7.2.4 container across gfx90a/gfx942/gfx950 (release-mi210/mi300x/mi350x presets), so a hip break or a gfx-specific divergence is caught without amd hardware. add a fast parity job that diffs the loops::xpu symbol set across backend/cuda.hxx and hip.hxx to catch "added to one, forgot the other". Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
a thread's first complete row can continue a row opened in the previous partition, whose head is emitted there as a remainder atomic; storing the first row directly raced with that atomic and intermittently dropped the head term. fold the first row in with atomicAdd onto the zero-initialized output, leaving interior rows as direct stores. Co-authored-by: Cursor <cursoragent@cursor.com>
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 pull request introduces comprehensive support for building and running the project on AMD GPUs using the HIP/ROCm backend, in addition to the existing CUDA/NVIDIA support. The changes include backend selection in the build system, new CMake presets for AMD architectures, CI integration for HIP builds, and code updates to ensure backend-agnostic usage of GPU streams and synchronization. Some CUDA-specific examples are now conditionally excluded from HIP builds.
Build system and configuration:
LOOPS_BACKENDCMake option to select between CUDA (default) and HIP backends, with corresponding logic for toolkit and architecture selection, dependency finding, compile flags, and target properties inCMakeLists.txt. This enables seamless building on either NVIDIA or AMD platforms. [1] [2] [3] [4] [5] [6] [7] [8] [9]CMakePresets.json, enabling easy configuration and build for MI210/MI250, MI300X/MI325X, and MI350X architectures. [1] [2]Continuous Integration:
.github/workflows/build.yml)Codebase portability and backend-agnostic changes:
xpu::stream_tandxpu::stream_synchronizeinstead of CUDA-specific types and functions, allowing the same code to work with both CUDA and HIP backends. [1] [2] [3] [4] [5] [6] [7] [8]Conditional example inclusion:
group_mapped.cuexample when building with HIP, as it relies on NVIDIA cooperative-groups, which are not available in HIP. [1] [2]These changes collectively enable robust, cross-vendor GPU support and improve maintainability and CI coverage for both CUDA and HIP backends.