cusbf::filter< Config > Class Template Reference

GPU Super Bloom filter (defined in filter.cuh). More...

Public Types
using	block_type = filter_block< Config >
	One 256-bit filter block stored as an array of Config::blockWordCount words.
using	Shard = block_type
	Alias for block_type (one 256-bit shard).

Public Member Functions
	filter (uint64_t requestedFilterBits)
	Constructs a Filter with at least requestedFilterBits bits of storage.
	filter (const filter &)=delete
	Non-copyable (owns device shard storage).
filter &	operator= (const filter &)=delete
	filter (filter &&) noexcept=default
	Move-constructs, transfers shard vectors and staging buffers.
filter &	operator= (filter &&) noexcept=default
	Move-assigns shard storage and staging buffers.
	~filter ()=default
	Destroys device allocations and releases staging scratch.
filter_ref< Config >	ref () const noexcept
	Non-owning device reference to this filter's shard storage.
Result< uint64_t >	insert_sequence (std::string_view sequence, cuda::stream_ref stream=cudaStream_t{})
	Inserts all valid k-mers from a host-resident sequence.
Result< uint64_t >	insert_sequence_async (device_span< const char > d_sequence, cuda::stream_ref stream=cudaStream_t{})
	Async insert of k-mers from a device-resident sequence.
Result< uint64_t >	insert_dense_packed_async (device_span< const uint64_t > d_words, uint64_t num_symbols, cuda::stream_ref stream=cudaStream_t{})
	Inserts all k-mers from a dense packed symbol buffer on the device.
Result< uint64_t >	insert_dense_packed (std::span< const uint64_t > words, uint64_t num_symbols, cuda::stream_ref stream=cudaStream_t{})
	Inserts all k-mers from a host-resident dense packed symbol buffer.
Result< void >	contains_dense_packed_async (device_span< const uint64_t > d_words, uint64_t num_symbols, device_span< uint8_t > d_output, cuda::stream_ref stream=cudaStream_t{}) const
	Async query of k-mers from a dense packed symbol buffer on the device.
Result< void >	contains_dense_packed (std::span< const uint64_t > words, uint64_t num_symbols, device_span< uint8_t > d_output, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all k-mers from a host-resident dense packed symbol buffer into device memory.
Result< std::vector< uint8_t > >	contains_dense_packed (std::span< const uint64_t > words, uint64_t num_symbols, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all k-mers from a host-resident dense packed symbol buffer.
constexpr uint64_t	dense_packed_kmer_count (uint64_t num_symbols) const
	Number of k-mer windows in a dense packed sequence of num_symbols symbols.
Result< FastxInsertReport >	insert_record_batch (RecordBatchView batch, cuda::stream_ref stream=cudaStream_t{})
	Inserts a dense host-resident record batch.
Result< uint64_t >	insert_record_batch_async (device_span< const char > d_sequence, std::span< const RecordRange > records, cuda::stream_ref stream=cudaStream_t{})
	Async insert of k-mers from a device-resident record batch.
Result< FastxInsertReport >	insert_fastx (std::istream &input, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{})
	Inserts all k-mers from a FASTA/FASTQ input stream.
Result< FastxInsertReport >	insert_fastx_file (const std::filesystem::path &path, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{})
	Inserts all k-mers from a FASTA/FASTQ file.
Result< void >	contains_sequence_async (device_span< const char > d_sequence, device_span< uint8_t > d_output, cuda::stream_ref stream=cudaStream_t{}) const
	Async query of k-mers from a device-resident sequence.
Result< void >	contains_sequence (std::string_view sequence, device_span< uint8_t > d_output, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all valid k-mers from a host-resident sequence into device memory.
Result< std::vector< uint8_t > >	contains_sequence (std::string_view sequence, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all valid k-mers from a host-resident sequence.
Result< FastxQueryReport >	query_record_batch (RecordBatchView batch, cuda::stream_ref stream=cudaStream_t{}) const
	Queries a dense host-resident record batch and returns aggregate counts.
template<RecordQueryConsumer Consumer>
Result< FastxQueryReport >	query_record_batch (RecordBatchView batch, Consumer &&consume, cuda::stream_ref stream=cudaStream_t{}) const
	Queries a dense host-resident record batch and streams per-record results.
Result< void >	contains_record_batch_async (device_span< const char > d_sequence, std::span< const RecordRange > records, device_span< uint8_t > d_output, cuda::stream_ref stream=cudaStream_t{}) const
	Async query of k-mers from a device-resident record batch.
Result< FastxQueryReport >	query_fastx (std::istream &input, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all k-mers from a FASTA/FASTQ input stream via chunked streaming.
Result< FastxQueryReport >	query_fastx_file (const std::filesystem::path &path, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all k-mers from a FASTA/FASTQ file.
template<FastxRecordConsumer Consumer>
Result< FastxQueryReport >	query_fastx_records (std::istream &input, Consumer &&consume, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{}) const
	Queries a FASTA/FASTQ stream and emits one record result per parsed record.
template<FastxRecordConsumer Consumer>
Result< FastxQueryReport >	query_fastx_file_records (const std::filesystem::path &path, Consumer &&consume, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{}) const
	Queries a FASTA/FASTQ file and emits one record result per parsed record.
Result< FastxDetailedQueryReport >	query_fastx_detailed (std::istream &input, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all k-mers from a FASTA/FASTQ input stream via chunked streaming and preserves per-record hit vectors.
Result< FastxDetailedQueryReport >	query_fastx_file_detailed (const std::filesystem::path &path, double fill_fraction=0.7, cuda::stream_ref stream=cudaStream_t{}) const
	Queries all k-mers from a FASTA/FASTQ file via chunked streaming and preserves per-record hit vectors.
Result< void >	clear (cuda::stream_ref stream=cudaStream_t{})
	Resets all filter bits to zero and synchronises the stream.
float	load_factor () const
	Computes the fraction of set bits in the filter.
uint64_t	filter_bits () const
	Returns the total allocated capacity of the filter in bits.
uint64_t	num_shards () const
	Returns the number of shards.

Static Public Member Functions
static constexpr uint64_t	dense_packed_word_count (uint64_t num_symbols)
	Number of uint64_t words for num_symbols dense packed symbols.

Detailed Description

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

GPU Super Bloom filter (defined in filter.cuh).

GPU-resident Super Bloom filter for batch k-mer insert and query.

Config fixes k-mer length, minimizer and s-mer widths, hash count, CUDA block size, and alphabet at compile time. Host bulk APIs synchronize before returning, device-span *_async methods do not.

FASTX and record-batch paths inject alphabet separator bytes between records so cross-record k-mers are never formed.

Template Parameters

Config

Compile-time filter configuration (cusbf::Config).

Definition at line 69 of file filter.cuh.

Member Typedef Documentation

block_type

template<typename Config >

using cusbf::filter< Config >::block_type = 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.

The struct is 32-byte aligned to enable vectorised loads via cusbf::detail::load256BitGlobalNC.

Definition at line 86 of file filter.cuh.

Shard

template<typename Config >

using cusbf::filter< Config >::Shard = block_type

Alias for block_type (one 256-bit shard).

Definition at line 88 of file filter.cuh.

Constructor & Destructor Documentation

filter() [1/3]

template<typename Config >

cusbf::filter< Config >::filter ( uint64_t  requestedFilterBits )

inlineexplicit

Constructs a Filter with at least requestedFilterBits bits of storage.

The actual allocated capacity is rounded up to the next power-of-two number of shards.

Parameters

requestedFilterBits

Desired filter capacity in bits.

Definition at line 108 of file filter.cuh.

        : num_shards_(
              cuda::std::bit_ceil(
                  std::max<uint64_t>(
                      1,
                      cuda::ceil_div(requestedFilterBits, Config::filterBlockBits)
                  )
              )
          ),
          filter_bits_(num_shards_ * Config::filterBlockBits),
          d_shards_(num_shards_) {
        CUSBF_UNWRAP(clear());
    }

Here is the call graph for this function:

filter() [2/3]

template<typename Config >

cusbf::filter< Config >::filter ( const filter< Config > &  )

delete

Non-copyable (owns device shard storage).

filter() [3/3]

template<typename Config >

cusbf::filter< Config >::filter ( filter< Config > &&  )

defaultnoexcept

Move-constructs, transfers shard vectors and staging buffers.

~filter()

template<typename Config >

cusbf::filter< Config >::~filter ( )

default

Destroys device allocations and releases staging scratch.

Member Function Documentation

clear()

template<typename Config >

Result< void > cusbf::filter< Config >::clear ( cuda::stream_ref  stream = cudaStream_t{} )

inline

Resets all filter bits to zero and synchronises the stream.

Parameters

stream

CUDA stream to use.

Definition at line 793 of file filter.cuh.

                                                                         {}) {
        CUSBF_CUDA_TRY(cudaMemsetAsync(
            thrust::raw_pointer_cast(d_shards_.data()),
            0,
            d_shards_.size() * sizeof(block_type),
            stream.get()
        ));
 
        CUSBF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
        return {};
    }

Here is the caller graph for this function:

contains_dense_packed() [1/2]

template<typename Config >

Result< std::vector< uint8_t > > cusbf::filter< Config >::contains_dense_packed ( std::span< const uint64_t >  words, uint64_t  num_symbols, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all k-mers from a host-resident dense packed symbol buffer.

Copies d_words to device, queries, copies results back, and synchronises.

Definition at line 289 of file filter.cuh.

                                            {}
    ) const {
        const uint64_t numKmers = dense_packed_kmer_count(num_symbols);
        if (numKmers == 0) {
            return std::vector<uint8_t>{};
        }
 
        std::vector<uint8_t> output(numKmers);
        fastx_state_.ensure_result_capacity(output.size());
        const device_span<uint8_t> d_output{
            thrust::raw_pointer_cast(fastx_state_.result_buffer_device().data()), output.size()
        };
        CUSBF_TRY(contains_dense_packed(words, num_symbols, d_output, stream));
        CUSBF_CUDA_TRY(cudaMemcpy(
            output.data(),
            thrust::raw_pointer_cast(fastx_state_.result_buffer_device().data()),
            output.size() * sizeof(uint8_t),
            cudaMemcpyDeviceToHost
        ));
        return output;
    }

contains_dense_packed() [2/2]

template<typename Config >

Result< void > cusbf::filter< Config >::contains_dense_packed ( std::span< const uint64_t >  words, uint64_t  num_symbols, device_span< uint8_t >  d_output, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all k-mers from a host-resident dense packed symbol buffer into device memory.

Copies words to device, launches the query kernel into d_output, and synchronises. Does not allocate the result buffer or copy hits back to the host.

Definition at line 267 of file filter.cuh.

                                            {}
    ) const {
        const uint64_t numKmers = dense_packed_kmer_count(num_symbols);
        if (numKmers == 0) {
            return {};
        }
 
        const auto staged = CUSBF_TRY(fastx_state_.staged_dense_packed_view(words, stream));
        CUSBF_TRY(launch_contains_dense_packed(staged, num_symbols, d_output, stream));
        CUSBF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
        return {};
    }

contains_dense_packed_async()

template<typename Config >

Result< void > cusbf::filter< Config >::contains_dense_packed_async ( device_span< const uint64_t >  d_words, uint64_t  num_symbols, device_span< uint8_t >  d_output, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Async query of k-mers from a dense packed symbol buffer on the device.

d_output receives one byte per k-mer (1 = present, 0 = absent). Does not synchronise the stream.

Definition at line 248 of file filter.cuh.

                                            {}
    ) const {
        if (dense_packed_kmer_count(num_symbols) == 0) {
            return {};
        }
 
        return launch_contains_dense_packed(d_words, num_symbols, d_output, stream);
    }

contains_record_batch_async()

template<typename Config >

Result< void > cusbf::filter< Config >::contains_record_batch_async ( device_span< const char >  d_sequence, std::span< const RecordRange >  records, device_span< uint8_t >  d_output, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Async query of k-mers from a device-resident record batch.

The device sequence is assumed to already contain alphabet separator bytes between records: the caller is responsible for embedding them. records provides ordered, non-overlapping byte ranges into d_sequence.

d_output receives one byte per k-mer across all records (1 = present, 0 = absent). The buffer must hold at least the total k-mer count.

Does not synchronise the stream.

Parameters

d_sequence	Device-resident sequence (records concatenated with separators already in place).
records	Ordered, non-overlapping record byte ranges.
d_output	Per-k-mer result buffer.
stream	CUDA stream to use.

Definition at line 603 of file filter.cuh.

                                            {}
    ) const {
        CUSBF_TRY(validate_device_record_batch(d_sequence, records));
        const uint64_t totalKmers = sequence_kmer_count(d_sequence);
        if (totalKmers == 0) {
            return {};
        }
        if (d_output.size() < totalKmers) {
            return Err(Error::invalid_argument("record batch query output span is too small"));
        }
        return launch_contains_sequence(d_sequence, d_output, stream);
    }

contains_sequence() [1/2]

template<typename Config >

Result< std::vector< uint8_t > > cusbf::filter< Config >::contains_sequence ( std::string_view  sequence, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all valid k-mers from a host-resident sequence.

Copies the sequence to device, queries, copies results back, and synchronises. The returned vector has one byte per k-mer: 1 = present, 0 = absent.

Parameters

sequence	Raw nucleotide sequence.
stream	CUDA stream to use.

Returns

Per-k-mer membership results (empty if sequence length < k).

Definition at line 515 of file filter.cuh.

                                                                                   {}) const {
        if (record_symbol_count(sequence.size()) < Config::k) {
            return std::vector<uint8_t>{};
        }
 
        std::vector<uint8_t> output(record_kmer_count(sequence.size()));
        fastx_state_.ensure_result_capacity(output.size());
        const device_span<uint8_t> d_output{
            thrust::raw_pointer_cast(fastx_state_.result_buffer_device().data()), output.size()
        };
        CUSBF_TRY(contains_sequence(sequence, d_output, stream));
        CUSBF_CUDA_TRY(cudaMemcpy(
            output.data(),
            thrust::raw_pointer_cast(fastx_state_.result_buffer_device().data()),
            output.size() * sizeof(uint8_t),
            cudaMemcpyDeviceToHost
        ));
        return output;
    }

contains_sequence() [2/2]

template<typename Config >

Result< void > cusbf::filter< Config >::contains_sequence ( std::string_view  sequence, device_span< uint8_t >  d_output, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all valid k-mers from a host-resident sequence into device memory.

Copies the sequence to device, launches the query kernel into d_output, and synchronises. Does not allocate the result buffer or copy hits back to the host.

Parameters

sequence	Raw nucleotide sequence.
d_output	Per-k-mer result buffer (must hold kmerCount() bytes).
stream	CUDA stream to use.

Definition at line 486 of file filter.cuh.

                                            {}
    ) const {
        if (record_symbol_count(sequence.size()) < Config::k) {
            return {};
        }
 
        const auto d_sequence = CUSBF_TRY(
            fastx_state_.staged_sequence_view({sequence.data(), sequence.size()}, stream)
        );
        CUSBF_TRY(launch_contains_sequence(d_sequence, d_output, stream));
        CUSBF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
        return {};
    }

contains_sequence_async()

template<typename Config >

Result< void > cusbf::filter< Config >::contains_sequence_async ( device_span< const char >  d_sequence, device_span< uint8_t >  d_output, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Async query of k-mers from a device-resident sequence.

Does not synchronise the stream. Results are written to d_output (one byte per k-mer: 1 = present, 0 = absent).

Parameters

d_sequence	Device-resident nucleotide sequence.
d_output	Per-k-mer result buffer (must hold kmerCount() bytes).
stream	CUDA stream to use.

Definition at line 464 of file filter.cuh.

                                            {}
    ) const {
        if (sequence_kmer_count(d_sequence) == 0) {
            return {};
        }
 
        return launch_contains_sequence(d_sequence, d_output, stream);
    }

dense_packed_kmer_count()

template<typename Config >

constexpr uint64_t cusbf::filter< Config >::dense_packed_kmer_count ( uint64_t  num_symbols ) const

inlineconstexpr

Number of k-mer windows in a dense packed sequence of num_symbols symbols.

Definition at line 320 of file filter.cuh.

                                                                                         {
        return detail::dense_packed_kmer_count<Config>(num_symbols);
    }

Here is the call graph for this function:

dense_packed_word_count()

template<typename Config >

static constexpr uint64_t cusbf::filter< Config >::dense_packed_word_count ( uint64_t  num_symbols )

inlinestaticconstexpr

Number of uint64_t words for num_symbols dense packed symbols.

Definition at line 315 of file filter.cuh.

                                                                                          {
        return detail::dense_packed_word_count<Config>(num_symbols);
    }

Here is the call graph for this function:

filter_bits()

template<typename Config >

uint64_t cusbf::filter< Config >::filter_bits ( ) const

inline

Returns the total allocated capacity of the filter in bits.

Definition at line 826 of file filter.cuh.

                                               {
        return filter_bits_;
    }

insert_dense_packed()

template<typename Config >

Result< uint64_t > cusbf::filter< Config >::insert_dense_packed ( std::span< const uint64_t >  words, uint64_t  num_symbols, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Inserts all k-mers from a host-resident dense packed symbol buffer.

Copies d_words to device staging, launches the insert kernel, and synchronises.

Definition at line 226 of file filter.cuh.

                                            {}
    ) {
        const uint64_t totalKmers = dense_packed_kmer_count(num_symbols);
        if (totalKmers == 0) {
            return 0;
        }
 
        const auto staged = CUSBF_TRY(fastx_state_.staged_dense_packed_view(words, stream));
        CUSBF_TRY(launch_insert_dense_packed(staged, num_symbols, stream));
        CUSBF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
        return totalKmers;
    }

insert_dense_packed_async()

template<typename Config >

Result< uint64_t > cusbf::filter< Config >::insert_dense_packed_async ( device_span< const uint64_t >  d_words, uint64_t  num_symbols, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Inserts all k-mers from a dense packed symbol buffer on the device.

d_words stores dense_packed_word_count(num_symbols) words using Config::symbolBits per encoded symbol. Adjacent k-mers overlap in the same uint64_t chunks; this path decodes a per-block symbol tile and slides packed k-mers like insert_sequence_async.

Does not synchronise the stream.

Parameters

d_words	Device-resident dense packed sequence.
num_symbols	Number of valid encoded symbols in d_words.
stream	CUDA stream to use.

Returns

Number of k-mers attempted.

Definition at line 207 of file filter.cuh.

                                            {}
    ) {
        const uint64_t totalKmers = dense_packed_kmer_count(num_symbols);
        if (totalKmers == 0) {
            return 0;
        }
 
        CUSBF_TRY(launch_insert_dense_packed(d_words, num_symbols, stream));
        return totalKmers;
    }

insert_fastx()

template<typename Config >

Result< FastxInsertReport > cusbf::filter< Config >::insert_fastx ( std::istream &  input, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Inserts all k-mers from a FASTA/FASTQ input stream.

Reads records in streaming fashion, accumulating them until the concatenated sequence approaches a GPU staging budget derived from free VRAM (fill_fraction of free memory minus a small slack reserve), then inserts each chunk independently. The budget accounts for normalized sequence device staging.

Parameters

input	Input stream containing FASTA or FASTQ records.
fill_fraction	Fraction of available GPU memory for per-chunk staging (default 0.7).
stream	CUDA stream to use.

Returns

Report summarising records indexed, bases processed, and k-mers inserted.

Definition at line 416 of file filter.cuh.

                                            {}
    ) {
        return insert_fastx_stream(input, "<stream>", fill_fraction, stream);
    }

insert_fastx_file()

template<typename Config >

Result< FastxInsertReport > cusbf::filter< Config >::insert_fastx_file ( const std::filesystem::path &  path, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Inserts all k-mers from a FASTA/FASTQ file.

Uses fill_fraction of free GPU memory to size chunks. When the file fits in a single chunk, reads via a lightweight stream (no mmap, no ping-pong). Larger files are mmap'd if they fit in host RAM, otherwise streamed, multi-chunk paths overlap host assembly with GPU work.

Parameters

path	Path to a FASTA or FASTQ file (optionally gzip-compressed).
fill_fraction	Fraction of free GPU memory for per-chunk staging (default 0.7).
stream	Optional CUDA stream, default uses an internal pipelined path.

Returns

Report summarising records indexed, bases processed, and k-mers inserted.

Definition at line 437 of file filter.cuh.

                                            {}
    ) {
        return detail::dispatch_fastx_file<Config>(
            path,
            detail::fastx_chunk_mode::insert,
            fill_fraction,
            [&](auto& reader, auto dispatch_path) {
                return insert_fastx_reader(
                    reader, std::string_view{path.native()}, fill_fraction, stream, dispatch_path
                );
            }
        );
    }

insert_record_batch()

template<typename Config >

Result< FastxInsertReport > cusbf::filter< Config >::insert_record_batch ( RecordBatchView  batch, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Inserts a dense host-resident record batch.

batch.sequence stores the raw record payloads back-to-back without separators. batch.records stores ordered, non-overlapping byte ranges into that dense buffer. The filter injects alphabet separators between records internally, so callers do not need to materialise separator bytes themselves.

Synchronises before returning.

Parameters

batch	Dense record batch to insert.
stream	CUDA stream to use.

Returns

Report summarising records indexed, bases processed, and k-mers inserted.

Definition at line 339 of file filter.cuh.

                                                                                   {}) {
        CUSBF_TRY(
            normalize_record_batch_into<Config>(
                batch,
                fastx_state_.normalized_sequence_scratch(),
                fastx_state_.normalized_records_scratch()
            )
        );
        FastxInsertReport report;
        report.recordsIndexed = fastx_state_.normalized_records_scratch().size();
        for (const NormalizedRecord& record : fastx_state_.normalized_records_scratch()) {
            report.indexedBases += record.size;
            report.insertedKmers += record.valid_kmers;
        }
        if (!fastx_state_.normalized_sequence_scratch().empty()) {
            const auto d_sequence = CUSBF_TRY(fastx_state_.staged_sequence_view(
                {fastx_state_.normalized_sequence_scratch().data(),
                 fastx_state_.normalized_sequence_scratch().size()},
                stream
            ));
            CUSBF_TRY(insert_sequence_async(d_sequence, stream));
            CUSBF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
        }
        fastx_state_.release_all();
        return report;
    }

insert_record_batch_async()

template<typename Config >

Result< uint64_t > cusbf::filter< Config >::insert_record_batch_async ( device_span< const char >  d_sequence, std::span< const RecordRange >  records, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Async insert of k-mers from a device-resident record batch.

The device sequence is assumed to already contain alphabet separator bytes between records, the caller is responsible for embedding them so that cross-record k-mers are not formed. records provides ordered, non-overlapping byte ranges into d_sequence.

The returned count is total attempted k-mers (maximum possible per record size). The kernel still skips k-mer windows that contain invalid symbols.

Does not synchronise the stream.

Parameters

d_sequence	Device-resident sequence (records concatenated with separators already in place).
records	Ordered, non-overlapping record byte ranges.
stream	CUDA stream to use.

Returns

Total attempted k-mers across all records (0 if none).

Definition at line 385 of file filter.cuh.

                                            {}
    ) {
        CUSBF_TRY(validate_device_record_batch(d_sequence, records));
        uint64_t totalKmers = 0;
        for (const RecordRange& record : records) {
            totalKmers += record_kmer_count(record.sequenceBytes);
        }
        if (totalKmers == 0) {
            return 0;
        }
        CUSBF_TRY(launch_insert_sequence(d_sequence, stream));
        return totalKmers;
    }

insert_sequence()

template<typename Config >

Result< uint64_t > cusbf::filter< Config >::insert_sequence ( std::string_view  sequence, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Inserts all valid k-mers from a host-resident sequence.

Copies the sequence to device, launches the insert kernel, and synchronises before returning. K-mer windows containing alphabet-invalid symbols are skipped.

Parameters

sequence	Host-resident sequence bytes (alphabet-encoded width per symbol).
stream	CUDA stream to use (default: null stream).

Returns

Number of k-mer windows attempted (sequences shorter than k yield 0).

Definition at line 155 of file filter.cuh.

                                                                                 {}) {
        if (record_symbol_count(sequence.size()) < Config::k) {
            return 0;
        }
 
        const uint64_t totalKmers = record_kmer_count(sequence.size());
        const auto d_sequence = CUSBF_TRY(
            fastx_state_.staged_sequence_view({sequence.data(), sequence.size()}, stream)
        );
        CUSBF_TRY(launch_insert_sequence(d_sequence, stream));
        CUSBF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
        return totalKmers;
    }

insert_sequence_async()

template<typename Config >

Result< uint64_t > cusbf::filter< Config >::insert_sequence_async ( device_span< const char >  d_sequence, cuda::stream_ref  stream = cudaStream_t{}  )

inline

Async insert of k-mers from a device-resident sequence.

Does not synchronise the stream, the caller is responsible for ordering relative to downstream operations.

Parameters

d_sequence	Device-resident nucleotide sequence.
stream	CUDA stream to use.

Returns

Number of k-mers attempted.

Definition at line 179 of file filter.cuh.

                                            {}
    ) {
        const uint64_t totalKmers = sequence_kmer_count(d_sequence);
        if (totalKmers == 0) {
            return 0;
        }
 
        CUSBF_TRY(launch_insert_sequence(d_sequence, stream));
        return totalKmers;
    }

load_factor()

template<typename Config >

float cusbf::filter< Config >::load_factor ( ) const

inline

Computes the fraction of set bits in the filter.

Returns

Load factor in [0, 1].

Definition at line 810 of file filter.cuh.

                                            {
        const auto* wordsBegin =
            reinterpret_cast<const uint64_t*>(thrust::raw_pointer_cast(d_shards_.data()));
        const uint64_t totalWords = num_shards_ * Config::blockWordCount;
        const uint64_t setBits = thrust::transform_reduce(
            thrust::device,
            wordsBegin,
            wordsBegin + totalWords,
            [] __device__(uint64_t w) -> uint64_t { return cuda::std::popcount(w); },
            uint64_t{0},
            cuda::std::plus<uint64_t>()
        );
        return static_cast<float>(setBits) / static_cast<float>(filter_bits_);
    }

num_shards()

template<typename Config >

uint64_t cusbf::filter< Config >::num_shards ( ) const

inline

Returns the number of shards.

Definition at line 831 of file filter.cuh.

                                              {
        return num_shards_;
    }

operator=() [1/2]

template<typename Config >

filter & cusbf::filter< Config >::operator= ( const filter< Config > &  )

delete

operator=() [2/2]

template<typename Config >

filter & cusbf::filter< Config >::operator= ( filter< Config > &&  )

defaultnoexcept

Move-assigns shard storage and staging buffers.

query_fastx()

template<typename Config >

Result< FastxQueryReport > cusbf::filter< Config >::query_fastx ( std::istream &  input, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all k-mers from a FASTA/FASTQ input stream via chunked streaming.

Returns aggregate counts only (no per-record hit vectors). For per-record callbacks use query_fastx_records, for owning per-record hits use query_fastx_detailed.

Parameters

input	Input stream containing FASTA or FASTQ records.
fill_fraction	Fraction of free GPU memory for per-chunk staging (default 0.7).
stream	CUDA stream to use.

Returns

Aggregate query summary for the whole stream.

See also

insert_fastx

Definition at line 632 of file filter.cuh.

                                            {}
    ) const {
        return query_fastx_stream(input, "<stream>", fill_fraction, stream);
    }

query_fastx_detailed()

template<typename Config >

Result< FastxDetailedQueryReport > cusbf::filter< Config >::query_fastx_detailed ( std::istream &  input, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all k-mers from a FASTA/FASTQ input stream via chunked streaming and preserves per-record hit vectors.

The returned report keeps aggregate counts plus one detailed record result in source order. Each detailed hit vector contains one byte per k-mer window: 1 = present, 0 = absent. Invalid-symbol windows remain in the vector as 0 and are excluded from queriedKmers.

Parameters

input	Input stream containing FASTA or FASTQ records.
fill_fraction	Fraction of free GPU memory for per-chunk staging (default 0.7).
stream	CUDA stream to use.

Returns

Aggregate and per-record query results.

See also

query_fastx

Definition at line 750 of file filter.cuh.

                                            {}
    ) const {
        detail::FastxReader reader(input, "<stream>");
        return query_fastx_detailed_stream(
            reader, "<stream>", fill_fraction, stream, detail::fastx_dispatch_path::chunked_stream
        );
    }

query_fastx_file()

template<typename Config >

Result< FastxQueryReport > cusbf::filter< Config >::query_fastx_file ( const std::filesystem::path &  path, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all k-mers from a FASTA/FASTQ file.

Returns aggregate counts only. See query_fastx_file_records and query_fastx_file_detailed for per-record results.

Parameters

path	Path to a FASTA or FASTQ file (optionally gzip-compressed).
fill_fraction	Fraction of free GPU memory for per-chunk staging (default 0.7).
stream	CUDA stream to use.

Returns

Aggregate query summary for the whole file.

See also

insert_fastx_file for dispatch and chunking behavior.

Definition at line 652 of file filter.cuh.

                                            {}
    ) const {
        return detail::dispatch_fastx_file<Config>(
            path,
            detail::fastx_chunk_mode::query,
            fill_fraction,
            [&](auto& reader, auto dispatch_path) {
                return query_fastx_reader(
                    reader, std::string_view{path.native()}, fill_fraction, stream, dispatch_path
                );
            }
        );
    }

query_fastx_file_detailed()

template<typename Config >

Result< FastxDetailedQueryReport > cusbf::filter< Config >::query_fastx_file_detailed ( const std::filesystem::path &  path, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries all k-mers from a FASTA/FASTQ file via chunked streaming and preserves per-record hit vectors.

Parameters

path	Path to a FASTA or FASTQ file (optionally gzip-compressed).
fill_fraction	Fraction of free GPU memory for per-chunk staging (default 0.7).
stream	CUDA stream to use.

Returns

Aggregate and per-record query results.

See also

query_fastx_detailed

Definition at line 771 of file filter.cuh.

                                            {}
    ) const {
        return detail::dispatch_fastx_file<Config>(
            path,
            detail::fastx_chunk_mode::query,
            fill_fraction,
            [&](auto& reader, auto dispatch_path) {
                return query_fastx_detailed_stream(
                    reader, std::string_view{path.native()}, fill_fraction, stream, dispatch_path
                );
            }
        );
    }

query_fastx_file_records()

template<typename Config >

template<FastxRecordConsumer Consumer>

Result< FastxQueryReport > cusbf::filter< Config >::query_fastx_file_records ( const std::filesystem::path &  path, Consumer &&  consume, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries a FASTA/FASTQ file and emits one record result per parsed record.

Parameters

path	Path to a FASTA or FASTQ file (optionally gzip-compressed).
consume	Per-record callback.
fill_fraction	Fraction of free GPU memory for per-chunk staging (default 0.7).
stream	CUDA stream to use.

Returns

Aggregate query summary for the whole file.

See also

query_fastx_records

Definition at line 712 of file filter.cuh.

                                            {}
    ) const {
        return detail::dispatch_fastx_file<Config>(
            path,
            detail::fastx_chunk_mode::query,
            fill_fraction,
            [&](auto& reader, auto dispatch_path) {
                return query_fastx_records_stream(
                    reader,
                    std::string_view{path.native()},
                    consume,
                    fill_fraction,
                    stream,
                    dispatch_path
                );
            }
        );
    }

query_fastx_records()

template<typename Config >

template<FastxRecordConsumer Consumer>

Result< FastxQueryReport > cusbf::filter< Config >::query_fastx_records ( std::istream &  input, Consumer &&  consume, double  fill_fraction = 0.7, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries a FASTA/FASTQ stream and emits one record result per parsed record.

The callback receives record headers, record sequences, aggregate counts, and the per-window hit span for each record as soon as its chunk has been processed. The hit span remains valid only for the duration of the callback.

Parameters

input	Input stream containing FASTA or FASTQ records.
consume	Per-record callback.
fill_fraction	Fraction of available GPU memory for per-chunk staging (default 0.7). Query mode reserves space for both sequence and per-k-mer hit buffers.
stream	CUDA stream to use.

Returns

Aggregate query summary for the whole stream.

Definition at line 684 of file filter.cuh.

                                            {}
    ) const {
        detail::FastxReader reader(input, "<stream>");
        return query_fastx_records_stream(
            reader,
            "<stream>",
            consume,
            fill_fraction,
            stream,
            detail::fastx_dispatch_path::chunked_stream
        );
    }

query_record_batch() [1/2]

template<typename Config >

template<RecordQueryConsumer Consumer>

Result< FastxQueryReport > cusbf::filter< Config >::query_record_batch ( RecordBatchView  batch, Consumer &&  consume, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries a dense host-resident record batch and streams per-record results.

The callback receives one RecordQueryView per input record in source order. The hit span remains valid only for the duration of the callback.

Synchronises before returning.

Parameters

batch	Dense record batch to query.
consume	Per-record callback.
stream	CUDA stream to use.

Returns

Aggregate query summary for the whole batch.

Definition at line 568 of file filter.cuh.

                                            {}
    ) const {
        CUSBF_TRY(
            normalize_record_batch_into<Config>(
                batch,
                fastx_state_.normalized_sequence_scratch(),
                fastx_state_.normalized_records_scratch()
            )
        );
        return query_normalized_record_batch_with_hits(
            batch.sequence, std::forward<Consumer>(consume), stream
        );
    }

query_record_batch() [2/2]

template<typename Config >

Result< FastxQueryReport > cusbf::filter< Config >::query_record_batch ( RecordBatchView  batch, cuda::stream_ref  stream = cudaStream_t{}  ) const

inline

Queries a dense host-resident record batch and returns aggregate counts.

batch.sequence stores raw record payloads back-to-back without separators. batch.records stores ordered, non-overlapping byte ranges into that dense buffer. The filter injects alphabet separators between records internally, so cross-record k-mers are never formed.

Synchronises before returning.

Parameters

batch	Dense record batch to query.
stream	CUDA stream to use.

Returns

Aggregate query summary for the whole batch.

Definition at line 550 of file filter.cuh.

                                                                                  {}) const {
        return query_record_batch_aggregate(batch, stream);
    }

ref()

template<typename Config >

filter_ref< Config > cusbf::filter< Config >::ref ( ) const

inlinenoexcept

Non-owning device reference to this filter's shard storage.

Trivially copyable, intended to be passed by value into CUDA kernels.

Definition at line 137 of file filter.cuh.

                                                          {
        return filter_ref<Config>{
            thrust::raw_pointer_cast(d_shards_.data()),
            num_shards_,
        };
    }

---

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

• config.cuh
• filter.cuh