Merge pull request #26 from psychocoderHPC/fix-hipWarpSize

support for any HIP compiler
kstppd · Oct 4, 2023 · b7f3f5a · b7f3f5a
2 parents 078890d + 5a5a160
commit b7f3f5a
Show file tree

Hide file tree

Showing 6 changed files with 21 additions and 17 deletions.
diff --git a/include/hashinator/defaults.h b/include/hashinator/defaults.h
@@ -25,13 +25,16 @@ namespace defaults {
 #ifdef __NVCC__
 constexpr int WARPSIZE = 32;
 constexpr int BUCKET_OVERFLOW = 32;
+#elif (__HIP__ && __AMDGCN_WAVEFRONT_SIZE)
+// AMDs GPU warp size depends on the GPU architecture
+constexpr int WARPSIZE = __AMDGCN_WAVEFRONT_SIZE;
+constexpr int BUCKET_OVERFLOW = __AMDGCN_WAVEFRONT_SIZE;
 #else
-constexpr int WARPSIZE = 64;
-constexpr int BUCKET_OVERFLOW = 64;
+#error "Warp size not known, please use a CUDA or HIP compiler."
 #endif
 constexpr int elementsPerWarp = 1;
 constexpr int MAX_BLOCKSIZE = 1024;
 template <typename T>
 using DefaultHashFunction = HashFunctions::Fibonacci<T>;
 } // namespace defaults
-} // namespace Hashinator
+} // namespace Hashinator
diff --git a/include/hashinator/hashers.h b/include/hashinator/hashers.h
@@ -612,7 +612,7 @@ __global__ void insert_index_kernel(KEY_TYPE* keys, hash_pair<KEY_TYPE, VAL_TYPE
 }
 
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
 
 /*
  * Resets all elements pointed by src to EMPTY in dst
@@ -1103,7 +1103,7 @@ __global__ void retrieve_kernel(KEY_TYPE* keys, VAL_TYPE* vals, hash_pair<KEY_TY
  VIRTUALWARP * subwarp_relative_index + VIRTUALWARP);
  }
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  uint64_t subwarp_relative_index = (wid) % (WARPSIZE / VIRTUALWARP);
  uint64_t submask;
  if constexpr (elementsPerWarp == 1) {
@@ -1168,7 +1168,7 @@ __global__ void retrieve_kernel(hash_pair<KEY_TYPE, VAL_TYPE>* src, hash_pair<KE
  VIRTUALWARP * subwarp_relative_index + VIRTUALWARP);
  }
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  uint64_t subwarp_relative_index = (wid) % (WARPSIZE / VIRTUALWARP);
  uint64_t submask;
  if constexpr (elementsPerWarp == 1) {
@@ -1239,7 +1239,7 @@ __global__ void delete_kernel(KEY_TYPE* keys, hash_pair<KEY_TYPE, VAL_TYPE>* buc
  VIRTUALWARP * subwarp_relative_index + VIRTUALWARP);
  }
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  uint64_t subwarp_relative_index = (wid) % (WARPSIZE / VIRTUALWARP);
  uint64_t submask;
  if constexpr (elementsPerWarp == 1) {

diff --git a/include/splitvector/archMacros.h b/include/splitvector/archMacros.h
@@ -92,7 +92,7 @@
 #define split_gpuMemoryAdvise cudaMemoryAdvise
 #define split_gpuMemAdvise cudaMemAdvise
 
-#elif __HIP_PLATFORM_HCC___
+#elif __HIP__
 
 #define split_gpuGetLastError hipGetLastError
 #define split_gpuGetErrorString hipGetErrorString

diff --git a/include/splitvector/gpu_wrappers.h b/include/splitvector/gpu_wrappers.h
@@ -23,7 +23,7 @@
 
 #ifdef __NVCC__
 #include <cuda_runtime_api.h>
-#else
+#elifdef __HIP__
 #include <hip/hip_runtime.h>
 #include <hip/hip_runtime_api.h>
 #endif
@@ -154,7 +154,7 @@ __device__ __forceinline__ T s_warpVote(bool predicate, T votingMask = T(-1)) no
  return __ballot_sync(votingMask, predicate);
 #endif
 
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  return __ballot(predicate);
 #endif
 }
@@ -172,7 +172,7 @@ __device__ __forceinline__ int s_findFirstSig(T mask) noexcept {
  return __ffs(mask);
 #endif
 
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  return __ffsll((unsigned long long)mask);
 #endif
 }
@@ -191,7 +191,7 @@ __device__ __forceinline__ int s_warpVoteAny(bool predicate, T votingMask = T(-1
  return __any_sync(votingMask, predicate);
 #endif
 
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  return __any(predicate);
 #endif
 }
@@ -208,7 +208,7 @@ __device__ __forceinline__ uint32_t s_pop_count(T mask) noexcept {
 #ifdef __NVCC__
  return __popc(mask);
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  if constexpr (sizeof(T) == 4) {
  return __popc(mask);
  } else if constexpr (sizeof(mask) == 8) {
@@ -236,7 +236,7 @@ __device__ __forceinline__ T s_shuffle(T variable, unsigned int source, U mask =
 #ifdef __NVCC__
  return __shfl_sync(mask, variable, source);
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  return __shfl(variable, source);
 #endif
 }
@@ -257,7 +257,7 @@ __device__ __forceinline__ T s_shuffle_down(T variable, unsigned int delta, U ma
 #ifdef __NVCC__
  return __shfl_down_sync(mask, variable, delta);
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
  return __shfl_down(variable, delta);
 #endif
 }

diff --git a/include/splitvector/split_allocators.h b/include/splitvector/split_allocators.h
@@ -22,6 +22,7 @@
  * */
 #pragma once
 #include "archMacros.h"
+#include "gpu_wrappers.h"
 #include <cassert>
 namespace split {
 
@@ -37,7 +38,7 @@ static void cuda_error(cudaError_t err, const char* file, int line) {
  }
 }
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
 /* Define the HIP error checking macro */
 #define SPLIT_CHECK_ERR(err) (split::hip_error(err, __FILE__, __LINE__))
 static void hip_error(hipError_t err, const char* file, int line) {

diff --git a/include/splitvector/split_tools.h b/include/splitvector/split_tools.h
@@ -38,7 +38,7 @@
 #define SPLIT_VOTING_MASK 0xFFFFFFFF // 32-bit wide for split_gpu warps
 #define WARPLENGTH 32
 #endif
-#ifdef __HIP_PLATFORM_HCC___
+#ifdef __HIP__
 #define SPLIT_VOTING_MASK 0xFFFFFFFFFFFFFFFFull // 64-bit wide for amd warps
 #define WARPLENGTH 64