|
| __host__ __device__ constexpr uint64_t | cuddl::detail::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 | cuddl::detail::top_k_heap_words (uint32_t k) noexcept |
| | Words of shared memory one heap takes: k keys, then their k packed counts.
|
| |
| __device__ void | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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 | cuddl::detail::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.
|
| |
| 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.
|
| |
| | cuddl::detail::if (query >=tile.query_count) return |
| |
| | cuddl::detail::top_k_warp_copy< true > (slots, row, heap) |
| |
| | cuddl::detail::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;}) |
| |
| | cuddl::detail::top_k_warp_copy< false > (slots, row, heap) |
| |
| 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.
|
| |
| | cuddl::detail::if (block >=static_cast< uint64_t >(row_words) *word_count) return |
| |
| | cuddl::detail::_Pragma ("unroll") for(uint32_t bit=0U |
| |
| std::vector< batch_search_result > | cuddl::select_top_k (std::vector< batch_search_result > rows, uint32_t k) |
| | The k best rows of each query among rows, as top_k_search ranks them, best first and query by query.
|
| |
