cuSBF
Loading...
Searching...
No Matches
Classes | Namespaces | Functions
sequence_kmer.cuh File Reference
#include <cuda_runtime.h>
#include <cub/warp/warp_reduce.cuh>
#include <cstdint>
#include <cusbf/config.cuh>
#include <cusbf/detail/filter_common.cuh>
#include <cusbf/detail/filter_impl.cuh>
#include <cusbf/device_span.cuh>
#include <cusbf/hashutil.cuh>
#include <cusbf/helpers.cuh>

Go to the source code of this file.

Classes

struct  cusbf::detail::SequenceKmerInput< Config >
 Device-side view of an encoded sequence for k-mer / s-mer counting. More...
 

Namespaces

namespace  cusbf
 
namespace  cusbf::detail
 

Functions

template<typename Config >
__device__ __forceinline__ uint64_t cusbf::detail::packed_kmer_minimizer_hash (uint64_t packed_kmer)
 Minimum minimizer hash over all m-mers in a packed k-mer.
 
template<typename Config >
__device__ __forceinline__ uint64_t cusbf::detail::packed_kmer_smer_hash (uint64_t packed_kmer, uint64_t start)
 Bloom hash for the s-mer at start within a packed k-mer.
 
template<typename Config >
__device__ __forceinline__ void cusbf::detail::load_shard_words4 (const filter_block< Config > *shards, uint64_t shard_index, uint64_t *w)
 Loads four 64-bit shard words with 256-bit (sm_100+) or 128-bit vector loads.
 
template<typename Config , uint64_t K>
__device__ __forceinline__ uint64_t cusbf::detail::pack_kmer_from_tile (const uint8_t *tile, uint64_t start)
 Packs K encoded symbols from a shared-memory tile starting at start.
 
template<typename Config , uint64_t K>
__device__ __forceinline__ uint64_t cusbf::detail::advance_packed_kmer (uint64_t packed, uint8_t new_base)
 Slides a packed k-mer window by one encoded base.
 
template<typename Config >
__device__ __forceinline__ bool cusbf::detail::kmer_is_valid (const uint8_t *tile, uint64_t start)
 True when no symbol in the k-mer window is the alphabet invalid sentinel.
 
template<typename Config >
__device__ __forceinline__ bool cusbf::detail::prepare_sequence_hash_tiles (const char *sequence, uint64_t block_start_kmer, uint64_t block_kmers, uint8_t *sequence_tile)
 Encodes a block's sequence slice into sequence_tile and reports global validity.
 
template<uint32_t k_stride, typename Config >
__device__ __forceinline__ uint32_t cusbf::detail::build_stride_kmer_valid_mask (uint64_t thread_offset, uint64_t block_kmers, bool block_all_valid, const uint8_t *sequence_tile)
 Builds the per-thread validity bitmask for strided query kernels.