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};