cusbf::filter_ref< Config > Class Template Reference

Non-owning device reference to sectorized filter storage. More...

Public Types
using	block_type = filter_block< Config >

Public Member Functions
__host__ __device__ constexpr	filter_ref () noexcept
	Default-constructed empty reference (no shard storage).
__host__ __device__ constexpr	filter_ref (block_type *blocks, uint64_t num_blocks) noexcept
	Constructs a reference over blocks spanning num_blocks shards.
__host__ __device__ block_type *	blocks () const noexcept
	Raw pointer to the first 256-bit filter block (shard array base).
__host__ __device__ uint64_t	num_blocks () const noexcept
	Number of power-of-two shards in this filter.
__device__ uint64_t	shard_index (uint64_t minimizer_hash) const noexcept
	Shard index for minimizer_hash within this filter.
__device__ void	apply_word_masks (block_type &block, uint64_t m0, uint64_t m1, uint64_t m2, uint64_t m3) const
	Atomically ORs non-zero per-word Bloom masks into block.

Static Public Member Functions
__host__ static __device__ uint64_t	shard_index (uint64_t minimizer_hash, uint64_t num_blocks) noexcept
	Maps a minimizer hash to a shard index via low bits.
static __device__ bool	sectorized_contains_packed_kmer (uint64_t packed_kmer, const uint64_t *shard_words)
	Tests sectorized membership for a packed k-mer against shard words.

Detailed Description

template<typename Config>
class cusbf::filter_ref< Config >

Non-owning device reference to sectorized filter storage.

Trivially copyable, intended to be passed by value into device code.

Definition at line 21 of file filter_ref.cuh.

Member Typedef Documentation

block_type

template<typename Config >

using cusbf::filter_ref< Config >::block_type = filter_block<Config>

Definition at line 23 of file filter_ref.cuh.

Constructor & Destructor Documentation

filter_ref() [1/2]

template<typename Config >

__host__ __device__ constexpr cusbf::filter_ref< Config >::filter_ref ( )

inlineconstexprnoexcept

Default-constructed empty reference (no shard storage).

Definition at line 26 of file filter_ref.cuh.

: blocks_{}, num_blocks_{} {}

filter_ref() [2/2]

template<typename Config >

__host__ __device__ constexpr cusbf::filter_ref< Config >::filter_ref ( block_type *  blocks, uint64_t  num_blocks  )

inlineconstexprnoexcept

Constructs a reference over blocks spanning num_blocks shards.

num_blocks must be a power of two, shard selection uses low bits of the minimizer hash.

Definition at line 33 of file filter_ref.cuh.

        : blocks_(blocks), num_blocks_(num_blocks) {}

Member Function Documentation

apply_word_masks()

template<typename Config >

__device__ void cusbf::filter_ref< Config >::apply_word_masks ( block_type &  block, uint64_t  m0, uint64_t  m1, uint64_t  m2, uint64_t  m3  ) const

inline

Atomically ORs non-zero per-word Bloom masks into block.

Used by the insert kernel after warp-local reduction of hash masks.

Definition at line 95 of file filter_ref.cuh.

                                                                                                  {
        if (m0 != 0) {
            detail::atomicOrWord(&block.words[0], m0);
        }
        if (m1 != 0) {
            detail::atomicOrWord(&block.words[1], m1);
        }
        if (m2 != 0) {
            detail::atomicOrWord(&block.words[2], m2);
        }
        if (m3 != 0) {
            detail::atomicOrWord(&block.words[3], m3);
        }
    }

Here is the call graph for this function:

Here is the caller graph for this function:

blocks()

template<typename Config >

__host__ __device__ block_type * cusbf::filter_ref< Config >::blocks ( ) const

inlinenoexcept

Raw pointer to the first 256-bit filter block (shard array base).

Definition at line 37 of file filter_ref.cuh.

                                                                          {
        return blocks_;
    }

num_blocks()

template<typename Config >

__host__ __device__ uint64_t cusbf::filter_ref< Config >::num_blocks ( ) const

inlinenoexcept

Number of power-of-two shards in this filter.

Definition at line 42 of file filter_ref.cuh.

                                                                           {
        return num_blocks_;
    }

Here is the caller graph for this function:

sectorized_contains_packed_kmer()

template<typename Config >

static __device__ bool cusbf::filter_ref< Config >::sectorized_contains_packed_kmer ( uint64_t  packed_kmer, const uint64_t *  shard_words  )

inlinestatic

Tests sectorized membership for a packed k-mer against shard words.

All findere s-mers must hit set bits in the shard, returns false if any s-mer misses.

Parameters

packed_kmer	K-mer packed with Config::symbolBits per symbol.
shard_words	Four 64-bit words of the target shard (unaligned load path).

Definition at line 71 of file filter_ref.cuh.

                                                                                       {
        bool present = true;
        _Pragma("unroll")
        for (uint64_t smer_offset = 0; smer_offset < Config::findereSpan; ++smer_offset) {
            const uint64_t smer_hash =
                detail::packed_kmer_smer_hash<Config>(packed_kmer, smer_offset);
            detail::forEachHashIndex<Config>(
                [&]<uint64_t HashIndex>(std::integral_constant<uint64_t, HashIndex>) {
                    constexpr uint64_t s = HashIndex % Config::blockWordCount;
                    const uint64_t bit_pos =
                        block_type::template sectorizedBitAddress<HashIndex>(smer_hash);
                    present &= ((shard_words[s] >> bit_pos) & 1) != 0;
                }
            );
        }
        return present;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

shard_index() [1/2]

template<typename Config >

__device__ uint64_t cusbf::filter_ref< Config >::shard_index ( uint64_t  minimizer_hash ) const

inlinenoexcept

Shard index for minimizer_hash within this filter.

Definition at line 58 of file filter_ref.cuh.

                                                                                          {
        return shard_index(minimizer_hash, num_blocks_);
    }

Here is the call graph for this function:

shard_index() [2/2]

template<typename Config >

__host__ static __device__ uint64_t cusbf::filter_ref< Config >::shard_index ( uint64_t  minimizer_hash, uint64_t  num_blocks  )

inlinestaticnoexcept

Maps a minimizer hash to a shard index via low bits.

Parameters

minimizer_hash	Minimizer hash for the k-mer window.
num_blocks	Power-of-two shard count.

Definition at line 53 of file filter_ref.cuh.

                                                                       {
        return minimizer_hash & (num_blocks - 1);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

---

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

• filter_ref.cuh