From af42affb77632acae0aef246c189653f9b679819 Mon Sep 17 00:00:00 2001 From: Igor Sfiligoi Date: Thu, 7 Aug 2025 07:53:26 -0700 Subject: [PATCH 1/2] Add HIP support for AMD GPUs --- README.rst | 6 ++++ src/Makefile | 39 +++++++++++++++++---- src/distance/permanova_dyn_impl.hpp | 54 ++++++++++++++++++++--------- src/util/skbb_accapi_impl.hpp | 43 ++++++++++++++++++----- 4 files changed, 111 insertions(+), 31 deletions(-) diff --git a/README.rst b/README.rst index c775745..bff1cdf 100644 --- a/README.rst +++ b/README.rst @@ -53,6 +53,12 @@ The CUDA build is not enabled by default. If you are interested in building the export NV_CUDA=Y +To build AMD-GPU-enabled code, you will also need the HIP compiler. + +The HIP build is not enabled by default. If you are interested in building the GPU-accelerated libraries, set:: + + export AMD_HIP=Y + There is also support for OpenMP Offload and OpenACC builds, by setting either ``AMD_CXX`` or ``NV_CXX``, but it is still experimental. Provided files and their usage diff --git a/src/Makefile b/src/Makefile index f930e73..97dc35d 100644 --- a/src/Makefile +++ b/src/Makefile @@ -58,7 +58,7 @@ ifneq ($(SKBB_ENABLE_ACC_NV),) # CUDA version SKBB_ENABLE_ACC_NV_BINS := 1 NV_CXX := nvcc - NV_CXXFLAGS := -DCUDA + NV_CXXFLAGS := -DSKBB_CUDA NV_CXXFLAGS += -O3 --use_fast_math -std=c++17 -I. $(OPT) -Xcompiler -fPIC NV_CXXFLAGS += -arch=all-major else @@ -75,14 +75,24 @@ ifneq ($(SKBB_ENABLE_ACC_NV),) endif endif -ifneq ($(AMD_CXX),) - ifneq ($(SKBB_ENABLE_ACC_AMD),) +ifneq ($(SKBB_ENABLE_ACC_AMD),) + ifneq ($(AMD_HIP),) + # HIP version + SKBB_ENABLE_ACC_AMD_BINS := 1 + AMD_CXX := hipcc + AMD_CXXFLAGS := -DSKBB_HIP + AMD_CXXFLAGS += -O3 -ffast-math -std=c++17 -I. $(OPT) -fPIC + AMD_CXXFLAGS += --offload-arch=gfx1100,gfx1101,gfx1102,gfx1103,gfx1030,gfx1031,gfx90a,gfx942 + AMD_SO_LDDFLAGS += -shared -fgpu-rdc --hip-link + else + ifneq ($(AMD_CXX),) SKBB_ENABLE_ACC_AMD_BINS := 1 AMD_CXXFLAGS += -fopenmp -fopenmp-offload-mandatory -DOMPGPU=1 AMD_CXXFLAGS += -O3 -ffast-math -std=c++17 -I. $(OPT) -fPIC AMD_LDFLAGS += -shared -fopenmp AMD_CXXFLAGS += --offload-arch=gfx1100,gfx1101,gfx1102,gfx1103,gfx1030,gfx1031,gfx90a,gfx942 AMD_LDFLAGS += --offload-arch=gfx1100,gfx1101,gfx1102,gfx1103,gfx1030,gfx1031,gfx90a,gfx942 + endif endif endif @@ -246,22 +256,32 @@ endif ifdef SKBB_ENABLE_ACC_AMD SKBB_OBJS += permanova_acc_amd.o - permanova_dyn_acc_amd.h: distance/permanova_dyn_impl.hpp ./tools/generate_permanova_dyn.py acc_amd api_h > $@ -permanova_dyn_acc_amd.cpp: distance/permanova_dyn_impl.hpp permanova_dyn_acc_amd.h - ./tools/generate_permanova_dyn.py acc_amd api > $@ permanova_acc_amd.cpp: distance/permanova_dyn_impl.hpp permanova_dyn_acc_amd.h ./tools/generate_permanova_dyn.py acc_amd indirect > $@ permanova_acc_amd.o: permanova_acc_amd.cpp permanova_dyn_acc_amd.h distance/permanova_dyn.hpp util/skbb_dl.cpp $(CXX) $(CXXFLAGS) -DSKBB_ACC_NM=skbb_acc_amd -c $< -o $@ +ifneq ($(AMD_HIP),) +# HIP version +permanova_dyn_acc_amd.hip: distance/permanova_dyn_impl.hpp permanova_dyn_acc_amd.h + ./tools/generate_permanova_dyn.py acc_amd api > $@ +permanova_dyn_acc_amd.o: permanova_dyn_acc_amd.hip distance/permanova_dyn.hpp distance/permanova_dyn_impl.hpp + $(AMD_CXX) $(AMD_CXXFLAGS) -DSKBB_ACC_NM=skbb_acc_amd -c $< -o $@ + +else +# OpenMP Target version +permanova_dyn_acc_amd.cpp: distance/permanova_dyn_impl.hpp permanova_dyn_acc_amd.h + ./tools/generate_permanova_dyn.py acc_amd api > $@ permanova_dyn_acc_amd.o: permanova_dyn_acc_amd.cpp distance/permanova_dyn.hpp distance/permanova_dyn_impl.hpp $(AMD_CXX) $(AMD_CXXFLAGS) -DSKBB_ACC_NM=skbb_acc_amd -c $< -o $@ endif +endif + @@ -317,8 +337,15 @@ ifdef SKBB_ENABLE_ACC_AMD_BINS SKBB_SHLIBS += libskbb_acc_amd.so +ifneq ($(AMD_HIP),) +#HIP version +libskbb_acc_amd.so: skbb_accapi_dyn_acc_amd.o permanova_dyn_acc_amd.o + $(AMD_CXX) $(AMD_SO_LDDFLAGS) -o $@ skbb_accapi_dyn_acc_amd.o permanova_dyn_acc_amd.o +else +#OpenMP Target version libskbb_acc_amd.so: skbb_accapi_dyn_acc_amd.o permanova_dyn_acc_amd.o $(AMD_CXX) $(AMD_LDFLAGS) $(SO_LDDFLAGS) -o $@ skbb_accapi_dyn_acc_amd.o permanova_dyn_acc_amd.o +endif endif diff --git a/src/distance/permanova_dyn_impl.hpp b/src/distance/permanova_dyn_impl.hpp index 3de133c..843c4f1 100644 --- a/src/distance/permanova_dyn_impl.hpp +++ b/src/distance/permanova_dyn_impl.hpp @@ -22,27 +22,41 @@ #include #include -#if defined(CUDA) +#if defined(SKBB_CUDA) #include #include -#elif !(defined(_OPENACC) || defined(OMPGPU)) +#elif defined(SKBB_HIP) -#include +#include +#include +#include #elif defined(_OPENACC) #include +#elif !(defined(_OPENACC) || defined(OMPGPU)) + +#include + +#define SKBB_CPU Y + #endif static inline int pmn_get_max_parallelism_T() { -#if defined(CUDA) +#if defined(SKBB_CPU) + // No good reason to do more than max threads + // (but use 2x to reduce thread spawning overhead) + // but we do use 16x blocking, so account for that, too + return 2*omp_get_max_threads()*16; + +#elif defined(SKBB_CUDA) int deviceID; cudaDeviceProp props; - cudaGetDevice(&deviceID); + if (cudaGetDevice(&deviceID)!=cudaSuccess) return 4000; // should never get in here, but just in case cudaGetDeviceProperties(&props, deviceID); // GPUs typically need at least 64 blocks per SM to be fully loaded @@ -50,11 +64,17 @@ static inline int pmn_get_max_parallelism_T() { // Most permanovas are multiple of 100, so double that makes a good constant return 200*props.multiProcessorCount; -#elif !(defined(_OPENACC) || defined(OMPGPU)) - // No good reason to do more than max threads - // (but use 2x to reduce thread spawning overhead) - // but we do use 16x blocking, so account for that, too - return 2*omp_get_max_threads()*16; +#elif defined(SKBB_HIP) + int deviceID = 0; + hipDeviceProp_t props; + + if (hipGetDevice(&deviceID)!=hipSuccess) return 4000; // should never get in here, but just in case + if (hipGetDeviceProperties(&props, deviceID)!=hipSuccess) throw std::runtime_error("hipGetDeviceProperties failed"); + + // GPUs typically need at least 64 blocks per SM to be fully loaded + // We want a few multiples of that to deal with unbalanced load + // Most permanovas are multiple of 100, so double that makes a good constant + return 200*props.multiProcessorCount; #else // 1k is enough for consumer-grade GPUs @@ -122,7 +142,7 @@ static inline void pmn_f_stat_sW_block( // inv_group_sizes is an array of size maxel(groupings) // Results in group_sWs, and array of size n_grouping_dims -#if !(defined(_OPENACC) || defined(OMPGPU) || defined(CUDA)) +#if defined(SKBB_CPU) template static inline void pmn_f_stat_sW_cpu( @@ -187,7 +207,7 @@ static inline void pmn_f_stat_sW_cpu( } } -#elif defined(CUDA) +#elif (defined(SKBB_CUDA) || defined(SKBB_HIP)) template __global__ void pmn_f_stat_sW_cuda_one( @@ -273,7 +293,7 @@ __global__ void pmn_f_stat_sW_cuda_one( } template -static inline void pmn_f_stat_sW_cuda( +static inline void pmn_f_stat_sW_gpu( const uint32_t n_dims, const TFloat * mat, const uint32_t n_grouping_dims, @@ -281,7 +301,11 @@ static inline void pmn_f_stat_sW_cuda( const TFloat *inv_group_sizes, TFloat *group_sWs) { pmn_f_stat_sW_cuda_one<<>>(n_dims,mat,n_grouping_dims,groupings,inv_group_sizes,group_sWs); +#if defined(SKBB_CUDA) cudaDeviceSynchronize(); +#else + if (hipDeviceSynchronize()!=hipSuccess) throw std::runtime_error("hipDeviceSynchronize failed"); +#endif } #else @@ -332,10 +356,8 @@ static inline void pmn_f_stat_sW_T( const uint32_t *groupings, const TFloat *inv_group_sizes, TFloat *group_sWs) { -#if !(defined(_OPENACC) || defined(OMPGPU) || defined(CUDA)) +#if defined(SKBB_CPU) pmn_f_stat_sW_cpu(n_dims, mat, n_grouping_dims, groupings, inv_group_sizes, group_sWs); -#elif defined(CUDA) - pmn_f_stat_sW_cuda(n_dims, mat, n_grouping_dims, groupings, inv_group_sizes, group_sWs); #else pmn_f_stat_sW_gpu(n_dims, mat, n_grouping_dims, groupings, inv_group_sizes, group_sWs); #endif diff --git a/src/util/skbb_accapi_impl.hpp b/src/util/skbb_accapi_impl.hpp index 165ff7a..a7e33ca 100644 --- a/src/util/skbb_accapi_impl.hpp +++ b/src/util/skbb_accapi_impl.hpp @@ -23,10 +23,16 @@ #include "util/skbb_accapi.hpp" #include -#if defined(CUDA) +#if defined(SKBB_CUDA) #include +#elif defined(SKBB_HIP) + +#include +#include +#include + #elif defined(OMPGPU) #include @@ -38,11 +44,16 @@ #endif static inline bool acc_found_gpu_T() { -#if defined(CUDA) +#if defined(SKBB_CUDA) int deviceCount; cudaError_t error = cudaGetDeviceCount(&deviceCount); if (error != cudaSuccess) return false; return deviceCount != 0; +#elif defined(SKBB_HIP) + int deviceCount; + hipError_t error = hipGetDeviceCount(&deviceCount); + if (error != hipSuccess) return false; + return deviceCount != 0; #elif defined(OMPGPU) return omp_get_num_devices() > 0; #elif defined(_OPENACC) @@ -54,7 +65,7 @@ static inline bool acc_found_gpu_T() { // is the implementation async, and need the alt structures? static inline bool acc_need_alt_T() { -#if defined(_OPENACC) || defined(OMPGPU) || defined(CUDA) +#if defined(_OPENACC) || defined(OMPGPU) || defined(SKBB_CUDA) || defined(SKBB_HIP) return true; #else return false; @@ -62,8 +73,10 @@ static inline bool acc_need_alt_T() { } static inline void acc_wait_T() { -#if defined(CUDA) +#if defined(SKBB_CUDA) cudaDeviceSynchronize(); +#elif defined(SKBB_HIP) + if (hipDeviceSynchronize()!=hipSuccess) throw std::runtime_error("hipDeviceSynchronize failed"); #elif defined(OMPGPU) // TODO: Change if we ever implement async in OMPGPU #elif defined(_OPENACC) @@ -76,8 +89,10 @@ static inline void acc_create_buf_T( TNum* buf_host, TNum** buf_device, uint64_t size) { -#if defined(CUDA) +#if defined(SKBB_CUDA) cudaMalloc((void**)buf_device, sizeof(TNum) * size); +#elif defined(SKBB_HIP) + if (hipMalloc((void**)buf_device, sizeof(TNum) * size)!=hipSuccess) throw std::runtime_error("hipMalloc failed"); #elif defined(OMPGPU) #pragma omp target enter data map(alloc:buf_host[0:size]) *buf_device = buf_host; @@ -94,9 +109,12 @@ static inline void acc_copyin_buf_T( TNum* buf_host, TNum** buf_device, uint64_t size) { -#if defined(CUDA) +#if defined(SKBB_CUDA) cudaMalloc((void**)buf_device, sizeof(TNum) * size); cudaMemcpy(*buf_device, buf_host, sizeof(TNum) * size, cudaMemcpyHostToDevice); +#elif defined(SKBB_HIP) + if (hipMalloc((void**)buf_device, sizeof(TNum) * size)!=hipSuccess) throw std::runtime_error("hipMalloc failed"); + if (hipMemcpy(*buf_device, buf_host, sizeof(TNum) * size, hipMemcpyHostToDevice)!=hipSuccess) throw std::runtime_error("hipMemcpy failed"); #elif defined(OMPGPU) #pragma omp target enter data map(to:buf_host[0:size]) *buf_device = buf_host; @@ -113,8 +131,10 @@ static inline void acc_update_device_T( TNum *buf_host, TNum *buf_device, uint64_t start, uint64_t end) { -#if defined(CUDA) +#if defined(SKBB_CUDA) cudaMemcpy(buf_device+start, buf_host+start, sizeof(TNum) * (end-start), cudaMemcpyHostToDevice); +#elif defined(SKBB_HIP) + if (hipMemcpy(buf_device+start, buf_host+start, sizeof(TNum) * (end-start), hipMemcpyHostToDevice)!=hipSuccess) throw std::runtime_error("hipMemcpy failed"); #elif defined(OMPGPU) // assert buf_host==buf_device #pragma omp target update to(buf_host[start:end]) @@ -129,9 +149,12 @@ static inline void acc_copyout_buf_T( TNum *buf_host, TNum *buf_device, uint64_t size) { -#if defined(CUDA) +#if defined(SKBB_CUDA) cudaMemcpy(buf_host, buf_device, sizeof(TNum) * size, cudaMemcpyDeviceToHost); cudaFree(buf_device); +#elif defined(SKBB_HIP) + if (hipMemcpy(buf_host, buf_device, sizeof(TNum) * size, hipMemcpyDeviceToHost)!=hipSuccess) throw std::runtime_error("hipMemcpy failed"); + if (hipFree(buf_device)!=hipSuccess) {} // ignore any errors, not critical #elif defined(OMPGPU) // assert buf_host==buf_device #pragma omp target exit data map(from:buf_device[0:size]) @@ -145,8 +168,10 @@ template static inline void acc_destroy_buf_T( TNum *buf_device, uint64_t size) { -#if defined(CUDA) +#if defined(SKBB_CUDA) cudaFree(buf_device); +#elif defined(SKBB_HIP) + if (hipFree(buf_device)!=hipSuccess) {} // ignore any errors, not critical #elif defined(OMPGPU) #pragma omp target exit data map(delete:buf_device[0:size]) #elif defined(_OPENACC) From 0a131fe5ab9465cafa7715cb1dd9ee7ca43d7366 Mon Sep 17 00:00:00 2001 From: Igor Sfiligoi Date: Thu, 7 Aug 2025 08:09:11 -0700 Subject: [PATCH 2/2] Improve error handling in CUDA code --- src/distance/permanova_dyn_impl.hpp | 3 ++- src/util/skbb_accapi_impl.hpp | 13 +++++++------ 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/src/distance/permanova_dyn_impl.hpp b/src/distance/permanova_dyn_impl.hpp index 843c4f1..a25dfe5 100644 --- a/src/distance/permanova_dyn_impl.hpp +++ b/src/distance/permanova_dyn_impl.hpp @@ -26,6 +26,7 @@ #include #include +#include #elif defined(SKBB_HIP) @@ -302,7 +303,7 @@ static inline void pmn_f_stat_sW_gpu( TFloat *group_sWs) { pmn_f_stat_sW_cuda_one<<>>(n_dims,mat,n_grouping_dims,groupings,inv_group_sizes,group_sWs); #if defined(SKBB_CUDA) - cudaDeviceSynchronize(); + if (cudaDeviceSynchronize()!=cudaSuccess) throw std::runtime_error("cudaDeviceSynchronize failed"); #else if (hipDeviceSynchronize()!=hipSuccess) throw std::runtime_error("hipDeviceSynchronize failed"); #endif diff --git a/src/util/skbb_accapi_impl.hpp b/src/util/skbb_accapi_impl.hpp index a7e33ca..a1bb317 100644 --- a/src/util/skbb_accapi_impl.hpp +++ b/src/util/skbb_accapi_impl.hpp @@ -26,6 +26,7 @@ #if defined(SKBB_CUDA) #include +#include #elif defined(SKBB_HIP) @@ -74,7 +75,7 @@ static inline bool acc_need_alt_T() { static inline void acc_wait_T() { #if defined(SKBB_CUDA) - cudaDeviceSynchronize(); + if (cudaDeviceSynchronize()!=cudaSuccess) throw std::runtime_error("cudaDeviceSynchronize failed"); #elif defined(SKBB_HIP) if (hipDeviceSynchronize()!=hipSuccess) throw std::runtime_error("hipDeviceSynchronize failed"); #elif defined(OMPGPU) @@ -90,7 +91,7 @@ static inline void acc_create_buf_T( TNum** buf_device, uint64_t size) { #if defined(SKBB_CUDA) - cudaMalloc((void**)buf_device, sizeof(TNum) * size); + if (cudaMalloc((void**)buf_device, sizeof(TNum) * size)!=cudaSuccess) throw std::runtime_error("cudaMalloc failed"); #elif defined(SKBB_HIP) if (hipMalloc((void**)buf_device, sizeof(TNum) * size)!=hipSuccess) throw std::runtime_error("hipMalloc failed"); #elif defined(OMPGPU) @@ -110,8 +111,8 @@ static inline void acc_copyin_buf_T( TNum** buf_device, uint64_t size) { #if defined(SKBB_CUDA) - cudaMalloc((void**)buf_device, sizeof(TNum) * size); - cudaMemcpy(*buf_device, buf_host, sizeof(TNum) * size, cudaMemcpyHostToDevice); + if (cudaMalloc((void**)buf_device, sizeof(TNum) * size)!=cudaSuccess) throw std::runtime_error("cudaMalloc failed"); + if (cudaMemcpy(*buf_device, buf_host, sizeof(TNum) * size, cudaMemcpyHostToDevice)!=cudaSuccess) throw std::runtime_error("cudaMemcpy failed"); #elif defined(SKBB_HIP) if (hipMalloc((void**)buf_device, sizeof(TNum) * size)!=hipSuccess) throw std::runtime_error("hipMalloc failed"); if (hipMemcpy(*buf_device, buf_host, sizeof(TNum) * size, hipMemcpyHostToDevice)!=hipSuccess) throw std::runtime_error("hipMemcpy failed"); @@ -132,7 +133,7 @@ static inline void acc_update_device_T( TNum *buf_device, uint64_t start, uint64_t end) { #if defined(SKBB_CUDA) - cudaMemcpy(buf_device+start, buf_host+start, sizeof(TNum) * (end-start), cudaMemcpyHostToDevice); + if (cudaMemcpy(buf_device+start, buf_host+start, sizeof(TNum) * (end-start), cudaMemcpyHostToDevice)!=cudaSuccess) throw std::runtime_error("cudaMemcpy failed"); #elif defined(SKBB_HIP) if (hipMemcpy(buf_device+start, buf_host+start, sizeof(TNum) * (end-start), hipMemcpyHostToDevice)!=hipSuccess) throw std::runtime_error("hipMemcpy failed"); #elif defined(OMPGPU) @@ -150,7 +151,7 @@ static inline void acc_copyout_buf_T( TNum *buf_device, uint64_t size) { #if defined(SKBB_CUDA) - cudaMemcpy(buf_host, buf_device, sizeof(TNum) * size, cudaMemcpyDeviceToHost); + if (cudaMemcpy(buf_host, buf_device, sizeof(TNum) * size, cudaMemcpyDeviceToHost)!=cudaSuccess) throw std::runtime_error("cudaMemcpy failed"); cudaFree(buf_device); #elif defined(SKBB_HIP) if (hipMemcpy(buf_host, buf_device, sizeof(TNum) * size, hipMemcpyDeviceToHost)!=hipSuccess) throw std::runtime_error("hipMemcpy failed");