cuddl::batch_result_tile Struct Reference

Non-owning view of one batch-search tile's results. More...

Collaboration diagram for cuddl::batch_result_tile:

Public Member Functions
__host__ __device__ constexpr uint32_t	words_per_query () const noexcept
__host__ __device__ uint64_t	slot_count () const noexcept
	Slots of the tile's layout, filled or not.
__host__ __device__ bool	passed (uint32_t query_id, uint32_t reference_id) const noexcept
	Whether the pair has a result in this tile, in either order for all-to-all tiles.
__host__ __device__ cuda::std::optional< batch_search_result >	find (uint32_t query_id, uint32_t reference_id) const noexcept
	The pair's result, or nothing when it did not pass.
__host__ __device__ uint32_t	match_count (uint32_t query_id, uint32_t reference_id) const noexcept
	The pair's index match count, in either order for all-to-all tiles; the pair must pass and counts must be present.
template<typename F >
__host__ __device__ void	for_each_passing (F &&f) const
	Calls f with every passing result, query-major and in reference order.
template<typename F >
__host__ __device__ void	for_each_passing (uint32_t begin, uint32_t end, F &&f) const
	Calls f with the passing results of the tile's query rows [begin, end), in the order of the whole-tile visit.
__host__ __device__ uint64_t	passing_before (uint32_t query) const noexcept
	Number of passing results in the tile's query rows before row query, which may be query_count.
__host__ __device__ uint64_t	count () const noexcept
	Number of passing results.

Public Attributes
packed_pairwise_counts const *	results {}
uint32_t const *	match_counts {}
	Index match counts parallel to results, when requested; otherwise null.
uint32_t const *	pass_bits {}
	Passing pairs, when the tile applied a threshold; null when every slot is filled.
uint32_t const *	word_offsets {}
	For packed results: the number of passing results before each pass-bit word.
uint32_t	first_query_id {}
uint32_t	query_count {}
uint32_t	reference_count {}
uint32_t	bucket_count {}
	Buckets per sketch, from which both-empty counts are recovered.
bool	upper_triangle {}

Detailed Description

Non-owning view of one batch-search tile's results.

Results are stored as packed_pairwise_counts; a result's position identifies its pair, so the accessors return full batch_search_result values. On the device, results sit at fixed slots: one row of reference_count slots per query, or for all-to-all tiles the packed strict upper triangle of pairs with query_id < reference_id. Exhaustive tiles fill every slot. Threshold (indexed) tiles fill only the passing pairs' slots and mark them in pass_bits, one bit per reference in rows of (reference_count + 31) / 32 words per query. Host copies (download) hold only the passing results, packed in order, and locate them through word_offsets.

All-to-all tiles store each pair once, in the row of its smaller ID. The lookups take the pair in either order: asked for (a, b) with b < a, they answer from the stored (b, a), with lower and higher swapped, as pairwise_counts::mirrored gives. The pair is then in the tile whose queries include b.

Definition at line 147 of file reference_database.cuh.

Member Function Documentation

count()

__host__ __device__ uint64_t cuddl::batch_result_tile::count ( ) const

inlinenoexcept

Number of passing results.

Definition at line 264 of file reference_database.cuh.

                                                                      {
        if (pass_bits == nullptr) return slot_count();
        auto const words = static_cast<size_t>(query_count) * words_per_query();
        if (words == 0U) return 0U;
        if (word_offsets != nullptr) {
            return static_cast<uint64_t>(word_offsets[words - 1U]) +
                   static_cast<uint64_t>(cuda::std::popcount(pass_bits[words - 1U]));
        }
        uint64_t total = 0U;
        for (size_t w = 0; w < words; ++w) {
            total += static_cast<uint64_t>(cuda::std::popcount(pass_bits[w]));
        }
        return total;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

find()

__host__ __device__ cuda::std::optional< batch_search_result > cuddl::batch_result_tile::find ( uint32_t  query_id, uint32_t  reference_id  ) const

inlinenoexcept

The pair's result, or nothing when it did not pass.

All-to-all tiles answer in either order, with the counts seen from query_id.

Definition at line 195 of file reference_database.cuh.

                                                                  {
        if (!passed(query_id, reference_id)) return cuda::std::nullopt;
        if (upper_triangle && reference_id < query_id) {
            auto const stored = results[position(reference_id, query_id)].unpack(bucket_count);
            return batch_search_result{query_id, reference_id, stored.mirrored()};
        }
        return result(query_id, reference_id, results[position(query_id, reference_id)]);
    }

Here is the call graph for this function:

for_each_passing() [1/2]

template<typename F >

__host__ __device__ void cuddl::batch_result_tile::for_each_passing ( F &&  f ) const

inline

Calls f with every passing result, query-major and in reference order.

Definition at line 216 of file reference_database.cuh.

                                                           {
        for_each_passing(0U, query_count, f);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

for_each_passing() [2/2]

template<typename F >

__host__ __device__ void cuddl::batch_result_tile::for_each_passing ( uint32_t  begin, uint32_t  end, F &&  f  ) const

inline

Calls f with the passing results of the tile's query rows [begin, end), in the order of the whole-tile visit.

Disjoint row ranges can be visited concurrently; the first result of row begin is passing result passing_before(begin) of the tile.

Definition at line 224 of file reference_database.cuh.

                                                                                         {
        auto next = begin < end ? stored_before(begin) : 0U;
        for (uint32_t query = begin; query < end; ++query) {
            auto const query_id = first_query_id + query;
            if (pass_bits == nullptr) {
                // Every slot passes, and slots are in this order.
                for (auto reference = upper_triangle ? query_id + 1U : 0U;
                     reference < reference_count;
                     ++reference) {
                    f(result(query_id, reference, results[next++]));
                }
                continue;
            }
            auto const words = words_per_query();
            for (uint32_t w = 0; w < words; ++w) {
                for (auto bits = pass_bits[static_cast<size_t>(query) * words + w]; bits != 0U;
                     bits &= bits - 1U) {
                    auto const reference =
                        w * 32U + static_cast<uint32_t>(cuda::std::countr_zero(bits));
                    auto const at = word_offsets != nullptr ? next++ : slot(query, reference);
                    f(result(query_id, reference, results[at]));
                }
            }
        }
    }

Here is the call graph for this function:

match_count()

__host__ __device__ uint32_t cuddl::batch_result_tile::match_count ( uint32_t  query_id, uint32_t  reference_id  ) const

inlinenoexcept

The pair's index match count, in either order for all-to-all tiles; the pair must pass and counts must be present.

Definition at line 207 of file reference_database.cuh.

                                                                         {
        if (upper_triangle && reference_id < query_id) {
            cuda::std::swap(query_id, reference_id);
        }
        return match_counts[position(query_id, reference_id)];
    }

passed()

__host__ __device__ bool cuddl::batch_result_tile::passed ( uint32_t  query_id, uint32_t  reference_id  ) const

inlinenoexcept

Whether the pair has a result in this tile, in either order for all-to-all tiles.

Definition at line 180 of file reference_database.cuh.

                                                                    {
        if (upper_triangle && reference_id < query_id) {
            cuda::std::swap(query_id, reference_id);
        }
        if (query_id < first_query_id || query_id - first_query_id >= query_count ||
            reference_id >= reference_count || (upper_triangle && reference_id <= query_id)) {
            return false;
        }
        if (pass_bits == nullptr) return true;
        return ((pass_bits[word(query_id, reference_id)] >> (reference_id % 32U)) & 1U) != 0U;
    }

Here is the caller graph for this function:

passing_before()

__host__ __device__ uint64_t cuddl::batch_result_tile::passing_before ( uint32_t  query ) const

inlinenoexcept

Number of passing results in the tile's query rows before row query, which may be query_count.

Definition at line 252 of file reference_database.cuh.

                                                                                             {
        if (query >= query_count) return count();
        if (pass_bits == nullptr || word_offsets != nullptr) return stored_before(query);
        auto const words = static_cast<size_t>(query) * words_per_query();
        uint64_t total = 0U;
        for (size_t w = 0; w < words; ++w) {
            total += static_cast<uint64_t>(cuda::std::popcount(pass_bits[w]));
        }
        return total;
    }

Here is the call graph for this function:

slot_count()

__host__ __device__ uint64_t cuddl::batch_result_tile::slot_count ( ) const

inlinenoexcept

Slots of the tile's layout, filled or not.

Definition at line 167 of file reference_database.cuh.

                                                                           {
        if (query_count == 0U) return 0U;
        auto const last = query_count - 1U;
        if (upper_triangle) {
            return first_query_id + last + 1U >= reference_count
                       ? slot(last, reference_count)
                       : slot(last, reference_count - 1U) + 1U;
        }
        return static_cast<uint64_t>(query_count) * reference_count;
    }

Here is the caller graph for this function:

words_per_query()

__host__ __device__ constexpr uint32_t cuddl::batch_result_tile::words_per_query ( ) const

inlineconstexprnoexcept

Definition at line 162 of file reference_database.cuh.

                                                                                          {
        return (reference_count + 31U) / 32U;
    }

Here is the caller graph for this function:

Member Data Documentation

bucket_count

uint32_t cuddl::batch_result_tile::bucket_count {}

Buckets per sketch, from which both-empty counts are recovered.

Definition at line 159 of file reference_database.cuh.

{};

first_query_id

uint32_t cuddl::batch_result_tile::first_query_id {}

Definition at line 155 of file reference_database.cuh.

{};

match_counts

uint32_t const* cuddl::batch_result_tile::match_counts {}

Index match counts parallel to results, when requested; otherwise null.

Definition at line 150 of file reference_database.cuh.

{};

pass_bits

uint32_t const* cuddl::batch_result_tile::pass_bits {}

Passing pairs, when the tile applied a threshold; null when every slot is filled.

Definition at line 152 of file reference_database.cuh.

{};

query_count

uint32_t cuddl::batch_result_tile::query_count {}

Definition at line 156 of file reference_database.cuh.

{};

reference_count

uint32_t cuddl::batch_result_tile::reference_count {}

Definition at line 157 of file reference_database.cuh.

{};

results

packed_pairwise_counts const* cuddl::batch_result_tile::results {}

Definition at line 148 of file reference_database.cuh.

{};

upper_triangle

bool cuddl::batch_result_tile::upper_triangle {}

Definition at line 160 of file reference_database.cuh.

{};

word_offsets

uint32_t const* cuddl::batch_result_tile::word_offsets {}

For packed results: the number of passing results before each pass-bit word.

Definition at line 154 of file reference_database.cuh.

{};

---

The documentation for this struct was generated from the following file:

• reference_database.cuh