cuSBF
Loading...
Searching...
No Matches
Classes | Namespaces | Functions | Variables
BloomFilter.cuh File Reference
#include <cuda/__cmath/ceil_div.h>
#include <cuda_runtime.h>
#include <cuda/std/bit>
#include <cuda/std/span>
#include <cuda/stream>
#include <cub/warp/warp_reduce.cuh>
#include <thrust/copy.h>
#include <thrust/detail/execution_policy.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/fill.h>
#include <thrust/transform_reduce.h>
#include <algorithm>
#include <concepts>
#include <cstddef>
#include <cstdint>
#include <limits>
#include <stdexcept>
#include <string_view>
#include <type_traits>
#include <utility>
#include <vector>
#include "Alphabet.cuh"
#include "device_span.cuh"
#include "Fastx.hpp"
#include "hashutil.cuh"
#include "helpers.cuh"

Go to the source code of this file.

Classes

struct  cusbf::Config< K_, S_, M_, HashCount_, CudaBlockSize_, Alphabet_ >
 Compile-time configuration for a cusbf::Filter. More...
 
struct  cusbf::detail::BitwiseOr< T >
 
struct  cusbf::detail::SaltLiteral< 0 >
 
struct  cusbf::detail::SaltLiteral< 1 >
 
struct  cusbf::detail::SaltLiteral< 2 >
 
struct  cusbf::detail::SaltLiteral< 3 >
 
struct  cusbf::detail::SaltLiteral< 4 >
 
struct  cusbf::detail::SaltLiteral< 5 >
 
struct  cusbf::detail::SaltLiteral< 6 >
 
struct  cusbf::detail::SaltLiteral< 7 >
 
struct  cusbf::detail::SaltLiteral< 8 >
 
struct  cusbf::detail::SaltLiteral< 9 >
 
struct  cusbf::detail::SaltLiteral< 10 >
 
struct  cusbf::detail::SaltLiteral< 11 >
 
struct  cusbf::detail::SaltLiteral< 12 >
 
struct  cusbf::detail::SaltLiteral< 13 >
 
struct  cusbf::detail::SaltLiteral< 14 >
 
struct  cusbf::detail::SaltLiteral< 15 >
 
class  cusbf::Filter< Config >
 cuSBF GPU-accelerated sectorized Bloom filter. More...
 
struct  cusbf::Filter< Config >::Shard
 One 256-bit filter block stored as an array of Config::blockWordCount words. More...
 
struct  cusbf::detail::SequenceKmerInput< Config >
 Kernel input descriptor for a sequence k-mer sweep. More...
 

Namespaces

namespace  cusbf
 
namespace  cusbf::detail
 

Functions

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.
 
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.
 
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.
 
template<uint64_t Index>
__host__ __device__ __forceinline__ constexpr uint64_t cusbf::detail::multiplicativeSaltLiteral ()
 Returns the multiplicative salt constant for hash function Index.
 
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).
 
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<typename Config , uint64_t Length>
__host__ __device__ __forceinline__ constexpr uint64_t cusbf::detail::packedWindowMask ()
 Returns a bitmask covering Length packed alphabet symbols.
 
template<typename Config , uint64_t WindowLength, uint64_t K>
__host__ __device__ __forceinline__ constexpr uint64_t cusbf::detail::extractPackedSubwindow (uint64_t packedKmer, uint64_t start)
 Extracts a packed sub-window from a packed k-mer.
 
__device__ __forceinline__ void cusbf::detail::atomicOrWord (uint64_t *ptr, uint64_t value)
 Atomically ORs value into the device word at ptr.
 
template<typename Config >
__device__ __forceinline__ uint64_t cusbf::detail::packedKmerMinimizerHash (uint64_t packedKmer)
 Computes the minimizer hash for a packed k-mer.
 
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<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.
 
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<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.
 
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.
 
template<typename Config >
__device__ __forceinline__ bool cusbf::detail::kmerIsValid (const uint8_t *tile, uint64_t start)
 

Variables

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