|
cuDDL
|
Namespaces | |
| namespace | compact_impl |
| namespace | device_gzip |
Classes | |
| struct | bases_packer |
| A base packer and the bases it takes per step, a multiple of sixteen up to 64. More... | |
| struct | blacklist_view |
| Exact membership within the candidate's destination bucket. More... | |
| struct | cardinality_payload |
| Per-thread accumulator feeding the CUB block reduction for cardinality. More... | |
| struct | cardinality_scratch |
| Device-visible scratch for the owning cardinality wrapper. More... | |
| struct | database_file_reader |
| struct | database_file_writer |
| class | database_stager |
| Encodes genome records into register rows through a device arena. More... | |
| struct | decoded_fasta |
| struct | decompression_source |
| Supplies storage for one file, or a null target when none is available. More... | |
| struct | decompression_target |
| Host-visible buffer a caller offers for decompressed bytes. More... | |
| class | device_fasta_pipeline |
| Normalises FASTA batches supplied compressed, uncompressed, or already inflated. More... | |
| struct | fasta_parse_result |
| Result of parsing one FASTA/FASTQ file into packed canonical k-mers. More... | |
| class | fastx_fastq_reader |
| FASTQ parser over a contiguous in-memory buffer. More... | |
| class | fastx_load_pool |
| class | fastx_mapped_file |
| Read-only contiguous file payload for in-memory FASTX parsing. More... | |
| struct | fastx_packed_record |
| One run of bases packed two bits each: its first word and its length in bases. More... | |
| struct | fastx_sequence_extent |
| A contiguous run of sequence bases (record headers and line breaks excluded). More... | |
| struct | fastx_sequence_file |
| struct | fastx_sequence_span |
| One parallel scan span of the combined FASTA sequence stream. More... | |
| class | file_writer |
| Writes a file from a thread of its own, so the caller produces the next bytes while earlier ones reach storage. More... | |
| struct | gzip_file_probe |
| What a gzip file's header and trailer say, or that a file is plain FASTA, read without inflating it. More... | |
| struct | inflated_genome |
| One genome the device inflated: its compacted bases, resident on the device. More... | |
| class | io_ring |
| Stand-in where the build has no io_uring headers: never available. More... | |
| struct | packed_identity |
| Weighted k-mer identity of one packed result, for the GPU conversion. More... | |
| class | path_loaders |
| Loader state a build fed by paths keeps for its whole run. More... | |
| class | pinned_pool_trim |
| Returns the free pages of a device_pinned_pool to the OS when it leaves scope. More... | |
| struct | plane_counts |
| Bit-sliced pairwise counts of one 32-bucket group, accumulated per lane. More... | |
| struct | posting_bitmap_cell |
| Selects posting lists whose bitmap uses no more storage than their reference IDs. More... | |
| class | reference_database_view |
| Non-owning, trivially copyable view of one selected reference-row backing. More... | |
| struct | reference_delta_view |
| Reference rows stored as a base reference plus the buckets that differ from it. More... | |
| struct | reference_deltas |
| struct | reference_file_assembly |
| Builds a reference_database_file from gathered parts; the file's members are private. More... | |
| struct | reference_index_digest |
| struct | refine_cells |
How a refinement-style block splits a tile: cell c holds query group c % groups and the reference words from first_word + (c / groups) * block_words. More... | |
| struct | sequence_batch_chunk |
| One piece of sequence of one genome, and where its bytes are. More... | |
| class | sequence_buffer_pool |
| struct | sequence_genome |
| One genome as its records, in order. No records keeps the genome's ID with a zero row. More... | |
| struct | sequence_record |
| One genome of bases the caller already holds. More... | |
| struct | shared_gzip_decompressor |
| libdeflate handle, one per thread because the object is single-threaded. More... | |
| struct | sketch_scratch |
| Persistent zero-copy scratch used by owning synchronous result wrappers. More... | |
| class | sketch_view |
| Non-owning implementation view of a DDL sketch. More... | |
| struct | span_result |
| One worker's rolling-window state and output. More... | |
| struct | sparse_index_key |
| struct | sparse_reference_id |
| struct | sparse_segment_offset |
| struct | staging_plan |
| Sizing a build resolves before its first genome. More... | |
| struct | summary_payload |
| Per-thread accumulator feeding a CUB block reduction for fused summaries. More... | |
| struct | top_k_row_offset |
Start of genome row i of a collector's heaps, rows of k slots. More... | |
| struct | top_k_slots |
| Each genome's k best neighbours so far: keys and packed counts in parallel rows of k, each row a min-heap of keys whose root is the row's worst key, zero while the row has room. More... | |
| struct | word_popcount |
| Popcount of one pass-bit word. More... | |
| class | worker_pool |
| Fixed threads that run indexed work with the caller, for work repeated per batch. More... | |
Typedefs | |
| using | fastx_line_end_fn = decltype(&fastx_line_end_scalar) |
| using | device_scope = ::cuda::__ensure_current_context |
Makes stream's device current for the enclosing scope and restores the previous context on exit. | |
Enumerations | |
| enum class | decompression_backend : uint8_t { automatic , hybrid , cpu , gpu , gpu_cores } |
| Where inputs are decompressed; sketch construction always runs on the GPU. More... | |
| enum class | transfer_mode : uint8_t { automatic , in_place , pinned , staged } |
| How a build gets CPU-parsed sequence bytes to the device. More... | |
| enum class | fastx_format : uint8_t { unknown , fasta , fastq } |
| Detected record format for a FASTX buffer. More... | |
| enum class | hybrid_variant : uint8_t { bbtools , paper } |
| Which HybridDDL blend the variant kernel computes. More... | |
| enum class | refine_candidates { bitmap , all , reference_mask } |
| Pairs refine_batch_bitmap_kernel compares. More... | |
| enum class | pinned_placement : uint8_t { device_node , host } |
| Where a device_pinned_pool takes its pages. More... | |
Functions | |
| constexpr bool | ablated (std::string_view) noexcept |
| Always false: this build compiles the ablation switches out. | |
| __host__ __device__ double | minimum_mle (double bucket_count, double empty_count, double sum_restored) noexcept |
| Estimates the number of distinct hashes from the minimum hash in each bucket. | |
| __host__ __device__ double | mean_m (double bucket_count, double filled, double sum_restored) noexcept |
| MeanM estimate retained as the full-sketch form of minimum_mle. | |
| __host__ __device__ double | cardinality (double bucket_count, double empty_count, double sum_restored) noexcept |
| Public DDL cardinality estimate. | |
| __host__ __device__ float | minimum_mle_f32 (float bucket_count, float empty_count, float sum_restored) noexcept |
| Single-precision variants of the cardinality estimators. | |
| __host__ __device__ float | mean_m_f32 (float bucket_count, float filled, float sum_restored) noexcept |
| Single-precision mean_m for the throughput-bound cardinality reduction. | |
| __host__ __device__ float | cardinality_f32 (float bucket_count, float empty_count, float sum_restored) noexcept |
| Single-precision cardinality for the throughput-bound cardinality reduction. | |
| __device__ void | classify (pairwise_counts &target, uint32_t this_score, uint32_t other_score) noexcept |
Classifies one score relative to this into target. | |
| template<size_t BucketCount, typename Layout = default_register_layout> | |
| __host__ Result< void > | launch_construction (device_span< uint64_t const > input, device_span< uint32_t > registers, cuda::stream_ref stream, cuda::std::optional< blacklist_view > blacklist=cuda::std::nullopt) |
Launches construction of a sketch through a CTA-local shared register array merged into registers (add_shared_kernel). | |
| bool | device_reads_pageable_memory (cuda::device_ref device) noexcept |
Whether device reads pageable host memory itself. | |
| bool | stages_in_place (transfer_mode mode, cuda::device_ref device) noexcept |
Whether a build in mode stages the caller's bytes in place. | |
| bool | pages_locked (transfer_mode mode, cuda::device_ref device) noexcept |
Whether a build in mode decompresses into page-locked buffers. | |
| Result< size_t > | available_device_bytes (cuda::stream_ref stream) |
Device bytes an allocation on stream can still get. | |
| template<uint32_t K, size_t BucketCount> | |
| Result< staging_plan > | plan_staging (size_t references, uint64_t staged_ceiling, std::optional< size_t > staging_bytes, cuda::stream_ref stream) |
| Sizes the arena and the row store from free device memory and the corpus. | |
| template<size_t BucketCount> | |
| Result< void > | download_store (device_span< uint32_t const > store, std::span< uint16_t > scores, cuda::stream_ref stream) |
| Copies the winner scores of device store rows into host score rows. | |
| decompression_target | acquire_sequence_target (void *context, size_t bytes) |
decompression_source trampoline for sequence_buffer_pool. | |
| template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout, typename Fill , typename Sink > | |
| Result< void > | stage_groups (size_t genomes, uint64_t staged_ceiling, std::optional< size_t > staging_bytes, cuda::stream_ref stream, reference_build_statistics *statistics, Fill &&fill, Sink &&sink, bool in_place=false, cuda::std::optional< blacklist_view > blacklist=cuda::std::nullopt) |
| Runs the staging loop over a collection, filling one group at a time. | |
| template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout, typename Sink > | |
| Result< void > | stage_paths (std::span< std::filesystem::path const > paths, cuda::stream_ref stream, std::optional< size_t > staging_bytes, unsigned parser_workers, transfer_mode transfer, decompression_backend decompression, bool direct_reads, reference_build_statistics *statistics, Sink &&sink, cuda::std::optional< blacklist_view > blacklist=cuda::std::nullopt) |
Stages a path collection, handing each group's device rows to sink. | |
| template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout, typename Sink > | |
| Result< void > | stage_sequences (std::span< sequence_genome const > genomes, cuda::stream_ref stream, std::optional< size_t > staging_bytes, reference_build_statistics *statistics, Sink &&sink, cuda::std::optional< blacklist_view > blacklist=cuda::std::nullopt) |
Stages bases the caller already holds, handing each group's device rows to sink. | |
| uint32_t | bgzf_block_size (unsigned char const *header) noexcept |
A BGZF block's total size from its 18-byte header, or 0 if it is not one: a gzip header with only FEXTRA set, holding just the BC subfield. | |
| uint32_t | crc32_combine (uint32_t a, uint32_t b, uint64_t length) noexcept |
| CRC32 of A followed by B, from CRC32(A), CRC32(B) and B's length, as zlib's crc32_combine: CRC32(A) times x^(8 * length) modulo the reflected polynomial, plus CRC32(B). | |
| template<typename Block > | |
| std::optional< std::pair< uint32_t, uint64_t > > | walk_bgzf (unsigned char const *bytes, uint64_t size, Block &&block) |
Calls block(deflate_at, deflate_bytes, inflated) for each non-empty block of the BGZF file in bytes. | |
| gzip_file_probe | probe_gzip_file (std::filesystem::path const &path) |
Reads the header and trailer of path, and of every block of a BGZF file. | |
| template<typename Work > | |
| void | parallel_for (size_t count, size_t workers, Work &&work) |
Runs work(index) for every index below count on workers threads. | |
| bool | has_gzip_signature_portable (char const *data, size_t size) |
True when data holds what could start another gzip member: 1f 8b 08, then a flag byte with its reserved bits clear. | |
| auto | resolve_gzip_signature_scanner () -> bool(*)(char const *, size_t) |
| Selects the widest signature scanner the host supports. | |
| constexpr uint8_t | encode_base (char base) noexcept |
Maps one DNA byte to its 2-bit symbol (A=0, C=1, T=2, G=3) or 0xff when invalid. | |
| constexpr uint64_t | reverse_complement (uint64_t packed, uint32_t k) noexcept |
| Reverse-complements a packed k-mer (2 bits per base) with bit-parallel reversal. | |
| std::vector< fastx_sequence_span > | fastx_split_sequence_spans (std::vector< fastx_sequence_extent > const &extents, uint32_t seed_bases, uint32_t count) |
Splits the combined FASTA sequence stream into count contiguous spans. | |
| void | consume_byte (char ch, uint32_t k, uint64_t mask, span_result &result, uint64_t &window, uint32_t &window_len) |
| Feeds one byte of sequence through the DDL rolling-window extraction. | |
| Result< fasta_parse_result > | parse_fasta (std::string const &path, uint32_t k, unsigned threads=std::thread::hardware_concurrency()) |
Parses a FASTA/FASTQ file's sequence bases into packed canonical k-mers of length k. | |
| constexpr bool | fastx_is_sequence_whitespace (char ch) noexcept |
| True for bytes that FASTA sequence consumers conventionally skip between bases. | |
| size_t | fastx_line_end_scalar (std::string_view data, size_t position) |
| fastx_line_end_fn | resolve_fastx_line_end () |
| Selects the fastest line-end scanner the host supports. | |
| size_t | fastx_line_end (std::string_view data, size_t position) |
First line-end position at or after position using the resolved scanner. | |
| uint32_t | pack_bases16_portable (char const *bases, size_t count, uint32_t &word) |
Packs up to sixteen bases into a two-bit word, base j in bits 2j and 2j+1. | |
| uint64_t | pack_bases_pieces (char const *bases, size_t count, uint32_t *words, size_t width) |
Packs width bases from sixteen-base pieces, for a tail shorter than a full step. | |
| uint64_t | pack_bases_portable (char const *bases, size_t count, uint32_t *words) |
| bases_packer | resolve_bases_packer () |
| Selects the widest base packer the host supports. | |
| void | pack_fastx_sequence_file (fastx_sequence_file &result, uint32_t k) |
| Packs a parsed file's records in place into two-bit words, run by run. | |
| libdeflate_result | decompress_gzip_member (libdeflate_decompressor *decoder, std::string_view input, char *output, size_t capacity, size_t &consumed, size_t &produced) |
| Decodes one gzip member, checking structure and ISIZE but not checksums. | |
| Result< void > | gunzip_members_into (std::string_view input, std::string &output, std::string const &path) |
Decompresses gzip members from a mapped buffer into output, appending. | |
| Result< size_t > | gunzip_members_direct (std::string_view input, decompression_target target, bool &insufficient, std::string const &path) |
| Decompresses gzip members straight into a caller's buffer. | |
| uint64_t | gzip_decompressed_size (std::string const &path) |
| Uncompressed size the gzip trailer advertises, or 0 when it cannot be read. | |
| size_t | gzip_size_hint (std::string_view input) |
| char * | compact_sequence_whitespace (char const *first, char const *last, char *out) |
Copies bases from [first, last) to out, dropping sequence whitespace. | |
| char const * | fasta_next_header (char const *first, char const *last) |
First '>' in [first, last) that starts a line, or last. | |
| char * | fasta_sequence_compact (char const *first, char const *last, char *out, char const *&stop) |
Compacts one FASTA record's sequence from first to out, stopping at the next header. | |
| std::vector< fastx_sequence_extent > | fastx_fasta_compact_extents (std::string_view data, char *out) |
Collects FASTA sequence extents with whitespace compacted out, in one pass over data. | |
| void | compact_fastx_sequence_extents (fastx_sequence_file &result) |
| Compacts sequence whitespace in place, leaving every extent a run of bases. | |
| Result< void > | parse_fastx_sequence_file (fastx_sequence_file &result, std::string const &path) |
| Result< void > | pack_fasta_chunks (fastx_sequence_file &result, std::string_view data, char *out, uint32_t k, std::string const &path) |
| Compact and pack one cache-sized block at a time, without a genome-sized ASCII write. | |
| Result< std::unique_ptr< fastx_sequence_file > > | finish_fastx_sequence_file (std::unique_ptr< fastx_sequence_file > result, std::string_view data, std::string const &path, bool defer_fasta, uint32_t pack_k) |
Parses, packs or defers the bytes data of a loaded file, which result holds. | |
| Result< std::optional< std::string_view > > | inflate_and_pack (fastx_sequence_file &result, std::string_view input, std::string const &path, decompression_source source, uint32_t k) |
Inflates gzip input in this thread's reusable heap buffer and, for FASTA, packs it there and moves only the packed words into a buffer from source. | |
| Result< std::unique_ptr< fastx_sequence_file > > | decode_fastx_sequence_file (std::unique_ptr< fastx_mapped_file > file, std::string const &path, decompression_source source, bool defer_fasta, uint32_t pack_k=0) |
| Result< std::unique_ptr< fastx_sequence_file > > | load_fastx_sequence_file (std::string const &path, decompression_source source={}, bool defer_fasta=false, bool read=false, uint32_t pack_k=0) |
| __host__ __device__ constexpr uint64_t | splitmix64 (uint64_t value) noexcept |
| SplitMix64 finalizer, applied to a packed k-mer XOR the domain seed. | |
| __host__ __device__ constexpr uint64_t | hash_kmer (uint64_t packed_kmer) noexcept |
| Hashes a canonical packed k-mer into the value space used by DDL registers. | |
| template<size_t BucketCount> | |
| __host__ __device__ constexpr size_t | bucket_of (uint64_t hash) noexcept |
Selects the register bucket for hash given a power-of-two BucketCount. | |
| constexpr __host__ __device__ __forceinline__ uint64_t | rotl64 (uint64_t x, int8_t r) |
Rotates x left by r bits. | |
| template<typename T > | |
| __host__ __device__ __forceinline__ T | load_chunk (const uint8_t *data, uint64_t index) |
Loads a chunk of type T from data at byte offset index*sizeof(T). | |
| constexpr __host__ __device__ __forceinline__ uint64_t | finalize (uint64_t h) |
| Applies the xxHash-64 final mixing (avalanche) step. | |
| __host__ __device__ uint64_t | xxhash64 (const uint8_t *bytes, uint64_t size, uint64_t seed) |
| Computes xxHash-64 over a byte range without allocating or copying it. | |
| template<typename T > | |
| __host__ __device__ uint64_t | xxhash64 (const T &key, uint64_t seed=0) |
| Computes xxHash-64 over the raw bytes of a value. | |
| __device__ double | interpolate_mean_m_cf (double estimate) noexcept |
Interpolates the MeanM correction factor for estimate from the calibration table. | |
| __device__ double | lc_min (uint32_t const *bins, double bucket_count) noexcept |
| Linear-counting estimate from the first occupied NLZ tier. | |
| __device__ double | dlc (uint32_t const *bins, double bucket_count, double lc) noexcept |
| Distributed linear-counting estimate, blended from per-tier estimates by precision. | |
| __device__ double | hybrid (double low_estimate, double zone_estimate, double corrected_mean, double bucket_count) noexcept |
| Log-blended low and corrected estimates across the hybrid transition zone. | |
| __device__ double | bbtools_mean_m (double bucket_count, double filled, double restored_sum) noexcept |
| BBTools occupancy-corrected MeanM estimate from restored magnitudes. | |
| __device__ hybrid_cardinality_estimates | hybrid_estimates (uint32_t const *bins, double bucket_count, double restored_sum) noexcept |
| Fused BBTools, paper, linear-counting, DLC, and raw MeanM estimates. | |
| __device__ float | interpolate_mean_m_cf_f32 (float estimate) noexcept |
| Single-precision variant of the hybrid estimator tail. | |
| __device__ float | lc_min_f32 (uint32_t const *bins, float bucket_count) noexcept |
| Single-precision lc_min for the latency-bound hybrid scan. | |
| __device__ float | dlc_f32 (uint32_t const *bins, float bucket_count, float lc) noexcept |
| Single-precision dlc for the latency-bound hybrid scan. | |
| __device__ float | hybrid_f32 (float low_estimate, float zone_estimate, float corrected_mean, float bucket_count) noexcept |
| Single-precision hybrid for the latency-bound hybrid scan. | |
| __device__ float | bbtools_mean_m_f32 (float bucket_count, float filled, float restored_sum) noexcept |
| Single-precision bbtools_mean_m for the latency-bound hybrid scan. | |
| __device__ hybrid_cardinality_estimates | hybrid_estimates_f32 (uint32_t const *bins, float bucket_count, float restored_sum) noexcept |
| Single-precision hybrid_estimates for the latency-bound hybrid scan. | |
| template<size_t BucketCount, typename Layout = default_register_layout, bool HasBlacklist = false> | |
| __global__ | __launch_bounds__ (shared_construction_block_size) void add_shared_kernel(uint64_t const *input |
| Constructs a sketch through a CTA-local shared-memory register array. | |
| __device__ uint2 | lane_posting_range (uint32_t const *offsets, uint16_t const *sorted_keys, uint32_t reference_count, uint32_t bucket, uint32_t key, uint32_t key_count) |
| One thread's posting range for a bucket/key cell of a dense or sparse index. | |
| __host__ __device__ constexpr uint32_t | sparse_directory_shift (uint32_t reference_count) noexcept |
Low key bits a sparse key directory cell spans for reference_count references. | |
| __host__ __device__ constexpr uint32_t | sparse_directory_entries (uint32_t reference_count) noexcept |
| Bounds per bucket in a sparse key directory. | |
| __global__ | __launch_bounds__ (block_size) static void build_sparse_directory_kernel(uint16_t const *sorted_keys |
Builds a sparse key directory: entry c of bucket b is the first position in b's sorted keys whose key is at least c << shift. | |
| for (auto i=static_cast< size_t >(blockIdx.x) *block_size+threadIdx.x;i< total;i+=static_cast< size_t >(gridDim.x) *block_size) | |
| if (local !=0U) | |
| __device__ uint2 | directory_posting_range (uint16_t const *sorted_keys, uint32_t const *directory, uint32_t reference_count, uint32_t bucket, uint32_t sparse_key) |
| One sparse bucket/key posting range, narrowed through the bucket's key directory. | |
| template<size_t BucketCount, typename QueryRow > | |
| __global__ | __launch_bounds__ (block_size) void count_batch_index_matches_kernel(QueryRow const *queries |
| Counts index matches for every query/reference pair in one tile. | |
| bool | uses_tiled_index_counts (uint32_t reference_count, uint32_t query_count, uint32_t indexed_bucket_count) noexcept |
| Whether a batch counts through count_batch_index_tile_kernel. | |
| uint32_t | index_tile_references (uint32_t reference_count, uint32_t query_count, uint32_t multiprocessors, uint32_t max_shared_bytes) noexcept |
| References per tile, bounded by the device's shared memory and small enough that the batch launches about four blocks per SM. | |
| template<size_t BucketCount> | |
| __global__ void | hybrid_posting_ranges_kernel (uint16_t const *queries, uint32_t const *offsets, uint32_t const *bitmap_cells, uint32_t bitmap_count, uint32_t words, uint32_t indexed_buckets, uint16_t key_mask, uint2 *ranges, uint32_t *bitmap_starts, uint32_t *bitmap_counts) |
| Splits each query's posting ranges into short lists and bitmap offsets. | |
| template<uint32_t CounterBits> | |
| __device__ uint32_t | sum_posting_bitmaps (uint32_t const *bitmaps, uint32_t const *starts, uint32_t count, uint32_t word, uint32_t minimum_matches, uint32_t passed, uint32_t(&counters)[CounterBits]) |
| Sums one query's posting bitmaps over one word of references in bit-sliced counters. | |
| template<size_t BucketCount, uint32_t CounterBits, typename QueryRow > | |
| __global__ | __maxnreg__ (48) void count_batch_index_tile_kernel(QueryRow const *queries |
| Counts index matches for one query against one reference tile in shared memory. | |
| if (candidate_bits !=nullptr &&all_to_all &&high<=query_id_offset+query_index+1U) | |
| if | constexpr (CounterBits !=0U) |
| __syncthreads () | |
| if (candidate_bits==nullptr) | |
| template<typename ReferenceScore > | |
| __device__ __forceinline__ void | classify_query_words (pairwise_counts &target, uint32_t const (&q)[8], ReferenceScore const *reference) noexcept |
| Classifies 16 buckets of an already loaded 16-bit query chunk against one load_256_alignment aligned reference chunk. | |
| template<typename SearchResult > | |
| __device__ void | refine_write_result (uint32_t index, uint32_t query_index, uint32_t reference_id, uint32_t query_id_offset, uint32_t const *match_counts, uint32_t pair_id, pairwise_counts local, uint32_t lane, SearchResult *results, uint32_t *result_match_counts) |
| Warp-reduces one refined candidate and writes its stable result slot. | |
| __host__ __device__ constexpr uint32_t | score_plane_index (uint32_t group, uint32_t plane) noexcept |
Position of plane plane of 32-bucket group group within a row's bit-planes. | |
| __device__ __forceinline__ void | write_score_planes (uint32_t score, uint32_t group, uint32_t lane, uint32_t *planes) |
| Writes the bit-planes of one 32-bucket group whose scores the warp holds, one per lane. | |
| template<size_t BucketCount> | |
| __global__ | __launch_bounds__ (block_size) void build_score_planes_kernel(uint16_t const *rows |
| Transposes compact score rows into the bit-plane layout of score_plane_index. | |
| __device__ __forceinline__ void | load_plane_group (uint4 const *row, uint32_t group, uint32_t(&words)[score_planes]) |
Loads group group of one bit-plane row (score_plane_index layout). | |
| __device__ __forceinline__ uint32_t | plane_any (uint32_t const (&planes)[score_planes]) noexcept |
| Buckets of a plane group holding a nonzero score. | |
| __device__ __forceinline__ uint32_t | plane_low_any (uint32_t const (&planes)[score_planes]) noexcept |
| Buckets of a plane group whose score is nonzero below the top plane. | |
| template<bool CountMatches, bool ReferenceFull = false> | |
| __device__ __forceinline__ void | compare_plane_group (uint32_t const (&q)[score_planes], uint32_t const (&r)[score_planes], uint32_t reference_any, uint32_t reference_low_any, uint16_t key_mask, uint32_t match_lanes, plane_counts &counts) noexcept |
Adds one group's comparison of query planes q with reference planes r, whose nonzero buckets reference_any the caller hoists out of per-query loops. | |
| uint32_t | warp_grid_blocks (size_t warps) noexcept |
Blocks for a grid-stride launch covering warps warps of block_size threads. | |
| __device__ __forceinline__ uint32_t | next_cell (uint32_t *cell_counter, uint32_t cell, uint32_t &claimed) |
The cell a refinement-style block processes after cell: claimed from cell_counter, a zeroed per-launch counter, so blocks balance uneven cells; without one, grid-stride. | |
| template<bool UpperTriangle> | |
| __device__ __forceinline__ refine_cells | make_refine_cells (uint32_t query_id_offset, uint32_t query_count, uint32_t group, uint32_t words_per_query, uint32_t max_block_words, uint32_t waves) |
Splits a tile into refine_cells, about waves cells per block when the tile allows, so that blocks claiming cells dynamically even out cells of uneven work. | |
| template<bool UpperTriangle> | |
| __host__ __device__ uint64_t | batch_result_slot (uint32_t query_index, uint32_t reference_id, uint32_t query_id_offset, uint32_t reference_count) noexcept |
Position of the pair (query_index of the tile, reference_id) in a tile's results: rows of reference_count references per query, or for all-to-all tiles the packed strict upper triangle of pairs with query_id < reference_id, where query IDs start at query_id_offset. | |
| template<size_t BucketCount, bool CountMatches, typename SearchResult , refine_candidates Candidates, bool UpperTriangle, uint32_t QueryGroup> | |
| __global__ | __launch_bounds__ (bitmap_refine_block_size) void refine_batch_bitmap_kernel(uint32_t const *query_planes |
| Exactly compares query bit-planes with references, sharing each reference row across a query group. | |
| __device__ __forceinline__ void | load_plane_chunk (uint4 const *row, uint32_t group, uint32_t chunk, uint32_t(&words)[4]) |
Loads plane chunk chunk (planes 4 chunk ... 4 chunk + 3) of group group of a bit-plane row (score_plane_index layout). | |
| template<size_t BucketCount, bool UpperTriangle, uint32_t QueryGroup> | |
| __global__ void | match_filter_kernel (uint32_t const *query_planes, uint32_t query_id_offset, uint32_t query_count, uint32_t const *reference_planes, uint32_t reference_count, uint32_t indexed_bucket_count, uint32_t minimum_matches, uint32_t *candidate_bits, uint32_t *cell_counter) |
Records in candidate_bits (candidate-bitmap layout) the pairs of the layout whose scores agree in their low match_filter_planes planes in at least minimum_matches buckets. | |
| template<size_t BucketCount, typename SearchResult , bool UpperTriangle, uint32_t QueryGroup> | |
| __global__ void | refine_all_pairs_kernel (uint32_t const *query_planes, uint32_t query_id_offset, uint32_t query_count, uint32_t const *reference_planes, uint32_t reference_count, SearchResult *results, uint32_t *cell_counter, uint32_t const *reference_mask) |
Exactly compares the layout pairs selected by reference_mask, without match counts. | |
| template<size_t BucketCount, typename SearchResult , bool UpperTriangle, bool Swar> | |
| __global__ | __launch_bounds__ (block_size) void row_major_all_pairs_kernel(uint16_t const *rows |
| The all-pairs comparison of refine_all_pairs_kernel over row-major 16-bit scores instead of bit-planes, one warp per pair, for ablation builds to compare with. | |
| template<size_t BucketCount, typename Layout = default_register_layout> | |
| __global__ void | cardinality_kernel (uint32_t const *registers, uint64_t *empty_out, double *estimate_out) |
| Computes the cardinality of a single constructed sketch. | |
| template<size_t BucketCount, typename Layout = default_register_layout> | |
| __global__ void | hybrid_cardinality_kernel (uint32_t const *const registers, hybrid_cardinality_estimates *const estimates) |
| Computes BBTools and paper-style HybridDDL estimates in one register scan. | |
| template<size_t BucketCount, hybrid_variant Variant, typename Layout = default_register_layout> | |
| __global__ void | hybrid_cardinality_variant_kernel (uint32_t const *const registers, double *const estimate) |
| Computes one HybridDDL variant estimate for estimator comparisons. | |
| template<size_t BucketCount, typename Layout = default_register_layout> | |
| __global__ void | batch_cardinality_kernel (uint32_t const *const registers, uint32_t const row_count, uint64_t *const empty_out, double *const estimates_out) |
| Computes the cardinality of every stored sketch in one launch. | |
| template<size_t BucketCount> | |
| __global__ void | batch_scores_kernel (uint32_t const *const registers, uint32_t const row_count, uint16_t *const scores_out) |
| Copies the winning score of every register into compact row-major scores. | |
| cuda::pinned_memory_pool_ref | device_pinned_pool (cuda::device_ref device, pinned_placement placement) |
A page-locked memory pool accessible from device only. | |
| template<typename Layout = default_register_layout> | |
| __host__ __device__ uint16_t | score (uint64_t hash) noexcept |
Computes the 16-bit DDL rarity score for hash. | |
| template<typename Layout = default_register_layout> | |
| __host__ __device__ constexpr uint64_t | restore (uint16_t stored) noexcept |
| Reconstructs the approximate original hash magnitude from a stored score. | |
| template<typename Layout = default_register_layout> | |
| __host__ __device__ constexpr uint64_t | restore_midpoint (uint16_t stored) noexcept |
| Reconstructs the interval midpoint for restore. | |
| __device__ __forceinline__ uint32_t | load_shared_relaxed (uint32_t const *address) |
| Relaxed CTA-scope load of a shared word that other threads update atomically. | |
| template<size_t BucketCount, typename Layout , bool HasBlacklist = false, bool Naive = false> | |
| __device__ __forceinline__ void | add_sequence_windows (char const *sequence, uint32_t windows, size_t first_tile, size_t tile_stride, uint32_t k, uint32_t *registers, bool packed_input=false, blacklist_view blacklist={}) |
| Encodes raw ASCII windows into registers using shared overlapping tiles. | |
| template<size_t BucketCount, typename Layout , bool HasBlacklist = false, bool Naive = false> | |
| __global__ void | add_sequence_batch_kernel (sequence_batch_chunk const *chunks, size_t chunk_count, size_t block_count, uint32_t k, uint32_t *registers, blacklist_view blacklist={}) |
| Batch kernel covering every staged chunk in one launch, including short records. | |
| template<size_t BucketCount, typename Layout , bool HasBlacklist = false> | |
| __global__ void | add_sequence_single_kernel (char const *sequence, uint32_t windows, uint32_t k, uint32_t *registers, blacklist_view blacklist={}) |
| Single-sequence kernel for the public raw-ASCII API. | |
| template<size_t BucketCount, typename Layout = default_register_layout> | |
| __host__ Result< void > | launch_sequence_add (device_span< char const > sequence, uint32_t k, device_span< uint32_t > registers, cuda::stream_ref stream, cuda::std::optional< blacklist_view > blacklist=cuda::std::nullopt) |
| Shared launch for one device-resident ASCII chunk. | |
| Result< void > | mpi_try (int status, char const *call) |
| Result< void > | agree (Result< void > local, MPI_Comm comm) |
Returns local on the rank that failed, and on every other rank an error naming the lowest failing rank, so that all ranks leave a collective step together. | |
| Result< std::vector< std::filesystem::path > > | share_listing (std::filesystem::path const &input, int rank, MPI_Comm comm) |
Rank 0's listing of input on every rank. | |
| size_t | slice_begin (std::span< uint64_t const > bytes, uint64_t total, int rank, int size) |
First genome of rank rank's slice: contiguous slices of about equal bytes. | |
| Result< std::vector< uint32_t > > | slice_counts (reference_database_file const &mine, genome_slice const &slice, MPI_Comm comm) |
Every rank's genome count, after checking collectively that mine holds the sketches of this rank's slice and that every rank built with rank 0's settings. | |
| Result< MPI_Datatype > | row_type (uint32_t buckets) |
One row of buckets scores as a single MPI element, so counts of rows stay within int. | |
| template<typename Search > | |
| Result< void > | shared_search (MPI_Comm comm, Search &&search) |
Runs search with a replica_search that shares tiles across the ranks of comm. | |
| template<typename T > | |
| T | try_unwrap_success (Result< T > &result) |
| void | try_unwrap_success (Result< void > &result) |
| cuda::std::unexpected< Error > | propagate_error (const Error &error) |
Copies error for propagation (avoids moving out of expected::error()). | |
| __host__ __device__ constexpr uint32_t | load_256_alignment () noexcept |
| Byte alignment load_256_global_nc requires on the compiled architecture. | |
| template<typename QueryScore , typename ReferenceScore > | |
| __host__ __device__ bool | wide_rows_aligned (QueryScore const *query, ReferenceScore const *reference) noexcept |
| True when both row bases meet load_256_alignment on the compiled architecture. | |
| template<typename T > | |
| __device__ __forceinline__ void | load_256_global_nc (const T *ptr, T *out) |
| Loads 256 bits from global memory using the non-coherent cache path. | |
| template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout> | |
| Result< void > | validate_score_compatibility (score_compatibility const &compatibility) |
| template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout> | |
| Result< void > | validate_non_indexed_score_compatibility (score_compatibility const &compatibility) |
| template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout> | |
| Result< void > | validate_indexed_score_compatibility (score_compatibility const &compatibility) |
| constexpr uint64_t | indexed_cell_count (score_compatibility const &compatibility) noexcept |
| constexpr uint64_t | indexed_posting_count (uint32_t reference_count, score_compatibility const &compatibility) noexcept |
| constexpr index_storage | smaller_index_storage (uint32_t reference_count, score_compatibility const &compatibility) noexcept |
Storage needing fewer device bytes for reference_count references; dense on ties. | |
| uintptr_t | align_up (uintptr_t address, size_t alignment) noexcept |
| constexpr uint32_t | batch_query_tile_size (uint32_t reference_count, uint32_t query_count, uint32_t tile_limit=batch_query_tile_count) noexcept |
| constexpr uint32_t | all_to_all_tile_queries (uint32_t reference_count, uint32_t parts=1U) noexcept |
| Query rows one all-to-all tile compares: about 2^25 pairs, and at least batch_query_tile_count. | |
| uint32_t | external_batch_query_limit (uint32_t reference_count, cuda::device_ref device) |
Query rows one external batch tile may hold on device. | |
| template<typename T > | |
| T | database_file_little_endian (T value) |
| template<typename IO > | |
| Result< void > | database_file_metadata (IO &io, reference_database_metadata &metadata, uint32_t &blacklist_count) |
| Result< uint64_t > | reference_database_digest (reference_database_file const &db) |
| __host__ __device__ constexpr uint64_t | top_k_key (pairwise_counts const &counts, uint32_t id) noexcept |
| Ranking key of a neighbour: its weighted k-mer identity above its complemented ID. | |
| constexpr size_t | top_k_heap_words (uint32_t k) noexcept |
| Words of shared memory one heap takes: k keys, then their k packed counts. | |
| __device__ void | top_k_replace_root (uint64_t *heap, uint32_t k, uint64_t key, uint64_t value) |
Replaces the root of a min-heap of k keys at heap, whose values follow the keys, with key and value, and sifts it down. | |
| template<bool Thresholded, bool UpperTriangle> | |
| __device__ uint32_t | top_k_word_bits (batch_result_tile const &tile, uint32_t query, uint32_t word, uint32_t words) |
The pass bits of word word of tile row query: the stored bits, or for exhaustive tiles every reference of the word. | |
| template<bool UpperTriangle> | |
| __device__ packed_pairwise_counts | top_k_stored (batch_result_tile const &tile, uint32_t query, uint32_t reference) |
The stored result of tile row query and reference, which must pass. | |
| __device__ uint32_t | top_k_nth_bit (uint32_t mask, uint32_t n) |
Position of set bit n, counting from zero, of mask, which has more than n set bits. | |
| template<typename Bits , typename Entry > | |
| __device__ void | top_k_warp_merge (uint64_t *heap, uint32_t k, uint32_t words, Bits bits_of, Entry entry_of) |
One warp merges the candidates of words bitmap words into the heap of k entries at heap, in shared memory, which only this warp updates, so it needs no lock. | |
| template<bool ToShared> | |
| __device__ void | top_k_warp_copy (top_k_slots const &slots, size_t row, uint64_t *heap) |
Copies the k entries of global heap row into, or for ToShared false out of, the shared heap heap with the warp's lanes. | |
| if (query >=tile.query_count) return | |
| top_k_warp_copy< true > (slots, row, heap) | |
| top_k_warp_merge (heap, slots.k, words - cuda::std::min(skipped, words), [&](uint32_t word) { return top_k_word_bits< Thresholded, UpperTriangle >(tile, query, skipped+word, words);}, [&](uint32_t candidate, uint64_t &key, uint64_t &value) { auto const reference=skipped *32U+candidate;auto const stored=top_k_stored< UpperTriangle >(tile, query, reference);key=top_k_key(stored.unpack(tile.bucket_count), reference);value=stored.bits;}) | |
| top_k_warp_copy< false > (slots, row, heap) | |
| if (block >=static_cast< uint64_t >(row_words) *word_count) return | |
| _Pragma ("unroll") for(uint32_t bit=0U | |
Variables | |
| constexpr std::array< std::string_view, 22 > | ablations |
Optimisations that the CUDDL_ABLATE environment variable can turn off in a build configured with -Dablations=true (CUDDL_ABLATIONS=1). | |
| constexpr bool | ablations_enabled = false |
| constexpr double | hash_range = 0x1p64 |
2^64, the range of the 64-bit hash used by the sketch. | |
| constexpr size_t | default_arena_ceiling = size_t{128} << 20 |
| Default arena ceiling. | |
| constexpr size_t | stager_arenas = 4 |
| Staging arenas the copy stream rotates through, so copies run ahead of the kernels that read earlier arenas while other work shares the device. | |
| constexpr size_t | stager_hold_slots = 16 |
| Host memory the device still has to read must stay alive: the transfer engine's read of a pageable buffer is not ordered with the host writes that follow it. | |
| constexpr size_t | packed_buffer_bytes = size_t{2} << 20 |
| Bytes reserved per loader buffer that holds packed bases: a bacterial genome's two-bit words, about 1.5 MB, fit. | |
| constexpr uint32_t | missing_cell = 0x00FFFFFFU |
| Eight ambiguous bases: an encoder cell past the end of a sequence. | |
| constexpr size_t | parallel_read_depth = 64 |
| Reads a parallel file system needs in flight to give one node its bandwidth: on MOGON's Lustre, 16 threads read 2.9 GB/s of ~1 MB files and 64 threads 6.7 GB/s. | |
| constexpr uint64_t | seed = 42 |
| Domain separation constant folded into every packed k-mer before hashing. | |
| constexpr uint64_t | PRIME64_1 = 11400714785074694791ULL |
| xxHash-64 round primes. | |
| constexpr uint64_t | PRIME64_2 = 14029467366897019727ULL |
| constexpr uint64_t | PRIME64_3 = 1609587929392839161ULL |
| constexpr uint64_t | PRIME64_4 = 9650029242287828579ULL |
| constexpr uint64_t | PRIME64_5 = 2870177450012600261ULL |
| constexpr int32_t | nlz_bins = 66 |
| NLZ histogram bins for the hybrid estimator, including the empty tier. | |
| constexpr double | dlc_blend_low = 0.088 |
| constexpr double | dlc_blend_high = 0.952 |
| constexpr double | dlc_min_fraction = 0.002 |
| constexpr double | dlc_info_power = 4.5 |
| constexpr double | hybrid_low = 0.2 |
| constexpr double | hybrid_high = 5.0 |
| constexpr uint32_t | ddl_cf_size = 13 |
| constexpr uint32_t | block_size = 256 |
| Threads per CTA for the single-pair and cardinality reduction kernels. | |
| constexpr size_t | maximum_batch_rows = (size_t{1} << 31) - 1U |
| Stored sketches one batch launch may index, bounded by the CUDA grid dimension. | |
| template<typename QueryScore , typename ReferenceScore > | |
| constexpr uint32_t | wide_chunk_buckets |
| Buckets one contiguous chunk covers, for a query/reference score pair. | |
| constexpr uint32_t | shared_construction_block_size = 768 |
| Threads per CTA for the CTA-local construction kernel. | |
| __global__ size_t | input_size |
| __global__ size_t uint32_t * | registers |
| __global__ size_t uint32_t bool | vector_input |
| __global__ size_t uint32_t bool blacklist_view | blacklist |
| __global__ size_t uint32_t bool blacklist_view bool | IncludeCardinality |
| __global__ size_t uint32_t bool blacklist_view bool | Layout |
| constexpr uint32_t | index_match_cells_per_warp = 32U |
| Cells one warp owns per iteration of count_batch_index_matches_kernel: one per lane. | |
| __global__ uint32_t | reference_count |
| __global__ uint32_t uint32_t | indexed_bucket_count |
| __global__ uint32_t uint32_t uint32_t * | directory |
| auto const | shift = sparse_directory_shift(reference_count) |
| auto const | total = static_cast<size_t>(indexed_bucket_count) * entries |
| __global__ uint16_t const * | sorted_keys |
| __global__ uint16_t const uint32_t uint32_t uint32_t | key_count = static_cast<uint32_t>(key_mask) + 1U |
| __global__ uint16_t const uint32_t uint32_t uint32_t unsigned long long * | work |
| unsigned long long | local = 0U |
| constexpr uint32_t | posting_range_block_size = 1024U |
| Threads of the single index_posting_ranges_kernel block. | |
| static __global__ uint32_t const * | offsets |
| static __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t | key_mask |
| static __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t uint32_t * | range_begin |
| static __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t uint32_t uint32_t * | range_end_scan |
| __shared__ block_scan::TempStorage | storage |
| uint32_t | carry = 0U |
| static __global__ uint32_t const uint32_t uint32_t const * | postings |
| static __global__ uint32_t const uint32_t uint32_t const uint32_t * | match_counts |
| auto const | stride = static_cast<uint32_t>(gridDim.x) * block_size |
| __global__ size_t | query_row_offset |
| __global__ size_t uint32_t | query_count |
| constexpr uint32_t | index_tile_lane_postings = 16U |
| Longest posting range one count_batch_index_tile_kernel lane walks alone. | |
| constexpr uint32_t | index_tile_long_capacity = 32U |
| Long posting ranges one count_batch_index_tile_kernel block walks together. | |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t | tile |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const * | ranges |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t | minimum_matches |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t | query_id_offset |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool | all_to_all |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t * | candidate_bits |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t uint32_t const * | bitmaps |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t uint32_t const uint32_t const * | bitmap_starts |
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t uint32_t const uint32_t const uint32_t const * | bitmap_counts |
| constexpr uint32_t | warps_per_block = index_tile_block_size / warp_width |
| __shared__ uint32_t | packed_counts [] |
| auto const | tiles = (reference_count + tile - 1U) / tile |
| auto const | query_index = static_cast<uint32_t>(blockIdx.x) / tiles |
| auto | low = static_cast<uint32_t>(blockIdx.x) % tiles * tile |
| auto const | high = min(low + tile, reference_count) |
| auto const | width = high - low |
| __shared__ uint2 | long_ranges [index_tile_long_capacity] |
| __shared__ uint32_t | long_count |
| auto const | lane = static_cast<uint32_t>(threadIdx.x) % warp_width |
| auto const | warp = static_cast<uint32_t>(threadIdx.x) / warp_width |
| else | |
| auto const * | query = queries + (query_row_offset + query_index) * BucketCount |
| auto const | count |
| auto const | long_ranges_used |
| auto *const | bits |
| constexpr size_t | refine_block_row_bytes = size_t{8} << 20U |
| Reference-row bytes one refine_batch_index_candidates_kernel block spans. | |
| constexpr uint32_t | score_planes = 16U |
| Bits per score, and so bit-planes per 32-bucket group. | |
| __global__ uint32_t | row_count |
| __global__ uint32_t uint32_t * | planes |
| constexpr uint32_t | groups_per_row = static_cast<uint32_t>(BucketCount / warp_width) |
| __global__ size_t uint16_t * | rows |
| __global__ uint32_t const * | reference_planes |
| __global__ uint32_t const uint32_t uint32_t const * | candidate_ids |
| __global__ uint32_t const uint32_t uint32_t const uint32_t const * | candidate_count |
| __global__ uint32_t const uint32_t uint32_t const uint32_t const SearchResult * | results |
| constexpr uint32_t | row_words = static_cast<uint32_t>(BucketCount / 2U) |
| auto const | items = candidate_ids == nullptr ? reference_count : *candidate_count |
| auto const | pairs = static_cast<uint64_t>(query_count) * reference_count |
| auto const | indexed_groups = indexed_bucket_count / warp_width |
| constexpr uint32_t | bitmap_refine_block_size = 768U |
| Threads per refine_batch_bitmap_kernel block. | |
| constexpr uint32_t | bitmap_refine_waves = 8U |
| Cells per refine_batch_bitmap_kernel block that a tile of bitmap-selected pairs is split into when it allows. | |
| constexpr size_t | bitmap_refine_shared_bytes = 96U * 1024U |
| Shared bytes refine_batch_bitmap_kernel may stage query planes in; one block per SM fits on every supported architecture. | |
| template<size_t BucketCount> | |
| constexpr bool | bitmap_refine_staged = bitmap_refine_shared_bytes >= BucketCount * sizeof(uint16_t) |
| Whether a query group's planes fit bitmap_refine_shared_bytes. | |
| template<size_t BucketCount> | |
| constexpr uint32_t | bitmap_refine_group |
| Queries one refine_batch_bitmap_kernel block refines together. | |
| template<size_t BucketCount> | |
| constexpr size_t | bitmap_refine_dynamic_bytes |
| Shared bytes one refine_batch_bitmap_kernel block stages. | |
| __global__ uint32_t const uint32_t const * | pass_bits |
| __global__ uint32_t const uint32_t const uint32_t const * | word_offsets |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t | first_query_id |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t | word_begin |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t | word_end |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t | packed_base |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t SearchResult * | packed |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t SearchResult uint32_t * | packed_match_counts |
| auto const | words_per_query = candidate_bit_words(reference_count) |
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const SearchResult uint32_t * | result_match_counts |
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const SearchResult uint32_t uint32_t uint32_t uint32_t * | cell_counter |
| constexpr uint32_t | group = QueryGroup |
| constexpr bool | staged = bitmap_refine_staged<BucketCount> |
| constexpr uint32_t | lane_groups = groups_per_row / warp_width |
| constexpr uint32_t | cached_groups = lane_groups < 2U ? lane_groups : 2U |
| constexpr uint32_t | max_block_words |
| constexpr uint32_t | word_slices = Candidates == refine_candidates::all ? 1U : 4U |
| __shared__ uint4 | query_plane_storage [] |
| __shared__ uint32_t | next_word |
| __shared__ uint32_t | claimed_cell |
| auto const | cells |
| auto const | groups = cells.groups |
| auto const | thresholded = CountMatches && pass_bits != nullptr |
| constexpr uint32_t | match_filter_block_size = 512U |
| Threads per match_filter_kernel block. | |
| template<size_t BucketCount> | |
| constexpr bool | match_filter_supported |
Whether match_filter_kernel supports BucketCount: a lane caches at most two plane groups of each of its references. | |
| constexpr uint32_t | match_filter_planes = 12U |
| Low planes match_filter_kernel compares. | |
| constexpr uint32_t | all_pairs_block_size = 512U |
| Threads per refine_all_pairs_kernel block. | |
| constexpr uint32_t | all_pairs_word_parts = 4U |
| Parts a refine_all_pairs_kernel warp claims each 32-reference word in. | |
| __global__ uint32_t uint32_t uint32_t SearchResult uint32_t const * | reference_mask |
| auto const | warps = uint64_t{gridDim.x} * (block_size / warp_width) |
| constexpr uint32_t | delta_change_limit = 160U |
| Most buckets in which a child reference may differ from its base reference. | |
| constexpr uint32_t | delta_signatures = 8U |
| Exact bucket signatures that propose base candidates, and buckets per signature. | |
| constexpr uint32_t | delta_signature_buckets = 8U |
| constexpr uint32_t | delta_signature_candidates = 4U |
| Highest-ID rows taken from each matching signature. | |
| constexpr uint32_t | delta_candidates = delta_signatures * delta_signature_candidates |
| constexpr uint32_t | delta_correction_warps = 24U |
| Warps per apply_reference_deltas_kernel block, and the most and fewest children one block cell takes. | |
| constexpr uint32_t | delta_child_tile = 2048U |
| constexpr uint32_t | delta_min_child_tile = 4U * delta_correction_warps |
| template<size_t BucketCount> | |
| constexpr uint32_t | delta_query_group = static_cast<uint32_t>(32768U / BucketCount) |
| Queries one correction cell stages bucket-major: 64 KiB of scores for every bucket count. | |
| template<size_t BucketCount> | |
| constexpr size_t | delta_correction_shared_bytes |
| __global__ uint32_t uint32_t | signature |
| __global__ uint32_t uint32_t uint64_t * | keys |
| __global__ uint32_t uint32_t uint64_t uint32_t * | ids |
| constexpr uint8_t | mantissa_bits = default_register_layout::mantissa_bits |
| Default score layout retained for source compatibility. | |
| constexpr uint32_t | mantissa_mask = (1U << mantissa_bits) - 1U |
| constexpr uint32_t | restore_shift_base = 63U - mantissa_bits |
| Default restoration constants retained for source compatibility. | |
| constexpr uint32_t | max_nlz = restore_shift_base |
| constexpr double | index_pair_fraction_limit = 0.03 |
| Index pair fraction (reference_index::pair_fraction) above which batch search compares bit-planes for every pair instead of counting through the index. | |
| constexpr uint32_t | batch_query_tile_count = 128U |
| Minimum query rows per batch tile. | |
| constexpr uint32_t | top_k_warps = block_size / 32U |
| Warps per merging block, each holding one heap; fewer when the heaps exceed shared memory. | |
| __global__ top_k_slots | slots |
| auto *const | heap = heaps + warp * top_k_heap_words(slots.k) |
| auto const | query_id = tile.first_query_id + query |
| auto const | row = static_cast<size_t>(query_id - slots.first_genome_id) |
| auto const | words = tile.words_per_query() |
| auto const | skipped = UpperTriangle ? (query_id + 1U) / 32U : 0U |
| __global__ uint32_t * | columns |
| __global__ uint32_t uint32_t | first_word |
| auto const | block = static_cast<uint64_t>(blockIdx.x) * (blockDim.x / 32U) + threadIdx.x / 32U |
| auto const | word_count = words - first_word |
| auto const | word = first_word + static_cast<uint32_t>(block % word_count) |
| auto const | row_word = static_cast<uint32_t>(block / word_count) |
| uint32_t | column = 0U |
| using cuddl::detail::device_scope = typedef ::cuda::__ensure_current_context |
Makes stream's device current for the enclosing scope and restores the previous context on exit.
Every public function that takes a stream opens one, so its kernel launches, CUB calls and function attributes target the stream's device, never the calling thread's current device. CCCL provides this only as an internal type.
Definition at line 55 of file reference_database.cuh.
| using cuddl::detail::fastx_line_end_fn = typedef decltype(&fastx_line_end_scalar) |
Definition at line 175 of file fastx_sequence_file.hpp.
|
strong |
Where inputs are decompressed; sketch construction always runs on the GPU.
Definition at line 46 of file database_staging.cuh.
|
strong |
Detected record format for a FASTX buffer.
| Enumerator | |
|---|---|
| unknown | Format not yet determined from the first header. |
| fasta | FASTA ( |
| fastq | FASTQ ( |
Definition at line 438 of file fastx_sequence_file.hpp.
|
strong |
Which HybridDDL blend the variant kernel computes.
| Enumerator | |
|---|---|
| bbtools | |
| paper | |
Definition at line 21 of file detail/hybrid_cardinality.cuh.
|
strong |
Where a device_pinned_pool takes its pages.
Definition at line 19 of file pinned_pool.cuh.
|
strong |
Pairs refine_batch_bitmap_kernel compares.
Definition at line 1737 of file kernels.cuh.
|
strong |
How a build gets CPU-parsed sequence bytes to the device.
Definition at line 60 of file database_staging.cuh.
| __global__ cuddl::detail::__launch_bounds__ | ( | bitmap_refine_block_size | ) | const |
Exactly compares query bit-planes with references, sharing each reference row across a query group.
Block cells are (reference block, group of bitmap_refine_group queries), reference-block major so resident blocks share reference rows through L2. A block stages its group's query bit-planes (score_plane_index) in shared memory when they fit; warps then claim 32-reference words and, for every reference any group query selects, compare the reference's bit-planes with each selecting query. A row thus crosses L2 once per group rather than once per pair.
Each written result lands at its batch_result_slot. Candidates selects which pairs are written: those in candidate_bits (one word per query and 32 references), or every pair of the layout. With pass_bits (counting kernels only), only the selected pairs with at least minimum_matches index matches are written and recorded in pass_bits, in the candidate-bitmap layout. In bitmap mode pass_bits must be candidate_bits, whose failing pairs' bits are cleared in place. With result_match_counts, results also get their index match counts, recomputed from the planes.
| __global__ cuddl::detail::__launch_bounds__ | ( | block_size | ) | const |
Builds a sparse key directory: entry c of bucket b is the first position in b's sorted keys whose key is at least c << shift.
Marks the first position of every run of equal sorted keys with its index.
Turns per-pair match counts into the query-major candidate bitmap.
Counts the query's posting matches for every reference, one posting per thread.
Adds the squared length of every nonzero-key posting list to work.
Dense indexes pass their cell offsets (key_count cells per bucket); sparse indexes pass their per-bucket sorted keys instead, whose runs are the posting lists. The sum is the number of (reference, reference, bucket) triples sharing an index key: an index walks about work / reference_count postings per database-like query.
| __global__ cuddl::detail::__launch_bounds__ | ( | block_size | ) | const |
Transposes compact score rows into the bit-plane layout of score_plane_index.
Transposes an all-to-all tile's pass bits from first_word on into per-genome columns: word w * reference_count + g of columns holds which of tile rows 32 w to 32 w + 31 pass with genome g.
Hashes signature signature of every row; rows are emitted in descending ID order so a stable sort keeps the highest IDs first within a signature.
Counts every query/reference pair's index matches from bit-planes, one warp per pair.
Recovers row-major scores from bit-plane rows, one warp per 32-bucket group.
Used when a search asks for index semantics without an index: the counts equal the ones the index postings would produce.
One warp per 32-by-32 bit block, which its lanes load a row each and turn with 32 ballots; neighbouring warps take neighbouring words of the same rows, so they share the rows' sectors, and every warp writes 32 adjacent words.
| __global__ cuddl::detail::__launch_bounds__ | ( | block_size | ) | const |
Counts index matches for every query/reference pair in one tile.
One warp per tile row merges the row's passing pairs into the row genome's heap.
Packs a threshold tile's passing results in order.
Compares one query's bit-planes with every reference, or with the selected candidates.
Resolves every query's sparse posting range for one indexed bucket per block.
Each lane of a warp resolves the posting range of its own (query, bucket) cell, so a warp keeps 32 independent range lookups in flight. The warp then walks every non-empty range with a lane stride, so hot keys with long lists (the dominant cost on skewed rows) are consumed 32 postings at a time.
Block b owns bucket b, so its lookups share the bucket's directory and keys through L1. ranges is query-major, one entry per (query, indexed bucket).
One warp compares one reference; lane l takes groups l, l + 32, ... With candidate_ids, item i refines reference candidate_ids[i] for *candidate_count items; otherwise item i is reference i.
One warp per pass-bit word in [word_begin, word_end); lane l handles reference 32 * word + l, so consecutive lanes read consecutive slots and write consecutive packed positions. word_offsets is the exclusive scan of the word popcounts; results land at their packed position minus packed_base.
| __global__ cuddl::detail::__launch_bounds__ | ( | block_size | ) | const |
The all-pairs comparison of refine_all_pairs_kernel over row-major 16-bit scores instead of bit-planes, one warp per pair, for ablation builds to compare with.
With Swar each lane compares 16 buckets per step from 256-bit loads, two packed scores per 32-bit subtraction (classify_query_words); without, it classifies one bucket at a time. rows holds every reference's row, and query q of the tile is row query_id_offset + q.
| __global__ cuddl::detail::__launch_bounds__ | ( | shared_construction_block_size | ) | const |
Constructs a sketch through a CTA-local shared-memory register array.
Each item costs a hash, a score, and one fire-and-forget shared atomic max; the CTA then merges its registers into the global sketch with one global atomic max per nonempty bucket.
The host launches two CTAs per SM and the kernel walks the input with a runtime grid-stride loop over 256-bit chunks, so the grid is a single balanced wave for every input size and the per-CTA merge traffic stays minimal.
| __global__ cuddl::detail::__maxnreg__ | ( | 48 | ) | const |
Counts index matches for one query against one reference tile in shared memory.
Block b owns query b / tiles and references [t * tile, (t + 1) * tile) for t = b % tiles, with two 16-bit counters per shared word. A query's tiles run as neighbouring blocks, so they share its posting-range lookups through the caches. Postings ascend by reference within each key, so each posting list is narrowed to the tile with two binary searches. The flush writes every pair of the tile, so match_counts needs no zeroing. With nonzero CounterBits, each lane first sums the query's posting bitmaps over one word of the tile in that many bit-sliced counters, seeding references that already reach minimum_matches with exactly that count; the posting lists then add the rest.
Launched with index_tile_block_size threads: the shared counters, not the thread count, bound how many blocks fit on an SM, so larger blocks keep more warps resident to hide the posting loads' latency. Capping registers at 48 leaves a quarter of the SM's registers to kernels of neighbouring tiles; a GH200 searched 115,223 genomes 12% faster than with 60.
| cuddl::detail::__syncthreads | ( | ) |
|
pure virtual |
|
constexprnoexcept |
Always false: this build compiles the ablation switches out.
Definition at line 82 of file ablation.hpp.
|
inline |
decompression_source trampoline for sequence_buffer_pool.
Definition at line 1012 of file database_staging.cuh.
| __global__ void cuddl::detail::add_sequence_batch_kernel | ( | sequence_batch_chunk const * | chunks, |
| size_t | chunk_count, | ||
| size_t | block_count, | ||
| uint32_t | k, | ||
| uint32_t * | registers, | ||
| blacklist_view | blacklist = {} |
||
| ) |
Batch kernel covering every staged chunk in one launch, including short records.
The register cap keeps the blacklist lookup from costing resident blocks: at 43 registers, a GH200 SM held five blocks instead of eight.
Definition at line 304 of file sequence_encode.cuh.
| __global__ void cuddl::detail::add_sequence_single_kernel | ( | char const * | sequence, |
| uint32_t | windows, | ||
| uint32_t | k, | ||
| uint32_t * | registers, | ||
| blacklist_view | blacklist = {} |
||
| ) |
Single-sequence kernel for the public raw-ASCII API.
There is no cross-chunk carry: the caller supplies any K-1 overlap explicitly.
Definition at line 350 of file sequence_encode.cuh.
| __device__ __forceinline__ void cuddl::detail::add_sequence_windows | ( | char const * | sequence, |
| uint32_t | windows, | ||
| size_t | first_tile, | ||
| size_t | tile_stride, | ||
| uint32_t | k, | ||
| uint32_t * | registers, | ||
| bool | packed_input = false, |
||
| blacklist_view | blacklist = {} |
||
| ) |
Encodes raw ASCII windows into registers using shared overlapping tiles.
Each shared word holds eight two-bit bases and eight ambiguity bits. Threads reuse neighboring words to construct eight overlapping windows on both strands. Naive instead builds every window base by base with a branch per base, for ablation studies only.
Definition at line 44 of file sequence_encode.cuh.
Returns local on the rank that failed, and on every other rank an error naming the lowest failing rank, so that all ranks leave a collective step together.
Definition at line 75 of file distributed.cuh.
|
inlinenoexcept |
Definition at line 944 of file reference_database.cuh.
|
constexprnoexcept |
Query rows one all-to-all tile compares: about 2^25 pairs, and at least batch_query_tile_count.
Enough query groups to fill the GPU for each launch, while result storage holds two tiles so one is searched as the other is consumed.
Split into parts shards, tiles shrink until each part gets about four, so the alternating assignment evens out the triangle. They shrink only when the whole search already spans several tiles: its two-slot result storage then holds the smaller tiles too.
Definition at line 977 of file reference_database.cuh.
|
inline |
Device bytes an allocation on stream can still get.
Storage the default pool caches from earlier allocations is reused by the next one, so it counts as available alongside what the device reports free.
Definition at line 127 of file database_staging.cuh.
| __global__ void cuddl::detail::batch_cardinality_kernel | ( | uint32_t const *const | registers, |
| uint32_t const | row_count, | ||
| uint64_t *const | empty_out, | ||
| double *const | estimates_out | ||
| ) |
Computes the cardinality of every stored sketch in one launch.
One block reduces one row of the row-major store of BucketCount registers per sketch.
Definition at line 2649 of file kernels.cuh.
|
constexprnoexcept |
Definition at line 957 of file reference_database.cuh.
|
inlinenoexcept |
Position of the pair (query_index of the tile, reference_id) in a tile's results: rows of reference_count references per query, or for all-to-all tiles the packed strict upper triangle of pairs with query_id < reference_id, where query IDs start at query_id_offset.
Definition at line 1751 of file kernels.cuh.
| __global__ void cuddl::detail::batch_scores_kernel | ( | uint32_t const *const | registers, |
| uint32_t const | row_count, | ||
| uint16_t *const | scores_out | ||
| ) |
Copies the winning score of every register into compact row-major scores.
Definition at line 2683 of file kernels.cuh.
|
inlinenoexcept |
BBTools occupancy-corrected MeanM estimate from restored magnitudes.
Definition at line 163 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
Single-precision bbtools_mean_m for the latency-bound hybrid scan.
Definition at line 305 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
A BGZF block's total size from its 18-byte header, or 0 if it is not one: a gzip header with only FEXTRA set, holding just the BC subfield.
Definition at line 76 of file device_gzip.cuh.
|
constexprnoexcept |
|
inlinenoexcept |
Public DDL cardinality estimate.
Minimum MLE is the current production estimator. Keeping this dispatch point makes changing the estimator explicit without exposing implementation-specific cardinality APIs.
Definition at line 59 of file cardinality.cuh.
|
inlinenoexcept |
Single-precision cardinality for the throughput-bound cardinality reduction.
Definition at line 91 of file cardinality.cuh.
| __global__ void cuddl::detail::cardinality_kernel | ( | uint32_t const * | registers, |
| uint64_t * | empty_out, | ||
| double * | estimate_out | ||
| ) |
Computes the cardinality of a single constructed sketch.
empty_out receives the empty-register count; estimate_out receives the estimate.
Definition at line 2551 of file kernels.cuh.
|
inlinenoexcept |
Classifies one score relative to this into target.
Takes 16-bit scores and 32-bit registers alike. Branchless: all four outcomes are computed as predicated arithmetic so a warp with mixed comparisons (random sketch data) never pays branch-reconvergence costs. Both-empty is the equality outcome restricted to zero scores.
Definition at line 16 of file comparison.cuh.
|
noexcept |
Classifies 16 buckets of an already loaded 16-bit query chunk against one load_256_alignment aligned reference chunk.
Definition at line 1241 of file kernels.cuh.
|
inline |
Compacts sequence whitespace in place, leaving every extent a run of bases.
Consumers stage extents verbatim and choose piece boundaries on any base, which needs contiguous bases. Compacting here moves the work onto the loader thread: each file is compacted while the other loaders are still decompressing, and a consumer that used to skip whitespace byte by byte now walks a dense span.
Definition at line 1413 of file fastx_sequence_file.hpp.
|
inline |
Copies bases from [first, last) to out, dropping sequence whitespace.
A sequence line is a long run of bases closed by a line ending, so a test per byte is the bulk of loading a corpus. Whole words are copied while none of their bytes is whitespace, and only a word that holds one falls back to bytes, which leaves the common case at one store per eight bases and keeps out no further along than the word being read.
Definition at line 1186 of file fastx_sequence_file.hpp.
|
noexcept |
Adds one group's comparison of query planes q with reference planes r, whose nonzero buckets reference_any the caller hoists out of per-query loops.
From the least significant plane up, each plane's verdict overrides the lower planes' unless its bits are equal: one 3-input logic operation per plane for lt, one for ne. With count_matches, buckets whose masked keys agree and whose scores are both nonzero are counted too, as the index counts them.
With CountMatches, matches are counted in the buckets of match_lanes; the two stored key masks come from the plane chain: a full key matches exactly the equal nonempty buckets, and the folded 15-bit key those whose low planes agree and that are nonempty in both rows, which needs the reference's low-plane nonzero mask reference_low_any. Only those two masks pass score_compatibility validation.
ReferenceFull asserts every reference score is nonzero, so no bucket is both empty.
Definition at line 1457 of file kernels.cuh.
|
constexpr |
Definition at line 968 of file kernels.cuh.
|
inline |
Feeds one byte of sequence through the DDL rolling-window extraction.
Definition at line 236 of file fasta_parser.hpp.
|
inlinenoexcept |
CRC32 of A followed by B, from CRC32(A), CRC32(B) and B's length, as zlib's crc32_combine: CRC32(A) times x^(8 * length) modulo the reflected polynomial, plus CRC32(B).
Definition at line 87 of file device_gzip.cuh.
|
inline |
Definition at line 40 of file reference_database_file.cuh.
| Result< void > cuddl::detail::database_file_metadata | ( | IO & | io, |
| reference_database_metadata & | metadata, | ||
| uint32_t & | blacklist_count | ||
| ) |
Definition at line 135 of file reference_database_file.cuh.
|
inline |
Definition at line 1670 of file fastx_sequence_file.hpp.
|
inline |
Decodes one gzip member, checking structure and ISIZE but not checksums.
Skipping the payload CRC avoids a second pass over every decompressed byte.
Definition at line 836 of file fastx_sequence_file.hpp.
|
inline |
A page-locked memory pool accessible from device only.
CCCL's default pinned pool is accessible from every visible device, so every allocation is mapped into all of them. With one MPI rank per GPU and all 8 GPUs of an A40 node visible to every rank, a 4-8 MB allocation took 27 ms instead of 1.3 ms, and sketch loaders spent most of their time waiting on the driver. Downloads into it with 8 ranks took 1.3 s, against 0.39 s from a device-node pool.
One pool per device and placement for the life of the process, like CUDA's default pools: never destroying it keeps it valid for buffers that live until process exit.
Definition at line 40 of file pinned_pool.cuh.
|
inlinenoexcept |
Whether device reads pageable host memory itself.
This is what makes a CPU/GPU system coherent, Grace Hopper and Grace Blackwell among them, and what lets a kernel read the loader's buffer instead of a copy of it.
Definition at line 80 of file database_staging.cuh.
|
inline |
One sparse bucket/key posting range, narrowed through the bucket's key directory.
Definition at line 505 of file kernels.cuh.
|
inlinenoexcept |
Distributed linear-counting estimate, blended from per-tier estimates by precision.
Definition at line 98 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
Single-precision dlc for the latency-bound hybrid scan.
Definition at line 239 of file detail/hybrid_cardinality.cuh.
| Result< void > cuddl::detail::download_store | ( | device_span< uint32_t const > | store, |
| std::span< uint16_t > | scores, | ||
| cuda::stream_ref | stream | ||
| ) |
Copies the winner scores of device store rows into host score rows.
Scores are extracted on the device, so only two bytes per bucket cross the bus. scores must stay alive until stream reaches this point.
Definition at line 844 of file database_staging.cuh.
|
inlineconstexprnoexcept |
Maps one DNA byte to its 2-bit symbol (A=0, C=1, T=2, G=3) or 0xff when invalid.
Case-insensitive; any ambiguous byte (including 'N') is invalid. The __host__ __device__ annotation only exists under the CUDA compiler so plain host translation units (e.g. via cuddl/fastx.hpp) see a normal inline function.
Definition at line 16 of file dna.hpp.
|
inline |
Query rows one external batch tile may hold on device.
A tile's results, match counts and counters cost about 20 bytes per pair, and a quarter of the device's memory goes to them, so a 16 GB device holds about 200 million pairs per tile. Larger tiles give the counting kernels more queries to overlap. The limit depends only on the device, so storage sized from the requirements always matches the search.
Definition at line 996 of file reference_database.cuh.
|
inline |
First '>' in [first, last) that starts a line, or last.
The byte before first must be a line ending.
Definition at line 1229 of file fastx_sequence_file.hpp.
|
inline |
Compacts one FASTA record's sequence from first to out, stopping at the next header.
The byte before first must be a line ending, and out may alias first. stop receives the next line-start '>' or last.
Definition at line 1366 of file fastx_sequence_file.hpp.
|
inline |
Collects FASTA sequence extents with whitespace compacted out, in one pass over data.
An extent spans the bases between one line-start '>' header line and the next header (or the end of data); bytes before the first header belong to no extent. Bases of consecutive records are packed back to back from out, which may alias data.
Definition at line 1384 of file fastx_sequence_file.hpp.
|
inlineconstexprnoexcept |
True for bytes that FASTA sequence consumers conventionally skip between bases.
Definition at line 61 of file fastx_sequence_file.hpp.
|
inline |
First line-end position at or after position using the resolved scanner.
Definition at line 195 of file fastx_sequence_file.hpp.
|
inline |
Definition at line 92 of file fastx_sequence_file.hpp.
|
inline |
Splits the combined FASTA sequence stream into count contiguous spans.
Each span carries the raw byte ranges of its slice of the sequence stream plus a materialised prefix of the preceding seed_bases valid bases (whitespace included, headers excluded, cut short at an invalid base or the buffer start). A rolling-window consumer that replays the prefix, then consumes the segments, reproduces the single-threaded byte stream exactly, so the per-span results can be concatenated in order. Whitespace between bases is preserved in both the prefix and the segments; consumers skip it with fastx_is_sequence_whitespace.
Definition at line 122 of file fasta_parser.hpp.
|
constexpr |
|
inline |
Parses, packs or defers the bytes data of a loaded file, which result holds.
Definition at line 1575 of file fastx_sequence_file.hpp.
| cuddl::detail::for | ( | ) |
Definition at line 444 of file kernels.cuh.
|
inline |
Decompresses gzip members straight into a caller's buffer.
Returns Error::resource when target is too small, which callers treat as a request to fall back to the growing-string path rather than a failure.
Definition at line 931 of file fastx_sequence_file.hpp.
|
inline |
Decompresses gzip members from a mapped buffer into output, appending.
Members are decompressed one at a time so concatenated streams (BGZF among them) work. Each member's structure, DEFLATE data and length trailer are checked, but CRC32 is not. output grows as needed, so pre-sizing storage saves reallocations without truncating data.
Definition at line 893 of file fastx_sequence_file.hpp.
|
inline |
Uncompressed size the gzip trailer advertises, or 0 when it cannot be read.
ISIZE holds one member's length modulo 2^32, so a concatenated stream is only sized correctly by its last member. A short buffer is refused later, never silently truncated.
Exact decompressed size of path, or 0 when it is not a readable gzip file.
A gzip member ends with a little-endian ISIZE, so this needs a 4-byte read instead of a full decompression pass. A corpus of these bounds how much a staged arena can ever hold.
Definition at line 977 of file fastx_sequence_file.hpp.
|
inline |
Definition at line 993 of file fastx_sequence_file.hpp.
|
inline |
True when data holds what could start another gzip member: 1f 8b 08, then a flag byte with its reserved bits clear.
A member with reserved flag bits set fails to decode, so it cannot be one the host path would inflate. Compressed data holds the three-byte signature alone in about one file in eleven, which would send those files back to the host.
Definition at line 221 of file device_gzip.cuh.
|
constexprnoexcept |
Hashes a canonical packed k-mer into the value space used by DDL registers.
Definition at line 21 of file hash.cuh.
|
inlinenoexcept |
Log-blended low and corrected estimates across the hybrid transition zone.
Definition at line 143 of file detail/hybrid_cardinality.cuh.
| __global__ void cuddl::detail::hybrid_cardinality_kernel | ( | uint32_t const *const | registers, |
| hybrid_cardinality_estimates *const | estimates | ||
| ) |
Computes BBTools and paper-style HybridDDL estimates in one register scan.
Definition at line 2575 of file kernels.cuh.
| __global__ void cuddl::detail::hybrid_cardinality_variant_kernel | ( | uint32_t const *const | registers, |
| double *const | estimate | ||
| ) |
Computes one HybridDDL variant estimate for estimator comparisons.
Definition at line 2610 of file kernels.cuh.
|
inlinenoexcept |
Fused BBTools, paper, linear-counting, DLC, and raw MeanM estimates.
Definition at line 174 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
Single-precision hybrid_estimates for the latency-bound hybrid scan.
Definition at line 316 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
Single-precision hybrid for the latency-bound hybrid scan.
Definition at line 285 of file detail/hybrid_cardinality.cuh.
| __global__ void cuddl::detail::hybrid_posting_ranges_kernel | ( | uint16_t const * | queries, |
| uint32_t const * | offsets, | ||
| uint32_t const * | bitmap_cells, | ||
| uint32_t | bitmap_count, | ||
| uint32_t | words, | ||
| uint32_t | indexed_buckets, | ||
| uint16_t | key_mask, | ||
| uint2 * | ranges, | ||
| uint32_t * | bitmap_starts, | ||
| uint32_t * | bitmap_counts | ||
| ) |
Splits each query's posting ranges into short lists and bitmap offsets.
With no dense offsets, consumes sparse ranges already staged in ranges. Packs bitmap offsets in bucket order so nearby queries walk their shared bitmaps together.
Definition at line 775 of file kernels.cuh.
| cuddl::detail::if | ( | candidate_bits ! | = nullptr && all_to_all && high <= query_id_offset + query_index + 1U | ) |
Definition at line 942 of file kernels.cuh.
| cuddl::detail::if | ( | candidate_bits | = = nullptr | ) |
Definition at line 1174 of file kernels.cuh.
| cuddl::detail::if | ( | local ! | = 0U | ) |
Definition at line 499 of file kernels.cuh.
| cuddl::detail::if | ( | query >=tile. | query_count | ) |
|
inlinenoexcept |
References per tile, bounded by the device's shared memory and small enough that the batch launches about four blocks per SM.
Tiles align to 64 references.
Definition at line 702 of file kernels.cuh.
|
inlineconstexprnoexcept |
Definition at line 921 of file reference_database.cuh.
|
inlineconstexprnoexcept |
|
inline |
Inflates gzip input in this thread's reusable heap buffer and, for FASTA, packs it there and moves only the packed words into a buffer from source.
The words are what crosses the bus, a quarter of the bases, so only they need page-locked memory. Pinning the whole inflated genome made pinning, which the driver runs one call at a time, what a fresh process with many loaders waited on. FASTQ is copied whole instead, and its bytes are returned for the parser; packed FASTA returns nothing.
Definition at line 1615 of file fastx_sequence_file.hpp.
|
inlinenoexcept |
Interpolates the MeanM correction factor for estimate from the calibration table.
Definition at line 59 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
Single-precision variant of the hybrid estimator tail.
The 2,048-register hybrid scan is memory/latency bound on every GPU, but the serial estimator tail in hybrid_estimates is FP64-heavy. Datacenter GPUs (H100/B200) execute those instructions at half FP32 rate, while consumer cards rate-limit FP64 to 1/64 of FP32. Keeping the scan and tail in FP32 avoids that consumer cliff; the ~1e-7 relative error is far below the ~2% standard error of a 2,048-register sketch.
Definition at line 200 of file detail/hybrid_cardinality.cuh.
|
inline |
One thread's posting range for a bucket/key cell of a dense or sparse index.
Definition at line 395 of file kernels.cuh.
|
inline |
Launches construction of a sketch through a CTA-local shared register array merged into registers (add_shared_kernel).
Definition at line 17 of file construction.cuh.
|
inline |
Shared launch for one device-resident ASCII chunk.
Windows derive as size >= k ? size - k + 1 : 0; short input is a no-op. Window counts above UINT32_MAX are rejected instead of narrowing, so size must not exceed UINT32_MAX + k - 1. No hidden carry: the caller supplies any K-1 overlap explicitly.
Definition at line 375 of file sequence_encode.cuh.
|
inlinenoexcept |
Linear-counting estimate from the first occupied NLZ tier.
Definition at line 81 of file detail/hybrid_cardinality.cuh.
|
inlinenoexcept |
Single-precision lc_min for the latency-bound hybrid scan.
Definition at line 222 of file detail/hybrid_cardinality.cuh.
|
constexprnoexcept |
Byte alignment load_256_global_nc requires on the compiled architecture.
sm_100 and later issue one 256-bit load per call, which needs 32-byte alignment. Earlier architectures issue two 128-bit loads, which need only 16 bytes, so a 16-byte-aligned row still takes the wide path there.
Definition at line 22 of file helpers.cuh.
| __device__ __forceinline__ void cuddl::detail::load_256_global_nc | ( | const T * | ptr, |
| T * | out | ||
| ) |
Loads 256 bits from global memory using the non-coherent cache path.
| T | Element type (uint32_t or uint64_t) |
| ptr | Source pointer (must be aligned to load_256_alignment) |
| out | Output array (4 elements for uint64_t, 8 for uint32_t) |
Definition at line 47 of file helpers.cuh.
| __host__ __device__ __forceinline__ T cuddl::detail::load_chunk | ( | const uint8_t * | data, |
| uint64_t | index | ||
| ) |
|
inline |
Definition at line 1738 of file fastx_sequence_file.hpp.
| __device__ __forceinline__ void cuddl::detail::load_plane_chunk | ( | uint4 const * | row, |
| uint32_t | group, | ||
| uint32_t | chunk, | ||
| uint32_t(&) | words[4] | ||
| ) |
Loads plane chunk chunk (planes 4 chunk ... 4 chunk + 3) of group group of a bit-plane row (score_plane_index layout).
Definition at line 2101 of file kernels.cuh.
| __device__ __forceinline__ void cuddl::detail::load_plane_group | ( | uint4 const * | row, |
| uint32_t | group, | ||
| uint32_t(&) | words[score_planes] | ||
| ) |
Loads group group of one bit-plane row (score_plane_index layout).
Definition at line 1380 of file kernels.cuh.
| __device__ __forceinline__ uint32_t cuddl::detail::load_shared_relaxed | ( | uint32_t const * | address | ) |
Relaxed CTA-scope load of a shared word that other threads update atomically.
cuda::atomic_ref emits a generic-address load here; the shared state space avoids the address conversion and the 64-bit address registers.
Definition at line 29 of file sequence_encode.cuh.
| __device__ __forceinline__ refine_cells cuddl::detail::make_refine_cells | ( | uint32_t | query_id_offset, |
| uint32_t | query_count, | ||
| uint32_t | group, | ||
| uint32_t | words_per_query, | ||
| uint32_t | max_block_words, | ||
| uint32_t | waves | ||
| ) |
Splits a tile into refine_cells, about waves cells per block when the tile allows, so that blocks claiming cells dynamically even out cells of uneven work.
Every query of an upper-triangle tile follows query_id_offset, so words wholly before it hold no pairs and get no cells; otherwise those empty cells would take most of the split for tiles far down the triangle.
Definition at line 1717 of file kernels.cuh.
| __global__ void cuddl::detail::match_filter_kernel | ( | uint32_t const * | query_planes, |
| uint32_t | query_id_offset, | ||
| uint32_t | query_count, | ||
| uint32_t const * | reference_planes, | ||
| uint32_t | reference_count, | ||
| uint32_t | indexed_bucket_count, | ||
| uint32_t | minimum_matches, | ||
| uint32_t * | candidate_bits, | ||
| uint32_t * | cell_counter | ||
| ) |
Records in candidate_bits (candidate-bitmap layout) the pairs of the layout whose scores agree in their low match_filter_planes planes in at least minimum_matches buckets.
Every validated key mask covers those planes, so the count bounds a pair's index matches from above and the candidates include every pair meeting the threshold. On RefSeq sketches the bound passes barely more pairs than the exact count; a thresholded bitmap refinement then counts matches exactly.
Cells, query staging, and word claiming follow refine_batch_bitmap_kernel. A warp takes a claimed word's references four at a time and keeps their low planes in registers, so each query plane chunk read from shared memory serves four pairs.
Definition at line 2125 of file kernels.cuh.
|
inlinenoexcept |
MeanM estimate retained as the full-sketch form of minimum_mle.
Definition at line 45 of file cardinality.cuh.
|
inlinenoexcept |
Single-precision mean_m for the throughput-bound cardinality reduction.
Definition at line 82 of file cardinality.cuh.
|
inlinenoexcept |
Estimates the number of distinct hashes from the minimum hash in each bucket.
Each distinct hash lands in one of bucket_count buckets. A bucket therefore receives about cardinality / bucket_count hashes. More hashes make its minimum smaller, so the observed minima tell us how many hashes must have produced them.
The denominator combines all observations after scaling hash values to [0, 1]:
1, meaning that its minimum lies beyond the hash range;Dividing bucket_count * filled_bucket_count by that total gives the maximum-likelihood cardinality estimate. If no bucket is empty, this reduces to the usual MeanM estimate.
| bucket_count | Total number of sketch buckets. |
| empty_count | Number of buckets that received no hashes. |
| sum_restored | Sum of the restored interval midpoint from every filled bucket. |
Definition at line 33 of file cardinality.cuh.
|
inlinenoexcept |
Single-precision variants of the cardinality estimators.
The reduction dominates the cardinality kernel. Keeping the restored-sum reduction and the scalar estimate in FP32 avoids the 1/64 FP32:FP64 throughput ratio on consumer Blackwell while preserving far more precision than the ~2% standard error of a 2,048-register sketch.
Definition at line 71 of file cardinality.cuh.
|
inline |
Definition at line 65 of file distributed.cuh.
| __device__ __forceinline__ uint32_t cuddl::detail::next_cell | ( | uint32_t * | cell_counter, |
| uint32_t | cell, | ||
| uint32_t & | claimed | ||
| ) |
The cell a refinement-style block processes after cell: claimed from cell_counter, a zeroed per-launch counter, so blocks balance uneven cells; without one, grid-stride.
Called by every thread of the block; claimed is the block's shared slot. The first cell is next_cell(cell_counter, blockIdx.x - gridDim.x, claimed).
Definition at line 1690 of file kernels.cuh.
|
inline |
Packs up to sixteen bases into a two-bit word, base j in bits 2j and 2j+1.
Codes are those of encode_base. A byte that is not a base gets a meaningless code.
j set when byte j is not a base, or is past count. Definition at line 591 of file fastx_sequence_file.hpp.
|
inline |
Packs width bases from sixteen-base pieces, for a tail shorter than a full step.
Definition at line 629 of file fastx_sequence_file.hpp.
|
inline |
Definition at line 638 of file fastx_sequence_file.hpp.
|
inline |
Compact and pack one cache-sized block at a time, without a genome-sized ASCII write.
Blocks end at line boundaries so header detection keeps its line-start invariant. Carrying k - 1 compacted bytes preserves every window without duplicating one.
Definition at line 1508 of file fastx_sequence_file.hpp.
|
inline |
Packs a parsed file's records in place into two-bit words, run by run.
A record splits at every byte that is not a base, and runs shorter than k are dropped: no k-mer holding such a byte is hashed, so the k-mers are those of the records. Output never passes unread input: words are written only after the packer has read their bytes, a full step writes a quarter of what it read, and a partial step writes only for a run of at least k >= 16 bases, whose words take at most half of the bytes it read.
Definition at line 761 of file fastx_sequence_file.hpp.
|
inlinenoexcept |
Whether a build in mode decompresses into page-locked buffers.
automatic keys on the device rather than the architecture: a device that reads pageable host memory is coherent, and there the host writes to page-locked memory are the slow part.
Definition at line 99 of file database_staging.cuh.
| void cuddl::detail::parallel_for | ( | size_t | count, |
| size_t | workers, | ||
| Work && | work | ||
| ) |
Runs work(index) for every index below count on workers threads.
Definition at line 203 of file device_gzip.cuh.
|
inline |
Parses a FASTA/FASTQ file's sequence bases into packed canonical k-mers of length k.
Record headers and FASTQ qualities are excluded, and every record's sequence is fed through the DDL rolling-window extraction as one combined genome. An invalid or ambiguous base breaks the rolling window, and no k-mer spanning it is emitted. Each k-mer is canonicalised to the larger packed orientation. Files above 1 MiB are parsed in parallel: the sequence stream is split into seeded spans (fastx_split_sequence_spans), each worker replays its window prefix and then consumes its segments, and the per-span results concatenate to the byte-identical single-threaded output.
| path | FASTA/FASTQ file to parse. |
| k | K-mer length. |
| threads | Worker count for parallel files, defaulting to std::thread::hardware_concurrency(); files at or below 1 MiB parse serially regardless. |
fasta_parse_result, or an error if the file cannot be opened, mapped, or contains bytes but no FASTA records, or malformed FASTQ records. Definition at line 288 of file fasta_parser.hpp.
|
inline |
Definition at line 1470 of file fastx_sequence_file.hpp.
|
inline |
Sizes the arena and the row store from free device memory and the corpus.
staged_ceiling bounds the bytes the corpus can stage, which keeps a small collection from reserving an arena it could never fill. An unset staging_bytes sizes the arena from what is free; a value is honoured, or rejected when it does not fit.
Definition at line 143 of file database_staging.cuh.
|
noexcept |
Buckets of a plane group holding a nonzero score.
Definition at line 1421 of file kernels.cuh.
|
noexcept |
Buckets of a plane group whose score is nonzero below the top plane.
Definition at line 1432 of file kernels.cuh.
|
inline |
Reads the header and trailer of path, and of every block of a BGZF file.
A candidate is a gzip file without FEXTRA whose trailer names a non-empty member, a BGZF file with a non-empty block, or an uncompressed file that opens with a FASTA header, any of them with a name that does not mark it FASTQ. For gzip without FEXTRA, additional member signatures are checked after reading the compressed bytes, before accepting the device output.
Definition at line 135 of file device_gzip.cuh.
|
inline |
Definition at line 42 of file reference_index_file.cuh.
| __global__ void cuddl::detail::refine_all_pairs_kernel | ( | uint32_t const * | query_planes, |
| uint32_t | query_id_offset, | ||
| uint32_t | query_count, | ||
| uint32_t const * | reference_planes, | ||
| uint32_t | reference_count, | ||
| SearchResult * | results, | ||
| uint32_t * | cell_counter, | ||
| uint32_t const * | reference_mask | ||
| ) |
Exactly compares the layout pairs selected by reference_mask, without match counts.
A null mask selects every reference.
Cells and query staging follow refine_batch_bitmap_kernel; warps claim all_pairs_word_parts parts of a 32-reference word at a time. A warp takes a claim's references two at a time and keeps both references' planes in registers, so each query plane group read from shared memory serves two pairs, and the two comparisons give every warp independent logic chains to interleave.
Definition at line 2316 of file kernels.cuh.
| __device__ void cuddl::detail::refine_write_result | ( | uint32_t | index, |
| uint32_t | query_index, | ||
| uint32_t | reference_id, | ||
| uint32_t | query_id_offset, | ||
| uint32_t const * | match_counts, | ||
| uint32_t | pair_id, | ||
| pairwise_counts | local, | ||
| uint32_t | lane, | ||
| SearchResult * | results, | ||
| uint32_t * | result_match_counts | ||
| ) |
Warp-reduces one refined candidate and writes its stable result slot.
Each counter reduces with one redux.sync instruction instead of a shuffle tree.
Definition at line 1290 of file kernels.cuh.
|
inline |
Selects the widest base packer the host supports.
Definition at line 739 of file fastx_sequence_file.hpp.
|
inline |
Selects the fastest line-end scanner the host supports.
Definition at line 178 of file fastx_sequence_file.hpp.
|
inline |
Selects the widest signature scanner the host supports.
Definition at line 305 of file device_gzip.cuh.
|
constexprnoexcept |
Reconstructs the approximate original hash magnitude from a stored score.
Inverts the score encoding: recovers NLZ from the exponent field, uninverts the mantissa bits and prepends the implicit leading one, then shifts the result back into 64-bit space.
Definition at line 68 of file register.cuh.
|
constexprnoexcept |
Reconstructs the interval midpoint for restore.
restore returns the lower bound of the hash interval represented by stored. The actual hash is uniformly distributed across that interval, so the midpoint is the minimum-variance scalar summary and reduces the small positive bias that lower-bound restoration introduces into hash-magnitude cardinality sums.
Definition at line 87 of file register.cuh.
|
inlineconstexprnoexcept |
Reverse-complements a packed k-mer (2 bits per base) with bit-parallel reversal.
Reverses the order of the 2-bit base digits with mask-swap rounds, shifts the reversed 2k bits down from the high end, then complements every base (^0b10). Byte-identical to the scalar loop formulation for every k <= 32; longer k-mers fall back to the scalar loop. The __host__ __device__ annotation only exists under the CUDA compiler so plain host translation units see a normal inline function, as with encode_base.
Definition at line 42 of file fasta_parser.hpp.
|
constexpr |
|
inline |
One row of buckets scores as a single MPI element, so counts of rows stay within int.
Definition at line 264 of file distributed.cuh.
|
inlinenoexcept |
Computes the 16-bit DDL rarity score for hash.
The exponent field stores the number of leading zeros (NLZ); the mantissa field stores the mantissa_bits bits immediately following the leading one, bitwise inverted. A larger score always encodes a rarer hash. Zero is reserved for the empty register, so a hash whose raw exponent/mantissa encoding is zero collapses to one.
Definition at line 43 of file register.cuh.
|
constexprnoexcept |
Position of plane plane of 32-bucket group group within a row's bit-planes.
A row of BucketCount 16-bit scores becomes BucketCount / 2 plane words: word p of group g holds bit p of the group's 32 scores. Lane g % 32 of a warp owns group g, and its 16 plane words form four consecutive uint4s at lane-contiguous positions, so a warp reads a row with coalesced 128-bit loads and shared copies with conflict-free 128-bit loads.
Definition at line 1334 of file kernels.cuh.
|
inline |
Rank 0's listing of input on every rank.
Definition at line 89 of file distributed.cuh.
| Result< void > cuddl::detail::shared_search | ( | MPI_Comm | comm, |
| Search && | search | ||
| ) |
Runs search with a replica_search that shares tiles across the ranks of comm.
Definition at line 449 of file distributed.cuh.
|
inline |
First genome of rank rank's slice: contiguous slices of about equal bytes.
Definition at line 120 of file distributed.cuh.
|
inline |
Every rank's genome count, after checking collectively that mine holds the sketches of this rank's slice and that every rank built with rank 0's settings.
Definition at line 235 of file distributed.cuh.
|
inlineconstexprnoexcept |
Storage needing fewer device bytes for reference_count references; dense on ties.
Both store one 32-bit posting per indexed row, so only the per-storage extras are compared: dense offsets per key cell against sparse 16-bit keys and their derived key directory.
Definition at line 936 of file reference_database.cuh.
|
constexprnoexcept |
Bounds per bucket in a sparse key directory.
Definition at line 429 of file kernels.cuh.
|
constexprnoexcept |
Low key bits a sparse key directory cell spans for reference_count references.
Cells average about eight keys per bucket for uniformly spread keys, so a lookup reads the cell's two bounds and then searches within about one memory sector.
Definition at line 420 of file kernels.cuh.
|
constexprnoexcept |
SplitMix64 finalizer, applied to a packed k-mer XOR the domain seed.
Definition at line 9 of file hash.cuh.
| Result< void > cuddl::detail::stage_groups | ( | size_t | genomes, |
| uint64_t | staged_ceiling, | ||
| std::optional< size_t > | staging_bytes, | ||
| cuda::stream_ref | stream, | ||
| reference_build_statistics * | statistics, | ||
| Fill && | fill, | ||
| Sink && | sink, | ||
| bool | in_place = false, |
||
| cuda::std::optional< blacklist_view > | blacklist = cuda::std::nullopt |
||
| ) |
Runs the staging loop over a collection, filling one group at a time.
fill receives the stager, a group's first genome and its size, and stages that whole group. Genomes are numbered within their group, so a fill stages genome genome - base. sink receives each finished group's device rows with its first genome and size, and must enqueue every read of them on stream: the next group reuses the rows.
Definition at line 1111 of file database_staging.cuh.
| Result< void > cuddl::detail::stage_paths | ( | std::span< std::filesystem::path const > | paths, |
| cuda::stream_ref | stream, | ||
| std::optional< size_t > | staging_bytes, | ||
| unsigned | parser_workers, | ||
| transfer_mode | transfer, | ||
| decompression_backend | decompression, | ||
| bool | direct_reads, | ||
| reference_build_statistics * | statistics, | ||
| Sink && | sink, | ||
| cuda::std::optional< blacklist_view > | blacklist = cuda::std::nullopt |
||
| ) |
Stages a path collection, handing each group's device rows to sink.
parser_workers is a value, not a sentinel: the build clamps it to what the inputs and the hardware allow, so asking for the default is asking for the default.
The arena ceiling is the decompressed size of every input: a bound from compressed sizes alone overestimates a corpus several times over, and an arena past what the corpus can hold is memory the device never needs.
Definition at line 1612 of file database_staging.cuh.
| Result< void > cuddl::detail::stage_sequences | ( | std::span< sequence_genome const > | genomes, |
| cuda::stream_ref | stream, | ||
| std::optional< size_t > | staging_bytes, | ||
| reference_build_statistics * | statistics, | ||
| Sink && | sink, | ||
| cuda::std::optional< blacklist_view > | blacklist = cuda::std::nullopt |
||
| ) |
Stages bases the caller already holds, handing each group's device rows to sink.
Definition at line 1761 of file database_staging.cuh.
|
inlinenoexcept |
Whether a build in mode stages the caller's bytes in place.
Automatic transfer reads CPU-decoded buffers in place on coherent devices.
Definition at line 90 of file database_staging.cuh.
| __device__ uint32_t cuddl::detail::sum_posting_bitmaps | ( | uint32_t const * | bitmaps, |
| uint32_t const * | starts, | ||
| uint32_t | count, | ||
| uint32_t | word, | ||
| uint32_t | minimum_matches, | ||
| uint32_t | passed, | ||
| uint32_t(&) | counters[CounterBits] | ||
| ) |
Sums one query's posting bitmaps over one word of references in bit-sliced counters.
Independent bitmap loads overlap in batches. Returns passed with every reference whose count reaches minimum_matches added; summing stops once all 32 have passed, so a passed reference's count is unused and may wrap.
Definition at line 844 of file kernels.cuh.
|
constexprnoexcept |
|
constexprnoexcept |
Ranking key of a neighbour: its weighted k-mer identity above its complemented ID.
Keys are unique, a larger key ranks higher, and equal identities rank the smaller ID higher, so any selection of the k largest keys is the same set. Zero means no candidate; a real key is never zero because IDs stay below UINT32_MAX.
Definition at line 29 of file top_k.cuh.
|
inline |
Position of set bit n, counting from zero, of mask, which has more than n set bits.
Definition at line 114 of file top_k.cuh.
|
inline |
Replaces the root of a min-heap of k keys at heap, whose values follow the keys, with key and value, and sifts it down.
The caller owns the heap and has checked that key beats the root.
Definition at line 66 of file top_k.cuh.
|
inline |
|
inline |
Copies the k entries of global heap row into, or for ToShared false out of, the shared heap heap with the warp's lanes.
| cuddl::detail::top_k_warp_copy< false > | ( | slots | , |
| row | , | ||
| heap | |||
| ) |
| cuddl::detail::top_k_warp_copy< true > | ( | slots | , |
| row | , | ||
| heap | |||
| ) |
| cuddl::detail::top_k_warp_merge | ( | heap | , |
| slots. | k, | ||
| words - | cuda::std::minskipped, words, | ||
| [&] (uint32_t word) { return top_k_word_bits< Thresholded, UpperTriangle >(tile, query, skipped+word, words);} | , | ||
| [&] (uint32_t candidate, uint64_t &key, uint64_t &value) { auto const reference=skipped *32U+candidate;auto const stored=top_k_stored< UpperTriangle >(tile, query, reference);key=top_k_key(stored.unpack(tile.bucket_count), reference);value=stored.bits;} | |||
| ) |
|
inline |
One warp merges the candidates of words bitmap words into the heap of k entries at heap, in shared memory, which only this warp updates, so it needs no lock.
bits_of gives word w's candidates and entry_of candidate w * 32 + bit's key and value. The warp takes 32 words at a time and hands their set bits out 32 per step, one per lane in order, so every step offers 32 candidates however the words are filled. Lanes test their candidates against the root; the warp then inserts the survivors one at a time.
Definition at line 131 of file top_k.cuh.
|
inline |
The pass bits of word word of tile row query: the stored bits, or for exhaustive tiles every reference of the word.
Pairs on or below the diagonal of all-to-all tiles are masked.
Definition at line 85 of file top_k.cuh.
| T cuddl::detail::try_unwrap_success | ( | Result< T > & | result | ) |
|
inline |
|
inlinenoexcept |
Whether a batch counts through count_batch_index_tile_kernel.
Its 16-bit counters hold at most 65535 matches. Tiny batches finish sooner on the global kernel, which needs no per-block counter initialization or flush.
Definition at line 691 of file kernels.cuh.
|
inline |
Definition at line 906 of file reference_database.cuh.
|
inline |
Definition at line 892 of file reference_database.cuh.
|
inline |
Definition at line 868 of file reference_database.cuh.
| std::optional< std::pair< uint32_t, uint64_t > > cuddl::detail::walk_bgzf | ( | unsigned char const * | bytes, |
| uint64_t | size, | ||
| Block && | block | ||
| ) |
Calls block(deflate_at, deflate_bytes, inflated) for each non-empty block of the BGZF file in bytes.
size bytes. Definition at line 111 of file device_gzip.cuh.
|
inlinenoexcept |
Blocks for a grid-stride launch covering warps warps of block_size threads.
Definition at line 1524 of file kernels.cuh.
|
inlinenoexcept |
True when both row bases meet load_256_alignment on the compiled architecture.
Definition at line 33 of file helpers.cuh.
| __device__ __forceinline__ void cuddl::detail::write_score_planes | ( | uint32_t | score, |
| uint32_t | group, | ||
| uint32_t | lane, | ||
| uint32_t * | planes | ||
| ) |
Writes the bit-planes of one 32-bucket group whose scores the warp holds, one per lane.
Definition at line 1340 of file kernels.cuh.
|
inline |
Computes xxHash-64 over the raw bytes of a value.
Definition at line 155 of file hash.cuh.
|
inline |
Computes xxHash-64 over a byte range without allocating or copying it.
Definition at line 63 of file hash.cuh.
|
inlineconstexpr |
Optimisations that the CUDDL_ABLATE environment variable can turn off in a build configured with -Dablations=true (CUDDL_ABLATIONS=1).
They exist to measure what each optimisation contributes (scripts/ablation.py); results are the same either way. Other builds compile every switch out.
Definition at line 19 of file ablation.hpp.
|
inlineconstexpr |
Definition at line 79 of file ablation.hpp.
|
constexpr |
Threads per refine_all_pairs_kernel block.
Two references' planes stay in registers, so a block holds fewer threads than a refine block.
Definition at line 2299 of file kernels.cuh.
|
constexpr |
Parts a refine_all_pairs_kernel warp claims each 32-reference word in.
Warps idle at the end of a cell until the slowest one finishes its last claim, and words differ in cost, so smaller claims shorten that wait. On a GPU running at its power limit the shorter wait only lowers the clock; with power headroom it shortens the search.
Definition at line 2305 of file kernels.cuh.
| static __global__ uint32_t uint32_t uint32_t uint32_t bool cuddl::detail::all_to_all |
Definition at line 928 of file kernels.cuh.
|
constexpr |
Minimum query rows per batch tile.
Definition at line 955 of file reference_database.cuh.
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t uint32_t const uint32_t const uint32_t const* cuddl::detail::bitmap_counts |
Definition at line 932 of file kernels.cuh.
|
constexpr |
Threads per refine_batch_bitmap_kernel block.
Leaves register space for query-plane loads alongside cached reference planes.
Definition at line 1649 of file kernels.cuh.
|
constexpr |
Shared bytes one refine_batch_bitmap_kernel block stages.
Definition at line 1680 of file kernels.cuh.
|
constexpr |
Queries one refine_batch_bitmap_kernel block refines together.
Each reference row is read once per group, so the group size divides the kernel's L2 traffic. Rows too large to stage in shared memory are read from global memory in groups of 8.
Definition at line 1670 of file kernels.cuh.
|
constexpr |
Shared bytes refine_batch_bitmap_kernel may stage query planes in; one block per SM fits on every supported architecture.
Definition at line 1659 of file kernels.cuh.
|
constexpr |
Whether a query group's planes fit bitmap_refine_shared_bytes.
Definition at line 1663 of file kernels.cuh.
|
constexpr |
Cells per refine_batch_bitmap_kernel block that a tile of bitmap-selected pairs is split into when it allows.
Selected pairs cluster around near-duplicate genomes, so cells differ in work by 20 times or more; with one cell per block the densest cell sets the kernel's time.
Definition at line 1655 of file kernels.cuh.
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t uint32_t const uint32_t const* cuddl::detail::bitmap_starts |
Definition at line 931 of file kernels.cuh.
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t uint32_t uint16_t const uint2 const uint32_t uint32_t bool uint32_t uint32_t const* cuddl::detail::bitmaps |
Definition at line 930 of file kernels.cuh.
| auto const cuddl::detail::bits |
Definition at line 1184 of file kernels.cuh.
| __global__ size_t uint32_t bool blacklist_view cuddl::detail::blacklist |
Definition at line 96 of file kernels.cuh.
| auto const cuddl::detail::block = static_cast<uint64_t>(blockIdx.x) * (blockDim.x / 32U) + threadIdx.x / 32U |
|
constexpr |
Threads per CTA for the single-pair and cardinality reduction kernels.
Definition at line 28 of file kernels.cuh.
|
constexpr |
Definition at line 1866 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const * cuddl::detail::candidate_bits |
Definition at line 929 of file kernels.cuh.
| __global__ uint32_t const uint32_t uint32_t const uint32_t const* cuddl::detail::candidate_count |
Definition at line 1542 of file kernels.cuh.
| __global__ uint32_t const uint32_t uint32_t const* cuddl::detail::candidate_ids |
Definition at line 1541 of file kernels.cuh.
| uint32_t cuddl::detail::carry = 0U |
Definition at line 547 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const SearchResult uint32_t uint32_t uint32_t uint32_t* cuddl::detail::cell_counter |
Definition at line 1855 of file kernels.cuh.
| auto const cuddl::detail::cells |
Definition at line 1878 of file kernels.cuh.
| __shared__ uint32_t cuddl::detail::claimed_cell |
Definition at line 1874 of file kernels.cuh.
| __global__ uint32_t cuddl::detail::count |
Definition at line 1031 of file kernels.cuh.
|
constexpr |
Definition at line 25 of file detail/hybrid_cardinality.cuh.
|
inlineconstexpr |
Default arena ceiling.
A batch is sketched only once its arena fills, so an arena that holds the whole corpus leaves the device idle until the last genome loads. Bounding it lets each batch's kernel overlap the loading of the next: on a GH200 building 2048 genomes this cut the build from 239 to 194 ms, and it cost nothing on a 24-thread RTX 5070 Ti host.
Definition at line 117 of file database_staging.cuh.
|
constexpr |
Definition at line 26 of file reference_deltas.cuh.
|
constexpr |
Most buckets in which a child reference may differ from its base reference.
Definition at line 20 of file reference_deltas.cuh.
|
constexpr |
Definition at line 30 of file reference_deltas.cuh.
|
constexpr |
Definition at line 39 of file reference_deltas.cuh.
|
constexpr |
Warps per apply_reference_deltas_kernel block, and the most and fewest children one block cell takes.
Definition at line 29 of file reference_deltas.cuh.
|
constexpr |
Definition at line 31 of file reference_deltas.cuh.
|
constexpr |
Queries one correction cell stages bucket-major: 64 KiB of scores for every bucket count.
Lane l serves query l % group and every (32 / group)-th change from l / group.
Definition at line 36 of file reference_deltas.cuh.
|
constexpr |
Definition at line 23 of file reference_deltas.cuh.
|
constexpr |
Highest-ID rows taken from each matching signature.
Definition at line 25 of file reference_deltas.cuh.
|
constexpr |
Exact bucket signatures that propose base candidates, and buckets per signature.
Definition at line 22 of file reference_deltas.cuh.
| __global__ size_t uint32_t uint16_t const uint32_t const * cuddl::detail::directory |
Definition at line 439 of file kernels.cuh.
|
constexpr |
Definition at line 14 of file detail/hybrid_cardinality.cuh.
|
constexpr |
Definition at line 13 of file detail/hybrid_cardinality.cuh.
|
constexpr |
Definition at line 16 of file detail/hybrid_cardinality.cuh.
|
constexpr |
Definition at line 15 of file detail/hybrid_cardinality.cuh.
| cuddl::detail::else |
Definition at line 1020 of file kernels.cuh.
| __global__ uint32_t const uint32_t const uint32_t const uint32_t cuddl::detail::first_query_id |
Definition at line 1786 of file kernels.cuh.
| __global__ uint32_t uint32_t cuddl::detail::first_word |
|
constexpr |
Definition at line 1858 of file kernels.cuh.
| auto const cuddl::detail::groups = cells.groups |
Definition at line 1886 of file kernels.cuh.
|
constexpr |
Definition at line 1360 of file kernels.cuh.
|
constexpr |
2^64, the range of the 64-bit hash used by the sketch.
Definition at line 10 of file cardinality.cuh.
| auto* const cuddl::detail::heap = heaps + warp * top_k_heap_words(slots.k) |
| auto const cuddl::detail::high = min(low + tile, reference_count) |
Definition at line 940 of file kernels.cuh.
|
constexpr |
Definition at line 18 of file detail/hybrid_cardinality.cuh.
|
constexpr |
Definition at line 17 of file detail/hybrid_cardinality.cuh.
| __global__ uint32_t uint32_t uint64_t uint32_t* cuddl::detail::ids |
Definition at line 85 of file reference_deltas.cuh.
| __global__ size_t uint32_t bool blacklist_view bool cuddl::detail::IncludeCardinality |
Definition at line 155 of file kernels.cuh.
|
constexpr |
Cells one warp owns per iteration of count_batch_index_matches_kernel: one per lane.
The host divides the grid by this so every launched warp stays busy.
Definition at line 392 of file kernels.cuh.
|
constexpr |
Index pair fraction (reference_index::pair_fraction) above which batch search compares bit-planes for every pair instead of counting through the index.
Definition at line 952 of file reference_database.cuh.
|
constexpr |
Longest posting range one count_batch_index_tile_kernel lane walks alone.
Definition at line 682 of file kernels.cuh.
|
constexpr |
Long posting ranges one count_batch_index_tile_kernel block walks together.
Definition at line 685 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t cuddl::detail::indexed_bucket_count |
Definition at line 438 of file kernels.cuh.
| auto const cuddl::detail::indexed_groups = indexed_bucket_count / warp_width |
Definition at line 1594 of file kernels.cuh.
| __global__ size_t cuddl::detail::input_size |
Definition at line 93 of file kernels.cuh.
| auto const cuddl::detail::items = candidate_ids == nullptr ? reference_count : *candidate_count |
Definition at line 1549 of file kernels.cuh.
| auto const cuddl::detail::key_count = static_cast<uint32_t>(key_mask) + 1U |
Definition at line 472 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t cuddl::detail::key_mask |
Definition at line 540 of file kernels.cuh.
| __global__ uint32_t uint32_t uint64_t* cuddl::detail::keys |
Definition at line 84 of file reference_deltas.cuh.
| auto const cuddl::detail::lane = static_cast<uint32_t>(threadIdx.x) % warp_width |
Definition at line 966 of file kernels.cuh.
|
constexpr |
Definition at line 1864 of file kernels.cuh.
| __global__ size_t uint32_t bool blacklist_view bool cuddl::detail::Layout |
Definition at line 155 of file kernels.cuh.
| unsigned long long cuddl::detail::local = 0U |
Definition at line 478 of file kernels.cuh.
| __shared__ uint32_t cuddl::detail::long_count |
Definition at line 965 of file kernels.cuh.
| __shared__ uint2 cuddl::detail::long_ranges[index_tile_long_capacity] |
Definition at line 964 of file kernels.cuh.
| auto const cuddl::detail::long_ranges_used |
Definition at line 1152 of file kernels.cuh.
Definition at line 939 of file kernels.cuh.
|
constexpr |
Default score layout retained for source compatibility.
Definition at line 27 of file register.cuh.
|
constexpr |
Definition at line 28 of file register.cuh.
| __global__ uint32_t const * cuddl::detail::match_counts |
Definition at line 575 of file kernels.cuh.
|
constexpr |
Threads per match_filter_kernel block.
Four references' low planes stay in registers, so a block holds fewer threads than a refine block.
Definition at line 2090 of file kernels.cuh.
|
constexpr |
Low planes match_filter_kernel compares.
Definition at line 2110 of file kernels.cuh.
|
constexpr |
Whether match_filter_kernel supports BucketCount: a lane caches at most two plane groups of each of its references.
Definition at line 2095 of file kernels.cuh.
|
constexpr |
Definition at line 1867 of file kernels.cuh.
|
constexpr |
Definition at line 32 of file register.cuh.
|
constexpr |
Stored sketches one batch launch may index, bounded by the CUDA grid dimension.
Definition at line 31 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const SearchResult uint32_t uint32_t cuddl::detail::minimum_matches |
Definition at line 926 of file kernels.cuh.
|
inlineconstexpr |
| __shared__ uint32_t cuddl::detail::next_word |
Definition at line 1873 of file kernels.cuh.
|
constexpr |
NLZ histogram bins for the hybrid estimator, including the empty tier.
Definition at line 12 of file detail/hybrid_cardinality.cuh.
| __global__ size_t uint32_t const * cuddl::detail::offsets |
Definition at line 536 of file kernels.cuh.
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t SearchResult* cuddl::detail::packed |
Definition at line 1791 of file kernels.cuh.
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t cuddl::detail::packed_base |
Definition at line 1790 of file kernels.cuh.
|
inlineconstexpr |
Bytes reserved per loader buffer that holds packed bases: a bacterial genome's two-bit words, about 1.5 MB, fit.
Definition at line 208 of file database_staging.cuh.
|
extern |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t SearchResult uint32_t* cuddl::detail::packed_match_counts |
Definition at line 1792 of file kernels.cuh.
| auto const cuddl::detail::pairs = static_cast<uint64_t>(query_count) * reference_count |
Definition at line 1592 of file kernels.cuh.
|
inlineconstexpr |
Reads a parallel file system needs in flight to give one node its bandwidth: on MOGON's Lustre, 16 threads read 2.9 GB/s of ~1 MB files and 64 threads 6.7 GB/s.
Definition at line 1784 of file fastx_sequence_file.hpp.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const SearchResult uint32_t uint32_t uint32_t * cuddl::detail::pass_bits |
Definition at line 1784 of file kernels.cuh.
| __global__ uint32_t uint32_t* cuddl::detail::planes |
Definition at line 1358 of file kernels.cuh.
|
constexpr |
Threads of the single index_posting_ranges_kernel block.
Definition at line 527 of file kernels.cuh.
| __global__ size_t uint32_t const uint32_t const * cuddl::detail::postings |
Definition at line 574 of file kernels.cuh.
|
constexpr |
|
constexpr |
|
constexpr |
|
constexpr |
|
constexpr |
| auto const cuddl::detail::query = queries + (query_row_offset + query_index) * BucketCount |
Definition at line 1030 of file kernels.cuh.
| __global__ uint32_t uint32_t cuddl::detail::query_count |
Definition at line 600 of file kernels.cuh.
| auto const cuddl::detail::query_id = tile.first_query_id + query |
| __global__ uint32_t cuddl::detail::query_id_offset |
Definition at line 927 of file kernels.cuh.
| auto const cuddl::detail::query_index = static_cast<uint32_t>(blockIdx.x) / tiles |
Definition at line 938 of file kernels.cuh.
|
extern |
| __global__ size_t cuddl::detail::query_row_offset |
Definition at line 599 of file kernels.cuh.
| __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t uint32_t* cuddl::detail::range_begin |
Definition at line 541 of file kernels.cuh.
| static __global__ uint32_t const * cuddl::detail::range_end_scan |
Definition at line 542 of file kernels.cuh.
| __global__ size_t uint32_t uint16_t const uint32_t const uint32_t uint32_t uint16_t uint2 * cuddl::detail::ranges |
Definition at line 925 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t cuddl::detail::reference_count |
Definition at line 437 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t SearchResult uint32_t const* cuddl::detail::reference_mask |
Definition at line 2491 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t const * cuddl::detail::reference_planes |
Definition at line 1539 of file kernels.cuh.
|
constexpr |
Reference-row bytes one refine_batch_index_candidates_kernel block spans.
Candidates are refined one (reference block, query) cell at a time, reference-block major, so the resident warps share a few blocks' rows through L2 instead of every query streaming every candidate reference row from DRAM.
Definition at line 1323 of file kernels.cuh.
| __global__ size_t uint32_t* cuddl::detail::registers |
Definition at line 94 of file kernels.cuh.
|
constexpr |
Default restoration constants retained for source compatibility.
Definition at line 31 of file register.cuh.
| __global__ uint32_t uint32_t uint32_t const uint32_t uint32_t uint16_t uint32_t const SearchResult uint32_t* cuddl::detail::result_match_counts |
Definition at line 1852 of file kernels.cuh.
| __global__ uint32_t uint32_t uint32_t SearchResult * cuddl::detail::results |
Definition at line 1543 of file kernels.cuh.
| auto const cuddl::detail::row = static_cast<size_t>(query_id - slots.first_genome_id) |
| __global__ size_t cuddl::detail::row_count |
Definition at line 1358 of file kernels.cuh.
| auto const cuddl::detail::row_word = static_cast<uint32_t>(block / word_count) |
|
constexpr |
Definition at line 1546 of file kernels.cuh.
| __global__ size_t uint16_t* cuddl::detail::rows |
Definition at line 1499 of file kernels.cuh.
|
constexpr |
Bits per score, and so bit-planes per 32-bucket group.
Definition at line 1326 of file kernels.cuh.
|
constexpr |
|
constexpr |
Threads per CTA for the CTA-local construction kernel.
Definition at line 78 of file kernels.cuh.
| auto const cuddl::detail::shift = sparse_directory_shift(reference_count) |
Definition at line 442 of file kernels.cuh.
| __global__ uint32_t uint32_t cuddl::detail::signature |
Definition at line 83 of file reference_deltas.cuh.
| auto const cuddl::detail::skipped = UpperTriangle ? (query_id + 1U) / 32U : 0U |
| __global__ top_k_slots cuddl::detail::slots |
| __global__ size_t uint32_t uint16_t const * cuddl::detail::sorted_keys |
Definition at line 469 of file kernels.cuh.
|
constexpr |
Definition at line 1860 of file kernels.cuh.
|
inlineconstexpr |
Staging arenas the copy stream rotates through, so copies run ahead of the kernels that read earlier arenas while other work shares the device.
Definition at line 121 of file database_staging.cuh.
|
inlineconstexpr |
Host memory the device still has to read must stay alive: the transfer engine's read of a pageable buffer is not ordered with the host writes that follow it.
Sixteen slots take the tail off that wait; at four the tail reached hundreds of milliseconds while dozens of loaders fed one consumer.
Definition at line 204 of file database_staging.cuh.
| __shared__ block_scan::TempStorage cuddl::detail::storage |
Definition at line 545 of file kernels.cuh.
| auto const cuddl::detail::stride = static_cast<uint32_t>(gridDim.x) * block_size |
Definition at line 578 of file kernels.cuh.
| auto const cuddl::detail::thresholded = CountMatches && pass_bits != nullptr |
Definition at line 1888 of file kernels.cuh.
| __global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t cuddl::detail::tile |
Definition at line 922 of file kernels.cuh.
| auto const cuddl::detail::tiles = (reference_count + tile - 1U) / tile |
Definition at line 937 of file kernels.cuh.
|
constexpr |
| auto const cuddl::detail::total = static_cast<size_t>(indexed_bucket_count) * entries |
Definition at line 443 of file kernels.cuh.
| __global__ size_t uint32_t bool cuddl::detail::vector_input |
Definition at line 95 of file kernels.cuh.
| auto const cuddl::detail::warp = static_cast<uint32_t>(threadIdx.x) / warp_width |
Definition at line 967 of file kernels.cuh.
| auto const cuddl::detail::warps = uint64_t{gridDim.x} * (block_size / warp_width) |
Definition at line 2495 of file kernels.cuh.
|
constexpr |
Definition at line 935 of file kernels.cuh.
|
constexpr |
Buckets one contiguous chunk covers, for a query/reference score pair.
One wide load covers 16 buckets whenever either row stores 16-bit scores, and 8 buckets when both store 32-bit registers. The scalar fallback covers half as many.
Definition at line 74 of file kernels.cuh.
Definition at line 963 of file kernels.cuh.
| auto const cuddl::detail::word = first_word + static_cast<uint32_t>(block % word_count) |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t cuddl::detail::word_begin |
Definition at line 1788 of file kernels.cuh.
| auto const cuddl::detail::word_count = words - first_word |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t cuddl::detail::word_end |
Definition at line 1789 of file kernels.cuh.
| __global__ uint32_t const uint32_t const uint32_t const* cuddl::detail::word_offsets |
Definition at line 1785 of file kernels.cuh.
|
constexpr |
Definition at line 1871 of file kernels.cuh.
| auto const cuddl::detail::words_per_query = candidate_bit_words(reference_count) |
Definition at line 1795 of file kernels.cuh.
| __global__ uint16_t const uint32_t uint32_t uint32_t unsigned long long* cuddl::detail::work |
Definition at line 473 of file kernels.cuh.