Offset-based alternate bucket policy from "Smaller and More Flexible Cuckoo Filters" by Schmitz, Zentgraf, and Rahmann. More...

Static Public Member Functions
template<typename H >
static __host__ __device__ uint64_t	hash64 (const H &key)
	Computes a 64-bit hash of the key.

static __host__ __device__ TagType	getPureFingerprint (TagType fp)
	Extracts the pure fingerprint (without choice bit).

static __host__ __device__ TagType	getChoiceBit (TagType fp)
	Extracts the choice bit (0 or 1).

static __host__ __device__ TagType	setChoiceBit (TagType fp, TagType choice)
	Sets the choice bit in a fingerprint.

static __host__ __device__ size_t	computeOffset (TagType pureFp, size_t numBuckets)
	Computes a non-zero offset from a fingerprint (uses pure fp).

static __host__ __device__ cuda::std::tuple< size_t, size_t, TagType, TagType >	getCandidateBucketsAndFPs (const KeyType &key, size_t numBuckets)
	Computes the primary and alternate bucket indices for a given key.

static __host__ __device__ size_t	getAlternateBucket (size_t bucket, TagType fp, size_t numBuckets)
	Computes the alternate bucket index using asymmetric offset.

static __host__ __device__ cuda::std::tuple< size_t, TagType >	getAlternateBucketWithNewFp (size_t bucket, TagType fp, size_t numBuckets)
	Computes alternate bucket AND updated fingerprint for eviction.

static size_t	calculateNumBuckets (size_t capacity)
	Calculates number of buckets without requiring power of two.

Static Public Attributes
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"

Detailed Description

template<typename KeyType, typename TagType, size_t bitsPerTag, size_t bucketSize>
struct cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >

Offset-based alternate bucket policy from "Smaller and More Flexible Cuckoo Filters" by Schmitz, Zentgraf, and Rahmann.

Uses an asymmetric offset with a choice bit:

b' = (b + offset(fp)) mod numBuckets (when choice=0, primary bucket)
b = (b' - offset(fp)) mod numBuckets (when choice=1, alternate bucket)

The choice bit is stored in the MSB of the fingerprint:

Lower (bitsPerTag-1) bits: actual fingerprint
MSB: choice bit (0 = primary, 1 = alternate)

This allows non-power-of-two bucket counts with proper symmetry.

Definition at line 201 of file bucket_policies.cuh.

Member Function Documentation

◆ calculateNumBuckets()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static size_t cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::calculateNumBuckets ( size_t capacity )

inlinestatic

Calculates number of buckets without requiring power of two.

Definition at line 348 of file bucket_policies.cuh.

                                                       {
        auto requiredBuckets = std::ceil(static_cast<double>(capacity) / bucketSize);
        return static_cast<size_t>(requiredBuckets);
    }

◆ computeOffset()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ size_t cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::computeOffset	(	TagType	pureFp,
		size_t	numBuckets
	)

inlinestatic

Computes a non-zero offset from a fingerprint (uses pure fp).

Parameters

pureFp	Pure fingerprint (without choice bit)
numBuckets	Total number of buckets

Returns: Non-zero offset value in [1, numBuckets-1]

Definition at line 250 of file bucket_policies.cuh.

                                                                                       {
        const uint64_t fpHash = hash64(pureFp);
        const uint64_t offset = fpHash % numBuckets;
        return offset == 0 ? 1 : offset;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ getAlternateBucket()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ size_t cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::getAlternateBucket	(	size_t	bucket,
		TagType	fp,
		size_t	numBuckets
	)

inlinestatic

Computes the alternate bucket index using asymmetric offset.

The choice bit determines direction:

choice=0: we're at primary bucket, add offset to get alternate
choice=1: we're at alternate bucket, subtract offset to get primary

Parameters

bucket	Current bucket index.
fp	Fingerprint (with choice bit).
numBuckets	Total number of buckets.

Returns: Alternate bucket index.

Definition at line 305 of file bucket_policies.cuh.

                                                                     {
        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;
        }
    }

Here is the call graph for this function:

◆ getAlternateBucketWithNewFp()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ cuda::std::tuple< size_t, TagType > cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::getAlternateBucketWithNewFp	(	size_t	bucket,
		TagType	fp,
		size_t	numBuckets
	)

inlinestatic

Computes alternate bucket AND updated fingerprint for eviction.

Returns what the alternate bucket would be AND what the new fp would be (with flipped choice bit). This is used during eviction when we need both.

Parameters

bucket	Current bucket index.
fp	Fingerprint (with choice bit).
numBuckets	Total number of buckets.

Returns: Tuple of (alternate bucket index, new fingerprint with flipped choice bit).

Definition at line 329 of file bucket_policies.cuh.

                                                                              {
        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};
    }

Here is the call graph for this function:

◆ getCandidateBucketsAndFPs()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ cuda::std::tuple< size_t, size_t, TagType, TagType > cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::getCandidateBucketsAndFPs	(	const KeyType &	key,
		size_t	numBuckets
	)

inlinestatic

Computes the primary and alternate bucket indices for a given key.

Returns both fingerprint variants to avoid redundant computations:

fp1: fingerprint with choice=0 for primary bucket
fp2: fingerprint with choice=1 for alternate bucket

Parameters

key	Key to hash
numBuckets	Number of buckets in the filter

Returns: A tuple of (bucket1, bucket2, fingerprintForBucket1, fingerprintForBucket2)

Definition at line 268 of file bucket_policies.cuh.

                                                                     {
        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};
    }

Here is the call graph for this function:

◆ getChoiceBit()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ TagType cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::getChoiceBit ( TagType fp )

inlinestatic

Extracts the choice bit (0 or 1).

Definition at line 232 of file bucket_policies.cuh.

                                                                {
        return (fp >> choiceBitPos) & 1;
    }

Here is the caller graph for this function:

◆ getPureFingerprint()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ TagType cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::getPureFingerprint ( TagType fp )

inlinestatic

Extracts the pure fingerprint (without choice bit).

Definition at line 225 of file bucket_policies.cuh.

                                                                      {
        return fp & pureFpMask;
    }

Here is the caller graph for this function:

◆ hash64()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

template<typename H >

static __host__ __device__ uint64_t cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::hash64 ( const H & key )

inlinestatic

Computes a 64-bit hash of the key.

Definition at line 218 of file bucket_policies.cuh.

                                                             {
        return xxhash::xxhash64(key);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ setChoiceBit()

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

static __host__ __device__ TagType cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::setChoiceBit	(	TagType	fp,
		TagType	choice
	)

inlinestatic

Sets the choice bit in a fingerprint.

Definition at line 239 of file bucket_policies.cuh.

                                                                                {
        return (fp & pureFpMask) | (choice << choiceBitPos);
    }

Here is the caller graph for this function:

Member Data Documentation

◆ choiceBitMask

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

constexpr TagType cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::choiceBitMask = TagType(1) << choiceBitPos

staticconstexpr

Definition at line 206 of file bucket_policies.cuh.

◆ choiceBitPos

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

constexpr size_t cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::choiceBitPos = bitsPerTag - 1

staticconstexpr

Definition at line 205 of file bucket_policies.cuh.

◆ fpMask

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

constexpr size_t cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::fpMask = (1ULL << bitsPerTag) - 1

staticconstexpr

Definition at line 208 of file bucket_policies.cuh.

◆ name

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

constexpr char cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::name[] = "OffsetAltBucketPolicy"

staticconstexpr

Definition at line 212 of file bucket_policies.cuh.

◆ pureFpMask

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

constexpr TagType cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::pureFpMask = (TagType(1) << choiceBitPos) - 1

staticconstexpr

Definition at line 207 of file bucket_policies.cuh.

◆ usesChoiceBit

template<typename KeyType , typename TagType , size_t bitsPerTag, size_t bucketSize>

constexpr bool cuckoogpu::OffsetAltBucketPolicy< KeyType, TagType, bitsPerTag, bucketSize >::usesChoiceBit = true

staticconstexpr

Definition at line 210 of file bucket_policies.cuh.

The documentation for this struct was generated from the following file:

bucket_policies.cuh

Static Public Member Functions

Static Public Attributes

Detailed Description

Member Function Documentation

◆ calculateNumBuckets()

◆ computeOffset()

◆ getAlternateBucket()

◆ getAlternateBucketWithNewFp()

◆ getCandidateBucketsAndFPs()

◆ getChoiceBit()

◆ getPureFingerprint()

◆ hash64()

◆ setChoiceBit()

Member Data Documentation

◆ choiceBitMask

◆ choiceBitPos

◆ fpMask

◆ name

◆ pureFpMask

◆ usesChoiceBit