A CUDA-accelerated Cuckoo Filter implementation. More...

Collaboration diagram for cuckoogpu::Filter< Config >:

Classes
struct	Bucket
	Bucket structure that holds the fingerprint and tags for a given bucket. More...

struct	PackedTag
	This is used by the sorted insert kernel to store the fingerprint and primary bucket index in a compact format that allows you to sort them directly since the bucket index lives in the upper bits. More...

Public Types
using	T = typename Config::KeyType

using	TagType = typename Config::TagType

using	AltBucketPolicy = typename Config::AltBucketPolicy

using	PackedTagType = typename std::conditional< bitsPerTag<=8, uint32_t, uint64_t >::type

Public Member Functions
	Filter (const Filter &)=delete

Filter &	operator= (const Filter &)=delete

	Filter (size_t capacity)
	Constructs a new Cuckoo Filter.

	~Filter ()
	Destroys the Cuckoo Filter.

size_t	insertMany (const T d_keys, const size_t n, bool d_output=nullptr, cudaStream_t stream={})
	Inserts a batch of keys into the filter.

size_t	insertManyWithEvictionCounts (const T d_keys, const size_t n, uint32_t d_evictionAttempts, bool *d_output=nullptr, cudaStream_t stream={})
	Inserts a batch of keys and records per-attempt eviction counts.

size_t	insertManySorted (const T d_keys, const size_t n, bool d_output=nullptr, cudaStream_t stream={})
	This pre-sorts the input keys based on the primary bucket index to allow for coalesced memory access when you later insert them into the filter.

size_t	insertManySortedWithEvictionCounts (const T d_keys, const size_t n, uint32_t d_evictionAttempts, bool *d_output=nullptr, cudaStream_t stream={})
	Inserts a pre-sorted batch of keys and records per-attempt eviction counts.

void	containsMany (const T d_keys, const size_t n, bool d_output, cudaStream_t stream={})
	Checks for the existence of a batch of keys.

size_t	deleteMany (const T d_keys, const size_t n, bool d_output=nullptr, cudaStream_t stream={})
	Tries to remove a set of keys from the filter.

size_t	insertMany (const thrust::device_vector< T > &d_keys, thrust::device_vector< bool > &d_output, cudaStream_t stream={})
	Inserts keys from a Thrust device vector.

size_t	insertMany (const thrust::device_vector< T > &d_keys, thrust::device_vector< uint8_t > &d_output, cudaStream_t stream={})
	Inserts keys from a Thrust device vector (uint8_t output).

size_t	insertMany (const thrust::device_vector< T > &d_keys, cudaStream_t stream={})
	Inserts keys from a Thrust device vector without outputting results.

size_t	insertManyWithEvictionCounts (const thrust::device_vector< T > &d_keys, thrust::device_vector< uint32_t > &d_evictionAttempts, thrust::device_vector< uint8_t > *d_output=nullptr, cudaStream_t stream={})
	Inserts keys from a Thrust device vector and records per-attempt eviction counts.

size_t	insertManySorted (const thrust::device_vector< T > &d_keys, thrust::device_vector< bool > &d_output, cudaStream_t stream={})
	Inserts keys from a Thrust device vector, sorting them first.

size_t	insertManySorted (const thrust::device_vector< T > &d_keys, thrust::device_vector< uint8_t > &d_output, cudaStream_t stream={})
	Inserts keys from a Thrust device vector, sorting them first (uint8_t output).

size_t	insertManySorted (const thrust::device_vector< T > &d_keys, cudaStream_t stream={})
	Inserts keys from a Thrust device vector, sorting them first, without outputting results.

size_t	insertManySortedWithEvictionCounts (const thrust::device_vector< T > &d_keys, thrust::device_vector< uint32_t > &d_evictionAttempts, thrust::device_vector< uint8_t > *d_output=nullptr, cudaStream_t stream={})
	Inserts keys from a Thrust device vector, sorting first, and records per-attempt eviction counts.

void	containsMany (const thrust::device_vector< T > &d_keys, thrust::device_vector< bool > &d_output, cudaStream_t stream={})
	Checks for existence of keys in a Thrust device vector.

void	containsMany (const thrust::device_vector< T > &d_keys, thrust::device_vector< uint8_t > &d_output, cudaStream_t stream={})
	Checks for existence of keys in a Thrust device vector (uint8_t output).

size_t	deleteMany (const thrust::device_vector< T > &d_keys, thrust::device_vector< bool > &d_output, cudaStream_t stream={})
	Deletes keys in a Thrust device vector.

size_t	deleteMany (const thrust::device_vector< T > &d_keys, thrust::device_vector< uint8_t > &d_output, cudaStream_t stream={})
	Deletes keys in a Thrust device vector (uint8_t output).

size_t	deleteMany (const thrust::device_vector< T > &d_keys, cudaStream_t stream={})
	Deletes keys in a Thrust device vector without outputting results.

void	clear ()
	Clears the filter, removing all items.

float	loadFactor ()
	Calculates the current load factor of the filter.

size_t	occupiedSlots ()
	Returns the total number of occupied slots.

size_t	evictionCount ()
	Returns the total number of evictions performed.

void	resetEvictionCount ()
	Resets the eviction counter to zero.

size_t	capacity ()
	Returns the total capacity of the filter.

size_t	getNumBuckets () const
	Returns the number of buckets in the filter.

size_t	sizeInBytes () const
	Returns the size of the filter in bytes.

size_t	countOccupiedSlots ()
	Counts occupied slots by iterating over all buckets on the host.

__device__ bool	tryRemoveAtBucket (size_t bucketIdx, TagType tag)
	Attempt to remove a single instance of a fingerprint from a bucket.

__device__ bool	tryInsertAtBucket (size_t bucketIdx, TagType tag)
	Attempts to insert a tag into a specific bucket.

__device__ bool	insertWithEvictionDFS (TagType fp, size_t startBucket, uint32_t *evictionAttempts=nullptr)
	Inserts a fingerprint into the filter by evicting existing fingerprints.

__device__ bool	insertWithEvictionBFS (TagType fp, size_t startBucket, uint32_t *evictionAttempts=nullptr)
	Inserts a fingerprint using repeated shallow breadth-first attempts.

__device__ bool	insert (const T &key, uint32_t *evictionAttempts=nullptr)
	Inserts a single key into the filter.

__device__ bool	contains (const T &key) const
	Checks if a key exists in the filter.

__device__ bool	remove (const T &key)
	Removes a key from the filter.

Static Public Member Functions
template<typename H >
static __host__ __device__ uint64_t	hash64 (const H &key)

static __host__ __device__ cuda::std::tuple< size_t, size_t, TagType, TagType >	getCandidateBucketsAndFPs (const T &key, size_t numBuckets)

static __host__ __device__ size_t	getAlternateBucket (size_t bucket, TagType fp, size_t numBuckets)
	Computes the alternate bucket for a fingerprint.

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 choice bit policies.

static size_t	calculateNumBuckets (size_t capacity)
	The number of buckets is enforced to be a power of two in order to allow for efficient modulo on the bucket indices.

Public Attributes
size_t	numBuckets
	Number of buckets in the filter.

Bucket *	d_buckets
	Pointer to the device memory for the buckets.

cuda::std::atomic< size_t > *	d_numOccupied {}
	Pointer to the device memory for the occupancy counter.

cuda::std::atomic< size_t > *	d_numEvictions {}
	Pointer to the device memory for the eviction counter.

size_t	h_numOccupied = 0
	Number of occupied buckets in the filter.

Static Public Attributes
static constexpr size_t	bitsPerTag = Config::bitsPerTag

static constexpr size_t	tagEntryBytes = sizeof(TagType)

static constexpr size_t	bucketSize = Config::bucketSize

static constexpr size_t	maxEvictions = Config::maxEvictions

static constexpr size_t	blockSize = Config::blockSize

static constexpr TagType	EMPTY = 0

static constexpr size_t	fpMask = (1ULL << bitsPerTag) - 1

Detailed Description

template<typename Config>
struct cuckoogpu::Filter< Config >

A CUDA-accelerated Cuckoo Filter implementation.

This class implements a Cuckoo Filter using CUDA for high-throughput insertion, lookup, and deletion. It supports concurrent operations and uses atomic operations for thread safety within buckets.

Template Parameters

Config The configuration structure defining filter parameters.

Definition at line 144 of file CuckooFilter.cuh.

Member Typedef Documentation

◆ AltBucketPolicy

template<typename Config >

using cuckoogpu::Filter< Config >::AltBucketPolicy = typename Config::AltBucketPolicy

Definition at line 149 of file CuckooFilter.cuh.

◆ PackedTagType

template<typename Config >

using cuckoogpu::Filter< Config >::PackedTagType = typename std::conditional<bitsPerTag <= 8, uint32_t, uint64_t>::type

Definition at line 163 of file CuckooFilter.cuh.

◆ T

template<typename Config >

using cuckoogpu::Filter< Config >::T = typename Config::KeyType

Definition at line 145 of file CuckooFilter.cuh.

◆ TagType

template<typename Config >

using cuckoogpu::Filter< Config >::TagType = typename Config::TagType

Definition at line 148 of file CuckooFilter.cuh.

Constructor & Destructor Documentation

◆ Filter() [1/2]

template<typename Config >

cuckoogpu::Filter< Config >::Filter ( const Filter< Config > & )

delete

◆ Filter() [2/2]

template<typename Config >

cuckoogpu::Filter< Config >::Filter ( size_t capacity )

inlineexplicit

Constructs a new Cuckoo Filter.

Allocates device memory for buckets and occupancy counters.

Parameters

capacity Desired capacity (number of items) for the filter.

Definition at line 408 of file CuckooFilter.cuh.

                                     : numBuckets(calculateNumBuckets(capacity)) {
        CUDA_CALL(cudaMalloc(&d_buckets, numBuckets * sizeof(Bucket)));
        CUDA_CALL(cudaMalloc(&d_numOccupied, sizeof(cuda::std::atomic<size_t>)));
#ifdef CUCKOO_FILTER_COUNT_EVICTIONS
        CUDA_CALL(cudaMalloc(&d_numEvictions, sizeof(cuda::std::atomic<size_t>)));
#endif
 
        clear();
    }

Here is the call graph for this function:

◆ ~Filter()

template<typename Config >

cuckoogpu::Filter< Config >::~Filter ( )

inline

Destroys the Cuckoo Filter.

Frees allocated device memory.

Definition at line 423 of file CuckooFilter.cuh.

              {
        if (d_buckets) {
            CUDA_CALL(cudaFree(d_buckets));
        }
        if (d_numOccupied) {
            CUDA_CALL(cudaFree(d_numOccupied));
        }
#ifdef CUCKOO_FILTER_COUNT_EVICTIONS
        if (d_numEvictions) {
            CUDA_CALL(cudaFree(d_numEvictions));
        }
#endif
    }

Member Function Documentation

◆ calculateNumBuckets()

template<typename Config >

static size_t cuckoogpu::Filter< Config >::calculateNumBuckets ( size_t capacity )

inlinestatic

The number of buckets is enforced to be a power of two in order to allow for efficient modulo on the bucket indices.

Definition at line 394 of file CuckooFilter.cuh.

                                                       {
        return AltBucketPolicy::calculateNumBuckets(capacity);
    }

Here is the call graph for this function:

◆ capacity()

template<typename Config >

size_t cuckoogpu::Filter< Config >::capacity ( )

inline

Returns the total capacity of the filter.

Returns: size_t Total number of slots.

Definition at line 1015 of file CuckooFilter.cuh.

                      {
        return numBuckets * bucketSize;
    }

Here is the caller graph for this function:

◆ clear()

template<typename Config >

void cuckoogpu::Filter< Config >::clear ( )

inline

Clears the filter, removing all items.

Definition at line 958 of file CuckooFilter.cuh.

                 {
        CUDA_CALL(cudaMemset(d_buckets, 0, numBuckets * sizeof(Bucket)));
        CUDA_CALL(cudaMemset(d_numOccupied, 0, sizeof(cuda::std::atomic<size_t>)));
#ifdef CUCKOO_FILTER_COUNT_EVICTIONS
        CUDA_CALL(cudaMemset(d_numEvictions, 0, sizeof(cuda::std::atomic<size_t>)));
#endif
        h_numOccupied = 0;
    }

Here is the caller graph for this function:

◆ contains()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::contains ( const T & key ) const

inline

Checks if a key exists in the filter.

Parameters

key	The key to check.

Returns: true if the key is found, false otherwise.

Definition at line 1396 of file CuckooFilter.cuh.

                                                 {
        auto [i1, i2, fp1, fp2] = getCandidateBucketsAndFPs(key, numBuckets);
 
        // fp1 is for bucket i1, fp2 is for bucket i2
        // For non-choice-bit policies, fp1 == fp2
        return d_buckets[i1].contains(fp1) || d_buckets[i2].contains(fp2);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ containsMany() [1/3]

template<typename Config >

void cuckoogpu::Filter< Config >::containsMany	(	const T *	d_keys,
		const size_t	n,
		bool *	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Checks for the existence of a batch of keys.

Parameters

d_keys	Pointer to device memory containing keys to check.
n	Number of keys to check.
d_output	Pointer to device memory to store results (true/false).
stream	CUDA stream to use for the operation.

Definition at line 631 of file CuckooFilter.cuh.

                                                                                             {}) {
        size_t numBlocks = SDIV(n, blockSize);
        detail::containsKernel<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_keys, d_output, n, this);
 
        CUDA_CALL(cudaStreamSynchronize(stream));
    }

Here is the caller graph for this function:

◆ containsMany() [2/3]

template<typename Config >

void cuckoogpu::Filter< Config >::containsMany	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< bool > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Checks for existence of keys in a Thrust device vector.

Parameters

d_keys	Vector of keys to check.
d_output	Vector to store results (bool). Resized if necessary.
stream	CUDA stream.

Definition at line 861 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        containsMany(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            thrust::raw_pointer_cast(d_output.data()),
            stream
        );
    }

◆ containsMany() [3/3]

template<typename Config >

void cuckoogpu::Filter< Config >::containsMany	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< uint8_t > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Checks for existence of keys in a Thrust device vector (uint8_t output).

Parameters

d_keys	Vector of keys to check.
d_output	Vector to store results (uint8_t). Resized if necessary.
stream	CUDA stream.

Definition at line 883 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        containsMany(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            reinterpret_cast<bool*>(thrust::raw_pointer_cast(d_output.data())),
            stream
        );
    }

◆ countOccupiedSlots()

template<typename Config >

size_t cuckoogpu::Filter< Config >::countOccupiedSlots ( )

inline

Counts occupied slots by iterating over all buckets on the host.

This is a slow operation used for verification/debugging.

Returns: size_t Actual number of occupied slots.

Definition at line 1042 of file CuckooFilter.cuh.

                                {
        std::vector<Bucket> h_buckets(numBuckets);
 
        CUDA_CALL(cudaMemcpy(
            h_buckets.data(), d_buckets, numBuckets * sizeof(Bucket), cudaMemcpyDeviceToHost
        ));
 
        size_t occupiedCount = 0;
 
        for (size_t bucketIdx = 0; bucketIdx < numBuckets; ++bucketIdx) {
            const Bucket& bucket = h_buckets[bucketIdx];
 
            for (size_t atomicIdx = 0; atomicIdx < Bucket::wordCount; ++atomicIdx) {
                uint64_t packed = reinterpret_cast<const uint64_t&>(bucket.packedTags[atomicIdx]);
 
                for (size_t tagIdx = 0; tagIdx < Bucket::tagsPerWord; ++tagIdx) {
                    auto tag = bucket.extractTag(packed, tagIdx);
 
                    if (tag != EMPTY) {
                        occupiedCount++;
                    }
                }
            }
        }
 
        return occupiedCount;
    }

Here is the call graph for this function:

◆ deleteMany() [1/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::deleteMany	(	const T *	d_keys,
		const size_t	n,
		bool *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

Tries to remove a set of keys from the filter.

If the key is not present, it is ignored.

Parameters

d_keys	Pointer to the array of keys to remove
n	Number of keys to remove
d_output	Optional pointer to an output array indicating the success of each key removal
stream	CUDA stream to use for the operation

Returns: size_t Updated number of occupied slots in the filter

Definition at line 650 of file CuckooFilter.cuh.

                              {}
    ) {
        size_t numBlocks = SDIV(n, blockSize);
        detail::deleteKernel<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_keys, d_output, n, this);
 
        CUDA_CALL(cudaStreamSynchronize(stream));
 
        return occupiedSlots();
    }

Here is the caller graph for this function:

◆ deleteMany() [2/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::deleteMany	(	const thrust::device_vector< T > &	d_keys,
		cudaStream_t	stream = `{}`
	)

inline

Deletes keys in a Thrust device vector without outputting results.

Parameters

d_keys	Vector of keys to delete.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 951 of file CuckooFilter.cuh.

                                                                                  {}) {
        return deleteMany(thrust::raw_pointer_cast(d_keys.data()), d_keys.size(), nullptr, stream);
    }

◆ deleteMany() [3/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::deleteMany	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< bool > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Deletes keys in a Thrust device vector.

Parameters

d_keys	Vector of keys to delete.
d_output	Vector to store results (bool). Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 906 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        return deleteMany(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            thrust::raw_pointer_cast(d_output.data()),
            stream
        );
    }

◆ deleteMany() [4/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::deleteMany	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< uint8_t > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Deletes keys in a Thrust device vector (uint8_t output).

Parameters

d_keys	Vector of keys to delete.
d_output	Vector to store results (uint8_t). Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 929 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        return deleteMany(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            reinterpret_cast<bool*>(thrust::raw_pointer_cast(d_output.data())),
            stream
        );
    }

◆ evictionCount()

template<typename Config >

size_t cuckoogpu::Filter< Config >::evictionCount ( )

inline

Returns the total number of evictions performed.

Only available when CUCKOO_FILTER_COUNT_EVICTIONS is defined.

Returns: size_t Number of evictions.

Definition at line 997 of file CuckooFilter.cuh.

                           {
        size_t count;
        CUDA_CALL(cudaMemcpy(&count, d_numEvictions, sizeof(size_t), cudaMemcpyDeviceToHost));
        return count;
    }

◆ getAlternateBucket()

template<typename Config >

static __host__ __device__ size_t cuckoogpu::Filter< Config >::getAlternateBucket	(	size_t	bucket,
		TagType	fp,
		size_t	numBuckets
	)

inlinestatic

Computes the alternate bucket for a fingerprint.

Definition at line 373 of file CuckooFilter.cuh.

                                                                     {
        return AltBucketPolicy::getAlternateBucket(bucket, fp, numBuckets);
    }

◆ getAlternateBucketWithNewFp()

template<typename Config >

static __host__ __device__ cuda::std::tuple< size_t, TagType > cuckoogpu::Filter< Config >::getAlternateBucketWithNewFp	(	size_t	bucket,
		TagType	fp,
		size_t	numBuckets
	)

inlinestatic

Computes alternate bucket AND updated fingerprint for choice bit policies.

For non-choice-bit policies, returns the original fingerprint unchanged.

Definition at line 382 of file CuckooFilter.cuh.

                                                                              {
        if constexpr (AltBucketPolicy::usesChoiceBit) {
            return AltBucketPolicy::getAlternateBucketWithNewFp(bucket, fp, numBuckets);
        } else {
            return {AltBucketPolicy::getAlternateBucket(bucket, fp, numBuckets), fp};
        }
    }

Here is the caller graph for this function:

◆ getCandidateBucketsAndFPs()

template<typename Config >

static __host__ __device__ cuda::std::tuple< size_t, size_t, TagType, TagType > cuckoogpu::Filter< Config >::getCandidateBucketsAndFPs	(	const T &	key,
		size_t	numBuckets
	)

inlinestatic

Definition at line 365 of file CuckooFilter.cuh.

                                                               {
        return AltBucketPolicy::getCandidateBucketsAndFPs(key, numBuckets);
    }

Here is the caller graph for this function:

◆ getNumBuckets()

template<typename Config >

size_t cuckoogpu::Filter< Config >::getNumBuckets ( ) const

inline

Returns the number of buckets in the filter.

Returns: size_t Number of buckets.

Definition at line 1023 of file CuckooFilter.cuh.

                                               {
        return numBuckets;
    }

◆ hash64()

template<typename Config >

template<typename H >

static __host__ __device__ uint64_t cuckoogpu::Filter< Config >::hash64 ( const H & key )

inlinestatic

Definition at line 360 of file CuckooFilter.cuh.

                                                             {
        return AltBucketPolicy::hash64(key);
    }

Here is the caller graph for this function:

◆ insert()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::insert	(	const T &	key,
		uint32_t *	evictionAttempts = `nullptr`
	)

inline

Inserts a single key into the filter.

Computes candidate buckets and attempts insertion, performing eviction if necessary.

Parameters

key	The key to insert.
evictionAttempts	Optional pointer to a counter for eviction attempts

Returns: true if insertion succeeded, false if the filter is too full (max evictions reached).

Definition at line 1358 of file CuckooFilter.cuh.

                                                                               {
        auto [i1, i2, fp1, fp2] = getCandidateBucketsAndFPs(key, numBuckets);
 
        // For all policies: fp1 is for bucket i1, fp2 is for bucket i2
        // For non-choice-bit policies, fp1 == fp2
        if (tryInsertAtBucket(i1, fp1) || tryInsertAtBucket(i2, fp2)) {
            return true;
        }
 
        // For eviction, use correct fingerprint for the starting bucket
        auto startBucket = (fp1 & 1) == 0 ? i1 : i2;
        TagType evictFp;
 
        if constexpr (AltBucketPolicy::usesChoiceBit) {
            evictFp = (fp1 & 1) == 0 ? fp1 : fp2;
        } else {
            evictFp = fp1;
        }
 
        if constexpr (Config::evictionPolicy == EvictionPolicy::BFS) {
            return insertWithEvictionBFS(evictFp, startBucket, evictionAttempts);
        } else if constexpr (Config::evictionPolicy == EvictionPolicy::DFS) {
            return insertWithEvictionDFS(evictFp, startBucket, evictionAttempts);
        } else {
            static_assert(
                Config::evictionPolicy == EvictionPolicy::DFS ||
                    Config::evictionPolicy == EvictionPolicy::BFS,
                "Unhandled eviction policy"
            );
        }
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ insertMany() [1/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertMany	(	const T *	d_keys,
		const size_t	n,
		bool *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

Inserts a batch of keys into the filter.

Parameters

d_keys	Pointer to device memory containing keys to insert.
n	Number of keys to insert.
d_output	Optional pointer to an output array indicating the success of each key insertion.
stream	CUDA stream to use for the operation.

Returns: size_t Total number of occupied slots after insertion.

Definition at line 447 of file CuckooFilter.cuh.

                              {}
    ) {
        size_t numBlocks = SDIV(n, blockSize);
        detail::insertKernel<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_keys, d_output, n, this, nullptr);
 
        CUDA_CALL(cudaStreamSynchronize(stream));
 
        return occupiedSlots();
    }

Here is the caller graph for this function:

◆ insertMany() [2/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertMany	(	const thrust::device_vector< T > &	d_keys,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector without outputting results.

Parameters

d_keys	Vector of keys to insert.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 717 of file CuckooFilter.cuh.

                                                                                  {}) {
        return insertMany(thrust::raw_pointer_cast(d_keys.data()), d_keys.size(), nullptr, stream);
    }

◆ insertMany() [3/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertMany	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< bool > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector.

Parameters

d_keys	Vector of keys to insert.
d_output	Vector to store results (bool). Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 672 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        return insertMany(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            thrust::raw_pointer_cast(d_output.data()),
            stream
        );
    }

◆ insertMany() [4/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertMany	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< uint8_t > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector (uint8_t output).

Parameters

d_keys	Vector of keys to insert.
d_output	Vector to store results (uint8_t). Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 695 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        return insertMany(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            reinterpret_cast<bool*>(thrust::raw_pointer_cast(d_output.data())),
            stream
        );
    }

◆ insertManySorted() [1/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManySorted	(	const T *	d_keys,
		const size_t	n,
		bool *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

This pre-sorts the input keys based on the primary bucket index to allow for coalesced memory access when you later insert them into the filter.

Parameters

d_keys	Pointer to device memory array of keys to insert
n	Number of keys to insert
d_output	Optional pointer to an output array indicating the success of each key insertion.
stream	CUDA stream to use for the operation.

Returns: size_t Updated number of occupied slots in the filter

Definition at line 502 of file CuckooFilter.cuh.

                              {}
    ) {
        PackedTagType* d_packedTags;
 
        CUDA_CALL(cudaMallocAsync(&d_packedTags, n * sizeof(PackedTagType), stream));
 
        size_t numBlocks = SDIV(n, blockSize);
 
        detail::computePackedTagsKernel<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_keys, d_packedTags, n, numBuckets);
 
        void* d_tempStorage = nullptr;
        size_t tempStorageBytes = 0;
 
        cub::DeviceRadixSort::SortKeys(
            d_tempStorage,
            tempStorageBytes,
            d_packedTags,
            d_packedTags,
            n,
            0,
            sizeof(PackedTagType) * 8,
            stream
        );
 
        CUDA_CALL(cudaMallocAsync(&d_tempStorage, tempStorageBytes, stream));
 
        cub::DeviceRadixSort::SortKeys(
            d_tempStorage,
            tempStorageBytes,
            d_packedTags,
            d_packedTags,
            n,
            0,
            sizeof(PackedTagType) * 8,
            stream
        );
 
        CUDA_CALL(cudaFreeAsync(d_tempStorage, stream));
 
        detail::insertKernelSorted<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_packedTags, d_output, n, this, nullptr);
 
        CUDA_CALL(cudaFreeAsync(d_packedTags, stream));
        CUDA_CALL(cudaStreamSynchronize(stream));
 
        return occupiedSlots();
    }

◆ insertManySorted() [2/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManySorted	(	const thrust::device_vector< T > &	d_keys,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector, sorting them first, without outputting results.

Parameters

d_keys	Vector of keys to insert.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 811 of file CuckooFilter.cuh.

                                                                                        {}) {
        return insertManySorted(
            thrust::raw_pointer_cast(d_keys.data()), d_keys.size(), nullptr, stream
        );
    }

◆ insertManySorted() [3/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManySorted	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< bool > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector, sorting them first.

Parameters

d_keys	Vector of keys to insert.
d_output	Vector to store results (bool). Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 765 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        return insertManySorted(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            thrust::raw_pointer_cast(d_output.data()),
            stream
        );
    }

◆ insertManySorted() [4/4]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManySorted	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< uint8_t > &	d_output,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector, sorting them first (uint8_t output).

Parameters

d_keys	Vector of keys to insert.
d_output	Vector to store results (uint8_t). Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 788 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_output.size() != d_keys.size()) {
            d_output.resize(d_keys.size());
        }
        return insertManySorted(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            reinterpret_cast<bool*>(thrust::raw_pointer_cast(d_output.data())),
            stream
        );
    }

◆ insertManySortedWithEvictionCounts() [1/2]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManySortedWithEvictionCounts	(	const T *	d_keys,
		const size_t	n,
		uint32_t *	d_evictionAttempts,
		bool *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

Inserts a pre-sorted batch of keys and records per-attempt eviction counts.

Parameters

d_keys	Pointer to device memory containing keys to insert.
n	Number of keys to insert.
d_evictionAttempts	Device output array of size n. Each entry stores the number of evictions performed by the corresponding insertion attempt.
d_output	Optional pointer to an output array indicating per-key insertion success.
stream	CUDA stream to use for the operation.

Returns: size_t Updated number of occupied slots in the filter.

Definition at line 567 of file CuckooFilter.cuh.

                              {}
    ) {
        PackedTagType* d_packedTags;
 
        CUDA_CALL(cudaMallocAsync(&d_packedTags, n * sizeof(PackedTagType), stream));
 
        size_t numBlocks = SDIV(n, blockSize);
 
        detail::computePackedTagsKernel<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_keys, d_packedTags, n, numBuckets);
 
        void* d_tempStorage = nullptr;
        size_t tempStorageBytes = 0;
 
        cub::DeviceRadixSort::SortKeys(
            d_tempStorage,
            tempStorageBytes,
            d_packedTags,
            d_packedTags,
            n,
            0,
            sizeof(PackedTagType) * 8,
            stream
        );
 
        CUDA_CALL(cudaMallocAsync(&d_tempStorage, tempStorageBytes, stream));
 
        cub::DeviceRadixSort::SortKeys(
            d_tempStorage,
            tempStorageBytes,
            d_packedTags,
            d_packedTags,
            n,
            0,
            sizeof(PackedTagType) * 8,
            stream
        );
 
        CUDA_CALL(cudaFreeAsync(d_tempStorage, stream));
 
        detail::insertKernelSorted<Config><<<numBlocks, blockSize, 0, stream>>>(
            d_packedTags, d_output, n, this, d_evictionAttempts
        );
 
        CUDA_CALL(cudaFreeAsync(d_packedTags, stream));
        CUDA_CALL(cudaStreamSynchronize(stream));
 
        return occupiedSlots();
    }

◆ insertManySortedWithEvictionCounts() [2/2]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManySortedWithEvictionCounts	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< uint32_t > &	d_evictionAttempts,
		thrust::device_vector< uint8_t > *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector, sorting first, and records per-attempt eviction counts.

Parameters

d_keys	Vector of keys to insert.
d_evictionAttempts	Vector to store per-key eviction counts. Resized if necessary.
d_output	Optional vector to store per-key insertion success. Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 827 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_evictionAttempts.size() != d_keys.size()) {
            d_evictionAttempts.resize(d_keys.size());
        }
 
        bool* d_outputPtr = nullptr;
        if (d_output != nullptr) {
            if (d_output->size() != d_keys.size()) {
                d_output->resize(d_keys.size());
            }
            d_outputPtr = reinterpret_cast<bool*>(thrust::raw_pointer_cast(d_output->data()));
        }
 
        return insertManySortedWithEvictionCounts(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            thrust::raw_pointer_cast(d_evictionAttempts.data()),
            d_outputPtr,
            stream
        );
    }

◆ insertManyWithEvictionCounts() [1/2]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManyWithEvictionCounts	(	const T *	d_keys,
		const size_t	n,
		uint32_t *	d_evictionAttempts,
		bool *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

Inserts a batch of keys and records per-attempt eviction counts.

Parameters

d_keys	Pointer to device memory containing keys to insert.
n	Number of keys to insert.
d_evictionAttempts	Device output array of size n. Each entry stores the number of evictions performed by the corresponding insertion attempt.
d_output	Optional pointer to an output array indicating per-key insertion success.
stream	CUDA stream to use for the operation.

Returns: size_t Total number of occupied slots after insertion.

Definition at line 474 of file CuckooFilter.cuh.

                              {}
    ) {
        size_t numBlocks = SDIV(n, blockSize);
        detail::insertKernel<Config>
            <<<numBlocks, blockSize, 0, stream>>>(d_keys, d_output, n, this, d_evictionAttempts);
 
        CUDA_CALL(cudaStreamSynchronize(stream));
 
        return occupiedSlots();
    }

◆ insertManyWithEvictionCounts() [2/2]

template<typename Config >

size_t cuckoogpu::Filter< Config >::insertManyWithEvictionCounts	(	const thrust::device_vector< T > &	d_keys,
		thrust::device_vector< uint32_t > &	d_evictionAttempts,
		thrust::device_vector< uint8_t > *	d_output = `nullptr`,
		cudaStream_t	stream = `{}`
	)

inline

Inserts keys from a Thrust device vector and records per-attempt eviction counts.

Parameters

d_keys	Vector of keys to insert.
d_evictionAttempts	Vector to store per-key eviction counts. Resized if necessary.
d_output	Optional vector to store per-key insertion success. Resized if necessary.
stream	CUDA stream.

Returns: size_t Total number of occupied slots.

Definition at line 730 of file CuckooFilter.cuh.

                              {}
    ) {
        if (d_evictionAttempts.size() != d_keys.size()) {
            d_evictionAttempts.resize(d_keys.size());
        }
 
        bool* d_outputPtr = nullptr;
        if (d_output != nullptr) {
            if (d_output->size() != d_keys.size()) {
                d_output->resize(d_keys.size());
            }
            d_outputPtr = reinterpret_cast<bool*>(thrust::raw_pointer_cast(d_output->data()));
        }
 
        return insertManyWithEvictionCounts(
            thrust::raw_pointer_cast(d_keys.data()),
            d_keys.size(),
            thrust::raw_pointer_cast(d_evictionAttempts.data()),
            d_outputPtr,
            stream
        );
    }

◆ insertWithEvictionBFS()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::insertWithEvictionBFS	(	TagType	fp,
		size_t	startBucket,
		uint32_t *	evictionAttempts = `nullptr`
	)

inline

Inserts a fingerprint using repeated shallow breadth-first attempts.

Each round scans a handful of candidate eviction slots in the current bucket and tries to place one candidate in its alternate bucket. If no shallow move succeeds, it evicts the last scanned candidate slot from the current bucket and restarts the same BFS round from that evicted tag's alternate bucket. This repeats until insertion succeeds or maxEvictions is reached.

Parameters

fp	Fingerprint to insert
evictionAttempts	Optional pointer to a counter for eviction attempts
startBucket	Index of the bucket to start the search from

Returns: true if the insertion was successful, false otherwise

Definition at line 1253 of file CuckooFilter.cuh.

                                                                                                {
        constexpr size_t numCandidates = std::max(1UL, bucketSize / 2);
 
        TagType currentFp = fp;
        size_t currentBucket = startBucket;
 
        size_t evictions = 0;
        while (evictions < maxEvictions) {
            Bucket& bucket = d_buckets[currentBucket];
            size_t restartWord = 0;
            size_t restartTagIdx = 0;
 
            for (size_t i = 0; i < numCandidates; ++i) {
                size_t evictSlot = (currentFp + i * 0x9E3779B1UL + (evictions + 1) * 0x85EBCA77) &
                                   (bucketSize - 1);
                size_t evictWord = evictSlot / Bucket::tagsPerWord;
                size_t evictTagIdx = evictSlot & (Bucket::tagsPerWord - 1);
                restartWord = evictWord;
                restartTagIdx = evictTagIdx;
 
                auto packed = bucket.packedTags[evictWord].load(cuda::memory_order_relaxed);
                TagType candidateFp = bucket.extractTag(packed, evictTagIdx);
 
                if (candidateFp == EMPTY) {
                    if (tryInsertAtBucket(currentBucket, currentFp)) {
                        return true;
                    }
                    continue;
                }
 
                auto [altBucket, altFp] =
                    getAlternateBucketWithNewFp(currentBucket, candidateFp, numBuckets);
                if (tryInsertAtBucket(altBucket, altFp)) {
                    // Successfully inserted the evicted tag at its alternate location
                    // Now atomically swap in our tag at the original location
                    auto expected = bucket.packedTags[evictWord].load(cuda::memory_order_relaxed);
 
                    // Verify the tag is still there and try to replace it
                    if (bucket.extractTag(expected, evictTagIdx) == candidateFp) {
                        auto desired = bucket.replaceTag(expected, evictTagIdx, currentFp);
 
                        if (bucket.packedTags[evictWord].compare_exchange_strong(
                                expected,
                                desired,
                                cuda::memory_order_relaxed,
                                cuda::memory_order_relaxed
                            )) {
#ifdef CUCKOO_FILTER_COUNT_EVICTIONS
                            d_numEvictions->fetch_add(1, cuda::memory_order_relaxed);
                            if (evictionAttempts != nullptr) {
                                (*evictionAttempts)++;
                            }
#endif
                            return true;
                        }
                    }
 
                    // Failed to swap, clean up the tag we inserted to avoid duplicates
                    tryRemoveAtBucket(altBucket, candidateFp);
                }
            }
 
            // Evict the last scanned candidate and continue from its alternate location.
            auto expected = bucket.packedTags[restartWord].load(cuda::memory_order_relaxed);
            typename Bucket::WordType desired;
            TagType evictedFp;
 
            do {
                evictedFp = bucket.extractTag(expected, restartTagIdx);
                desired = bucket.replaceTag(expected, restartTagIdx, currentFp);
            } while (!bucket.packedTags[restartWord].compare_exchange_strong(
                expected, desired, cuda::memory_order_relaxed, cuda::memory_order_relaxed
            ));
 
            if (evictedFp == EMPTY) {
                return true;
            }
 
#ifdef CUCKOO_FILTER_COUNT_EVICTIONS
            d_numEvictions->fetch_add(1, cuda::memory_order_relaxed);
            if (evictionAttempts != nullptr) {
                (*evictionAttempts)++;
            }
#endif
 
            evictions++;
            auto [altBucket, altFp] =
                getAlternateBucketWithNewFp(currentBucket, evictedFp, numBuckets);
            currentBucket = altBucket;
            currentFp = altFp;
        }
 
        return false;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ insertWithEvictionDFS()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::insertWithEvictionDFS	(	TagType	fp,
		size_t	startBucket,
		uint32_t *	evictionAttempts = `nullptr`
	)

inline

Inserts a fingerprint into the filter by evicting existing fingerprints.

The thread first picks a pseudo-random target to replace with the new fingerprint. Then it tries to insert the evicted fingerprint into its alternate bucket. This process is repeated until either a fingerprint is inserted into an empty slot or the maximum number of evictions is reached.

Parameters

fp	Fingerprint to insert
startBucket	Index of the bucket to start the search from
evictionAttempts	Optional pointer to a counter for eviction attempts

Returns: true if the insertion was successful, false otherwise

Definition at line 1196 of file CuckooFilter.cuh.

                                                                                                {
        TagType currentFp = fp;
        size_t currentBucket = startBucket;
 
        for (size_t evictions = 0; evictions < maxEvictions; ++evictions) {
            auto evictSlot = (currentFp + (evictions + 1) * 0x9E3779B1UL) & (bucketSize - 1);
 
            size_t evictWord = evictSlot / Bucket::tagsPerWord;
            size_t evictTagIdx = evictSlot & (Bucket::tagsPerWord - 1);
 
            Bucket& bucket = d_buckets[currentBucket];
            auto expected = bucket.packedTags[evictWord].load(cuda::memory_order_relaxed);
            typename Bucket::WordType desired;
            TagType evictedFp;
 
            do {
                evictedFp = bucket.extractTag(expected, evictTagIdx);
                desired = bucket.replaceTag(expected, evictTagIdx, currentFp);
            } while (!bucket.packedTags[evictWord].compare_exchange_strong(
                expected, desired, cuda::memory_order_relaxed, cuda::memory_order_relaxed
            ));
 
#ifdef CUCKOO_FILTER_COUNT_EVICTIONS
            d_numEvictions->fetch_add(1, cuda::memory_order_relaxed);
            if (evictionAttempts != nullptr) {
                (*evictionAttempts)++;
            }
#endif
 
            currentFp = evictedFp;
            auto [altBucket, newFp] =
                getAlternateBucketWithNewFp(currentBucket, evictedFp, numBuckets);
            currentBucket = altBucket;
            currentFp = newFp;
 
            if (tryInsertAtBucket(currentBucket, currentFp)) {
                return true;
            }
        }
        return false;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ loadFactor()

template<typename Config >

float cuckoogpu::Filter< Config >::loadFactor ( )

inline

Calculates the current load factor of the filter.

Returns: float Load factor (occupied slots / total capacity).

Definition at line 971 of file CuckooFilter.cuh.

                                     {
        return static_cast<float>(occupiedSlots()) / (numBuckets * bucketSize);
    }

Here is the call graph for this function:

◆ occupiedSlots()

template<typename Config >

size_t cuckoogpu::Filter< Config >::occupiedSlots ( )

inline

Returns the total number of occupied slots.

Retrieves the value from the device counter.

Returns: size_t Number of occupied slots.

Definition at line 982 of file CuckooFilter.cuh.

                           {
        CUDA_CALL(
            cudaMemcpy(&h_numOccupied, d_numOccupied, sizeof(size_t), cudaMemcpyDeviceToHost)
        );
        return h_numOccupied;
    }

Here is the caller graph for this function:

◆ operator=()

template<typename Config >

Filter & cuckoogpu::Filter< Config >::operator= ( const Filter< Config > & )

delete

◆ remove()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::remove ( const T & key )

inline

Removes a key from the filter.

Parameters

key	The key to remove.

Returns: true if the key was found and removed, false otherwise.

Definition at line 1410 of file CuckooFilter.cuh.

                                         {
        auto [i1, i2, fp1, fp2] = getCandidateBucketsAndFPs(key, numBuckets);
 
        // fp1 is for bucket i1, fp2 is for bucket i2
        // For non-choice-bit policies, fp1 == fp2
        return tryRemoveAtBucket(i1, fp1) || tryRemoveAtBucket(i2, fp2);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ resetEvictionCount()

template<typename Config >

void cuckoogpu::Filter< Config >::resetEvictionCount ( )

inline

Resets the eviction counter to zero.

Definition at line 1006 of file CuckooFilter.cuh.

                              {
        CUDA_CALL(cudaMemset(d_numEvictions, 0, sizeof(cuda::std::atomic<size_t>)));
    }

◆ sizeInBytes()

template<typename Config >

size_t cuckoogpu::Filter< Config >::sizeInBytes ( ) const

inline

Returns the size of the filter in bytes.

Returns: size_t Size in bytes.

Definition at line 1031 of file CuckooFilter.cuh.

                                             {
        return numBuckets * sizeof(Bucket);
    }

◆ tryInsertAtBucket()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::tryInsertAtBucket	(	size_t	bucketIdx,
		TagType	tag
	)

inline

Attempts to insert a tag into a specific bucket.

Scans the bucket for an empty slot and attempts to atomically place the tag.

Parameters

bucketIdx	Index of the bucket.
tag	Tag to insert.

Returns: true if insertion succeeded, false if the bucket is full.

Definition at line 1143 of file CuckooFilter.cuh.

                                                                     {
        Bucket& bucket = d_buckets[bucketIdx];
        const uint32_t startIdx = tag & (bucketSize - 1);
        const size_t startWord = startIdx / Bucket::tagsPerWord;
 
        using WordType = typename Bucket::WordType;
 
        for (size_t i = 0; i < Bucket::wordCount; ++i) {
            const size_t currWord = (startWord + i) & (Bucket::wordCount - 1);
            auto expected = bucket.packedTags[currWord].load(cuda::memory_order_relaxed);
 
            while (true) {
                WordType zeroMask = detail::getZeroMask<TagType, WordType>(expected);
 
                if (zeroMask == 0) {
                    // No empty slots in this word, move to next
                    break;
                }
 
                // Find position of first empty slot (returns 1-indexed bit position)
                int bitPos;
                if constexpr (sizeof(WordType) == 4) {
                    bitPos = __ffs(static_cast<int>(zeroMask)) - 1;
                } else {
                    bitPos = __ffsll(static_cast<long long>(zeroMask)) - 1;
                }
                size_t j = bitPos / bitsPerTag;
 
                auto desired = bucket.replaceTag(expected, j, tag);
 
                if (bucket.packedTags[currWord].compare_exchange_strong(
                        expected, desired, cuda::memory_order_relaxed, cuda::memory_order_relaxed
                    )) {
                    return true;
                }
            }
        }
        return false;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

◆ tryRemoveAtBucket()

template<typename Config >

__device__ bool cuckoogpu::Filter< Config >::tryRemoveAtBucket	(	size_t	bucketIdx,
		TagType	tag
	)

inline

Attempt to remove a single instance of a fingerprint from a bucket.

Scans the atomic words that make up the bucket and attempts a CAS on each matching tag position until one removal succeeds. This allows multiple concurrent deleters to remove distinct copies when a bucket contains several identical fingerprints instead of all trying to clear the same slot

The function is lock-free and uses per-word compare-and-swap operations. It does NOT update any global occupancy counter, the caller is responsible for decrementing d_numOccupied if appropriate

Parameters

bucketIdx	Index of the bucket to search.
tag	Fingerprint value to remove (must not be EMPTY).

Returns: true if a single instance of tag was removed from the bucket; false if no matching tag remained (or another thread removed the last matching instance before this call could succeed).

Definition at line 1088 of file CuckooFilter.cuh.

                                                                     {
        Bucket& bucket = d_buckets[bucketIdx];
 
        const uint32_t startSlot = tag & (bucketSize - 1);
        const size_t startWord = startSlot / Bucket::tagsPerWord;
 
        using WordType = typename Bucket::WordType;
 
        for (size_t i = 0; i < Bucket::wordCount; ++i) {
            const size_t currIdx = (startWord + i) & (Bucket::wordCount - 1);
 
            while (true) {
                auto expected = bucket.packedTags[currIdx].load(cuda::memory_order_relaxed);
 
                WordType matchMask = detail::getZeroMask<TagType, WordType>(
                    expected ^ detail::replicateTag<TagType, WordType>(tag)
                );
 
                if (matchMask == 0) {
                    // No matching tags in this word
                    break;
                }
 
                // Find position of first matching tag
                int bitPos;
                if constexpr (sizeof(WordType) == 4) {
                    bitPos = __ffs(static_cast<int>(matchMask)) - 1;
                } else {
                    bitPos = __ffsll(static_cast<long long>(matchMask)) - 1;
                }
                size_t tagIdx = bitPos / bitsPerTag;
 
                auto desired = bucket.replaceTag(expected, tagIdx, EMPTY);
 
                if (bucket.packedTags[currIdx].compare_exchange_weak(
                        expected, desired, cuda::memory_order_relaxed, cuda::memory_order_relaxed
                    )) {
                    return true;
                }
                // CAS failed, retry with updated expected value
            }
        }
 
        return false;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

Member Data Documentation

◆ bitsPerTag

template<typename Config >

constexpr size_t cuckoogpu::Filter< Config >::bitsPerTag = Config::bitsPerTag

staticconstexpr

Definition at line 146 of file CuckooFilter.cuh.

◆ blockSize

template<typename Config >

constexpr size_t cuckoogpu::Filter< Config >::blockSize = Config::blockSize

staticconstexpr

Definition at line 155 of file CuckooFilter.cuh.

◆ bucketSize

template<typename Config >

constexpr size_t cuckoogpu::Filter< Config >::bucketSize = Config::bucketSize

staticconstexpr

Definition at line 152 of file CuckooFilter.cuh.

◆ d_buckets

template<typename Config >

Bucket* cuckoogpu::Filter< Config >::d_buckets

Pointer to the device memory for the buckets.

Definition at line 348 of file CuckooFilter.cuh.

◆ d_numEvictions

template<typename Config >

cuda::std::atomic<size_t>* cuckoogpu::Filter< Config >::d_numEvictions {}

Pointer to the device memory for the eviction counter.

Definition at line 354 of file CuckooFilter.cuh.

354{};

◆ d_numOccupied

template<typename Config >

cuda::std::atomic<size_t>* cuckoogpu::Filter< Config >::d_numOccupied {}

Pointer to the device memory for the occupancy counter.

Definition at line 350 of file CuckooFilter.cuh.

350{};

◆ EMPTY

template<typename Config >

constexpr TagType cuckoogpu::Filter< Config >::EMPTY = 0

staticconstexpr

Definition at line 213 of file CuckooFilter.cuh.

◆ fpMask

template<typename Config >

constexpr size_t cuckoogpu::Filter< Config >::fpMask = (1ULL << bitsPerTag) - 1

staticconstexpr

Definition at line 214 of file CuckooFilter.cuh.

◆ h_numOccupied

template<typename Config >

size_t cuckoogpu::Filter< Config >::h_numOccupied = 0

Number of occupied buckets in the filter.

Definition at line 357 of file CuckooFilter.cuh.

◆ maxEvictions

template<typename Config >

constexpr size_t cuckoogpu::Filter< Config >::maxEvictions = Config::maxEvictions

staticconstexpr

Definition at line 154 of file CuckooFilter.cuh.

◆ numBuckets

template<typename Config >

size_t cuckoogpu::Filter< Config >::numBuckets

Number of buckets in the filter.

Definition at line 347 of file CuckooFilter.cuh.

◆ tagEntryBytes

template<typename Config >

constexpr size_t cuckoogpu::Filter< Config >::tagEntryBytes = sizeof(TagType)

staticconstexpr

Definition at line 151 of file CuckooFilter.cuh.

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

CuckooFilter.cuh

Classes

Public Types

Public Member Functions

Static Public Member Functions

Public Attributes

Static Public Attributes

Detailed Description

Member Typedef Documentation

◆ AltBucketPolicy

◆ PackedTagType

◆ T

◆ TagType

Constructor & Destructor Documentation

◆ Filter() [1/2]

◆ Filter() [2/2]

◆ ~Filter()

Member Function Documentation

◆ calculateNumBuckets()

◆ capacity()

◆ clear()

◆ contains()

◆ containsMany() [1/3]

◆ containsMany() [2/3]

◆ containsMany() [3/3]

◆ countOccupiedSlots()

◆ deleteMany() [1/4]

◆ deleteMany() [2/4]

◆ deleteMany() [3/4]

◆ deleteMany() [4/4]

◆ evictionCount()

◆ getAlternateBucket()

◆ getAlternateBucketWithNewFp()

◆ getCandidateBucketsAndFPs()

◆ getNumBuckets()

◆ hash64()

◆ insert()

◆ insertMany() [1/4]

◆ insertMany() [2/4]

◆ insertMany() [3/4]

◆ insertMany() [4/4]

◆ insertManySorted() [1/4]

◆ insertManySorted() [2/4]

◆ insertManySorted() [3/4]

◆ insertManySorted() [4/4]

◆ insertManySortedWithEvictionCounts() [1/2]

◆ insertManySortedWithEvictionCounts() [2/2]

◆ insertManyWithEvictionCounts() [1/2]

◆ insertManyWithEvictionCounts() [2/2]

◆ insertWithEvictionBFS()

◆ insertWithEvictionDFS()

◆ loadFactor()

◆ occupiedSlots()

◆ operator=()

◆ remove()

◆ resetEvictionCount()

◆ sizeInBytes()

◆ tryInsertAtBucket()

◆ tryRemoveAtBucket()

Member Data Documentation

◆ bitsPerTag

◆ blockSize

◆ bucketSize

◆ d_buckets

◆ d_numEvictions

◆ d_numOccupied

◆ EMPTY

◆ fpMask

◆ h_numOccupied

◆ maxEvictions

◆ numBuckets

◆ tagEntryBytes