Namespaces
namespace	xxhash
	XXHash_64 implementation from.

Classes
struct	BitwiseOr

struct	SaltLiteral
	Compile-time golden-ratio-derived multiplicative salt constants. More...

struct	SaltLiteral< 0 >

struct	SaltLiteral< 1 >

struct	SaltLiteral< 10 >

struct	SaltLiteral< 11 >

struct	SaltLiteral< 12 >

struct	SaltLiteral< 13 >

struct	SaltLiteral< 14 >

struct	SaltLiteral< 15 >

struct	SaltLiteral< 2 >

struct	SaltLiteral< 3 >

struct	SaltLiteral< 4 >

struct	SaltLiteral< 5 >

struct	SaltLiteral< 6 >

struct	SaltLiteral< 7 >

struct	SaltLiteral< 8 >

struct	SaltLiteral< 9 >

struct	SequenceKmerInput
	Kernel input descriptor for a sequence k-mer sweep. More...

Functions
template<typename T >
consteval uint64_t	validByteCount ()

template<typename T >
consteval bool	separatorPositionAlwaysEncodesInvalid (char *input, uint64_t separatorPosition, uint64_t index)
	Recursively tests whether placing the separator byte at any position in an input of valid bytes always results in an invalid encoding.

template<typename T >
consteval bool	separatorByteAlwaysEncodesInvalid ()
	Tests that for every position in the input, placing the separator byte at that position always results in an invalid encoding.

template<typename Config >
__global__ void	containsSequenceKmersKernel (SequenceKmerInput< Config > input, device_span< const typename Filter< Config >::Shard > shards, device_span< uint8_t > output)
	CUDA kernel: queries k-mers from a sequence against the filter.

template<typename Config >
__device__ __forceinline__ bool	prepareSequenceHashTiles (const char sequence, uint64_t blockStartKmer, uint64_t blockKmers, uint8_t sequenceTile)
	Cooperatively loads and encodes a tile of symbols into shared memory.

template<typename Config >
__global__ void	insertSequenceKmersKernel (SequenceKmerInput< Config > input, device_span< typename Filter< Config >::Shard > shards)
	CUDA kernel: inserts k-mers from a sequence into the filter.

template<uint64_t Index>
__host__ __device__ __forceinline__ constexpr uint64_t	multiplicativeSaltLiteral ()
	Returns the multiplicative salt constant for hash function `Index`.

template<typename Config , typename Fn , uint64_t... HashIndices>
__host__ __device__ __forceinline__ void	forEachHashIndexImpl (Fn &&fn, std::index_sequence< HashIndices... >)
	Implementation helper for forEachHashIndex (fold-expression over an index sequence).

template<typename Config , typename Fn >
__host__ __device__ __forceinline__ void	forEachHashIndex (Fn &&fn)
	Invokes `fn` for each hash index in [0, Config::hashCount) at compile time.

template<typename Config , uint64_t Length>
__host__ __device__ __forceinline__ constexpr uint64_t	packedWindowMask ()
	Returns a bitmask covering `Length` packed alphabet symbols.

template<typename Config , uint64_t WindowLength, uint64_t K>
__host__ __device__ __forceinline__ constexpr uint64_t	extractPackedSubwindow (uint64_t packedKmer, uint64_t start)
	Extracts a packed sub-window from a packed k-mer.

__device__ __forceinline__ void	atomicOrWord (uint64_t *ptr, uint64_t value)
	Atomically ORs `value` into the device word at `ptr`.

template<typename Config >
__device__ __forceinline__ uint64_t	packedKmerMinimizerHash (uint64_t packedKmer)
	Computes the minimizer hash for a packed k-mer.

template<typename Config >
__device__ __forceinline__ uint64_t	packedKmerSmerHash (uint64_t packedKmer, uint64_t start)
	Computes the hash for the s-mer at position `start` within a packed k-mer.

template<typename Config >
__device__ __forceinline__ void	loadShardWords4 (const typename Filter< Config >::Shard shards, uint64_t shardIndex, uint64_t w)
	Loads all four 64-bit words of a shard into a local array.

template<typename Config , uint64_t K>
__device__ __forceinline__ uint64_t	packKmerFromTile (const uint8_t *tile, uint64_t start)
	Packs `K` symbols from a shared-memory tile into an integer.

template<typename Config , uint64_t K>
__device__ __forceinline__ uint64_t	advancePackedKmer (uint64_t packed, uint8_t newBase)
	Slides the packed k-mer window forward by one symbol.

template<typename Config >
__device__ __forceinline__ bool	sectorizedContainsPackedKmer (uint64_t packedKmer, const uint64_t *w)
	Tests whether a packed k-mer is present in a pre-loaded shard.

template<typename Config >
__device__ __forceinline__ bool	kmerIsValid (const uint8_t *tile, uint64_t start)

constexpr __host__ __device__ __forceinline__ uint64_t	hash64 (uint64_t key)
	Fast 64-bit integer hash (non-cryptographic).

constexpr __host__ __device__ __forceinline__ uint64_t	minimizerHash64 (uint64_t key)
	Fast 64-bit hash sufficient for uniform minimizer selection.

Variables
constexpr uint32_t	kContainsSequenceStride = 4

constexpr uint64_t	kInvalidHash = std::numeric_limits<uint64_t>::max()
	Sentinel hash value indicating "no valid minimizer found".

Function Documentation

◆ advancePackedKmer()

template<typename Config , uint64_t K>

__device__ __forceinline__ uint64_t cusbf::detail::advancePackedKmer	(	uint64_t	packed,
		uint8_t	newBase
	)

Slides the packed k-mer window forward by one symbol.

Shifts the existing packed representation left by one symbol, inserts the new symbol in the least-significant position, and masks to K symbols.

Template Parameters

K	k-mer length.

Parameters

packed	Current packed k-mer.
newBase	Pre-encoded new symbol.

Returns: Updated packed k-mer.

Definition at line 1411 of file BloomFilter.cuh.

                                                                                        {
    return ((packed << Config::symbolBits) | (newBase & Config::symbolMask)) &
           packedWindowMask<Config, K>();
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ atomicOrWord()

__device__ __forceinline__ void cusbf::detail::atomicOrWord	(	uint64_t *	ptr,
		uint64_t	value
	)

Atomically ORs value into the device word at ptr.

Parameters

ptr	Target device word.
value	Value to OR in.

Definition at line 319 of file BloomFilter.cuh.

                                                                            {
    atomicOr(reinterpret_cast<unsigned long long*>(ptr), static_cast<unsigned long long>(value));
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ containsSequenceKmersKernel()

template<typename Config >

__global__ void cusbf::detail::containsSequenceKmersKernel	(	SequenceKmerInput< Config >	input,
		device_span< const typename Filter< Config >::Shard >	shards,
		device_span< uint8_t >	output
	)

CUDA kernel: queries k-mers from a sequence against the filter.

Each thread processes kStride consecutive k-mers to amortise packing and shard loads. Threads sharing the same shard collaborate via __match_any_sync to load the shard words once and broadcast them.

Template Parameters

Config Filter configuration.

Parameters

input	Sequence descriptor (device span + k-mer count).
shards	Device-resident shard array (read-only).
output	Per-k-mer result buffer (1 = present, 0 = absent).

Definition at line 1503 of file BloomFilter.cuh.

  {
    // Each thread handles this many consecutive k-mers to amortise packing
    constexpr uint32_t kStride = kContainsSequenceStride;
    constexpr uint64_t sequenceTileBases = Config::cudaBlockSize * kStride + Config::k - 1;
 
    __shared__ uint8_t sequenceTile[sequenceTileBases];
 
    const uint64_t numKmers = input.kmerCount();
    const uint64_t blockStartKmer =
        static_cast<uint64_t>(blockIdx.x) * Config::cudaBlockSize * kStride;
    if (blockStartKmer >= numKmers) {
        return;
    }
 
    const uint64_t blockKmers = min(Config::cudaBlockSize * kStride, numKmers - blockStartKmer);
 
    const bool blockAllValid = prepareSequenceHashTiles<Config>(
        input.sequence.data(), blockStartKmer, blockKmers, sequenceTile
    );
 
    const uint64_t threadOffset = static_cast<uint64_t>(threadIdx.x) * kStride;
    if (threadOffset >= blockKmers) {
        return;
    }
 
    // Bitmask: bit s set = k-mer at offset s is valid.
    uint32_t kmerValidMask = 0;
    _Pragma("unroll")
    for (uint32_t s = 0; s < kStride; ++s) {
        if ((threadOffset + s) < blockKmers) {
            kmerValidMask |= (1u << s);
        }
    }
 
    if (!blockAllValid) {
        _Pragma("unroll")
        for (uint32_t s = 0; s < kStride; ++s) {
            if (!(kmerValidMask & (1u << s))) {
                continue;
            }
            const uint64_t localIdx = threadOffset + s;
            if (!kmerIsValid<Config>(sequenceTile, localIdx)) {
                kmerValidMask &= ~(1u << s);
            }
        }
    }
 
    // Always pack from position 0.  Sliding propagates the packed value forward
    // invalid bases from earlier k-mers are simply shifted out.
    uint64_t packedKmer = packKmerFromTile<Config, Config::k>(sequenceTile, threadOffset);
 
    for (uint32_t s = 0; s < kStride; ++s) {
        const uint64_t localIdx = threadOffset + s;
        if (localIdx >= blockKmers) {
            break;
        }
 
        const uint64_t kmerIndex = blockStartKmer + localIdx;
 
        if (s > 0) {
            packedKmer = advancePackedKmer<Config, Config::k>(
                packedKmer, sequenceTile[localIdx + Config::k - 1]
            );
        }
 
        if (!(kmerValidMask & (1u << s))) {
            output[kmerIndex] = 0;
            continue;
        }
 
        const uint64_t minimizerHash = packedKmerMinimizerHash<Config>(packedKmer);
 
        // Warp-level shard sharing.
        const auto shardIdx = static_cast<uint32_t>(minimizerHash & (shards.size() - 1));
        const uint32_t peers = __match_any_sync(0xFFFFFFFFu, shardIdx);
        const int leader = __ffs(static_cast<int>(peers)) - 1;
 
        uint64_t w[4];
        if (static_cast<int>(threadIdx.x & 31u) == leader) {
            loadShardWords4<Config>(shards.data(), shardIdx, w);
        }
        w[0] = __shfl_sync(peers, w[0], leader);
        w[1] = __shfl_sync(peers, w[1], leader);
        w[2] = __shfl_sync(peers, w[2], leader);
        w[3] = __shfl_sync(peers, w[3], leader);
 
        const bool present = sectorizedContainsPackedKmer<Config>(packedKmer, w);
        output[kmerIndex] = present;
    }
}

◆ extractPackedSubwindow()

template<typename Config , uint64_t WindowLength, uint64_t K>

__host__ __device__ __forceinline__ constexpr uint64_t cusbf::detail::extractPackedSubwindow	(	uint64_t	packedKmer,
		uint64_t	start
	)

constexpr

Extracts a packed sub-window from a packed k-mer.

Extracts WindowLength consecutive bases starting at start from a packed k-mer of length K (MSB = first base).

Template Parameters

WindowLength	Length of the sub-window to extract.
K	Length of the full k-mer.

Parameters

packedKmer	Packed k-mer (MSB = first base).
start	Zero-based start position.

Returns: Packed sub-window.

Definition at line 307 of file BloomFilter.cuh.

                                                            {
    static_assert(WindowLength <= K, "WindowLength must not exceed K");
    return (packedKmer >> (Config::symbolBits * (K - (start + WindowLength)))) &
           packedWindowMask<Config, WindowLength>();
}

Here is the call graph for this function:

◆ forEachHashIndex()

template<typename Config , typename Fn >

__host__ __device__ __forceinline__ void cusbf::detail::forEachHashIndex ( Fn && fn )

Invokes fn for each hash index in [0, Config::hashCount) at compile time.

Template Parameters

Config	Filter configuration.
Fn	Callable with signature `void(std::integral_constant<uint64_t, I>)`.

Parameters

fn	Callable to invoke for each index.

Definition at line 271 of file BloomFilter.cuh.

                                                                   {
    forEachHashIndexImpl<Config>(
        static_cast<Fn&&>(fn), std::make_index_sequence<Config::hashCount>{}
    );
}

Here is the call graph for this function:

◆ forEachHashIndexImpl()

template<typename Config , typename Fn , uint64_t... HashIndices>

__host__ __device__ __forceinline__ void cusbf::detail::forEachHashIndexImpl	(	Fn &&	fn,
		std::index_sequence< HashIndices... >
	)

Implementation helper for forEachHashIndex (fold-expression over an index sequence).

Definition at line 259 of file BloomFilter.cuh.

                                                                 {
    (fn(std::integral_constant<uint64_t, HashIndices>{}), ...);
}

Here is the call graph for this function:

◆ hash64()

constexpr __host__ __device__ __forceinline__ uint64_t cusbf::detail::hash64 ( uint64_t key )

constexpr

Fast 64-bit integer hash (non-cryptographic).

One multiplicative step followed by an xorshift. Used to hash s-mer packed representations for Bloom bit-position selection.

Parameters

key	Input value.

Returns: Hashed value.

Definition at line 192 of file hashutil.cuh.

                                                                            {
    key *= 0x9e3779b97f4a7c15ULL;
    key ^= key >> 33;
    return key;
}

Here is the call graph for this function:

◆ insertSequenceKmersKernel()

template<typename Config >

__global__ void cusbf::detail::insertSequenceKmersKernel	(	SequenceKmerInput< Config >	input,
		device_span< typename Filter< Config >::Shard >	shards
	)

CUDA kernel: inserts k-mers from a sequence into the filter.

Each thread processes one k-mer. Consecutive threads targeting the same shard use cub::WarpReduce::HeadSegmentedReduce to merge bitmasks before the run head issues a minimal number of atomicOr operations.

Template Parameters

Config Filter configuration.

Parameters

input	Sequence descriptor.
shards	Device-resident shard array (modified in place).

Definition at line 1610 of file BloomFilter.cuh.

  {
    constexpr uint64_t sequenceTileBases = Config::cudaBlockSize + Config::k - 1;
    constexpr uint32_t warpSize = 32;
    constexpr uint32_t warpsPerBlock = Config::cudaBlockSize / warpSize;
 
    using WarpReduceWord = cub::WarpReduce<uint64_t>;
 
    __shared__ uint8_t sequenceTile[sequenceTileBases];
    __shared__ typename WarpReduceWord::TempStorage reduceStorage[warpsPerBlock][4];
 
    const uint64_t numKmers = input.kmerCount();
    const uint64_t blockStartKmer = static_cast<uint64_t>(blockIdx.x) * Config::cudaBlockSize;
    if (blockStartKmer >= numKmers) {
        return;
    }
 
    const uint64_t blockKmers = min(Config::cudaBlockSize, numKmers - blockStartKmer);
    const auto localKmerIndex = static_cast<uint64_t>(threadIdx.x);
    const bool inRange = localKmerIndex < blockKmers;
 
    const bool blockAllValid = prepareSequenceHashTiles<Config>(
        input.sequence.data(), blockStartKmer, blockKmers, sequenceTile
    );
 
    // Avoid early returns so all warp lanes can participate in the segmented
    // warp reductions below.
    bool active = inRange;
 
    if (active && !blockAllValid) {
        active = kmerIsValid<Config>(sequenceTile, localKmerIndex);
    }
 
    // Inactive threads keep zero masks and a per-lane sentinel shard index so
    // contiguous run detection naturally splits around them.
    uint64_t minimizerHash = 0;
    uint64_t wordMask0 = 0;
    uint64_t wordMask1 = 0;
    uint64_t wordMask2 = 0;
    uint64_t wordMask3 = 0;
 
    if (active) {
        const uint64_t packedKmer =
            packKmerFromTile<Config, Config::k>(sequenceTile, localKmerIndex);
        minimizerHash = packedKmerMinimizerHash<Config>(packedKmer);
 
        uint64_t h_s = packedKmerSmerHash<Config>(packedKmer, 0);
        Filter<Config>::Shard::sectorizedHashToMasks(
            h_s, wordMask0, wordMask1, wordMask2, wordMask3
        );
        _Pragma("unroll")
        for (uint64_t smerOffset = 1; smerOffset < Config::findereSpan; ++smerOffset) {
            h_s = packedKmerSmerHash<Config>(packedKmer, smerOffset);
            Filter<Config>::Shard::sectorizedHashToMasks(
                h_s, wordMask0, wordMask1, wordMask2, wordMask3
            );
        }
    }
 
    // Warp-local segmented reductions: contiguous threads sharing the same
    // shard merge their masks so only the run head issues the atomicOrs.
    const auto shardIdx =
        static_cast<uint32_t>(active ? (minimizerHash & (shards.size() - 1)) : ~threadIdx.x);
 
    const uint32_t lane = threadIdx.x & (warpSize - 1);
    const uint32_t warpIdx = threadIdx.x / warpSize;
    const uint32_t prevShardIdx = __shfl_up_sync(0xffffffff, shardIdx, 1);
    const bool runHead = (lane == 0) || (shardIdx != prevShardIdx);
    const BitwiseOr<uint64_t> bitwiseOr{};
 
    wordMask0 = WarpReduceWord(reduceStorage[warpIdx][0])
                    .HeadSegmentedReduce(wordMask0, runHead, bitwiseOr);
    wordMask1 = WarpReduceWord(reduceStorage[warpIdx][1])
                    .HeadSegmentedReduce(wordMask1, runHead, bitwiseOr);
    wordMask2 = WarpReduceWord(reduceStorage[warpIdx][2])
                    .HeadSegmentedReduce(wordMask2, runHead, bitwiseOr);
    wordMask3 = WarpReduceWord(reduceStorage[warpIdx][3])
                    .HeadSegmentedReduce(wordMask3, runHead, bitwiseOr);
 
    if (runHead && active) {
        auto& shard = shards[shardIdx];
        if (wordMask0 != 0) {
            atomicOrWord(&shard.words[0], wordMask0);
        }
        if (wordMask1 != 0) {
            atomicOrWord(&shard.words[1], wordMask1);
        }
        if (wordMask2 != 0) {
            atomicOrWord(&shard.words[2], wordMask2);
        }
        if (wordMask3 != 0) {
            atomicOrWord(&shard.words[3], wordMask3);
        }
    }
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ kmerIsValid()

template<typename Config >

__device__ __forceinline__ bool cusbf::detail::kmerIsValid	(	const uint8_t *	tile,
		uint64_t	start
	)

Definition at line 1447 of file BloomFilter.cuh.

                                                                                 {
    _Pragma("unroll")
    for (uint64_t i = 0; i < Config::k; ++i) {
        if (tile[start + i] == Config::Alphabet::invalidSymbol) {
            return false;
        }
    }
    return true;
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ loadShardWords4()

template<typename Config >

__device__ __forceinline__ void cusbf::detail::loadShardWords4	(	const typename Filter< Config >::Shard *	shards,
		uint64_t	shardIndex,
		uint64_t *	w
	)

Loads all four 64-bit words of a shard into a local array.

On sm_100+ issues a single 256-bit non-coherent global load, on older architectures falls back to two 128-bit loads.

Template Parameters

Config Filter configuration.

Parameters

shards	Pointer to the device shard array.
shardIndex	Index of the shard to load.
w	Output array of (at least) four words.

Definition at line 1372 of file BloomFilter.cuh.

                                                                                              {
#if __CUDA_ARCH__ >= 1000
    detail::load256BitGlobalNC(shards[shardIndex].words, w[0], w[1], w[2], w[3]);
#else
    detail::load128BitGlobalNC(shards[shardIndex].words + 0, w[0], w[1]);
    detail::load128BitGlobalNC(shards[shardIndex].words + 2, w[2], w[3]);
#endif
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ minimizerHash64()

constexpr __host__ __device__ __forceinline__ uint64_t cusbf::detail::minimizerHash64 ( uint64_t key )

constexpr

Fast 64-bit hash sufficient for uniform minimizer selection.

A single Knuth multiplicative step — provides enough uniformity for shard selection without the full avalanche quality of hash64.

Parameters

key	Packed m-mer input.

Returns: Hash value used to select the minimum (minimizer).

Definition at line 209 of file hashutil.cuh.

                                                                                     {
    return key * 0x9E3779B97F4A7C15ULL;
}

Here is the call graph for this function:

◆ multiplicativeSaltLiteral()

template<uint64_t Index>

__host__ __device__ __forceinline__ constexpr uint64_t cusbf::detail::multiplicativeSaltLiteral ( )

constexpr

Returns the multiplicative salt constant for hash function Index.

Template Parameters

Index Salt index in [0, 15].

Returns: Salt value.

Definition at line 251 of file BloomFilter.cuh.

                                                                                                 {
    static_assert(Index < 16, "Salt index out of range");
    return SaltLiteral<Index>::value;
}

Here is the call graph for this function:

◆ packedKmerMinimizerHash()

template<typename Config >

__device__ __forceinline__ uint64_t cusbf::detail::packedKmerMinimizerHash ( uint64_t packedKmer )

Computes the minimizer hash for a packed k-mer.

Iterates over all m-mers within the k-mer and returns the minimum hash value, which is used to select the target shard.

Template Parameters

Config Filter configuration.

Parameters

packedKmer 2-bit packed k-mer.

Returns: Minimizer hash value.

Definition at line 1332 of file BloomFilter.cuh.

                                                                                               {
    uint64_t minimizerHash = kInvalidHash;
    _Pragma("unroll")
    for (uint64_t offset = 0; offset < Config::minimizerSpan; ++offset) {
        const uint64_t packedMmer =
            extractPackedSubwindow<Config, Config::m, Config::k>(packedKmer, offset);
        minimizerHash = min(minimizerHash, detail::minimizerHash64(packedMmer));
    }
    return minimizerHash;
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ packedKmerSmerHash()

template<typename Config >

__device__ __forceinline__ uint64_t cusbf::detail::packedKmerSmerHash	(	uint64_t	packedKmer,
		uint64_t	start
	)

Computes the hash for the s-mer at position start within a packed k-mer.

Template Parameters

Config Filter configuration.

Parameters

packedKmer	2-bit packed k-mer.
start	Zero-based start position of the s-mer within the k-mer.

Returns: Hash of the s-mer.

Definition at line 1353 of file BloomFilter.cuh.

                                                        {
    const uint64_t packedSmer =
        extractPackedSubwindow<Config, Config::s, Config::k>(packedKmer, start);
    return detail::hash64(packedSmer);
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ packedWindowMask()

template<typename Config , uint64_t Length>

__host__ __device__ __forceinline__ constexpr uint64_t cusbf::detail::packedWindowMask ( )

constexpr

Returns a bitmask covering Length packed alphabet symbols.

Returns UINT64_MAX when the packed window consumes all 64 bits.

Template Parameters

Length Number of symbols.

Definition at line 285 of file BloomFilter.cuh.

                                                                                        {
    if constexpr (Length * Config::symbolBits >= 64) {
        return std::numeric_limits<uint64_t>::max();
    } else {
        return (uint64_t{1} << (Config::symbolBits * Length)) - 1;
    }
}

Here is the call graph for this function:

◆ packKmerFromTile()

template<typename Config , uint64_t K>

__device__ __forceinline__ uint64_t cusbf::detail::packKmerFromTile	(	const uint8_t *	tile,
		uint64_t	start
	)

Packs K symbols from a shared-memory tile into an integer.

Template Parameters

K	k-mer length.

Parameters

tile	Encoded symbol tile in shared memory.
start	Start position within the tile.

Returns: Packed k-mer.

Definition at line 1390 of file BloomFilter.cuh.

                                                                                          {
    uint64_t packed = 0;
    _Pragma("unroll")
    for (uint64_t i = 0; i < K; ++i) {
        packed = (packed << Config::symbolBits) | (tile[start + i] & Config::symbolMask);
    }
    return packed;
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ prepareSequenceHashTiles()

template<typename Config >

__device__ __forceinline__ bool cusbf::detail::prepareSequenceHashTiles	(	const char *	sequence,
		uint64_t	blockStartKmer,
		uint64_t	blockKmers,
		uint8_t *	sequenceTile
	)

Cooperatively loads and encodes a tile of symbols into shared memory.

All threads in the block participate. The return value (via __syncthreads_count) is true only if every base in the tile is a valid alphabet symbol.

Template Parameters

Config Filter configuration.

Parameters

sequence	Device-resident sequence pointer.
blockStartKmer	Index of the first k-mer assigned to this block.
blockKmers	Number of k-mers handled by this block.
sequenceTile	Shared-memory output buffer (blockKmers + k - 1 bytes).

Returns: true if no invalid symbols are present in the tile.

Definition at line 1472 of file BloomFilter.cuh.

  {
    const uint64_t tileBases = blockKmers + Config::k - 1;
 
    bool localInvalidBase = false;
    for (uint64_t idx = threadIdx.x; idx < tileBases; idx += Config::cudaBlockSize) {
        const uint8_t encodedBase =
            Config::Alphabet::encode(sequence + (blockStartKmer + idx) * Config::symbolWidth);
        sequenceTile[idx] = encodedBase;
        localInvalidBase |= (encodedBase == Config::Alphabet::invalidSymbol);
    }
    return __syncthreads_count(localInvalidBase) == 0;
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ sectorizedContainsPackedKmer()

template<typename Config >

__device__ __forceinline__ bool cusbf::detail::sectorizedContainsPackedKmer	(	uint64_t	packedKmer,
		const uint64_t *	w
	)

Tests whether a packed k-mer is present in a pre-loaded shard.

Checks all s-mer hashes across the k-mer against the four shard words. Returns false as soon as any required bit is absent.

Template Parameters

Config Filter configuration.

Parameters

packedKmer	Packed k-mer to query.
w	The four pre-loaded shard words.

Returns: true if all required bits are set.

Definition at line 1429 of file BloomFilter.cuh.

                                                                     {
    bool present = true;
    _Pragma("unroll")
    for (uint64_t smerOffset = 0; smerOffset < Config::findereSpan; ++smerOffset) {
        const uint64_t smerHash = packedKmerSmerHash<Config>(packedKmer, smerOffset);
        detail::forEachHashIndex<Config>(
            [&]<uint64_t HashIndex>(std::integral_constant<uint64_t, HashIndex>) {
                constexpr uint64_t s = HashIndex % Config::blockWordCount;
                const uint64_t bitPos =
                    Filter<Config>::Shard::template sectorizedBitAddress<HashIndex>(smerHash);
                present &= ((w[s] >> bitPos) & 1) != 0;
            }
        );
    }
    return present;
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ separatorByteAlwaysEncodesInvalid()

template<typename T >

consteval bool cusbf::detail::separatorByteAlwaysEncodesInvalid ( )

Tests that for every position in the input, placing the separator byte at that position always results in an invalid encoding.

This is a necessary condition for the separator to function correctly when concatenating sequences, as it prevents the creation of valid symbols that span across sequence boundaries.

Template Parameters

T	Alphabet type to test.

Returns: bool True if the separator byte always produces an invalid encoding at every position, false if any position allows the separator to be part of a valid encoding.

Definition at line 67 of file Alphabet.cuh.

                                                   {
    for (uint64_t separatorPosition = 0; separatorPosition < T::symbolWidth; ++separatorPosition) {
        char input[T::symbolWidth]{};
        if (!separatorPositionAlwaysEncodesInvalid<T>(input, separatorPosition, 0)) {
            return false;
        }
    }
    return true;
}

Here is the call graph for this function:

◆ separatorPositionAlwaysEncodesInvalid()

template<typename T >

consteval bool cusbf::detail::separatorPositionAlwaysEncodesInvalid	(	char *	input,
		uint64_t	separatorPosition,
		uint64_t	index
	)

Recursively tests whether placing the separator byte at any position in an input of valid bytes always results in an invalid encoding.

This ensures that the separator cannot be confused with valid symbols when concatenating sequences.

Template Parameters

T	Alphabet type to test.

Parameters

input	Buffer to construct input strings for encoding. Must have length at least `T::symbolWidth`.
separatorPosition	Position at which to place the separator byte in the input.
index	Current index being set in the input. Should be called with 0 initially.

Returns: bool True if the separator byte always produces an invalid encoding, false if any combination of valid bytes with the separator produces a valid encoding.

Definition at line 37 of file Alphabet.cuh.

                                                                                               {
    if (index == T::symbolWidth) {
        return T::encode(input) == static_cast<uint8_t>(T::invalidSymbol);
    }
 
    if (index == separatorPosition) {
        input[index] = static_cast<char>(T::separator);
        return separatorPositionAlwaysEncodesInvalid<T>(input, separatorPosition, index + 1);
    }
 
    for (uint64_t byteIndex = 0; byteIndex < validByteCount<T>(); ++byteIndex) {
        input[index] = T::validBytes[byteIndex];
        if (!separatorPositionAlwaysEncodesInvalid<T>(input, separatorPosition, index + 1)) {
            return false;
        }
    }
    return true;
}

Here is the call graph for this function:

Here is the caller graph for this function:

◆ validByteCount()

template<typename T >

consteval uint64_t cusbf::detail::validByteCount ( )

Definition at line 14 of file Alphabet.cuh.

                                    {
    uint64_t count = 0;
    while (T::validBytes[count] != '\0') {
        ++count;
    }
    return count;
}

Variable Documentation

◆ kContainsSequenceStride

constexpr uint32_t cusbf::detail::kContainsSequenceStride = 4

inlineconstexpr

Definition at line 143 of file BloomFilter.cuh.

◆ kInvalidHash

constexpr uint64_t cusbf::detail::kInvalidHash = std::numeric_limits<uint64_t>::max()

inlineconstexpr

Sentinel hash value indicating "no valid minimizer found".

Definition at line 167 of file BloomFilter.cuh.

Namespaces

Classes

Functions

Variables

Function Documentation

◆ advancePackedKmer()

◆ atomicOrWord()

◆ containsSequenceKmersKernel()

◆ extractPackedSubwindow()

◆ forEachHashIndex()

◆ forEachHashIndexImpl()

◆ hash64()

◆ insertSequenceKmersKernel()

◆ kmerIsValid()

◆ loadShardWords4()

◆ minimizerHash64()

◆ multiplicativeSaltLiteral()

◆ packedKmerMinimizerHash()

◆ packedKmerSmerHash()

◆ packedWindowMask()

◆ packKmerFromTile()

◆ prepareSequenceHashTiles()

◆ sectorizedContainsPackedKmer()

◆ separatorByteAlwaysEncodesInvalid()

◆ separatorPositionAlwaysEncodesInvalid()

◆ validByteCount()

Variable Documentation

◆ kContainsSequenceStride

◆ kInvalidHash