Update master to most recent dev

include/common.h

@@ -50,12 +50,34 @@ constexpr inline size_t nextPow2(size_t v) noexcept {
  * @brief Computes the next optimal overflow for the hasher kernels
  */
 HASHINATOR_HOSTDEVICE
-[[nodiscard]]
+//[[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);
 }
 
+inline bool isDeviceAccessible(void* ptr){
+#ifdef __NVCC__
+ cudaPointerAttributes attributes;
+ cudaPointerGetAttributes(&attributes, ptr);
+ if (attributes.type != cudaMemoryType::cudaMemoryTypeManaged &&
+ attributes.type != cudaMemoryType::cudaMemoryTypeDevice) {
+ return false;
+ }
+ return true;
+#endif
+
+#ifdef __HIP__
+ hipPointerAttribute_t attributes;
+ hipPointerGetAttributes(&attributes, ptr);
+ if (attributes.type != hipMemoryType::hipMemoryTypeManaged &&
+ attributes.type != hipMemoryType::hipMemoryTypeDevice) {
+ return false;
+ }
+ return true;
+#endif
+}
+
 /**
  * @brief Enum for error checking in Hahsinator.
  */

include/hashinator/defaults.h

@@ -40,4 +40,17 @@ constexpr int MAX_BLOCKSIZE = 1024;
 template <typename T>
 using DefaultHashFunction = HashFunctions::Fibonacci<T>;
 } // namespace defaults
+
+struct Info {
+ Info(){};
+ Info(int sz)
+ : sizePower(sz), fill(0), currentMaxBucketOverflow(defaults::BUCKET_OVERFLOW), tombstoneCounter(0),
+ err(status::invalid) {}
+ int sizePower;
+ size_t fill;
+ size_t currentMaxBucketOverflow;
+ size_t tombstoneCounter;
+ status err;
+};
+
 } // namespace Hashinator

include/hashinator/hashers.h

@@ -143,20 +143,22 @@ class Hasher {
 
  // Reset wrapper
  static void reset(hash_pair<KEY_TYPE, VAL_TYPE>* src, hash_pair<KEY_TYPE, VAL_TYPE>* dst, const int sizePower,
- size_t maxoverflow, size_t len, split_gpuStream_t s = 0) {
+ size_t maxoverflow, Hashinator::Info* info, size_t len, split_gpuStream_t s = 0) {
  size_t blocks, blockSize;
  launchParams(len, blocks, blockSize);
  Hashinator::Hashers::reset_to_empty<KEY_TYPE, VAL_TYPE, EMPTYBUCKET, HashFunction, defaults::WARPSIZE,
  elementsPerWarp>
- <<<blocks, blockSize, 0, s>>>(src, dst, sizePower, maxoverflow, len);
+ <<<blocks, blockSize, 0, s>>>(src, dst, sizePower, maxoverflow, info, len);
  SPLIT_CHECK_ERR(split_gpuStreamSynchronize(s));
  }
 
  // Reset wrapper for all elements
- static void reset_all(hash_pair<KEY_TYPE, VAL_TYPE>* dst, size_t len, split_gpuStream_t s = 0) {
+ static void reset_all(hash_pair<KEY_TYPE, VAL_TYPE>* dst, Hashinator::Info* info, size_t len,
+ split_gpuStream_t s = 0) {
  size_t blocksNeeded = len / defaults::MAX_BLOCKSIZE;
  blocksNeeded = blocksNeeded + (blocksNeeded == 0);
- reset_all_to_empty<KEY_TYPE, VAL_TYPE, EMPTYBUCKET><<<blocksNeeded, defaults::MAX_BLOCKSIZE, 0, s>>>(dst, len);
+ reset_all_to_empty<KEY_TYPE, VAL_TYPE, EMPTYBUCKET>
+ <<<blocksNeeded, defaults::MAX_BLOCKSIZE, 0, s>>>(dst, info, len);
  SPLIT_CHECK_ERR(split_gpuStreamSynchronize(s));
  }