Merge pull request #48 from kstppd/fasterCompactions

Faster compactions for up to 1024 elements

include/hashinator/hashinator.h

@@ -1110,7 +1110,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, elements, rule, nBlocks, mPool, s);

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