|
cuDDL
|
#include <cooperative_groups.h>#include <cooperative_groups/reduce.h>#include <cuda_runtime.h>#include <cub/block/block_histogram.cuh>#include <cub/block/block_reduce.cuh>#include <cub/block/block_scan.cuh>#include <cuda/std/algorithm>#include <cuda/std/bit>#include <cuda/std/cstddef>#include <cuda/std/cstdint>#include <cuddl/blacklist.cuh>#include <cuddl/detail/cardinality.cuh>#include <cuddl/detail/comparison.cuh>#include <cuddl/detail/hash.cuh>#include <cuddl/detail/hybrid_cardinality.cuh>#include <cuddl/detail/register.cuh>#include <cuddl/helpers.cuh>#include <cuddl/pairwise_counts.cuh>Go to the source code of this file.
Classes | |
| struct | cuddl::detail::summary_payload |
| Per-thread accumulator feeding a CUB block reduction for fused summaries. More... | |
| struct | cuddl::detail::posting_bitmap_cell |
| Selects posting lists whose bitmap uses no more storage than their reference IDs. More... | |
| struct | cuddl::detail::plane_counts |
| Bit-sliced pairwise counts of one 32-bucket group, accumulated per lane. More... | |
| struct | cuddl::detail::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 | cuddl::detail::word_popcount |
| Popcount of one pass-bit word. More... | |
| struct | cuddl::detail::cardinality_payload |
| Per-thread accumulator feeding the CUB block reduction for cardinality. More... | |
Namespaces | |
| namespace | cuddl |
| Batch operations over contiguous device rows. | |
| namespace | cuddl::detail |
Enumerations | |
| enum class | cuddl::detail::refine_candidates { cuddl::detail::bitmap , cuddl::detail::all , cuddl::detail::reference_mask } |
| Pairs refine_batch_bitmap_kernel compares. More... | |
Functions | |
| template<size_t BucketCount, typename Layout = default_register_layout, bool HasBlacklist = false> | |
| __global__ | cuddl::detail::__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 | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::sparse_directory_entries (uint32_t reference_count) noexcept |
| Bounds per bucket in a sparse key directory. | |
| __global__ | cuddl::detail::__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. | |
| cuddl::detail::for (auto i=static_cast< size_t >(blockIdx.x) *block_size+threadIdx.x;i< total;i+=static_cast< size_t >(gridDim.x) *block_size) | |
| cuddl::detail::if (local !=0U) | |
| __device__ uint2 | cuddl::detail::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__ | cuddl::detail::__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 | cuddl::detail::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 | cuddl::detail::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 | 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. | |
| template<uint32_t CounterBits> | |
| __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. | |
| template<size_t BucketCount, uint32_t CounterBits, typename QueryRow > | |
| __global__ | cuddl::detail::__maxnreg__ (48) void count_batch_index_tile_kernel(QueryRow const *queries |
| Counts index matches for one query against one reference tile in shared memory. | |
| cuddl::detail::if (candidate_bits !=nullptr &&all_to_all &&high<=query_id_offset+query_index+1U) | |
| if | cuddl::detail::constexpr (CounterBits !=0U) |
| cuddl::detail::__syncthreads () | |
| cuddl::detail::if (candidate_bits==nullptr) | |
| template<typename ReferenceScore > | |
| __device__ __forceinline__ void | cuddl::detail::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 | 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. | |
| __host__ __device__ constexpr uint32_t | cuddl::detail::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 | 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. | |
| template<size_t BucketCount> | |
| __global__ | cuddl::detail::__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 | 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). | |
| __device__ __forceinline__ uint32_t | cuddl::detail::plane_any (uint32_t const (&planes)[score_planes]) noexcept |
| Buckets of a plane group holding a nonzero score. | |
| __device__ __forceinline__ uint32_t | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::warp_grid_blocks (size_t warps) noexcept |
Blocks for a grid-stride launch covering warps warps of block_size threads. | |
| __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. | |
| template<bool UpperTriangle> | |
| __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. | |
| template<bool UpperTriangle> | |
| __host__ __device__ uint64_t | cuddl::detail::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__ | cuddl::detail::__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 | 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). | |
| template<size_t BucketCount, bool UpperTriangle, uint32_t QueryGroup> | |
| __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. | |
| template<size_t BucketCount, typename SearchResult , bool UpperTriangle, uint32_t QueryGroup> | |
| __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. | |
| template<size_t BucketCount, typename SearchResult , bool UpperTriangle, bool Swar> | |
| __global__ | cuddl::detail::__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 | cuddl::detail::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 | 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. | |
| template<size_t BucketCount, hybrid_variant Variant, typename Layout = default_register_layout> | |
| __global__ void | cuddl::detail::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 | 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. | |
| template<size_t BucketCount> | |
| __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. | |
Variables | |
| constexpr uint32_t | cuddl::detail::block_size = 256 |
| Threads per CTA for the single-pair and cardinality reduction kernels. | |
| constexpr size_t | cuddl::detail::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 | cuddl::detail::wide_chunk_buckets |
| Buckets one contiguous chunk covers, for a query/reference score pair. | |
| constexpr uint32_t | cuddl::detail::shared_construction_block_size = 768 |
| Threads per CTA for the CTA-local construction kernel. | |
| __global__ size_t | cuddl::detail::input_size |
| __global__ size_t uint32_t * | cuddl::detail::registers |
| __global__ size_t uint32_t bool | cuddl::detail::vector_input |
| __global__ size_t uint32_t bool blacklist_view | cuddl::detail::blacklist |
| __global__ size_t uint32_t bool blacklist_view bool | cuddl::detail::IncludeCardinality |
| __global__ size_t uint32_t bool blacklist_view bool | cuddl::detail::Layout |
| constexpr uint32_t | cuddl::detail::index_match_cells_per_warp = 32U |
| Cells one warp owns per iteration of count_batch_index_matches_kernel: one per lane. | |
| __global__ uint32_t | cuddl::detail::reference_count |
| __global__ uint32_t uint32_t | cuddl::detail::indexed_bucket_count |
| __global__ uint32_t uint32_t uint32_t * | cuddl::detail::directory |
| auto const | cuddl::detail::shift = sparse_directory_shift(reference_count) |
| auto const | cuddl::detail::total = static_cast<size_t>(indexed_bucket_count) * entries |
| __global__ uint16_t const * | cuddl::detail::sorted_keys |
| __global__ uint16_t const uint32_t uint32_t uint32_t | cuddl::detail::key_count = static_cast<uint32_t>(key_mask) + 1U |
| __global__ uint16_t const uint32_t uint32_t uint32_t unsigned long long * | cuddl::detail::work |
| unsigned long long | cuddl::detail::local = 0U |
| constexpr uint32_t | cuddl::detail::posting_range_block_size = 1024U |
| Threads of the single index_posting_ranges_kernel block. | |
| static __global__ uint32_t const * | cuddl::detail::offsets |
| static __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t | cuddl::detail::key_mask |
| static __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t uint32_t * | cuddl::detail::range_begin |
| static __global__ uint32_t const uint16_t const uint32_t uint32_t uint16_t uint32_t uint32_t * | cuddl::detail::range_end_scan |
| __shared__ block_scan::TempStorage | cuddl::detail::storage |
| uint32_t | cuddl::detail::carry = 0U |
| static __global__ uint32_t const uint32_t uint32_t const * | cuddl::detail::postings |
| static __global__ uint32_t const uint32_t uint32_t const uint32_t * | cuddl::detail::match_counts |
| auto const | cuddl::detail::stride = static_cast<uint32_t>(gridDim.x) * block_size |
| __global__ size_t | cuddl::detail::query_row_offset |
| __global__ size_t uint32_t | cuddl::detail::query_count |
| constexpr uint32_t | cuddl::detail::index_tile_lane_postings = 16U |
| Longest posting range one count_batch_index_tile_kernel lane walks alone. | |
| constexpr uint32_t | cuddl::detail::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 | cuddl::detail::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 * | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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 * | cuddl::detail::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 * | cuddl::detail::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 * | cuddl::detail::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 * | cuddl::detail::bitmap_counts |
| constexpr uint32_t | cuddl::detail::warps_per_block = index_tile_block_size / warp_width |
| __shared__ uint32_t | cuddl::detail::packed_counts [] |
| auto const | cuddl::detail::tiles = (reference_count + tile - 1U) / tile |
| auto const | cuddl::detail::query_index = static_cast<uint32_t>(blockIdx.x) / tiles |
| auto | cuddl::detail::low = static_cast<uint32_t>(blockIdx.x) % tiles * tile |
| auto const | cuddl::detail::high = min(low + tile, reference_count) |
| auto const | cuddl::detail::width = high - low |
| __shared__ uint2 | cuddl::detail::long_ranges [index_tile_long_capacity] |
| __shared__ uint32_t | cuddl::detail::long_count |
| auto const | cuddl::detail::lane = static_cast<uint32_t>(threadIdx.x) % warp_width |
| auto const | cuddl::detail::warp = static_cast<uint32_t>(threadIdx.x) / warp_width |
| cuddl::detail::else | |
| auto const * | cuddl::detail::query = queries + (query_row_offset + query_index) * BucketCount |
| auto const | cuddl::detail::count |
| auto const | cuddl::detail::long_ranges_used |
| auto *const | cuddl::detail::bits |
| constexpr size_t | cuddl::detail::refine_block_row_bytes = size_t{8} << 20U |
| Reference-row bytes one refine_batch_index_candidates_kernel block spans. | |
| constexpr uint32_t | cuddl::detail::score_planes = 16U |
| Bits per score, and so bit-planes per 32-bucket group. | |
| __global__ uint32_t | cuddl::detail::row_count |
| __global__ uint32_t uint32_t * | cuddl::detail::planes |
| constexpr uint32_t | cuddl::detail::groups_per_row = static_cast<uint32_t>(BucketCount / warp_width) |
| __global__ size_t uint16_t * | cuddl::detail::rows |
| __global__ uint32_t const * | cuddl::detail::reference_planes |
| __global__ uint32_t const uint32_t uint32_t const * | cuddl::detail::candidate_ids |
| __global__ uint32_t const uint32_t uint32_t const uint32_t const * | cuddl::detail::candidate_count |
| __global__ uint32_t const uint32_t uint32_t const uint32_t const SearchResult * | cuddl::detail::results |
| constexpr uint32_t | cuddl::detail::row_words = static_cast<uint32_t>(BucketCount / 2U) |
| auto const | cuddl::detail::items = candidate_ids == nullptr ? reference_count : *candidate_count |
| auto const | cuddl::detail::pairs = static_cast<uint64_t>(query_count) * reference_count |
| auto const | cuddl::detail::indexed_groups = indexed_bucket_count / warp_width |
| constexpr uint32_t | cuddl::detail::bitmap_refine_block_size = 768U |
| Threads per refine_batch_bitmap_kernel block. | |
| constexpr uint32_t | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::bitmap_refine_group |
| Queries one refine_batch_bitmap_kernel block refines together. | |
| template<size_t BucketCount> | |
| constexpr size_t | cuddl::detail::bitmap_refine_dynamic_bytes |
| Shared bytes one refine_batch_bitmap_kernel block stages. | |
| __global__ uint32_t const uint32_t const * | cuddl::detail::pass_bits |
| __global__ uint32_t const uint32_t const uint32_t const * | cuddl::detail::word_offsets |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t | cuddl::detail::first_query_id |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t | cuddl::detail::word_begin |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t | cuddl::detail::word_end |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t | cuddl::detail::packed_base |
| __global__ uint32_t const uint32_t const uint32_t const uint32_t uint32_t size_t size_t uint32_t SearchResult * | cuddl::detail::packed |
| __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 |
| auto const | cuddl::detail::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 * | cuddl::detail::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 * | cuddl::detail::cell_counter |
| constexpr uint32_t | cuddl::detail::group = QueryGroup |
| constexpr bool | cuddl::detail::staged = bitmap_refine_staged<BucketCount> |
| constexpr uint32_t | cuddl::detail::lane_groups = groups_per_row / warp_width |
| constexpr uint32_t | cuddl::detail::cached_groups = lane_groups < 2U ? lane_groups : 2U |
| constexpr uint32_t | cuddl::detail::max_block_words |
| constexpr uint32_t | cuddl::detail::word_slices = Candidates == refine_candidates::all ? 1U : 4U |
| __shared__ uint4 | cuddl::detail::query_plane_storage [] |
| __shared__ uint32_t | cuddl::detail::next_word |
| __shared__ uint32_t | cuddl::detail::claimed_cell |
| auto const | cuddl::detail::cells |
| auto const | cuddl::detail::groups = cells.groups |
| auto const | cuddl::detail::thresholded = CountMatches && pass_bits != nullptr |
| constexpr uint32_t | cuddl::detail::match_filter_block_size = 512U |
| Threads per match_filter_kernel block. | |
| template<size_t BucketCount> | |
| constexpr bool | cuddl::detail::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 | cuddl::detail::match_filter_planes = 12U |
| Low planes match_filter_kernel compares. | |
| constexpr uint32_t | cuddl::detail::all_pairs_block_size = 512U |
| Threads per refine_all_pairs_kernel block. | |
| constexpr uint32_t | cuddl::detail::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 * | cuddl::detail::reference_mask |
| auto const | cuddl::detail::warps = uint64_t{gridDim.x} * (block_size / warp_width) |
1.9.8