Merge dev

kstppd · Apr 9, 2024 · cd1ea1c · cd1ea1c
2 parents 8725a4b + 3f214f5
commit cd1ea1c
Show file tree

Hide file tree

Showing 12 changed files with 528 additions and 166 deletions.
diff --git a/include/common.h b/include/common.h
@@ -49,8 +49,8 @@ constexpr inline size_t nextPow2(size_t v) noexcept {
 /**
  * @brief Computes the next optimal overflow for the hasher kernels
  */
-[[nodiscard]]
 HASHINATOR_HOSTDEVICE
+//[[nodiscard]]
 constexpr inline size_t nextOverflow(size_t currentOverflow, size_t virtualWarp) noexcept {
  size_t remainder = currentOverflow % virtualWarp;
  return ((remainder)==0)?currentOverflow: currentOverflow + (virtualWarp - remainder);

diff --git a/include/hashinator/hashmap/hashmap.h b/include/hashinator/hashmap/hashmap.h
@@ -416,6 +416,12 @@ class Hashmap {
  HASHINATOR_HOSTDEVICE
  size_t bucket_count() const { return buckets.size(); }
 
+ HASHINATOR_HOSTDEVICE
+ constexpr KEY_TYPE get_emptybucket() const { return EMPTYBUCKET; }
+
+ HASHINATOR_HOSTDEVICE
+ constexpr KEY_TYPE get_tombstone() const { return TOMBSTONE; }
+
  HASHINATOR_HOSTDEVICE
  float load_factor() const { return (float)size() / bucket_count(); }
 
@@ -1110,7 +1116,7 @@ class Hashmap {
 
  // Allocate with Mempool
  const size_t memory_for_pool = 8 * nBlocks * sizeof(uint32_t);
- split::tools::Cuda_mempool mPool(memory_for_pool, s);
+ split::tools::splitStackArena mPool(memory_for_pool, s);
  size_t retval =
  split::tools::copy_if_raw<hash_pair<KEY_TYPE, VAL_TYPE>, Rule, defaults::MAX_BLOCKSIZE, defaults::WARPSIZE>(
  buckets.data(), elements, buckets.size(), rule, nBlocks, mPool, s);
@@ -1276,29 +1282,35 @@ class Hashmap {
  if (neededPowerSize > _mapInfo->sizePower) {
  resize(neededPowerSize, targets::device, s);
  }
- buckets.optimizeGPU(s);
  DeviceHasher::insert(src, buckets.data(), _mapInfo->sizePower, _mapInfo->currentMaxBucketOverflow,
  &_mapInfo->currentMaxBucketOverflow, &_mapInfo->fill, len, &_mapInfo->err, s);
  return;
  }
 
  // Uses Hasher's retrieve_kernel to read all elements
- void retrieve(KEY_TYPE* keys, VAL_TYPE* vals, size_t len, split_gpuStream_t s = 0) {
- buckets.optimizeGPU(s);
+ void retrieve(KEY_TYPE* keys, VAL_TYPE* vals, size_t len, split_gpuStream_t s = 0,bool prefetches=true) {
+ if (prefetches){
+ buckets.optimizeGPU(s);
+ }
  DeviceHasher::retrieve(keys, vals, buckets.data(), _mapInfo->sizePower, _mapInfo->currentMaxBucketOverflow, len,
  s);
  return;
  }
 
  // Uses Hasher's retrieve_kernel to read all elements
- void retrieve(hash_pair<KEY_TYPE, VAL_TYPE>* src, size_t len, split_gpuStream_t s = 0) {
- buckets.optimizeGPU(s);
+ void retrieve(hash_pair<KEY_TYPE, VAL_TYPE>* src, size_t len, split_gpuStream_t s = 0, bool prefetches=true) {
+ if (prefetches){
+ buckets.optimizeGPU(s);
+ }
  DeviceHasher::retrieve(src, buckets.data(), _mapInfo->sizePower, _mapInfo->currentMaxBucketOverflow, len, s);
  return;
  }
 
  // Uses Hasher's erase_kernel to delete all elements
- void erase(KEY_TYPE* keys, size_t len, split_gpuStream_t s = 0) {
+ void erase(KEY_TYPE* keys, size_t len, split_gpuStream_t s = 0,bool prefetches=true) {
+ if (prefetches){
+ buckets.optimizeGPU(s);
+ }
  // Remember the last numeber of tombstones
  size_t tbStore = tombstone_count();
  DeviceHasher::erase(keys, buckets.data(), &_mapInfo->tombstoneCounter, _mapInfo->sizePower,

diff --git a/include/hashinator/unordered_set/unordered_set.h b/include/hashinator/unordered_set/unordered_set.h
@@ -330,7 +330,7 @@ class Unordered_Set {
 
  // Allocate with Mempool
  const size_t memory_for_pool = 8 * nBlocks * sizeof(uint32_t);
- split::tools::Cuda_mempool mPool(memory_for_pool, s);
+ split::tools::splitStackArena mPool(memory_for_pool, s);
  size_t retval =
  split::tools::copy_if_raw<KEY_TYPE, Rule, defaults::MAX_BLOCKSIZE, defaults::WARPSIZE>(
  buckets.data(), elements, buckets.size(), rule, nBlocks, mPool, s);

diff --git a/include/splitvector/gpu_wrappers.h b/include/splitvector/gpu_wrappers.h
@@ -261,5 +261,26 @@ __device__ __forceinline__ T s_shuffle_down(T variable, unsigned int delta, U ma
  return __shfl_down(variable, delta);
 #endif
 }
+
+/**
+ * @brief Wrapper for performing an up register shuffle operation.
+ *
+ * @tparam T The data type of the variable.
+ * @tparam U The data type of the mask.
+ * @param variable The variable to shuffle.
+ * @param delta The offset.
+ * @param mask Voting mask.
+ * @return The shuffled variable.
+ */
+template <typename T, typename U>
+__device__ __forceinline__ T s_shuffle_up(T variable, unsigned int delta, U mask = 0) noexcept {
+ static_assert(std::is_integral<T>::value && "Only integers supported");
+#ifdef __NVCC__
+ return __shfl_up_sync(mask, variable, delta);
+#endif
+#ifdef __HIP__
+ return __shfl_up(variable, delta);
+#endif
+}
 } // namespace split
 #endif