bucket_policies.cuh

Go to the documentation of this file.

#pragma once
 
#include <cmath>
#include <cstdint>
#include <cuda/std/tuple>
#include "hashutil.cuh"
#include "helpers.cuh"
 
namespace cuckoogpu {
 
template <typename KeyType, typename TagType, size_t bitsPerTag, size_t bucketSize>
struct XorAltBucketPolicy {
    static constexpr size_t fpMask = (1ULL << bitsPerTag) - 1;
 
    static constexpr bool usesChoiceBit = false;
 
    static constexpr char name[] = "XorAltBucketPolicy";
 
    template <typename H>
    static __host__ __device__ uint64_t hash64(const H& key) {
        return xxhash::xxhash64(key);
    }
 
    static __host__ __device__ cuda::std::tuple<size_t, size_t, TagType, TagType>
    getCandidateBucketsAndFPs(const KeyType& key, size_t numBuckets) {
        const uint64_t h = hash64(key);
 
        // Upper 32 bits for the fingerprint
        const uint32_t h_fp = h >> 32;
        const uint32_t masked = h_fp & fpMask;
 
        const auto fp = TagType(masked == 0 ? 1 : masked);
 
        // Lower 32 bits for the bucket indices
        const uint32_t h_bucket = h & 0xFFFFFFFF;
        const size_t i1 = h_bucket & (numBuckets - 1);
        const size_t i2 = getAlternateBucket(i1, fp, numBuckets);
 
        return {i1, i2, fp, fp};
    }
 
    static __host__ __device__ size_t
    getAlternateBucket(size_t bucket, TagType fp, size_t numBuckets) {
        return bucket ^ (hash64(fp) & (numBuckets - 1));
    }
 
    static size_t calculateNumBuckets(size_t capacity) {
        auto requiredBuckets = std::ceil(static_cast<double>(capacity) / bucketSize);
        return detail::nextPowerOfTwo(requiredBuckets);
    }
};
 
template <typename KeyType, typename TagType, size_t bitsPerTag, size_t bucketSize>
struct AddSubAltBucketPolicy {
    static constexpr size_t fpMask = (1ULL << bitsPerTag) - 1;
 
    static constexpr bool usesChoiceBit = false;
 
    static constexpr char name[] = "AddSubAltBucketPolicy";
 
    template <typename H>
    static __host__ __device__ uint64_t hash64(const H& key) {
        return xxhash::xxhash64(key);
    }
 
    static __host__ __device__ cuda::std::tuple<size_t, size_t, TagType, TagType>
    getCandidateBucketsAndFPs(const KeyType& key, size_t numBuckets) {
        const uint64_t h = hash64(key);
 
        // Upper 32 bits for the fingerprint
        const uint32_t h_fp = h >> 32;
        const uint32_t masked = h_fp & fpMask;
        const auto fp = TagType(masked == 0 ? 1 : masked);
 
        // Lower 32 bits for bucket index in block 0
        const uint32_t h_bucket = h & 0xFFFFFFFF;
        const size_t bucketsPerBlock = numBuckets / 2;
 
        const size_t i1 = h_bucket % bucketsPerBlock;
        const size_t i2 = getAlternateBucket(i1, fp, numBuckets);
 
        return {i1, i2, fp, fp};
    }
 
    static __host__ __device__ size_t
    getAlternateBucket(size_t bucket, TagType fp, size_t numBuckets) {
        const size_t bucketsPerBlock = numBuckets / 2;
        const uint64_t fpHash = hash64(fp) % bucketsPerBlock;
 
        if (bucket < bucketsPerBlock) {
            return ((bucket + fpHash) % bucketsPerBlock) + bucketsPerBlock;
        } else {
            return (bucket - fpHash) % bucketsPerBlock;
        }
    }
 
    static size_t calculateNumBuckets(size_t capacity) {
        auto requiredBuckets = std::ceil(static_cast<double>(capacity) / bucketSize);
 
        // Round up to next even number to ensure equal block sizes
        if (static_cast<size_t>(requiredBuckets) % 2 != 0) {
            requiredBuckets += 1;
        }
 
        return static_cast<size_t>(requiredBuckets);
    }
};
 
template <typename KeyType, typename TagType, size_t bitsPerTag, size_t bucketSize>
struct OffsetAltBucketPolicy {
    static_assert(bitsPerTag >= 2, "bitsPerTag must be at least 2 for choice bit");
 
    // Choice bit is MSB, fingerprint uses remaining bits
    static constexpr size_t choiceBitPos = bitsPerTag - 1;
    static constexpr TagType choiceBitMask = TagType(1) << choiceBitPos;
    static constexpr TagType pureFpMask = (TagType(1) << choiceBitPos) - 1;
    static constexpr size_t fpMask = (1ULL << bitsPerTag) - 1;
 
    static constexpr bool usesChoiceBit = true;
 
    static constexpr char name[] = "OffsetAltBucketPolicy";
 
    template <typename H>
    static __host__ __device__ uint64_t hash64(const H& key) {
        return xxhash::xxhash64(key);
    }
 
    static __host__ __device__ TagType getPureFingerprint(TagType fp) {
        return fp & pureFpMask;
    }
 
    static __host__ __device__ TagType getChoiceBit(TagType fp) {
        return (fp >> choiceBitPos) & 1;
    }
 
    static __host__ __device__ TagType setChoiceBit(TagType fp, TagType choice) {
        return (fp & pureFpMask) | (choice << choiceBitPos);
    }
 
    static __host__ __device__ size_t computeOffset(TagType pureFp, size_t numBuckets) {
        const uint64_t fpHash = hash64(pureFp);
        const uint64_t offset = fpHash % numBuckets;
        return offset == 0 ? 1 : offset;
    }
 
    static __host__ __device__ cuda::std::tuple<size_t, size_t, TagType, TagType>
    getCandidateBucketsAndFPs(const KeyType& key, size_t numBuckets) {
        const uint64_t h = hash64(key);
 
        // Upper 32 bits for the fingerprint (use bitsPerTag-1 bits)
        const uint32_t h_fp = h >> 32;
        const uint32_t masked = h_fp & pureFpMask;
        // Ensure non-zero (0 is reserved for empty)
        const TagType pureFp = TagType(masked == 0 ? 1 : masked);
 
        // Create both fingerprint variants (choice=0 and choice=1)
        const TagType fp1 = setChoiceBit(pureFp, 0);
        const TagType fp2 = setChoiceBit(pureFp, 1);
 
        // Lower 32 bits for the bucket index
        const uint32_t h_bucket = h & 0xFFFFFFFF;
        const size_t i1 = h_bucket % numBuckets;
 
        // Compute i2 using the fingerprint offset
        const size_t offset = computeOffset(pureFp, numBuckets);
        const size_t i2 = (i1 + offset) % numBuckets;
 
        return {i1, i2, fp1, fp2};
    }
 
    static __host__ __device__ size_t
    getAlternateBucket(size_t bucket, TagType fp, size_t numBuckets) {
        const TagType pureFp = getPureFingerprint(fp);
        const TagType choice = getChoiceBit(fp);
        const size_t offset = computeOffset(pureFp, numBuckets);
 
        if (choice == 0) {
            return (bucket + offset) % numBuckets;
        } else {
            return (bucket + numBuckets - offset) % numBuckets;
        }
    }
 
    static __host__ __device__ cuda::std::tuple<size_t, TagType>
    getAlternateBucketWithNewFp(size_t bucket, TagType fp, size_t numBuckets) {
        const TagType pureFp = getPureFingerprint(fp);
        const TagType choice = getChoiceBit(fp);
        const size_t offset = computeOffset(pureFp, numBuckets);
 
        size_t alt;
        if (choice == 0) {
            alt = (bucket + offset) % numBuckets;
        } else {
            alt = (bucket + numBuckets - offset) % numBuckets;
        }
 
        TagType newFp = setChoiceBit(pureFp, 1 - choice);
        return {alt, newFp};
    }
 
    static size_t calculateNumBuckets(size_t capacity) {
        auto requiredBuckets = std::ceil(static_cast<double>(capacity) / bucketSize);
        return static_cast<size_t>(requiredBuckets);
    }
};
 
}  // namespace cuckoogpu