cusbf::filter_block< Config > Struct Template 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)
	Bit position within shard word HashIndex % blockWordCount for baseHash.
__device__ static __forceinline__ void	sectorizedHashToMasks (uint64_t baseHash, uint64_t &mask0, uint64_t &mask1, uint64_t &mask2, uint64_t &mask3)
	Accumulates Bloom bit masks for all hash indices into four shard words.

Public Attributes
uint64_t	words [wordCount]
	Sector words, hash i maps to word i % blockWordCount.

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_block< Config >

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.

Definition at line 19 of file filter_impl.cuh.

Member Function Documentation

sectorizedBitAddress()

template<typename Config >

template<uint64_t HashIndex>

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

inlinestaticconstexpr

Bit position within shard word HashIndex % blockWordCount for baseHash.

Uses bit-slicing when sliceWidth >= wordBitsLog2, otherwise a salted multiply-shift.

Definition at line 37 of file filter_impl.cuh.

      {
        static_assert(HashIndex < Config::hashCount, "Hash index out of range");
        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_block< Config >::sectorizedHashToMasks ( uint64_t  baseHash, uint64_t &  mask0, uint64_t &  mask1, uint64_t &  mask2, uint64_t &  mask3  )

inlinestatic

Accumulates Bloom bit masks for all hash indices into four shard words.

ORs into the output masks in place, used during insert to build per-k-mer updates.

Definition at line 54 of file filter_impl.cuh.

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

Here is the call graph for this function:

Here is the caller graph for this function:

Member Data Documentation

hashShift

template<typename Config >

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

staticconstexpr

Definition at line 24 of file filter_impl.cuh.

sliceWidth

template<typename Config >

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

staticconstexpr

Definition at line 25 of file filter_impl.cuh.

useBitSlicing

template<typename Config >

constexpr bool cusbf::filter_block< Config >::useBitSlicing = sliceWidth >= wordBitsLog2

staticconstexpr

Definition at line 26 of file filter_impl.cuh.

wordBits

template<typename Config >

constexpr uint64_t cusbf::filter_block< Config >::wordBits = Config::wordBits

staticconstexpr

Definition at line 21 of file filter_impl.cuh.

wordBitsLog2

template<typename Config >

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

staticconstexpr

Definition at line 22 of file filter_impl.cuh.

wordCount

template<typename Config >

constexpr uint64_t cusbf::filter_block< Config >::wordCount = Config::blockWordCount

staticconstexpr

Definition at line 20 of file filter_impl.cuh.

wordMask

template<typename Config >

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

staticconstexpr

Definition at line 23 of file filter_impl.cuh.

words

template<typename Config >

uint64_t cusbf::filter_block< Config >::words[wordCount]

Sector words, hash i maps to word i % blockWordCount.

Definition at line 29 of file filter_impl.cuh.

---

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

• filter_impl.cuh