From 45df0148c07231f2a2b8b7ed76f4cebd413c5d53 Mon Sep 17 00:00:00 2001 From: Kerelos Gayed Date: Thu, 1 Oct 2026 07:31:35 -0400 Subject: [PATCH] Add AMD HIP/ROCm backend Add HIP build selection, runtime operations, hipCUB algorithms, hipRAND generators, and hipSOLVER linear algebra. Share the GPU test suite across CUDA and HIP and document the toolchain requirements and opt-in ROCm CI. --- .github/workflows/test.yml | 13 +++ CMakeLists.txt | 103 +++++++++++++++++++++++- README.md | 96 ++++++++++++++++------ scripts/test.sh | 63 +++++++++++---- tests/CMakeLists.txt | 20 ++++- tests/buffer/cases.hpp | 2 +- tests/buffer/cuda.cu | 9 ++- tests/checked/cases.hpp | 2 +- tests/config/cpu.cpp | 4 +- tests/config/cuda.cu | 4 +- tests/launch/cuda.cu | 7 +- tests/math/cuda.cu | 9 ++- tests/memory/cuda.cu | 9 ++- tests/random/cuda.cu | 3 +- tests/support/device.hpp | 29 +++++++ xpu/algorithm.hpp | 27 +++++++ xpu/config.hpp | 45 +++++++++-- xpu/detail/checked.hpp | 3 + xpu/launch.hpp | 93 ++++++++++++++++++---- xpu/linear_algebra.hpp | 158 ++++++++++++++++++++++--------------- xpu/math.hpp | 10 +-- xpu/memory.hpp | 9 +++ xpu/random.hpp | 23 +++++- 23 files changed, 581 insertions(+), 160 deletions(-) create mode 100644 tests/support/device.hpp diff --git a/.github/workflows/test.yml b/.github/workflows/test.yml index 93ef72e..9dbdedc 100644 --- a/.github/workflows/test.yml +++ b/.github/workflows/test.yml @@ -45,3 +45,16 @@ jobs: - name: Test run: ./scripts/test.sh --cu --clean + + hip: + name: HIP + if: ${{ vars.XPU_HIP_CI == 'true' && github.event_name != 'pull_request' }} + runs-on: [self-hosted, linux, x64, rocm] + timeout-minutes: 30 + + steps: + - name: Checkout + uses: actions/checkout@v6 + + - name: Test + run: ./scripts/test.sh --hip --clean diff --git a/CMakeLists.txt b/CMakeLists.txt index 50c6876..168f272 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -14,6 +14,7 @@ endif() project(xpu LANGUAGES CXX) option(XPU_ENABLE_CUDA "Build the CUDA backend if a CUDA compiler is available" ON) +option(XPU_ENABLE_HIP "Build the HIP backend if a HIP compiler is available and CUDA is not used" ON) option(XPU_ENABLE_LINALG "Build the optional linear-algebra component" OFF) option(XPU_ARCH_NATIVE "Tune for the host architecture (-march=native)" ON) option(XPU_BUILD_TESTS "Build the xpu tests" ${PROJECT_IS_TOP_LEVEL}) @@ -64,10 +65,89 @@ if(XPU_ENABLE_CUDA) "xpu: CUDA ${CMAKE_CUDA_COMPILER_VERSION}, arch ${CMAKE_CUDA_ARCHITECTURES}, " "host ${CMAKE_CXX_COMPILER_ID} ${CMAKE_CXX_COMPILER_VERSION}") else() - message(STATUS "xpu: no CUDA compiler found, building the CPU-only path") + message(STATUS "xpu: no CUDA compiler found") endif() endif() +# HIP +set(XPU_HIP OFF) + +if(XPU_ENABLE_HIP AND NOT XPU_CUDA) + include(CheckLanguage) + check_language(HIP) + + if(CMAKE_HIP_COMPILER) + enable_language(HIP) + + if(DEFINED CMAKE_HIP_PLATFORM AND NOT CMAKE_HIP_PLATFORM STREQUAL "amd") + message(FATAL_ERROR + "xpu's HIP backend targets AMD GPUs, but CMAKE_HIP_PLATFORM is ${CMAKE_HIP_PLATFORM}. " + "Use the CUDA backend on NVIDIA, or configure with -DXPU_ENABLE_HIP=OFF.") + endif() + + set(CMAKE_HIP_EXTENSIONS OFF) + + # Let the ROCm packages, and the dependencies they look up, find ROCm + list(APPEND CMAKE_PREFIX_PATH ${CMAKE_HIP_COMPILER_ROCM_ROOT} $ENV{ROCM_PATH} /opt/rocm) + + # hip-config probes for GPUs unless GPU_TARGETS is set, so reuse CMake's list + if(NOT DEFINED GPU_TARGETS) + set(GPU_TARGETS ${CMAKE_HIP_ARCHITECTURES}) + endif() + + foreach(package IN ITEMS hip rocprim hipcub hiprand) + find_package(${package} CONFIG) + + if(NOT ${package}_FOUND) + message(FATAL_ERROR + "xpu: found a HIP compiler but not the ROCm ${package} package. Point ROCM_PATH or " + "CMAKE_PREFIX_PATH at ROCm, or configure with -DXPU_ENABLE_HIP=OFF to build the CPU-only path.") + endif() + endforeach() + + if(hip_VERSION VERSION_LESS 7.0) + message(FATAL_ERROR + "xpu needs ROCm 7.0 or newer for the HIP backend (found HIP ${hip_VERSION}). " + "Configure with -DXPU_ENABLE_HIP=OFF to build the CPU-only path.") + endif() + + # hipCUB's C++23 path uses std::extents, which needs from the standard library + include(CheckSourceCompiles) + block() + # try_compile adds a -std of its own from the CMAKE_HIP_* standard settings after + # CMAKE_REQUIRED_FLAGS, so pin that one to C++23 as well + set(CMAKE_HIP_STANDARD 23) + set(CMAKE_REQUIRED_FLAGS -std=c++23) + set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY) + check_source_compiles(HIP [=[ +#include +int main() { return static_cast(std::extents::rank()) - 1; } +]=] XPU_HIP_HAS_MDSPAN) + endblock() + + if(NOT XPU_HIP_HAS_MDSPAN) + message(FATAL_ERROR + "The HIP compiler's C++ standard library has no std::extents in , which hipCUB " + "needs under C++23. Install GCC 15 (ROCm's clang uses the newest GCC it finds) or use " + "libc++, then reconfigure in a fresh build directory.") + endif() + + set(XPU_HIP ON) + message(STATUS + "xpu: HIP ${hip_VERSION}, arch ${CMAKE_HIP_ARCHITECTURES}, " + "compiler ${CMAKE_HIP_COMPILER_ID} ${CMAKE_HIP_COMPILER_VERSION}") + else() + message(STATUS "xpu: no HIP compiler found") + endif() +endif() + +if(XPU_CUDA OR XPU_HIP) + set(XPU_GPU ON) +else() + set(XPU_GPU OFF) + message(STATUS "xpu: building the CPU-only path") +endif() + # Targets add_library(xpu INTERFACE) add_library(xpu::xpu ALIAS xpu) @@ -84,6 +164,16 @@ if(XPU_CUDA) target_link_libraries(xpu INTERFACE CUDA::cudart) target_compile_options(xpu INTERFACE $<$:-std=c++23>) +elseif(XPU_HIP) + target_compile_definitions(xpu INTERFACE XPU_HIP) + # hip::hipcub pulls in hip::device, whose -x hip and --offload-arch would land on every CXX + # source of a consumer, so take the hipCUB and rocPRIM headers without it + target_link_libraries(xpu INTERFACE hip::host roc::rocprim hip::hiprand) + target_include_directories(xpu SYSTEM INTERFACE + $ + ) + + target_compile_options(xpu INTERFACE $<$:-std=c++23>) else() find_package(OpenMP QUIET COMPONENTS CXX) if(OpenMP_CXX_FOUND) @@ -103,6 +193,16 @@ if(XPU_ENABLE_LINALG) CUDA::cusolver CUDA::cublas ) + elseif(XPU_HIP) + find_package(hipsolver CONFIG) + + if(NOT hipsolver_FOUND) + message(FATAL_ERROR + "xpu: XPU_ENABLE_LINALG on the HIP backend needs the ROCm hipsolver package. " + "Install it, or configure with -DXPU_ENABLE_LINALG=OFF.") + endif() + + target_link_libraries(xpu_linalg INTERFACE roc::hipsolver) else() find_package(LAPACK REQUIRED) find_path(XPU_LAPACKE_INCLUDE_DIR NAMES lapacke.h REQUIRED) @@ -122,6 +222,7 @@ if(XPU_ARCH_NATIVE AND NOT MSVC) target_compile_options(xpu INTERFACE $<$:-march=native> $<$:-Xcompiler=-march=native> + $<$:-march=native> ) endif() diff --git a/README.md b/README.md index 76c303a..6627589 100644 --- a/README.md +++ b/README.md @@ -1,8 +1,8 @@ # xpu -`xpu` is a small, header-only C++ library for code that runs on a CPU or NVIDIA -CUDA. It provides backend-aware allocation, contiguous buffers, -structure-of-arrays storage, math helpers, and basic CUDA launch utilities. +`xpu` is a small, header-only C++ library for code that runs on a CPU, NVIDIA +CUDA, or AMD HIP. It provides backend-aware allocation, contiguous buffers, +structure-of-arrays storage, math helpers, and basic GPU launch utilities. The project is in early development. The API may change. @@ -12,8 +12,13 @@ The project is in early development. The API may change. - CMake 3.25 or newer - CUDA 13.3 or newer for the CUDA backend - A compatible CUDA host compiler, with GCC 15 or newer when GCC is used +- ROCm 7.0 or newer for the HIP backend, with hipCUB, rocPRIM and hipRAND +- A C++ standard library with `` for HIP builds: libstdc++ from GCC 15 + or newer, or libc++ - LAPACKE for the optional CPU linear-algebra component +- hipSOLVER for the optional HIP linear-algebra component - Linux, or Windows through WSL2, for CUDA builds +- Linux for HIP builds The base CPU-only library has no external dependencies. CPU builds use OpenMP when CMake finds it; otherwise execution is serial. @@ -45,12 +50,30 @@ cmake -S . -B build -G Ninja \ cmake --build build ``` +HIP is enabled by default when CMake finds ROCm's `clang++` and CUDA is not in +use. CUDA takes precedence when both toolchains are available. Point CMake at +ROCm's compiler and set the target GPU architecture explicitly: + +```bash +cmake -S . -B build-hip -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ + -DCMAKE_HIP_ARCHITECTURES=gfx942 \ + -DXPU_ENABLE_CUDA=OFF + +cmake --build build-hip +``` + +Without `CMAKE_HIP_ARCHITECTURES`, CMake targets the GPUs that +`rocm_agent_enumerator` reports, or the compiler's default when it finds none. + For a CPU-only build: ```bash cmake -S . -B build-cpu -G Ninja \ -DCMAKE_BUILD_TYPE=Release \ - -DXPU_ENABLE_CUDA=OFF + -DXPU_ENABLE_CUDA=OFF \ + -DXPU_ENABLE_HIP=OFF cmake --build build-cpu ``` @@ -63,6 +86,7 @@ sudo apt install liblapacke-dev cmake -S . -B build-cpu-linalg -G Ninja \ -DCMAKE_BUILD_TYPE=Release \ -DXPU_ENABLE_CUDA=OFF \ + -DXPU_ENABLE_HIP=OFF \ -DXPU_ENABLE_LINALG=ON cmake --build build-cpu-linalg @@ -77,8 +101,8 @@ target_link_libraries(my_target PRIVATE xpu::xpu) Link `xpu::linalg` instead when using ``. -Set `XPU_ENABLE_CUDA` and `XPU_ENABLE_LINALG` before `add_subdirectory` when -you need to select them explicitly. +Set `XPU_ENABLE_CUDA`, `XPU_ENABLE_HIP` and `XPU_ENABLE_LINALG` before +`add_subdirectory` when you need to select them explicitly. ## Buffers @@ -98,8 +122,8 @@ const auto count{values.count()}; const auto capacity{values.capacity()}; ``` -`count()` is the requested number of elements. `capacity()` includes any CPU -padding. +`count()` is the requested number of elements. `capacity()` includes any +alignment padding. ## Structure of arrays @@ -150,31 +174,40 @@ is valid only while the original `soa` owns the allocation. ## Backend and memory model -`XPU_CUDA` selects the allocation backend: +`XPU_CUDA` or `XPU_HIP` selects the allocation backend: | Backend | Allocation | |---|---| | CPU | aligned `operator new` | | CUDA | `cudaMalloc` | +| HIP | `hipMalloc` | -The `xpu::xpu` CMake target sets `XPU_CUDA` for CUDA builds. Do not set it on -individual source files. It must have the same value in every translation unit -linked into a program. +The `xpu::xpu` CMake target sets `XPU_CUDA` for CUDA builds and `XPU_HIP` for +HIP builds. Do not set them on individual source files. They must have the same +value in every translation unit linked into a program. `` +defines `XPU_GPU` for either GPU backend, and `xpu::xpu_cuda`, `xpu::xpu_hip` +and `xpu::xpu_gpu` expose the same choice as constants. When CUDA is enabled, every translation unit that includes an xpu header must be compiled by nvcc. These files normally use the `.cu` extension. -CUDA allocations are device memory. Pointers returned by `buffer` and `soa` -cannot be dereferenced by host code. The library does not currently wrap memory -transfers, so use the CUDA runtime directly when transfers are required. +When HIP is enabled, every translation unit that includes an xpu header must be +compiled as HIP. Use the `.hip` extension or set the `LANGUAGE HIP` source file +property. The test suite does the latter to reuse its `.cu` sources. + +GPU allocations are device memory. Pointers returned by `buffer` and `soa` +cannot be dereferenced by host code. Use `xpu::copy_n` or `xpu::memcpy` to move +data between host and device memory; the runtime infers the direction. Allocation failure terminates the process with `std::abort`. ## Padding -CPU allocations are aligned to at least `xpu::simd_bytes`. Capacities and SoA -strides are padded to SIMD-lane multiples when the element type is smaller than -the SIMD width. CUDA uses a tight layout with no row padding. +CPU allocations are aligned to at least `xpu::simd_bytes`. GPU builds use +`xpu::cuda_align_bytes`, which is 128 bytes, instead. When the element type is +smaller than that alignment, capacities and SoA strides are padded to a multiple +of `alignment / sizeof(T)` elements (integer division), which fills whole +alignment blocks when `sizeof(T)` is a power of two. The default SIMD width is 64 bytes with AVX-512, 32 bytes with AVX or AVX2, and 16 bytes otherwise. Pin it when layout must remain stable across machines: @@ -203,32 +236,43 @@ if ( Use `solve()` for linear systems. Use `invert()` only when the inverse itself is required. Matrices use row-major layout, and strides are measured in -elements. +elements. The component uses LAPACKE on the CPU, cuSOLVER on CUDA, and +hipSOLVER on HIP. ## Testing ```bash -./scripts/test.sh # CPU and CUDA test suites +./scripts/test.sh # every available suite: CPU, CUDA, then HIP ./scripts/test.sh --sanitize # CUDA Compute Sanitizer ./scripts/test.sh --cpp # CPU test suite only ./scripts/test.sh --cu # CUDA test suite only +./scripts/test.sh --hip # HIP test suite only ``` -Behavioral tests are grouped by component under `tests/`. CPU and CUDA entry -points are kept separate, while backend-neutral cases live beside them in a +The script finds ROCm's `clang++` through `HIPCXX` or +`$ROCM_PATH/llvm/bin/clang++`, where `ROCM_PATH` defaults to `/opt/rocm`. It +skips a GPU suite whose compiler is missing unless that suite was requested +explicitly. `--sanitize` applies to the CUDA suite only. + +Behavioral tests are grouped by component under `tests/`. CPU and GPU entry +points are kept separate in `cpu.cpp` and `cuda.cu`, and HIP builds compile the +`cuda.cu` entry points as HIP. Backend-neutral cases live beside them in a shared `cases.hpp`. Common test support lives in `tests/support`, umbrella-header coverage lives in `tests/integration`, and every exported header also gets a compile-only self-containment check. GitHub Actions runs the CPU suite on Ubuntu 26.04 with GCC 15. CUDA runtime testing is enabled when the repository variable `XPU_CUDA_CI` is `true` and a -self-hosted Linux x64 runner with the `gpu` label is available. +self-hosted Linux x64 runner with the `gpu` label is available. HIP runtime +testing works the same way with the `XPU_HIP_CI` variable and a runner with the +`rocm` label. ## Current limitations -- NVIDIA CUDA is the only GPU backend. -- CPU and CUDA backends cannot be mixed in one linked program. -- CUDA memory transfers are not wrapped. +- A build uses one backend. CPU, CUDA and HIP cannot be mixed in one linked + program. +- Random sequences are backend-specific. The same seed need not produce the + same values on the CPU, CUDA and HIP. - Accessors do not perform bounds checking. - Multidimensional launch configuration is still a work in progress. - Installed CMake package metadata is incomplete. diff --git a/scripts/test.sh b/scripts/test.sh index 69f1566..e9098be 100755 --- a/scripts/test.sh +++ b/scripts/test.sh @@ -1,8 +1,9 @@ #!/usr/bin/env bash # -# ./scripts/test.sh both paths: tests/smoke.cpp then tests/smoke.cu +# ./scripts/test.sh every available path: CPU, then CUDA, then HIP # ./scripts/test.sh --cpp CPU path only # ./scripts/test.sh --cu CUDA path only +# ./scripts/test.sh --hip HIP path only # ./scripts/test.sh --sanitize also run compute-sanitizer on the CUDA path # ./scripts/test.sh --clean wipe the build dirs first @@ -18,6 +19,7 @@ for arg in "$@"; do case $arg in --cpp) paths+=(cpp) ;; --cu) paths+=(cu) ;; + --hip) paths+=(hip) ;; --sanitize) sanitize=1 ;; --clean) clean=1 ;; *) echo "unknown option: $arg" >&2; exit 2 ;; @@ -27,7 +29,7 @@ done explicit=1 if [[ ${#paths[@]} == 0 ]]; then explicit=0 - paths=(cpp cu) + paths=(cpp cu hip) fi wants() { @@ -38,11 +40,24 @@ wants() { return 1 } +skip() { + local unwanted=$1 p kept=() + for p in "${paths[@]}"; do + [[ $p == "$unwanted" ]] || kept+=("$p") + done + paths=("${kept[@]}") +} + nvcc="${CUDA_HOME:-/usr/local/cuda}/bin/nvcc" if ! [[ -x $nvcc ]]; then nvcc=$(command -v nvcc || true) fi +hipcxx=${HIPCXX:-${ROCM_PATH:-/opt/rocm}/llvm/bin/clang++} +if [[ $hipcxx != /* ]]; then + hipcxx=$(command -v "$hipcxx" || true) +fi + cxx=${CXX:-} if [[ -z $cxx ]]; then cxx=$(command -v g++-15 || command -v g++ || true) @@ -55,15 +70,24 @@ if [[ -z $cxx || ! -x $cxx ]]; then exit 2 fi -# a CUDA-less configure silently falls back to smoke.cpp, which would look like -# the .cu path passing when it never got compiled +# a configure without the GPU compiler silently falls back to the CPU tests, +# which would look like the GPU path passing when it never got compiled if wants cu && [[ -z $nvcc ]]; then if [[ $explicit == 1 ]]; then echo "--cu needs nvcc on PATH or at \$CUDA_HOME/bin/nvcc" >&2 exit 2 fi echo "xpu: no nvcc found, skipping the CUDA path" - paths=(cpp) + skip cu +fi + +if wants hip && [[ -z $hipcxx || ! -x $hipcxx ]]; then + if [[ $explicit == 1 ]]; then + echo "--hip needs ROCm's clang++ at \$HIPCXX or \$ROCM_PATH/llvm/bin/clang++" >&2 + exit 2 + fi + echo "xpu: no ROCm clang++ found, skipping the HIP path" + skip hip fi if [[ $sanitize == 1 ]] && ! wants cu; then @@ -72,7 +96,8 @@ if [[ $sanitize == 1 ]] && ! wants cu; then fi run() { - local label=$1 cuda=$2 build=$3 + local label=$1 build=$2 + shift 2 echo echo "=== $label path ($build) ===" @@ -85,28 +110,27 @@ run() { -S . -B "$build" -G Ninja -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_COMPILER="$cxx" - -DXPU_ENABLE_CUDA=$cuda -DXPU_ENABLE_LINALG=ON + "$@" ) - if [[ $cuda == ON ]]; then - args+=( - -DCMAKE_CUDA_COMPILER="$nvcc" - -DCMAKE_CUDA_HOST_COMPILER="$cxx" - ) - fi - cmake "${args[@]}" cmake --build "$build" ctest --test-dir "$build" --output-on-failure --no-tests=ignore } if wants cpp; then - run CPU OFF build-cpp + run CPU build-cpp \ + -DXPU_ENABLE_CUDA=OFF \ + -DXPU_ENABLE_HIP=OFF fi if wants cu; then - run CUDA ON build-cuda + run CUDA build-cuda \ + -DXPU_ENABLE_CUDA=ON \ + -DXPU_ENABLE_HIP=OFF \ + -DCMAKE_CUDA_COMPILER="$nvcc" \ + -DCMAKE_CUDA_HOST_COMPILER="$cxx" if [[ $sanitize == 1 ]]; then echo @@ -120,3 +144,10 @@ if wants cu; then done fi fi + +if wants hip; then + run HIP build-hip \ + -DXPU_ENABLE_CUDA=OFF \ + -DXPU_ENABLE_HIP=ON \ + -DCMAKE_HIP_COMPILER="$hipcxx" +fi diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index a3c6d84..9e4c943 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -6,6 +6,11 @@ function(xpu_configure_test_target target) CUDA_STANDARD 23 CUDA_STANDARD_REQUIRED ON ) + elseif(XPU_HIP) + set_target_properties(${target} PROPERTIES + HIP_STANDARD 23 + HIP_STANDARD_REQUIRED ON + ) else() set_target_properties(${target} PROPERTIES CXX_STANDARD 23 @@ -19,7 +24,7 @@ function(xpu_add_backend_test module) return() endif() - if(XPU_CUDA) + if(XPU_GPU) set(source ${module}/cuda.cu) else() set(source ${module}/cpu.cpp) @@ -29,14 +34,19 @@ function(xpu_add_backend_test module) return() endif() + # The GPU tests are shared, so HIP builds compile the .cu sources as HIP + if(XPU_HIP) + set_source_files_properties(${source} PROPERTIES LANGUAGE HIP) + endif() + set(target xpu_test_${module}) add_executable(${target} ${source}) xpu_configure_test_target(${target}) - if(module STREQUAL "math" AND NOT XPU_CUDA) + if(module STREQUAL "math" AND NOT XPU_GPU) find_package(Threads REQUIRED) target_link_libraries(${target} PRIVATE Threads::Threads) endif() - if(module STREQUAL "reduce" AND NOT XPU_CUDA) + if(module STREQUAL "reduce" AND NOT XPU_GPU) find_package(OpenMP QUIET) if(OpenMP_CXX_FOUND) target_link_libraries(${target} PRIVATE OpenMP::OpenMP_CXX) @@ -46,7 +56,7 @@ function(xpu_add_backend_test module) target_link_libraries(${target} PRIVATE xpu::linalg) endif() add_test(NAME ${target} COMMAND ${target}) - if(module STREQUAL "reduce" AND NOT XPU_CUDA) + if(module STREQUAL "reduce" AND NOT XPU_GPU) set_tests_properties(${target} PROPERTIES ENVIRONMENT "OMP_NUM_THREADS=4") endif() endfunction() @@ -86,6 +96,8 @@ foreach(header IN ITEMS if(XPU_CUDA) set(extension cu) + elseif(XPU_HIP) + set(extension hip) else() set(extension cpp) endif() diff --git a/tests/buffer/cases.hpp b/tests/buffer/cases.hpp index 491ed3e..52964f6 100644 --- a/tests/buffer/cases.hpp +++ b/tests/buffer/cases.hpp @@ -78,7 +78,7 @@ inline int run_buffer_cases() { return test::fail("non-empty buffer has a null data pointer"); } - if constexpr (!xpu::xpu_cuda) { + if constexpr (!xpu::xpu_gpu) { view[0uz] = 7.0f; if (&view[0uz] != values.data() || readonly[0uz] != 7.0f) { return test::fail("buffer owner view does not refer to its storage"); diff --git a/tests/buffer/cuda.cu b/tests/buffer/cuda.cu index 2d01b52..a365223 100644 --- a/tests/buffer/cuda.cu +++ b/tests/buffer/cuda.cu @@ -1,4 +1,5 @@ #include "cases.hpp" +#include "../support/device.hpp" namespace { @@ -24,16 +25,16 @@ auto run_device_buffer_view_cases() -> int { auto values{xpu::buffer{count}}; auto output{xpu::buffer{count}}; write_buffer_view<<<1, 32>>>(values.view()); - xpu::cu_check(cudaGetLastError()); + test::check_launch(); read_buffer_view<<<1, 32>>>(std::as_const(values).view(), output.view()); - xpu::cu_check(cudaGetLastError()); - xpu::cu_check(cudaDeviceSynchronize()); + test::check_launch(); + test::synchronize(); int result[count]{}; xpu::copy_n(result, output.data(), count); for (auto i{0uz}; i < count; ++i) { if (result[i] != static_cast(2uz * (i + 1uz))) { - return test::fail("CUDA buffer view indexing is incorrect"); + return test::fail("GPU buffer view indexing is incorrect"); } } return 0; diff --git a/tests/checked/cases.hpp b/tests/checked/cases.hpp index 5fba676..9633dbd 100644 --- a/tests/checked/cases.hpp +++ b/tests/checked/cases.hpp @@ -98,7 +98,7 @@ inline auto run_checked_cases() -> int { return failure; } -#if !defined(XPU_CUDA) +#if !defined(XPU_GPU) if (const auto failure{test::check_abort([] { const auto range = xpu::range<2uz>{ {0uz, 0uz}, diff --git a/tests/config/cpu.cpp b/tests/config/cpu.cpp index 0dcba66..3139bf5 100644 --- a/tests/config/cpu.cpp +++ b/tests/config/cpu.cpp @@ -4,8 +4,8 @@ int main() { if (const auto status{check_simd_width()}; status != 0) { return status; } - if (xpu::xpu_cuda) { - return test::fail("CPU configuration flag is incorrect"); + if (xpu::xpu_cuda || xpu::xpu_hip || xpu::xpu_gpu) { + return test::fail("CPU configuration flags are incorrect"); } return 0; diff --git a/tests/config/cuda.cu b/tests/config/cuda.cu index 2eddc0a..0af777a 100644 --- a/tests/config/cuda.cu +++ b/tests/config/cuda.cu @@ -4,8 +4,8 @@ int main() { if (const auto status{check_simd_width()}; status != 0) { return status; } - if (!xpu::xpu_cuda) { - return test::fail("CUDA configuration flag is incorrect"); + if (!xpu::xpu_gpu || xpu::xpu_cuda == xpu::xpu_hip) { + return test::fail("GPU configuration flags are incorrect"); } return 0; diff --git a/tests/launch/cuda.cu b/tests/launch/cuda.cu index 39a7eee..75587c7 100644 --- a/tests/launch/cuda.cu +++ b/tests/launch/cuda.cu @@ -1,5 +1,6 @@ #include "../support/check.hpp" #include "../support/aborts.hpp" +#include "../support/device.hpp" #include @@ -57,8 +58,8 @@ int main() { write_launch_coordinates<<>>( indices.data(), strides.data(), test::count ); - xpu::cu_check(cudaGetLastError()); - xpu::cu_check(cudaDeviceSynchronize()); + test::check_launch(); + test::synchronize(); auto host_indices{std::make_unique(test::count)}; auto host_strides{std::make_unique(test::count)}; @@ -67,7 +68,7 @@ int main() { for (auto i{0uz}; i < test::count; ++i) { if (host_indices[i] != i || host_strides[i] != expected_stride) { - return test::fail("CUDA launch coordinates are incorrect"); + return test::fail("GPU launch coordinates are incorrect"); } } diff --git a/tests/math/cuda.cu b/tests/math/cuda.cu index 43ad8d3..0f9ce65 100644 --- a/tests/math/cuda.cu +++ b/tests/math/cuda.cu @@ -1,4 +1,5 @@ #include "cases.hpp" +#include "../support/device.hpp" #include #include @@ -47,8 +48,8 @@ int main() { add_under_contention<<>>( atomic_result.data(), atomic_add_count ); - xpu::cu_check(cudaGetLastError()); - xpu::cu_check(cudaDeviceSynchronize()); + test::check_launch(); + test::synchronize(); auto host_atomic_result{0}; xpu::copy_n(&host_atomic_result, atomic_result.data(), 1uz); @@ -67,8 +68,8 @@ int main() { compute_inverse_norms<<>>( result.data(), a.data(), b.data(), result.count() ); - xpu::cu_check(cudaGetLastError()); - xpu::cu_check(cudaDeviceSynchronize()); + test::check_launch(); + test::synchronize(); auto host_result{std::make_unique(test::count)}; xpu::copy_n(host_result.get(), result.data(), test::count); diff --git a/tests/memory/cuda.cu b/tests/memory/cuda.cu index 721191f..d0aabd9 100644 --- a/tests/memory/cuda.cu +++ b/tests/memory/cuda.cu @@ -1,4 +1,5 @@ #include "cases.hpp" +#include "../support/device.hpp" namespace { @@ -25,10 +26,10 @@ auto run_device_soa_view_cases() -> int { auto values{xpu::soa{count}}; auto output{xpu::soa{count}}; write_soa_view<<<1, 32>>>(values.view()); - xpu::cu_check(cudaGetLastError()); + test::check_launch(); read_soa_view<<<1, 32>>>(values.view(), output.view()); - xpu::cu_check(cudaGetLastError()); - xpu::cu_check(cudaDeviceSynchronize()); + test::check_launch(); + test::synchronize(); int result[2][count]{}; xpu::copy_n(result[0], output.view()[0], count); @@ -36,7 +37,7 @@ auto run_device_soa_view_cases() -> int { for (auto i{0uz}; i < count; ++i) { if (result[0][i] != static_cast(2uz * (i + 1uz)) || result[1][i] != static_cast((2uz * (i + 1uz)) + 1uz)) { - return test::fail("CUDA SoA View Indexing is incorrect"); + return test::fail("GPU SoA View Indexing is incorrect"); } } return 0; diff --git a/tests/random/cuda.cu b/tests/random/cuda.cu index f0eaf75..883c896 100644 --- a/tests/random/cuda.cu +++ b/tests/random/cuda.cu @@ -1,4 +1,5 @@ #include "cases.hpp" +#include "../support/device.hpp" #include #include @@ -30,7 +31,7 @@ int main() { auto samples{xpu::buffer{random_sample_count}}; generate_random_samples<<<1u, 1u>>>(samples.data()); - xpu::cu_check(cudaGetLastError()); + test::check_launch(); auto host_samples = std::array{}; xpu::copy_n( diff --git a/tests/support/device.hpp b/tests/support/device.hpp new file mode 100644 index 0000000..08914b7 --- /dev/null +++ b/tests/support/device.hpp @@ -0,0 +1,29 @@ +#pragma once + +#include + +#include + +namespace test { + +inline auto check_launch( + std::source_location loc = std::source_location::current() +) -> void { +#if defined(XPU_CUDA) + xpu::cu_check(cudaGetLastError(), cudaSuccess, loc); +#else + xpu::cu_check(hipGetLastError(), hipSuccess, loc); +#endif +} + +inline auto synchronize( + std::source_location loc = std::source_location::current() +) -> void { +#if defined(XPU_CUDA) + xpu::cu_check(cudaDeviceSynchronize(), cudaSuccess, loc); +#else + xpu::cu_check(hipDeviceSynchronize(), hipSuccess, loc); +#endif +} + +} // namespace test diff --git a/xpu/algorithm.hpp b/xpu/algorithm.hpp index 60af72c..49a6f9f 100644 --- a/xpu/algorithm.hpp +++ b/xpu/algorithm.hpp @@ -8,6 +8,11 @@ #include #include #include +#elif defined(XPU_HIP) + // hipCUB's C++23 path names std::extents without including + #include + #include + #include #else #include #include @@ -17,6 +22,23 @@ namespace xpu { +#if defined(XPU_HIP) +namespace detail { + +template +struct fill_function { + T* ptr; + T value; + + XPU_DEVICE_ONLY + auto operator()(std::size_t linear) const -> void { + ptr[linear] = value; + } +}; + +} // namespace xpu::detail +#endif + template inline auto fill_n( T* XPU_RESTRICT ptr, @@ -25,6 +47,11 @@ inline auto fill_n( ) -> void { #if defined(XPU_CUDA) xpu::cu_check(cub::DeviceTransform::Fill(ptr, count, value)); +#elif defined(XPU_HIP) + if (count == 0uz) { return; } + + // hipCUB has no DeviceTransform::Fill + xpu::cu_check(hipcub::DeviceFor::Bulk(count, detail::fill_function{ptr, value})); #else std::fill_n(ptr, count, value); #endif diff --git a/xpu/config.hpp b/xpu/config.hpp index aea5c96..c895536 100644 --- a/xpu/config.hpp +++ b/xpu/config.hpp @@ -1,12 +1,30 @@ #pragma once +#if defined(XPU_CUDA) && defined(XPU_HIP) + #error "ERROR: XPU_CUDA and XPU_HIP cannot both be defined." +#endif + #if defined(XPU_CUDA) && !defined(__CUDACC__) #error "ERROR: must use nvcc when compiling with the XPU_CUDA flag." #endif +#if defined(XPU_HIP) && !defined(__HIP__) + #error "ERROR: must compile as HIP when compiling with the XPU_HIP flag." +#endif + +#if defined(XPU_CUDA) || defined(XPU_HIP) + #define XPU_GPU 1 +#endif + +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + #define XPU_DEVICE_COMPILE 1 +#endif + #if defined(XPU_CUDA) #include #include +#elif defined(XPU_HIP) + #include #endif #include @@ -24,7 +42,7 @@ namespace xpu { -#if defined(XPU_CUDA) +#if defined(XPU_GPU) template inline auto cu_check( @@ -34,9 +52,16 @@ inline auto cu_check( ) noexcept -> void { if (result == success) { return; } +#if defined(XPU_CUDA) + constexpr auto backend{"CUDA"}; +#else + constexpr auto backend{"HIP"}; +#endif + std::fprintf( stderr, - "CUDA error at %s:%u in %s\n status code: %d\n", + "%s error at %s:%u in %s\n status code: %d\n", + backend, loc.file_name(), loc.line(), loc.function_name(), @@ -56,13 +81,13 @@ inline auto cu_check( #define XPU_RESTRICT #endif -#if defined(XPU_CUDA) +#if defined(XPU_GPU) #define XPU_CUDA_CALLABLE __host__ __device__ #else #define XPU_CUDA_CALLABLE #endif -#if defined(XPU_CUDA) +#if defined(XPU_GPU) #define XPU_DEVICE_ONLY __device__ #else #define XPU_DEVICE_ONLY @@ -101,8 +126,18 @@ inline constexpr auto xpu_cuda{ #endif }; +inline constexpr auto xpu_hip{ +#if defined(XPU_HIP) + true +#else + false +#endif +}; + +inline constexpr auto xpu_gpu{xpu_cuda || xpu_hip}; + inline constexpr auto alignment_bytes{ - xpu_cuda ? cuda_align_bytes : simd_bytes + xpu_gpu ? cuda_align_bytes : simd_bytes }; } // namespace xpu diff --git a/xpu/detail/checked.hpp b/xpu/detail/checked.hpp index 5121ba4..456b83f 100644 --- a/xpu/detail/checked.hpp +++ b/xpu/detail/checked.hpp @@ -34,6 +34,9 @@ inline auto checked_error(const char* message) noexcept -> void { #if defined(__CUDA_ARCH__) printf("xpu: %s\n", message); __trap(); +#elif defined(__HIP_DEVICE_COMPILE__) + printf("xpu: %s\n", message); + __builtin_trap(); #else std::fprintf(stderr, "xpu: %s\n", message); std::abort(); diff --git a/xpu/launch.hpp b/xpu/launch.hpp index 2e283af..a9d284e 100644 --- a/xpu/launch.hpp +++ b/xpu/launch.hpp @@ -15,13 +15,21 @@ #include #include #include +#elif defined(XPU_HIP) + // hipCUB's C++23 path names std::extents without including + #include + #include + #include + #include + #include + #include #else #include #endif namespace xpu { -#if defined(XPU_CUDA) && defined(__CUDACC__) +#if defined(XPU_GPU) namespace detail { @@ -29,10 +37,17 @@ namespace detail { inline auto device_SMs() -> unsigned int { static const auto cached{[] { auto device{0}, sms{0}; +#if defined(XPU_CUDA) xpu::cu_check(cudaGetDevice(&device)); xpu::cu_check(cudaDeviceGetAttribute( &sms, cudaDevAttrMultiProcessorCount, device )); +#else + xpu::cu_check(hipGetDevice(&device)); + xpu::cu_check(hipDeviceGetAttribute( + &sms, hipDeviceAttributeMultiprocessorCount, device + )); +#endif const auto multiprocessors{xpu::detail::checked_cast(sms)}; return multiprocessors; @@ -45,13 +60,22 @@ inline auto device_SMs() -> unsigned int { inline auto retain_default_pool() -> void { static const auto once{[] { auto device{0}; - auto pool{cudaMemPool_t{}}; auto threshold{std::numeric_limits::max()}; +#if defined(XPU_CUDA) + auto pool{cudaMemPool_t{}}; xpu::cu_check(cudaGetDevice(&device)); xpu::cu_check(cudaDeviceGetDefaultMemPool(&pool, device)); xpu::cu_check(cudaMemPoolSetAttribute( pool, cudaMemPoolAttrReleaseThreshold, &threshold )); +#else + auto pool{hipMemPool_t{}}; + xpu::cu_check(hipGetDevice(&device)); + xpu::cu_check(hipDeviceGetDefaultMemPool(&pool, device)); + xpu::cu_check(hipMemPoolSetAttribute( + pool, hipMemPoolAttrReleaseThreshold, &threshold + )); +#endif return true; }() @@ -67,9 +91,15 @@ inline auto wave_blocks(Kernel kernel, dim3 threads, std::size_t smem = 0uz) -> )}; auto blocks_per_SM{0}; +#if defined(XPU_CUDA) cu_check(cudaOccupancyMaxActiveBlocksPerMultiprocessor( &blocks_per_SM, kernel, xpu::detail::checked_cast(thread_budget), smem )); +#else + cu_check(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &blocks_per_SM, kernel, xpu::detail::checked_cast(thread_budget), smem + )); +#endif const auto blocks{xpu::detail::checked_mul( xpu::detail::device_SMs(), @@ -128,7 +158,7 @@ template <> struct Coord<1> { std::size_t x{}; }; template <> struct Coord<2> { std::size_t x{}, y{}; }; template <> struct Coord<3> { std::size_t x{}, y{}, z{}; }; -template __device__ [[nodiscard]] +template [[nodiscard]] XPU_DEVICE_ONLY inline auto global_index() noexcept -> Coord { auto id{Coord{}}; @@ -151,7 +181,7 @@ inline constexpr auto block_per_dim(std::size_t size, unsigned int dim_threads) return blocks; } -template __device__ [[nodiscard]] +template [[nodiscard]] XPU_DEVICE_ONLY inline auto global_stride() noexcept -> Coord { auto stride{Coord{}}; @@ -229,13 +259,20 @@ inline constexpr auto itr_index( return index; } +// HIP-Clang checks a requires-clause as the host caller, which rejects __device__ contributions. +// A constexpr function is host+device there, so the call is checked from that context instead template -concept sum_contribution = requires( - const std::remove_reference_t& contribution, - const xpu::array& index -) { - { contribution(index) } -> std::same_as; -}; +inline constexpr auto is_sum_contribution() -> bool { + return requires( + const std::remove_reference_t& contribution, + const xpu::array& index + ) { + { contribution(index) } -> std::same_as; + }; +} + +template +concept sum_contribution = is_sum_contribution(); template struct reduction_contribution { @@ -251,7 +288,7 @@ struct reduction_contribution { } }; -#if defined(XPU_CUDA) +#if defined(XPU_GPU) template inline auto reduction_input( const xpu::range& range, @@ -263,8 +300,13 @@ inline auto reduction_input( range, function_t{std::forward(contribution)} }}; +#if defined(XPU_CUDA) const auto indices{thrust::counting_iterator{0uz}}; const auto input{thrust::make_transform_iterator(indices, transform)}; +#else + const auto indices{rocprim::counting_iterator{0uz}}; + const auto input{rocprim::make_transform_iterator(indices, transform)}; +#endif return input; } @@ -283,7 +325,7 @@ struct parallel_for_function { }; #endif -#if !defined(XPU_CUDA) +#if !defined(XPU_GPU) template inline auto parallel_for_coordinate( const xpu::range& range, @@ -358,7 +400,7 @@ inline auto parallel_for_impl( const xpu::range& range, F&& fcn ) -> void { -#if defined(XPU_CUDA) +#if defined(XPU_GPU) const auto total{xpu::detail::num_itrs(range)}; if (total == 0uz) { return; } @@ -369,7 +411,11 @@ inline auto parallel_for_impl( function_t{std::forward(fcn)} }}; +#if defined(XPU_CUDA) xpu::cu_check(cub::DeviceFor::Bulk(total, operation)); +#else + xpu::cu_check(hipcub::DeviceFor::Bulk(total, operation)); +#endif #else auto dense{true}; for (auto dimension{0uz}; dimension < dims; ++dimension) { @@ -399,7 +445,7 @@ inline auto parallel_for_impl( #endif } -#if !defined(XPU_CUDA) +#if !defined(XPU_GPU) template XPU_FORCE_INLINE auto parallel_reduce_nested( const xpu::range& range, @@ -487,12 +533,27 @@ inline auto parallel_reduce_sum_impl( T* output_ptr, F&& contribution ) -> void { -#if defined(XPU_CUDA) +#if defined(XPU_GPU) const auto input{reduction_input(range, std::forward(contribution))}; const auto count{xpu::detail::checked_cast(total)}; xpu::detail::retain_default_pool(); +#if defined(XPU_CUDA) xpu::cu_check(cub::DeviceReduce::Sum(input, output_ptr, count)); +#else + // hipCUB still needs explicit scratch; keep its lifetime on the default stream. + auto scratch_bytes{0uz}; + xpu::cu_check(hipcub::DeviceReduce::Sum( + nullptr, scratch_bytes, input, output_ptr, count + )); + + auto scratch{static_cast(nullptr)}; + xpu::cu_check(hipMallocAsync(&scratch, scratch_bytes, nullptr)); + xpu::cu_check(hipcub::DeviceReduce::Sum( + scratch, scratch_bytes, input, output_ptr, count + )); + xpu::cu_check(hipFreeAsync(scratch, nullptr)); +#endif #else if (range.step[dims - 1uz] == 1uz) { *output_ptr = parallel_reduce_cpu(range, total, contribution); @@ -529,6 +590,8 @@ inline auto parallel_reduce_sum( if (empty) { #if defined(XPU_CUDA) xpu::cu_check(cudaMemsetAsync(output_ptr, 0, sizeof(T))); +#elif defined(XPU_HIP) + xpu::cu_check(hipMemsetAsync(output_ptr, 0, sizeof(T))); #else *output_ptr = T{}; #endif diff --git a/xpu/linear_algebra.hpp b/xpu/linear_algebra.hpp index d59e883..ed24b60 100644 --- a/xpu/linear_algebra.hpp +++ b/xpu/linear_algebra.hpp @@ -9,6 +9,8 @@ #if defined(XPU_CUDA) #include #include +#elif defined(XPU_HIP) + #include #else #include #endif @@ -18,6 +20,15 @@ #include #include +// cuSOLVER and hipSOLVER share the dense API, so one GPU path calls either +#if defined(XPU_CUDA) + #define XPU_SOLVER(name) cusolverDn##name + #define XPU_SOLVER_NAME "cuSOLVER" +#elif defined(XPU_HIP) + #define XPU_SOLVER(name) hipsolverDn##name + #define XPU_SOLVER_NAME "hipSOLVER" +#endif + namespace xpu { namespace linalg { @@ -40,7 +51,23 @@ inline auto linalg_error(const char* message) noexcept -> void { std::abort(); } +#if defined(XPU_GPU) + #if defined(XPU_CUDA) +using solver_handle = cusolverDnHandle_t; +using solver_operation = cublasOperation_t; + +inline constexpr auto solver_fill_upper{CUBLAS_FILL_MODE_UPPER}; +inline constexpr auto solver_op_n{CUBLAS_OP_N}; +inline constexpr auto solver_op_t{CUBLAS_OP_T}; +#else +using solver_handle = hipsolverDnHandle_t; +using solver_operation = hipblasOperation_t; + +inline constexpr auto solver_fill_upper{HIPBLAS_FILL_MODE_UPPER}; +inline constexpr auto solver_op_n{HIPBLAS_OP_N}; +inline constexpr auto solver_op_t{HIPBLAS_OP_T}; +#endif template struct build_identity { @@ -85,15 +112,15 @@ struct transpose_square { } }; -inline auto create_cusolver_handle() -> cusolverDnHandle_t { - auto handle{cusolverDnHandle_t{}}; - xpu::cu_check(cusolverDnCreate(&handle)); +inline auto create_solver_handle() -> solver_handle { + auto handle{solver_handle{}}; + xpu::cu_check(XPU_SOLVER(Create)(&handle)); return handle; } template inline auto getrf_workspace_size( - cusolverDnHandle_t handle, + solver_handle handle, std::size_t order, std::size_t stride ) -> std::size_t { @@ -103,14 +130,14 @@ inline auto getrf_workspace_size( auto size{0}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSgetrf_bufferSize( + xpu::cu_check(XPU_SOLVER(Sgetrf_bufferSize)( handle, vendor_order, vendor_order, nullptr, vendor_stride, &size )); } else { - xpu::cu_check(cusolverDnDgetrf_bufferSize( + xpu::cu_check(XPU_SOLVER(Dgetrf_bufferSize)( handle, vendor_order, vendor_order, nullptr, vendor_stride, @@ -125,7 +152,7 @@ inline auto getrf_workspace_size( template inline auto potrf_workspace_size( - cusolverDnHandle_t handle, + solver_handle handle, std::size_t order, std::size_t stride ) -> std::size_t { @@ -135,15 +162,15 @@ inline auto potrf_workspace_size( auto size{0}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSpotrf_bufferSize( - handle, CUBLAS_FILL_MODE_UPPER, + xpu::cu_check(XPU_SOLVER(Spotrf_bufferSize)( + handle, solver_fill_upper, vendor_order, nullptr, vendor_stride, &size )); } else { - xpu::cu_check(cusolverDnDpotrf_bufferSize( - handle, CUBLAS_FILL_MODE_UPPER, + xpu::cu_check(XPU_SOLVER(Dpotrf_bufferSize)( + handle, solver_fill_upper, vendor_order, nullptr, vendor_stride, &size @@ -156,8 +183,8 @@ inline auto potrf_workspace_size( } template -inline auto cusolver_getrf( - cusolverDnHandle_t handle, +inline auto solver_getrf( + solver_handle handle, std::size_t order, std::size_t stride, T* matrix, @@ -169,14 +196,14 @@ inline auto cusolver_getrf( const auto vendor_stride{xpu::detail::checked_cast(stride)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSgetrf( + xpu::cu_check(XPU_SOLVER(Sgetrf)( handle, vendor_order, vendor_order, matrix, vendor_stride, workspace, pivot, info )); } else { - xpu::cu_check(cusolverDnDgetrf( + xpu::cu_check(XPU_SOLVER(Dgetrf)( handle, vendor_order, vendor_order, matrix, vendor_stride, @@ -186,8 +213,8 @@ inline auto cusolver_getrf( } template -inline auto cusolver_potrf( - cusolverDnHandle_t handle, +inline auto solver_potrf( + solver_handle handle, std::size_t order, std::size_t stride, T* matrix, @@ -199,16 +226,16 @@ inline auto cusolver_potrf( const auto vendor_workspace_size{xpu::detail::checked_cast(workspace.count())}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSpotrf( - handle, CUBLAS_FILL_MODE_UPPER, + xpu::cu_check(XPU_SOLVER(Spotrf)( + handle, solver_fill_upper, vendor_order, matrix, vendor_stride, workspace.data(), vendor_workspace_size, info )); } else { - xpu::cu_check(cusolverDnDpotrf( - handle, CUBLAS_FILL_MODE_UPPER, + xpu::cu_check(XPU_SOLVER(Dpotrf)( + handle, solver_fill_upper, vendor_order, matrix, vendor_stride, workspace.data(), vendor_workspace_size, @@ -218,9 +245,9 @@ inline auto cusolver_potrf( } template -inline auto cusolver_getrs( - cusolverDnHandle_t handle, - cublasOperation_t operation, +inline auto solver_getrs( + solver_handle handle, + solver_operation operation, std::size_t order, std::size_t right_hand_sides, const T* lower_upper, @@ -236,7 +263,7 @@ inline auto cusolver_getrs( const auto vendor_solution_stride{xpu::detail::checked_cast(solution_stride)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSgetrs( + xpu::cu_check(XPU_SOLVER(Sgetrs)( handle, operation, vendor_order, vendor_right_hand_sides, lower_upper, vendor_lower_upper_stride, @@ -245,7 +272,7 @@ inline auto cusolver_getrs( info )); } else { - xpu::cu_check(cusolverDnDgetrs( + xpu::cu_check(XPU_SOLVER(Dgetrs)( handle, operation, vendor_order, vendor_right_hand_sides, lower_upper, vendor_lower_upper_stride, @@ -257,8 +284,8 @@ inline auto cusolver_getrs( } template -inline auto cusolver_potrs( - cusolverDnHandle_t handle, +inline auto solver_potrs( + solver_handle handle, std::size_t order, std::size_t right_hand_sides, const T* factor, @@ -273,16 +300,16 @@ inline auto cusolver_potrs( const auto vendor_solution_stride{xpu::detail::checked_cast(solution_stride)}; if constexpr (std::same_as) { - xpu::cu_check(cusolverDnSpotrs( - handle, CUBLAS_FILL_MODE_UPPER, + xpu::cu_check(XPU_SOLVER(Spotrs)( + handle, solver_fill_upper, vendor_order, vendor_right_hand_sides, factor, vendor_factor_stride, solution, vendor_solution_stride, info )); } else { - xpu::cu_check(cusolverDnDpotrs( - handle, CUBLAS_FILL_MODE_UPPER, + xpu::cu_check(XPU_SOLVER(Dpotrs)( + handle, solver_fill_upper, vendor_order, vendor_right_hand_sides, factor, vendor_factor_stride, solution, vendor_solution_stride, @@ -455,7 +482,7 @@ inline auto transpose_square( static_cast(xpu::detail::checked_bytes(matrix_size)); -#if defined(XPU_CUDA) +#if defined(XPU_GPU) const auto range = xpu::range<2uz>{ {0uz, 0uz}, {order, order}, @@ -490,14 +517,14 @@ class cholesky_factorization { std::size_t stride_; T* factor_{}; -#if defined(XPU_CUDA) +#if defined(XPU_GPU) using dimension_type = int; #else using dimension_type = lapack_int; #endif -#if defined(XPU_CUDA) - cusolverDnHandle_t handle_; +#if defined(XPU_GPU) + detail::solver_handle handle_; xpu::buffer workspace_; xpu::buffer info_; #endif @@ -524,16 +551,16 @@ class cholesky_factorization { cholesky_factorization(std::size_t order, std::size_t stride) : order_{checked_order(order, stride)} , stride_{stride} -#if defined(XPU_CUDA) - , handle_{detail::create_cusolver_handle()} +#if defined(XPU_GPU) + , handle_{detail::create_solver_handle()} , workspace_{detail::potrf_workspace_size(handle_, order, stride)} , info_{1uz} #endif { } ~cholesky_factorization() { -#if defined(XPU_CUDA) - xpu::cu_check(cusolverDnDestroy(handle_)); +#if defined(XPU_GPU) + xpu::cu_check(XPU_SOLVER(Destroy)(handle_)); #endif } @@ -551,8 +578,8 @@ class cholesky_factorization { auto factorize(T* XPU_RESTRICT matrix) noexcept -> status { factor_ = nullptr; -#if defined(XPU_CUDA) - detail::cusolver_potrf( +#if defined(XPU_GPU) + detail::solver_potrf( handle_, order_, stride_, matrix, workspace_.view(), info_.data() ); @@ -560,7 +587,7 @@ class cholesky_factorization { auto info{0}; xpu::copy_n(&info, info_.data(), 1uz); if (info < 0) { - detail::linalg_error("cuSOLVER potrf received an invalid argument"); + detail::linalg_error(XPU_SOLVER_NAME " potrf received an invalid argument"); } if (info > 0) { return status::not_pd; } #else @@ -596,8 +623,8 @@ class cholesky_factorization { xpu::copy_n(solution, rhs, order_); -#if defined(XPU_CUDA) - detail::cusolver_potrs( +#if defined(XPU_GPU) + detail::solver_potrs( handle_, order_, 1uz, factor, stride_, solution, order_, @@ -607,7 +634,7 @@ class cholesky_factorization { auto info{0}; xpu::copy_n(&info, info_.data(), 1uz); if (info != 0) { - detail::linalg_error("cuSOLVER potrs received an invalid argument"); + detail::linalg_error(XPU_SOLVER_NAME " potrs received an invalid argument"); } #else const auto info{ @@ -632,15 +659,15 @@ class lu_factorization { std::size_t stride_; T* lower_upper_{}; -#if defined(XPU_CUDA) +#if defined(XPU_GPU) using dimension_type = int; #else using dimension_type = lapack_int; #endif -#if defined(XPU_CUDA) +#if defined(XPU_GPU) xpu::buffer pivot_; - cusolverDnHandle_t handle_; + detail::solver_handle handle_; xpu::buffer workspace_; xpu::buffer info_; #else @@ -669,9 +696,9 @@ class lu_factorization { lu_factorization(std::size_t order, std::size_t stride) : order_{checked_order(order, stride)} , stride_{stride} -#if defined(XPU_CUDA) +#if defined(XPU_GPU) , pivot_{order} - , handle_{detail::create_cusolver_handle()} + , handle_{detail::create_solver_handle()} , workspace_{detail::getrf_workspace_size(handle_, order, stride)} , info_{1uz} #else @@ -680,8 +707,8 @@ class lu_factorization { { } ~lu_factorization() { -#if defined(XPU_CUDA) - xpu::cu_check(cusolverDnDestroy(handle_)); +#if defined(XPU_GPU) + xpu::cu_check(XPU_SOLVER(Destroy)(handle_)); #endif } @@ -699,8 +726,8 @@ class lu_factorization { auto factorize(T* XPU_RESTRICT matrix) noexcept -> status { lower_upper_ = nullptr; -#if defined(XPU_CUDA) - detail::cusolver_getrf( +#if defined(XPU_GPU) + detail::solver_getrf( handle_, order_, stride_, matrix, workspace_.data(), pivot_.data(), info_.data() ); @@ -708,7 +735,7 @@ class lu_factorization { auto info{0}; xpu::copy_n(&info, info_.data(), 1uz); if (info < 0) { - detail::linalg_error("cuSOLVER getrf received an invalid argument"); + detail::linalg_error(XPU_SOLVER_NAME " getrf received an invalid argument"); } if (info > 0) { return status::singular; } #else @@ -743,9 +770,9 @@ class lu_factorization { xpu::copy_n(solution, rhs, order_); -#if defined(XPU_CUDA) - detail::cusolver_getrs( - handle_, CUBLAS_OP_T, +#if defined(XPU_GPU) + detail::solver_getrs( + handle_, detail::solver_op_t, order_, 1uz, lower_upper, stride_, pivot_.data(), @@ -756,7 +783,7 @@ class lu_factorization { auto info{0}; xpu::copy_n(&info, info_.data(), 1uz); if (info != 0) { - detail::linalg_error("cuSOLVER getrs received an invalid argument"); + detail::linalg_error(XPU_SOLVER_NAME " getrs received an invalid argument"); } #else const auto info{ @@ -784,7 +811,7 @@ class lu_factorization { detail::linalg_error("lower_upper and inverse must not alias"); } -#if defined(XPU_CUDA) +#if defined(XPU_GPU) const auto range = xpu::range<2uz>{ {0uz, 0uz}, {order_, order_}, @@ -798,8 +825,8 @@ class lu_factorization { xpu::parallel_for(range, initialize); - detail::cusolver_getrs( - handle_, CUBLAS_OP_N, + detail::solver_getrs( + handle_, detail::solver_op_n, order_, order_, lower_upper, stride_, pivot_.data(), @@ -810,7 +837,7 @@ class lu_factorization { auto info{0}; xpu::copy_n(&info, info_.data(), 1uz); if (info != 0) { - detail::linalg_error("cuSOLVER getrs received an invalid argument"); + detail::linalg_error(XPU_SOLVER_NAME " getrs received an invalid argument"); } #else for (auto row{0uz}; row < order_; ++row) { @@ -841,3 +868,6 @@ class lu_factorization { } // namespace xpu::linalg } // namespace xpu + +#undef XPU_SOLVER +#undef XPU_SOLVER_NAME diff --git a/xpu/math.hpp b/xpu/math.hpp index c83a829..17f693b 100644 --- a/xpu/math.hpp +++ b/xpu/math.hpp @@ -81,7 +81,7 @@ inline constexpr auto ceiling_div(T num, T den) noexcept -> T { template XPU_CUDA_CALLABLE inline auto sincos(T arg, T* XPU_RESTRICT s, T* XPU_RESTRICT c) noexcept -> void { -#if defined(__CUDA_ARCH__) +#if defined(XPU_DEVICE_COMPILE) if constexpr (std::is_same_v) { ::sincosf(arg, s, c); } else { ::sincos(arg, s, c); } #else @@ -91,7 +91,7 @@ inline auto sincos(T arg, T* XPU_RESTRICT s, T* XPU_RESTRICT c) noexcept -> void template [[nodiscard]] XPU_CUDA_CALLABLE inline auto rsqrt(T x) noexcept -> T { -#if defined(__CUDA_ARCH__) +#if defined(XPU_DEVICE_COMPILE) if constexpr (std::is_same_v) { return ::rsqrtf(x); } else { return ::rsqrt(x); } #else return T{1} / xpu::sqrt(x); @@ -100,7 +100,7 @@ inline auto rsqrt(T x) noexcept -> T { template [[nodiscard]] XPU_CUDA_CALLABLE inline auto norm3d(T x, T y, T z) noexcept -> T { -#if defined(__CUDA_ARCH__) +#if defined(XPU_DEVICE_COMPILE) if constexpr (std::is_same_v) { return ::norm3df(x, y, z); } else { return ::norm3d(x, y, z); } #else @@ -110,7 +110,7 @@ inline auto norm3d(T x, T y, T z) noexcept -> T { template [[nodiscard]] XPU_CUDA_CALLABLE inline auto rnorm3d(T x, T y, T z) noexcept -> T { -#if defined(__CUDA_ARCH__) +#if defined(XPU_DEVICE_COMPILE) if constexpr (std::is_same_v) { return ::rnorm3df(x, y, z); } else { return ::rnorm3d(x, y, z); } #else @@ -120,7 +120,7 @@ inline auto rnorm3d(T x, T y, T z) noexcept -> T { template XPU_CUDA_CALLABLE inline auto atomic_add(T* ptr, T value) noexcept -> void { -#if defined(__CUDA_ARCH__) +#if defined(XPU_DEVICE_COMPILE) atomicAdd(ptr, value); #else std::atomic_ref(*ptr).fetch_add(value, std::memory_order_relaxed); diff --git a/xpu/memory.hpp b/xpu/memory.hpp index 4ce95cb..0328e5e 100644 --- a/xpu/memory.hpp +++ b/xpu/memory.hpp @@ -49,6 +49,9 @@ inline auto alloc(std::size_t count) -> T* { #if defined(XPU_CUDA) auto ptr{static_cast(nullptr)}; if(cudaMalloc(&ptr, bytes) != cudaSuccess) { ptr = nullptr; } +#elif defined(XPU_HIP) + auto ptr{static_cast(nullptr)}; + if(hipMalloc(&ptr, bytes) != hipSuccess) { ptr = nullptr; } #else auto ptr{::operator new(bytes, std::align_val_t{default_align}, std::nothrow)}; #endif @@ -71,6 +74,8 @@ inline auto free(T* ptr) noexcept -> void { #if defined(XPU_CUDA) cudaFree(ptr); +#elif defined(XPU_HIP) + static_cast(hipFree(ptr)); #else ::operator delete(ptr, std::align_val_t(default_align)); #endif @@ -109,6 +114,8 @@ inline auto memset( #if defined(XPU_CUDA) xpu::cu_check(cudaMemset(dst, value, bytes)); +#elif defined(XPU_HIP) + xpu::cu_check(hipMemset(dst, value, bytes)); #else std::memset(dst, value, bytes); #endif @@ -123,6 +130,8 @@ inline auto memcpy( #if defined(XPU_CUDA) xpu::cu_check(cudaMemcpy(dst, src, bytes, cudaMemcpyDefault)); +#elif defined(XPU_HIP) + xpu::cu_check(hipMemcpy(dst, src, bytes, hipMemcpyDefault)); #else std::memcpy(dst, src, bytes); #endif diff --git a/xpu/random.hpp b/xpu/random.hpp index 189fd6e..f5325e9 100644 --- a/xpu/random.hpp +++ b/xpu/random.hpp @@ -5,6 +5,8 @@ #if defined(XPU_CUDA) #include +#elif defined(XPU_HIP) + #include #else #include #endif @@ -21,12 +23,20 @@ namespace random { class generator { private: +#if defined(XPU_GPU) #if defined(XPU_CUDA) curandStatePhilox4_32_10_t engine_; +#else + hiprandStatePhilox4_32_10_t engine_; +#endif [[nodiscard]] XPU_DEVICE_ONLY auto next_u32() -> std::uint32_t { +#if defined(XPU_CUDA) return curand(&engine_); +#else + return hiprand(&engine_); +#endif } [[nodiscard]] XPU_DEVICE_ONLY @@ -48,6 +58,8 @@ class generator { ) -> void { #if defined(XPU_CUDA) curand_init(master_seed, stream_id, offset, &engine_); +#elif defined(XPU_HIP) + hiprand_init(master_seed, stream_id, offset, &engine_); #else auto seed{std::seed_seq{ static_cast(master_seed), @@ -68,6 +80,13 @@ class generator { } else { return T{1} - curand_uniform_double(&engine_); } +#elif defined(XPU_HIP) + if constexpr (std::same_as) { + // 1 - hiprand_uniform rounds to 1 for the smallest draws, so take [0, 1) from the top 24 bits + return static_cast(next_u32() >> 8u) * 0x1p-24f; + } else { + return T{1} - hiprand_uniform_double(&engine_); + } #else return std::generate_canonical::digits>(engine_); #endif @@ -79,7 +98,7 @@ class generator { const auto value{minimum + (maximum - minimum) * uniform()}; if (value < maximum) { return value; } -#if defined(XPU_CUDA) +#if defined(XPU_GPU) if constexpr (std::same_as) { return ::nextafterf(maximum, minimum); } else { @@ -94,7 +113,7 @@ class generator { auto uniform_index(std::size_t upper_bound) -> std::size_t { assert(upper_bound != 0uz); -#if defined(XPU_CUDA) +#if defined(XPU_GPU) const auto bound{static_cast(upper_bound)}; const auto threshold{(std::uint64_t{0} - bound) % bound};