EBGeometry_GPURuntime.hpp

Go to the documentation of this file.

// SPDX-FileCopyrightText: 2026 Robert Marskar <robert.marskar@sintef.no>
//
// SPDX-License-Identifier: GPL-3.0-or-later
 
#ifndef EBGEOMETRY_GPURUNTIME_HPP
#define EBGEOMETRY_GPURUNTIME_HPP
 
// Our includes
#include "EBGeometry_GPU.hpp"
 
#if defined(EBGEOMETRY_CUDA) || defined(EBGEOMETRY_HIP) || defined(EBGEOMETRY_DOXYGEN)
 
// Std includes (EBGEOMETRY_GPU_CHECK needs fprintf/abort; the wrappers need size_t).
#include <cstddef>
#include <cstdio>
#include <cstdlib>
 
// Backend runtime header. HIP is checked FIRST: hipcc on the NVIDIA platform defines both
// __HIPCC__ and __CUDACC__, and there the hip* entry points exist and forward to cuda*, so
// mapping to hip* whenever __HIPCC__ is defined is correct on both HIP platforms.
#if defined(EBGEOMETRY_HIP)
#include <hip/hip_runtime.h> // NOT pre-included by the compiler driver; must be explicit.
#elif defined(EBGEOMETRY_CUDA)
#include <cuda_runtime.h> // nvcc pre-includes this; spelled out for hygiene.
#endif
 
namespace EBGeometry {
 
namespace GPU {
 
#if defined(EBGEOMETRY_DOXYGEN)
 
using Error = int;
 
using MemcpyKind = int;
 
inline constexpr Error Success = 0;
 
inline constexpr MemcpyKind MemcpyHostToDevice = 1;
 
inline constexpr MemcpyKind MemcpyDeviceToHost = 2;
 
[[nodiscard]] inline Error
memAlloc(void** a_devicePtr, size_t a_numBytes) noexcept;
 
[[nodiscard]] inline Error
memFree(void* a_devicePtr) noexcept;
 
[[nodiscard]] inline Error
memAllocManaged(void** a_ptr, size_t a_numBytes) noexcept;
 
[[nodiscard]] inline Error
memAllocHost(void** a_ptr, size_t a_numBytes) noexcept;
 
[[nodiscard]] inline Error
memAllocHostMapped(void** a_ptr, size_t a_numBytes) noexcept;
 
[[nodiscard]] inline Error
memFreeHost(void* a_ptr) noexcept;
 
[[nodiscard]] inline Error
memcpy(void* a_dst, const void* a_src, size_t a_numBytes, MemcpyKind a_kind) noexcept;
 
[[nodiscard]] inline Error
getDeviceCount(int* a_count) noexcept;
 
[[nodiscard]] inline Error
getLastError() noexcept;
 
[[nodiscard]] inline Error
deviceSynchronize() noexcept;
 
[[nodiscard]] inline const char*
getErrorString(Error a_error) noexcept;
 
[[nodiscard]] inline bool
isShutdownError(Error a_error) noexcept;
 
[[nodiscard]] inline const char*
getErrorHint(Error a_error) noexcept;
 
#elif defined(EBGEOMETRY_HIP)
 
using Error      = hipError_t;
using MemcpyKind = hipMemcpyKind;
 
inline constexpr Error      Success            = hipSuccess;
inline constexpr MemcpyKind MemcpyHostToDevice = hipMemcpyHostToDevice;
inline constexpr MemcpyKind MemcpyDeviceToHost = hipMemcpyDeviceToHost;
 
[[nodiscard]] inline Error
memAlloc(void** a_devicePtr, size_t a_numBytes) noexcept
{
  return hipMalloc(a_devicePtr, a_numBytes);
}
 
[[nodiscard]] inline Error
memFree(void* a_devicePtr) noexcept
{
  return hipFree(a_devicePtr);
}
 
[[nodiscard]] inline Error
memAllocManaged(void** a_ptr, size_t a_numBytes) noexcept
{
  return hipMallocManaged(a_ptr, a_numBytes);
}
 
[[nodiscard]] inline Error
memAllocHost(void** a_ptr, size_t a_numBytes) noexcept
{
  return hipHostMalloc(a_ptr, a_numBytes, 0u);
}
 
[[nodiscard]] inline Error
memAllocHostMapped(void** a_ptr, size_t a_numBytes) noexcept
{
  return hipHostMalloc(a_ptr, a_numBytes, hipHostMallocMapped);
}
 
[[nodiscard]] inline Error
memFreeHost(void* a_ptr) noexcept
{
  return hipHostFree(a_ptr);
}
 
[[nodiscard]] inline Error
memcpy(void* a_dst, const void* a_src, size_t a_numBytes, MemcpyKind a_kind) noexcept
{
  return hipMemcpy(a_dst, a_src, a_numBytes, a_kind);
}
 
[[nodiscard]] inline Error
getDeviceCount(int* a_count) noexcept
{
  return hipGetDeviceCount(a_count);
}
 
[[nodiscard]] inline Error
getLastError() noexcept
{
  return hipGetLastError();
}
 
[[nodiscard]] inline Error
deviceSynchronize() noexcept
{
  return hipDeviceSynchronize();
}
 
[[nodiscard]] inline const char*
getErrorString(Error a_error) noexcept
{
  return hipGetErrorString(a_error);
}
 
[[nodiscard]] inline bool
isShutdownError(Error a_error) noexcept
{
  return a_error == hipErrorDeinitialized || a_error == hipErrorContextIsDestroyed;
}
 
[[nodiscard]] inline const char*
getErrorHint(Error a_error) noexcept
{
  return (a_error == hipErrorNoBinaryForGpu)
           ? "\n  hint: the binary has no code for this GPU; build for its architecture (--offload-arch, "
             "HIP_ARCH or CMAKE_HIP_ARCHITECTURES)"
           : "";
}
 
#elif defined(EBGEOMETRY_CUDA)
 
using Error      = cudaError_t;
using MemcpyKind = cudaMemcpyKind;
 
inline constexpr Error      Success            = cudaSuccess;
inline constexpr MemcpyKind MemcpyHostToDevice = cudaMemcpyHostToDevice;
inline constexpr MemcpyKind MemcpyDeviceToHost = cudaMemcpyDeviceToHost;
 
[[nodiscard]] inline Error
memAlloc(void** a_devicePtr, size_t a_numBytes) noexcept
{
  return cudaMalloc(a_devicePtr, a_numBytes);
}
 
[[nodiscard]] inline Error
memFree(void* a_devicePtr) noexcept
{
  return cudaFree(a_devicePtr);
}
 
[[nodiscard]] inline Error
memAllocManaged(void** a_ptr, size_t a_numBytes) noexcept
{
  return cudaMallocManaged(a_ptr, a_numBytes, cudaMemAttachGlobal);
}
 
[[nodiscard]] inline Error
memAllocHost(void** a_ptr, size_t a_numBytes) noexcept
{
  return cudaHostAlloc(a_ptr, a_numBytes, cudaHostAllocDefault);
}
 
[[nodiscard]] inline Error
memAllocHostMapped(void** a_ptr, size_t a_numBytes) noexcept
{
  return cudaHostAlloc(a_ptr, a_numBytes, cudaHostAllocMapped);
}
 
[[nodiscard]] inline Error
memFreeHost(void* a_ptr) noexcept
{
  return cudaFreeHost(a_ptr);
}
 
[[nodiscard]] inline Error
memcpy(void* a_dst, const void* a_src, size_t a_numBytes, MemcpyKind a_kind) noexcept
{
  return cudaMemcpy(a_dst, a_src, a_numBytes, a_kind);
}
 
[[nodiscard]] inline Error
getDeviceCount(int* a_count) noexcept
{
  return cudaGetDeviceCount(a_count);
}
 
[[nodiscard]] inline Error
getLastError() noexcept
{
  return cudaGetLastError();
}
 
[[nodiscard]] inline Error
deviceSynchronize() noexcept
{
  return cudaDeviceSynchronize();
}
 
[[nodiscard]] inline const char*
getErrorString(Error a_error) noexcept
{
  return cudaGetErrorString(a_error);
}
 
[[nodiscard]] inline bool
isShutdownError(Error a_error) noexcept
{
  return a_error == cudaErrorCudartUnloading || a_error == cudaErrorContextIsDestroyed;
}
 
[[nodiscard]] inline const char*
getErrorHint(Error a_error) noexcept
{
  return (a_error == cudaErrorUnsupportedPtxVersion || a_error == cudaErrorNoKernelImageForDevice)
           ? "\n  hint: the binary has no machine code for this GPU, and the driver cannot compile its PTX; "
             "build for the GPU's architecture (-arch, CUDA_ARCH or CMAKE_CUDA_ARCHITECTURES), or update "
             "the driver"
           : "";
}
 
#endif
 
} // namespace GPU
} // namespace EBGeometry
 
#define EBGEOMETRY_GPU_CHECK(x)                                                        \
  do {                                                                                 \
    const EBGeometry::GPU::Error ebgeometry_gpu_err = (x);                             \
    if (ebgeometry_gpu_err != EBGeometry::GPU::Success) {                              \
      std::fprintf(stderr,                                                             \
                   "EBGeometry GPU error: %s\n  call: %s\n  file: %s\n  line: %d%s\n", \
                   EBGeometry::GPU::getErrorString(ebgeometry_gpu_err),                \
                   #x,                                                                 \
                   __FILE__,                                                           \
                   __LINE__,                                                           \
                   EBGeometry::GPU::getErrorHint(ebgeometry_gpu_err));                 \
      std::abort();                                                                    \
    }                                                                                  \
  } while (0)
 
#endif // EBGEOMETRY_CUDA || EBGEOMETRY_HIP || EBGEOMETRY_DOXYGEN
 
#endif // EBGEOMETRY_GPURUNTIME_HPP