GPU-Accelerated Cuckoo Filter
Loading...
Searching...
No Matches
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | Static Public Attributes | List of all members
cuckoogpu::Filter< Config >::Bucket Struct Reference

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

Public Types

using WordType = typename Config::WordType
 

Public Member Functions

__host__ __device__ __forceinline__ TagType extractTag (WordType packed, size_t tagIdx) const
 
__host__ __device__ __forceinline__ WordType replaceTag (WordType packed, size_t tagIdx, TagType newTag) const
 
template<size_t N>
__device__ __forceinline__ void load128Bit (size_t startIdx, WordType(&out)[N]) const
 Loads words using 128-bit vectorized loads into a fixed-size array.
 
__device__ bool contains (TagType tag) const
 Checks if a tag is present in the bucket using vectorized loads.
 

Static Public Member Functions

template<size_t N>
__device__ static __forceinline__ bool checkWords (const WordType(&loaded)[N], WordType replicatedTag)
 Checks an array of loaded words for a matching tag using SWAR.
 

Public Attributes

cuda::std::atomic< WordType > packedTags [wordCount]
 

Static Public Attributes

static constexpr size_t tagsPerWord = sizeof(WordType) / sizeof(TagType)
 
static constexpr size_t wordCount = bucketSize / tagsPerWord
 

Detailed Description

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

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

The bucket is divided into words, where each word contains one or more fingerprints depending on tag size.

This optimisation allows us to avoid having to perform atomic operations on every fingerprint in the bucket, the extra computational overhead is negligible.

For efficiency reasons, the number of fingerprints per word is enforced to be a power of 2, same goes for the total number of fingerprints in the bucket.

Definition at line 225 of file CuckooFilter.cuh.

Member Typedef Documentation

◆ WordType

template<typename Config >
using cuckoogpu::Filter< Config >::Bucket::WordType = typename Config::WordType

Definition at line 228 of file CuckooFilter.cuh.

Member Function Documentation

◆ checkWords()

template<typename Config >
template<size_t N>
__device__ static __forceinline__ bool cuckoogpu::Filter< Config >::Bucket::checkWords ( const WordType(&)  loaded[N],
WordType  replicatedTag 
)
inlinestatic

Checks an array of loaded words for a matching tag using SWAR.

Definition at line 257 of file CuckooFilter.cuh.

257 {
258 _Pragma("unroll")
259 for (size_t j = 0; j < N; ++j) {
260 if (detail::hasZero<TagType, WordType>(loaded[j] ^ replicatedTag)) {
261 return true;
262 }
263 }
264 return false;
265 }
__host__ __device__ __forceinline__ constexpr bool hasZero(WordType v)
Checks if a packed word contains a zero slot.
Definition helpers.cuh:96
Here is the call graph for this function:
Here is the caller graph for this function:

◆ contains()

template<typename Config >
__device__ bool cuckoogpu::Filter< Config >::Bucket::contains ( TagType  tag) const
inline

Checks if a tag is present in the bucket using vectorized loads.

Automatically selects the best load width based on bucket size and architecture.

Definition at line 291 of file CuckooFilter.cuh.

291 {
292 const WordType replicatedTag = detail::replicateTag<TagType, WordType>(tag);
293
294 // Scalar path
295 if constexpr (wordCount == 1) {
296 const auto packed = reinterpret_cast<const WordType&>(packedTags[0]);
297 return detail::hasZero<TagType, WordType>(packed ^ replicatedTag);
298 }
299
300 const uint32_t startSlot = tag & (bucketSize - 1);
301 const size_t startWordIdx = startSlot / tagsPerWord;
302
303#if __CUDA_ARCH__ >= 1000 && !defined(CUCKOO_FILTER_DISABLE_256BIT_LOADS)
304 // 256-bit load path
305 constexpr size_t wordsPerLoad256 = (sizeof(WordType) == 4) ? 8 : 4;
306 if constexpr (wordCount >= wordsPerLoad256) {
307 constexpr size_t alignMask = wordsPerLoad256 - 1;
308 const size_t startAlignedIdx = startWordIdx & ~alignMask;
309
310 WordType loaded[wordsPerLoad256];
311 for (size_t i = 0; i < wordCount / wordsPerLoad256; i++) {
312 const size_t idx = (startAlignedIdx + i * wordsPerLoad256) & (wordCount - 1);
314 reinterpret_cast<const WordType*>(&packedTags[idx]), loaded
315 );
316 if (checkWords(loaded, replicatedTag)) {
317 return true;
318 }
319 }
320 return false;
321 }
322#endif
323 // 128-bit load path
324 constexpr size_t wordsPerLoad128 = (sizeof(WordType) == 4) ? 4 : 2;
325 if constexpr (wordCount >= wordsPerLoad128) {
326 constexpr size_t alignMask = wordsPerLoad128 - 1;
327 const size_t startAlignedIdx = startWordIdx & ~alignMask;
328
329 WordType loaded[wordsPerLoad128];
330 for (size_t i = 0; i < wordCount / wordsPerLoad128; i++) {
331 const size_t idx = (startAlignedIdx + i * wordsPerLoad128) & (wordCount - 1);
332 load128Bit(idx, loaded);
333 if (checkWords(loaded, replicatedTag)) {
334 return true;
335 }
336 }
337 return false;
338 }
339
340 return false;
341 }
static constexpr size_t bucketSize
typename Config::WordType WordType
static constexpr size_t tagsPerWord
__device__ __forceinline__ void load128Bit(size_t startIdx, WordType(&out)[N]) const
Loads words using 128-bit vectorized loads into a fixed-size array.
static constexpr size_t wordCount
__device__ static __forceinline__ bool checkWords(const WordType(&loaded)[N], WordType replicatedTag)
Checks an array of loaded words for a matching tag using SWAR.
cuda::std::atomic< WordType > packedTags[wordCount]
Here is the call graph for this function:
Here is the caller graph for this function:

◆ extractTag()

template<typename Config >
__host__ __device__ __forceinline__ TagType cuckoogpu::Filter< Config >::Bucket::extractTag ( WordType  packed,
size_t  tagIdx 
) const
inline

Definition at line 241 of file CuckooFilter.cuh.

241 {
242 return static_cast<TagType>((packed >> (tagIdx * bitsPerTag)) & fpMask);
243 }
static constexpr size_t fpMask
typename Config::TagType TagType
static constexpr size_t bitsPerTag
Here is the caller graph for this function:

◆ load128Bit()

template<typename Config >
template<size_t N>
__device__ __forceinline__ void cuckoogpu::Filter< Config >::Bucket::load128Bit ( size_t  startIdx,
WordType(&)  out[N] 
) const
inline

Loads words using 128-bit vectorized loads into a fixed-size array.

Definition at line 271 of file CuckooFilter.cuh.

271 {
272 static_assert(N == 2 || N == 4, "128-bit loads support 2 or 4 words");
273 if constexpr (sizeof(WordType) == 4) {
274 auto vec = __ldg(reinterpret_cast<const uint4*>(&packedTags[startIdx]));
275 out[0] = vec.x;
276 out[1] = vec.y;
277 out[2] = vec.z;
278 out[3] = vec.w;
279 } else {
280 auto vec = __ldg(reinterpret_cast<const ulonglong2*>(&packedTags[startIdx]));
281 out[0] = vec.x;
282 out[1] = vec.y;
283 }
284 }
Here is the caller graph for this function:

◆ replaceTag()

template<typename Config >
__host__ __device__ __forceinline__ WordType cuckoogpu::Filter< Config >::Bucket::replaceTag ( WordType  packed,
size_t  tagIdx,
TagType  newTag 
) const
inline

Definition at line 246 of file CuckooFilter.cuh.

246 {
247 size_t shift = tagIdx * bitsPerTag;
248 WordType cleared = packed & ~(static_cast<WordType>(fpMask) << shift);
249 return cleared | (static_cast<WordType>(newTag) << shift);
250 }
Here is the caller graph for this function:

Member Data Documentation

◆ packedTags

template<typename Config >
cuda::std::atomic<WordType> cuckoogpu::Filter< Config >::Bucket::packedTags[wordCount]

Definition at line 238 of file CuckooFilter.cuh.

◆ tagsPerWord

template<typename Config >
constexpr size_t cuckoogpu::Filter< Config >::Bucket::tagsPerWord = sizeof(WordType) / sizeof(TagType)
staticconstexpr

Definition at line 230 of file CuckooFilter.cuh.

◆ wordCount

template<typename Config >
constexpr size_t cuckoogpu::Filter< Config >::Bucket::wordCount = bucketSize / tagsPerWord
staticconstexpr

Definition at line 235 of file CuckooFilter.cuh.


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