cusbf::Filter< Config >::Shard Struct Reference

One 256-bit filter block stored as an array of Config::blockWordCount words. More...

Static Public Member Functions
template<uint64_t HashIndex>
constexpr __host__ static __device__ uint64_t	sectorizedBitAddress (uint64_t baseHash)
	Maps a base hash to a bit position within word sector HashIndex.
__device__ static __forceinline__ void	sectorizedHashToMasks (uint64_t baseHash, uint64_t &mask0, uint64_t &mask1, uint64_t &mask2, uint64_t &mask3)
	Computes four word bitmasks from a single base hash.

Public Attributes
uint64_t	words [wordCount]

Static Public Attributes
static constexpr uint64_t	wordCount = Config::blockWordCount
static constexpr uint64_t	wordBits = Config::wordBits
static constexpr int	wordBitsLog2 = cuda::std::bit_width(wordBits) - 1
static constexpr uint64_t	wordMask = (1ULL << wordBitsLog2) - 1
static constexpr int	hashShift = 64 - wordBitsLog2
static constexpr uint64_t	sliceWidth = 64 / Config::hashCount
static constexpr bool	useBitSlicing = sliceWidth >= wordBitsLog2

Detailed Description

template<typename Config>
struct cusbf::Filter< Config >::Shard

One 256-bit filter block stored as an array of Config::blockWordCount words.

Each shard is addressed as a unit: a minimizer hash selects the shard, and the s-mer-derived hashes set/test bits within it.

The struct is 32-byte aligned to enable vectorised loads via cusbf::detail::load256BitGlobalNC.

Definition at line 397 of file BloomFilter.cuh.

Member Function Documentation

sectorizedBitAddress()

template<typename Config >

template<uint64_t HashIndex>

constexpr __host__ static __device__ uint64_t cusbf::Filter< Config >::Shard::sectorizedBitAddress ( uint64_t  baseHash )

inlinestaticconstexpr

Maps a base hash to a bit position within word sector HashIndex.

Uses bit-slicing when the hash has enough entropy per slice; otherwise applies a multiplicative salt to redistribute bits.

Template Parameters

HashIndex

Hash function index in [0, Config::hashCount).

Parameters

baseHash

The s-mer-derived hash value.

Returns

Bit position (0-based) within the word at sector HashIndex % blockWordCount.

Definition at line 419 of file BloomFilter.cuh.

          {
            static_assert(HashIndex < Config::hashCount, "Hash index out of range");
            // When there are enough bits in a 64-bit hash to give each hash
            // index its own slice, avoid the extra multiply and use
            // bit-slicing instead.
            if constexpr (useBitSlicing) {
                return (baseHash >> (sliceWidth * HashIndex)) & wordMask;
            } else {
                const uint64_t mixed = baseHash * detail::multiplicativeSaltLiteral<HashIndex>();
                return mixed >> hashShift;
            }
        }

Here is the call graph for this function:

sectorizedHashToMasks()

template<typename Config >

__device__ static __forceinline__ void cusbf::Filter< Config >::Shard::sectorizedHashToMasks ( uint64_t  baseHash, uint64_t &  mask0, uint64_t &  mask1, uint64_t &  mask2, uint64_t &  mask3  )

inlinestatic

Computes four word bitmasks from a single base hash.

Iterates over all Config::hashCount hash functions and ORs the corresponding bit into one of the four output word masks (sectors 0-3).

Parameters

baseHash	s-mer-derived hash value.
mask0	Accumulated bits for word 0 (in/out).
mask1	Accumulated bits for word 1 (in/out).
mask2	Accumulated bits for word 2 (in/out).
mask3	Accumulated bits for word 3 (in/out).

Definition at line 446 of file BloomFilter.cuh.

          {
            detail::forEachHashIndex<Config>(
                [&]<uint64_t HashIndex>(std::integral_constant<uint64_t, HashIndex>) {
                    constexpr uint64_t s = HashIndex % Config::blockWordCount;
                    const uint64_t bitPos = sectorizedBitAddress<HashIndex>(baseHash);
                    const uint64_t bit = uint64_t{1} << bitPos;
                    // clang-format off
                    if      constexpr (s == 0) mask0 |= bit;
                    else if constexpr (s == 1) mask1 |= bit;
                    else if constexpr (s == 2) mask2 |= bit;
                    else                       mask3 |= bit;
                    // clang-format on
                }
            );
        }

Here is the call graph for this function:

Member Data Documentation

hashShift

template<typename Config >

constexpr int cusbf::Filter< Config >::Shard::hashShift = 64 - wordBitsLog2

staticconstexpr

Definition at line 402 of file BloomFilter.cuh.

sliceWidth

template<typename Config >

constexpr uint64_t cusbf::Filter< Config >::Shard::sliceWidth = 64 / Config::hashCount

staticconstexpr

Definition at line 403 of file BloomFilter.cuh.

useBitSlicing

template<typename Config >

constexpr bool cusbf::Filter< Config >::Shard::useBitSlicing = sliceWidth >= wordBitsLog2

staticconstexpr

Definition at line 404 of file BloomFilter.cuh.

wordBits

template<typename Config >

constexpr uint64_t cusbf::Filter< Config >::Shard::wordBits = Config::wordBits

staticconstexpr

Definition at line 399 of file BloomFilter.cuh.

wordBitsLog2

template<typename Config >

constexpr int cusbf::Filter< Config >::Shard::wordBitsLog2 = cuda::std::bit_width(wordBits) - 1

staticconstexpr

Definition at line 400 of file BloomFilter.cuh.

wordCount

template<typename Config >

constexpr uint64_t cusbf::Filter< Config >::Shard::wordCount = Config::blockWordCount

staticconstexpr

Definition at line 398 of file BloomFilter.cuh.

wordMask

template<typename Config >

constexpr uint64_t cusbf::Filter< Config >::Shard::wordMask = (1ULL << wordBitsLog2) - 1

staticconstexpr

Definition at line 401 of file BloomFilter.cuh.

words

template<typename Config >

uint64_t cusbf::Filter< Config >::Shard::words[wordCount]

Definition at line 406 of file BloomFilter.cuh.

---

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

• BloomFilter.cuh