cuckoo_hash.cuh

Go to the documentation of this file.

// SPDX-License-Identifier: Apache-2.0
 
#pragma once
#include <cuda_runtime.h>
#include <cmath>
#include <limits>
#include <span>
#include <utility>
#include <random>
#include <fss/prp.cuh>
#include <stdexcept>
 
namespace fss::cuckoo_hash {

constexpr int ChBucket(int t, int lambda);
 
namespace detail {
 
// Candidate count that callers without access to the schemes' private constant
// use. The three-choice construction fixes the count; the schemes pass theirs.
inline constexpr int kDefaultCuckooChoices = 3;

constexpr int kRejectedBucket = std::numeric_limits<int>::max();

struct LocateResult {
  int bucket;  
  __uint128_t index;  
};

template <typename Prp, typename In>
  requires Permutable<Prp>
struct PrpHash {
  Prp prp;  

  __host__ __device__ LocateResult Locate(int4 sigma, In x, int k, __uint128_t n,
      __uint128_t b_size, int choices = kDefaultCuckooChoices);
};

template <typename Prp, typename In>
  requires Permutable<Prp>
struct Compact {
  Prp prp;  

  int Run(std::span<const In> as, int m, int4 sigma, __uint128_t n, __uint128_t b_size, int max_evictions,
      std::span<std::pair<int, int>> table);
};
 
}  // namespace detail
 
// -----------------------------------------------------------------------------
 
namespace detail {
 
constexpr double Log2(double x) {
  if (x <= 0) return -1e308;
  int e = 0;
  double m = x;
  while (m >= 2.0) {
    m /= 2.0;
    ++e;
  }
  while (m < 1.0) {
    m *= 2.0;
    --e;
  }
  double y = (m - 1.0) / (m + 1.0);
  double y2 = y * y, sum = 0, term = y;
  for (int k = 0; k < 40; ++k) {
    sum += term / (2 * k + 1);
    term *= y2;
  }
  return e + 2.0 * sum / 0.6931471805599453;
}
 
constexpr double Ceil(double x) {
  auto i = static_cast<long long>(x);
  return (x > static_cast<double>(i)) ? static_cast<double>(i + 1) : static_cast<double>(i);
}

__host__ __device__ constexpr __uint128_t CeilDiv(__uint128_t num, __uint128_t den) {
  return (num - 1) / den + 1;
}

__host__ __device__ constexpr __uint128_t BucketIndexBound(__uint128_t n, int choices, __uint128_t bucket_size) {
  return CeilDiv(n * choices, bucket_size);
}

__host__ __device__ constexpr __uint128_t MinBucketSize(__uint128_t n, int choices, __uint128_t bucket_count) {
  return CeilDiv(n * choices, bucket_count);
}

__host__ __device__ constexpr bool ValidBucketSize(__uint128_t b_size_rt,
    __uint128_t bucket_size_min, int bucket_bits) {
  return b_size_rt >= bucket_size_min && b_size_rt <= (__uint128_t{1} << bucket_bits);
}
 
}  // namespace detail
 
constexpr int ChBucket(int t, int lambda) {
  // Remark 1 applies for t >= 30.
  // For smaller t, the full Lemma 5 formula with Normal CDF factors is needed,
  // but the paper considers t >= 30 to capture nearly all practical use cases.
  if (t < 30)
    throw std::invalid_argument("cuckoo_hash::ChBucket: t must be at least 30");
  double td = static_cast<double>(t);
  double e = (static_cast<double>(lambda) + 130.0 + detail::Log2(td)) / 123.5;
  return static_cast<int>(detail::Ceil(e * td));
}
 
namespace detail {
 
template <typename Prp, typename In>
  requires Permutable<Prp>
inline __host__ __device__ LocateResult PrpHash<Prp, In>::Locate(
    int4 sigma, In x, int k, __uint128_t n, __uint128_t b_size, int choices) {
  __uint128_t val = static_cast<__uint128_t>(x) + n * k;
  __uint128_t domain_size = n * static_cast<__uint128_t>(choices);
  __uint128_t y = prp.Permu(sigma, val, domain_size);
 
  // Compare the quotient in __uint128_t: a zero b_size has no bucket at all,
  // and an over-int quotient must saturate rather than wrap negative through
  // the cast, because the callers filter buckets with `bucket >= bound` and
  // would accept a negative value before indexing the per-bucket key arrays.
  if (b_size == 0) {
    return {kRejectedBucket, 0};
  }
  __uint128_t bucket = y / b_size;
  if (bucket > static_cast<__uint128_t>(kRejectedBucket)) {
    return {kRejectedBucket, 0};
  }
  return {static_cast<int>(bucket), y % b_size};
}
 
template <typename Prp, typename In>
  requires Permutable<Prp>
inline int Compact<Prp, In>::Run(std::span<const In> as, int m, int4 sigma, __uint128_t n,
    __uint128_t b_size, int max_evictions, std::span<std::pair<int, int>> table) {
  PrpHash<Prp, In> hasher{prp};
  int t = static_cast<int>(as.size());
 
  // Initialize table to empty.
  for (int i = 0; i < m; ++i) {
    table[i] = {-1, -1};
  }
 
  std::mt19937 rng(42);
 
  for (int omega = 0; omega < t; ++omega) {
    int cur_idx = omega;
    int cur_k = static_cast<int>(rng() % 3);
    int evictions = 0;
 
    for (;;) {
      auto [bucket, _index] = hasher.Locate(sigma, as[cur_idx], cur_k, n, b_size);
      // Clamp bucket to [0, m).
      bucket = bucket % m;
 
      if (table[bucket].first == -1) {
        // Empty slot found.
        table[bucket] = {cur_idx, cur_k};
        break;
      }
 
      // Evict the existing entry and take its place.
      int evicted_idx = table[bucket].first;
      table[bucket] = {cur_idx, cur_k};
 
      // The evicted element becomes the current element with a new random hash.
      cur_idx = evicted_idx;
      cur_k = static_cast<int>(rng() % 3);
 
      ++evictions;
      if (evictions > max_evictions) {
        return 1;
      }
    }
  }
 
  return 0;
}
 
}  // namespace detail
 
}  // namespace fss::cuckoo_hash