cuddl::detail::reference_database_view< K, BucketCount, Layout > Class Template Reference

Non-owning, trivially copyable view of one selected reference-row backing. More...

Public Types
using	score_type = uint16_t
using	layout_type = Layout
using	result_type = reference_search_result
using	batch_result_type = packed_pairwise_counts

Public Member Functions
__host__ __device__ constexpr	reference_database_view (device_span< uint32_t const > planes, reference_database_metadata metadata, device_span< uint32_t const > index_offsets={}, device_span< uint32_t const > index_postings={}, bool indexed=false, device_span< uint16_t const > index_keys={}, device_span< uint32_t const > key_directory={}, double index_pair_fraction=0.0, device_span< uint32_t const > bitmap_cells={}, device_span< uint32_t const > posting_bitmaps={}, reference_delta_view deltas={}) noexcept
__host__ __device__ constexpr device_span< uint32_t const >	planes () const noexcept
	Reference bit-plane rows (detail::score_plane_index layout), in reference-ID order.
constexpr reference_database_metadata	metadata () const noexcept
	Compatibility metadata and reference count.
constexpr uint32_t	reference_count () const noexcept
	Reference count.
constexpr bool	has_index () const noexcept
	True when an index backs this view.
constexpr size_t	persistent_row_bytes () const noexcept
	Bytes this view's own row store needs.
constexpr size_t	persistent_index_bytes () const noexcept
	Bytes this view's index needs for offsets, postings, and keys.
constexpr size_t	single_query_workspace_bytes () const noexcept
	Caller-owned bytes for one exhaustive query on this view.
constexpr uint32_t	single_query_result_count () const noexcept
	Results one exhaustive query writes: one per reference.
Result< size_t >	indexed_single_query_workspace_bytes (cuda::stream_ref stream) const
	Caller-owned bytes required by one positive-threshold indexed query.
Result< cuddl::batch_search_requirements >	batch_search_requirements (uint32_t query_count, cuda::stream_ref stream) const
	Storage reused while exhaustively searching query_count compact rows.
Result< cuddl::batch_search_requirements >	indexed_batch_search_requirements (uint32_t query_count, cuda::stream_ref stream) const
	Storage reused while searching query_count compact rows through the index.
Result< cuddl::batch_search_requirements >	all_to_all_search_requirements () const
	Storage reused while exhaustively searching every unique database-row pair.
Result< cuddl::batch_search_requirements >	indexed_all_to_all_search_requirements () const
	Storage for every unique database-row pair through the index, or with a threshold but no index; it also covers the exhaustive all-to-all search.
Result< void >	search_async (device_span< score_type const > query, score_compatibility const &query_compatibility, device_span< uint8_t > workspace, device_span< result_type > results, cuda::stream_ref stream) const
	Compares one compact query row with every reference row on stream.
Result< void >	search_indexed_async (device_span< score_type const > query, score_compatibility const &query_compatibility, device_span< uint8_t > workspace, device_span< result_type > results, device_span< uint32_t > result_count, search_options options, cuda::stream_ref stream) const
	Finds and exactly refines references meeting options on stream.
template<typename OnTile >
Result< void >	search_batch_async (device_span< score_type const > queries, score_compatibility const &query_compatibility, uint32_t query_id_offset, device_span< uint8_t > workspace, device_span< batch_result_type > results, OnTile &&on_tile, device_span< uint32_t > result_match_counts, cuda::stream_ref stream) const
	Exhaustively searches compact queries using bounded reusable storage.
template<typename OnTile >
Result< void >	search_batch_indexed_async (device_span< score_type const > queries, score_compatibility const &query_compatibility, uint32_t query_id_offset, device_span< uint8_t > workspace, device_span< batch_result_type > results, OnTile &&on_tile, device_span< uint32_t > result_match_counts, search_options options, cuda::stream_ref stream) const
	Indexed search for compact queries using bounded reusable storage.
template<typename OnTile >
Result< void >	search_all_to_all_async (device_span< uint8_t > workspace, device_span< batch_result_type > results, OnTile &&on_tile, device_span< uint32_t > result_match_counts, cuda::stream_ref stream, all_to_all_shard shard={}) const
	Exhaustively searches every unique database-row pair using bounded storage.
template<typename OnTile >
Result< void >	search_all_to_all_async (device_span< uint8_t > workspace, device_span< batch_result_type > results, OnTile &&on_tile, device_span< uint32_t > result_match_counts, cuda::stream_ref stream, all_to_all_tile_queue &tiles) const
	Exhaustive all-to-all search of the tiles claimed from tiles, which other searches, for example on other devices' replicas, may claim from concurrently.
template<typename OnTile >
Result< void >	search_all_to_all_indexed_async (device_span< uint8_t > workspace, device_span< batch_result_type > results, OnTile &&on_tile, device_span< uint32_t > result_match_counts, search_options options, cuda::stream_ref stream, all_to_all_shard shard={}) const
	Indexed search of every unique database-row pair using bounded storage.
template<typename OnTile >
Result< void >	search_all_to_all_indexed_async (device_span< uint8_t > workspace, device_span< batch_result_type > results, OnTile &&on_tile, device_span< uint32_t > result_match_counts, search_options options, cuda::stream_ref stream, all_to_all_tile_queue &tiles) const
	Indexed all-to-all search of the tiles claimed from tiles, which other searches may claim from concurrently.

Static Public Member Functions
static constexpr size_t	persistent_row_bytes (uint32_t reference_count) noexcept
	Bytes one compact row store needs for reference_count references.
static constexpr size_t	single_query_workspace_bytes (uint32_t reference_count) noexcept
	Caller-owned bytes for one exhaustive query: the query's bit-planes.
static constexpr uint32_t	single_query_result_count (uint32_t reference_count) noexcept
static constexpr uint32_t	all_to_all_result_capacity (uint32_t reference_count) noexcept

Detailed Description

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
class cuddl::detail::reference_database_view< K, BucketCount, Layout >

Non-owning, trivially copyable view of one selected reference-row backing.

Inputs, database rows, workspace, and results must remain valid until the supplied stream completes.

Definition at line 1129 of file reference_database.cuh.

Member Typedef Documentation

batch_result_type

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

using cuddl::detail::reference_database_view< K, BucketCount, Layout >::batch_result_type = packed_pairwise_counts

Definition at line 1138 of file reference_database.cuh.

layout_type

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

using cuddl::detail::reference_database_view< K, BucketCount, Layout >::layout_type = Layout

Definition at line 1136 of file reference_database.cuh.

result_type

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

using cuddl::detail::reference_database_view< K, BucketCount, Layout >::result_type = reference_search_result

Definition at line 1137 of file reference_database.cuh.

score_type

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

using cuddl::detail::reference_database_view< K, BucketCount, Layout >::score_type = uint16_t

Definition at line 1135 of file reference_database.cuh.

Constructor & Destructor Documentation

reference_database_view()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

__host__ __device__ constexpr cuddl::detail::reference_database_view< K, BucketCount, Layout >::reference_database_view ( device_span< uint32_t const >  planes, reference_database_metadata  metadata, device_span< uint32_t const >  index_offsets = {}, device_span< uint32_t const >  index_postings = {}, bool  indexed = false, device_span< uint16_t const >  index_keys = {}, device_span< uint32_t const >  key_directory = {}, double  index_pair_fraction = 0.0, device_span< uint32_t const >  bitmap_cells = {}, device_span< uint32_t const >  posting_bitmaps = {}, reference_delta_view  deltas = {}  )

inlineconstexprnoexcept

Definition at line 1140 of file reference_database.cuh.

                                                    {},
        device_span<uint32_t const> index_postings = {},
        bool indexed = false,
        device_span<uint16_t const> index_keys = {},
        device_span<uint32_t const> key_directory = {},
        double index_pair_fraction = 0.0,
        device_span<uint32_t const> bitmap_cells = {},
        device_span<uint32_t const> posting_bitmaps = {},
        reference_delta_view deltas = {}
    ) noexcept
        : planes_(planes),
          metadata_(metadata),
          index_offsets_(index_offsets),
          index_postings_(index_postings),
          index_keys_(index_keys),
          key_directory_(key_directory),
          bitmap_cells_(bitmap_cells),
          posting_bitmaps_(posting_bitmaps),
          deltas_(deltas),
          index_pair_fraction_(index_pair_fraction),
          indexed_(indexed) {}

Member Function Documentation

all_to_all_result_capacity()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

static constexpr uint32_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::all_to_all_result_capacity ( uint32_t  reference_count )

inlinestaticconstexprnoexcept

Definition at line 1300 of file reference_database.cuh.

               {
        if (reference_count < 2U) {
            return 0U;
        }
        auto const query_count = detail::all_to_all_tile_queries(reference_count);
        auto const pair_count = static_cast<uint64_t>(query_count) *
                                (2ULL * reference_count - query_count - 1ULL) / 2ULL;
        auto const tiles = (reference_count + query_count - 1U) / query_count;
        return static_cast<uint32_t>(pair_count * pipeline_slots(tiles));
    }

Here is the call graph for this function:

Here is the caller graph for this function:

all_to_all_search_requirements()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< cuddl::batch_search_requirements > cuddl::detail::reference_database_view< K, BucketCount, Layout >::all_to_all_search_requirements ( ) const

inline

Storage reused while exhaustively searching every unique database-row pair.

Definition at line 1290 of file reference_database.cuh.

                                                                                                {
        auto const slot = CUDDL_TRY(all_to_all_slot_pairs());
        if (slot > std::numeric_limits<uint32_t>::max() / 2U) {
            return Err(Error::resource("all-to-all tile exceeds 32-bit result counts"));
        }
        return make_batch_requirements(
            0U, pipeline_slots(all_to_all_tile_count()) * slot, false, 0U
        );
    }

Here is the call graph for this function:

Here is the caller graph for this function:

batch_search_requirements()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< cuddl::batch_search_requirements > cuddl::detail::reference_database_view< K, BucketCount, Layout >::batch_search_requirements ( uint32_t  query_count, cuda::stream_ref  stream  ) const

inline

Storage reused while exhaustively searching query_count compact rows.

Definition at line 1270 of file reference_database.cuh.

                                                                                 {
        auto const pair_count =
            CUDDL_TRY(dense_batch_pair_count(external_batch_tile(query_count, stream)));
        return make_batch_requirements(
            pair_count, pair_count, false, pair_count / reference_divisor()
        );
    }

has_index()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr bool cuddl::detail::reference_database_view< K, BucketCount, Layout >::has_index ( ) const

inlineconstexprnoexcept

True when an index backs this view.

Definition at line 1182 of file reference_database.cuh.

                                                            {
        return indexed_;
    }

indexed_all_to_all_search_requirements()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< cuddl::batch_search_requirements > cuddl::detail::reference_database_view< K, BucketCount, Layout >::indexed_all_to_all_search_requirements ( ) const

inline

Storage for every unique database-row pair through the index, or with a threshold but no index; it also covers the exhaustive all-to-all search.

Definition at line 1316 of file reference_database.cuh.

                                                   {
        auto requirements = CUDDL_TRY(indexed_all_to_all_tile_search_requirements(
            0U, detail::all_to_all_tile_queries(metadata_.reference_count)
        ));
        auto const exhaustive = CUDDL_TRY(all_to_all_search_requirements());
        requirements.maximum_pair_count =
            std::max(requirements.maximum_pair_count, exhaustive.maximum_pair_count);
        requirements.result_bytes = std::max(requirements.result_bytes, exhaustive.result_bytes);
        requirements.match_count_bytes =
            std::max(requirements.match_count_bytes, exhaustive.match_count_bytes);
        return requirements;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

indexed_batch_search_requirements()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< cuddl::batch_search_requirements > cuddl::detail::reference_database_view< K, BucketCount, Layout >::indexed_batch_search_requirements ( uint32_t  query_count, cuda::stream_ref  stream  ) const

inline

Storage reused while searching query_count compact rows through the index.

Definition at line 1280 of file reference_database.cuh.

                                                                                         {
        CUDDL_TRY(validate_index_storage());
        auto const pair_count =
            CUDDL_TRY(dense_batch_pair_count(external_batch_tile(query_count, stream)));
        return make_batch_requirements(
            pair_count, pair_count, true, pair_count / reference_divisor()
        );
    }

Here is the caller graph for this function:

indexed_single_query_workspace_bytes()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< size_t > cuddl::detail::reference_database_view< K, BucketCount, Layout >::indexed_single_query_workspace_bytes ( cuda::stream_ref  stream ) const

inline

Caller-owned bytes required by one positive-threshold indexed query.

Definition at line 1226 of file reference_database.cuh.

            {
        CUDDL_TRY(validate_index_storage());
        if (metadata_.reference_count == 0U) {
            return size_t{0};
        }
 
        size_t selection_bytes = 0;
        auto const ids = cuda::make_counting_iterator(uint32_t{0});
        auto const selection = cuda_try(
            cub::DeviceSelect::If(
                nullptr,
                selection_bytes,
                ids,
                static_cast<uint32_t*>(nullptr),
                static_cast<uint32_t*>(nullptr),
                static_cast<int64_t>(metadata_.reference_count),
                detail::minimum_match_predicate{nullptr, 1U},
                stream.get()
            )
        );
        if (!selection) {
            return Err(selection.error());
        }
        // The counting pass stages two words per indexed bucket in the selection storage.
        selection_bytes = std::max(
            selection_bytes,
            static_cast<size_t>(metadata_.compatibility.indexed_bucket_count) * 2U *
                sizeof(uint32_t)
        );
 
        constexpr size_t alignment_slack = plane_alignment - 1U + 3U + 255U;
        auto const arrays_bytes =
            query_plane_bytes(1U) +
            static_cast<size_t>(metadata_.reference_count) * 2U * sizeof(uint32_t);
        if (selection_bytes > std::numeric_limits<size_t>::max() - arrays_bytes - alignment_slack) {
            return Err(Error::resource("indexed single-query workspace size overflows"));
        }
        return arrays_bytes + alignment_slack + selection_bytes;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

metadata()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr reference_database_metadata cuddl::detail::reference_database_view< K, BucketCount, Layout >::metadata ( ) const

inlineconstexprnoexcept

Compatibility metadata and reference count.

Definition at line 1172 of file reference_database.cuh.

                                                                                  {
        return metadata_;
    }

persistent_index_bytes()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr size_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::persistent_index_bytes ( ) const

inlineconstexprnoexcept

Bytes this view's index needs for offsets, postings, and keys.

Definition at line 1197 of file reference_database.cuh.

                                                                           {
        return index_offsets_.size_bytes() + index_postings_.size_bytes() +
               index_keys_.size_bytes();
    }

persistent_row_bytes() [1/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr size_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::persistent_row_bytes ( ) const

inlineconstexprnoexcept

Bytes this view's own row store needs.

Definition at line 1192 of file reference_database.cuh.

                                                                         {
        return planes_.size_bytes();
    }

Here is the caller graph for this function:

persistent_row_bytes() [2/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

static constexpr size_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::persistent_row_bytes ( uint32_t  reference_count )

inlinestaticconstexprnoexcept

Bytes one compact row store needs for reference_count references.

Definition at line 1187 of file reference_database.cuh.

                                                                                                  {
        return static_cast<size_t>(reference_count) * BucketCount * sizeof(score_type);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

planes()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

__host__ __device__ constexpr device_span< uint32_t const > cuddl::detail::reference_database_view< K, BucketCount, Layout >::planes ( ) const

inlineconstexprnoexcept

Reference bit-plane rows (detail::score_plane_index layout), in reference-ID order.

Definition at line 1167 of file reference_database.cuh.

                            {
        return planes_;
    }

reference_count()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr uint32_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::reference_count ( ) const

inlineconstexprnoexcept

Reference count.

Definition at line 1177 of file reference_database.cuh.

                                                                      {
        return metadata_.reference_count;
    }

Here is the caller graph for this function:

search_all_to_all_async() [1/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

template<typename OnTile >

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_all_to_all_async ( device_span< uint8_t >  workspace, device_span< batch_result_type >  results, OnTile &&  on_tile, device_span< uint32_t >  result_match_counts, cuda::stream_ref  stream, all_to_all_shard  shard = {}  ) const

inline

Exhaustively searches every unique database-row pair using bounded storage.

on_tile receives each tile's batch_result_tile. Before returning, it must consume the tile synchronously or enqueue dependent work on stream because the supplied storage is reused by the next tile. shard restricts the search to one part of the pairs.

Definition at line 1670 of file reference_database.cuh.

                                 {}
    ) const {
        auto tiles = CUDDL_TRY(all_to_all_tile_queue::make(metadata_.reference_count, shard));
        return search_all_to_all_async(
            workspace, results, std::forward<OnTile>(on_tile), result_match_counts, stream, tiles
        );
    }

Here is the caller graph for this function:

search_all_to_all_async() [2/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

template<typename OnTile >

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_all_to_all_async ( device_span< uint8_t >  workspace, device_span< batch_result_type >  results, OnTile &&  on_tile, device_span< uint32_t >  result_match_counts, cuda::stream_ref  stream, all_to_all_tile_queue &  tiles  ) const

inline

Exhaustive all-to-all search of the tiles claimed from tiles, which other searches, for example on other devices' replicas, may claim from concurrently.

Definition at line 1687 of file reference_database.cuh.

            {
        device_scope const scope{stream};
        CUDDL_TRY(validate_stored_rows());
        CUDDL_TRY(validate_tile_queue(tiles));
        auto const requirements = CUDDL_TRY(all_to_all_search_requirements());
        CUDDL_TRY(validate_batch_outputs(requirements, results, result_match_counts));
        std::vector<all_to_all_tile_queue::tile> claimed;
        return pipeline_tiles(
            [&]() -> std::optional<uint32_t> {
                auto const tile = tiles.claim();
                if (!tile) return std::nullopt;
                claimed.push_back(*tile);
                return static_cast<uint32_t>(claimed.size() - 1U);
            },
            CUDDL_TRY(all_to_all_slot_pairs()),
            results,
            result_match_counts,
            [&](uint32_t index,
                device_span<batch_result_type> slot_results,
                device_span<uint32_t> slot_matches,
                cuda::stream_ref work) {
                auto const tile = claimed[index];
                return launch_batch_exhaustive<true>(
                    {},
                    tile.first_query_id,
                    tile.query_count,
                    tile.first_query_id,
                    workspace,
                    slot_results,
                    slot_matches,
                    work
                );
            },
            on_tile,
            stream
        );
    }

Here is the call graph for this function:

search_all_to_all_indexed_async() [1/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

template<typename OnTile >

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_all_to_all_indexed_async ( device_span< uint8_t >  workspace, device_span< batch_result_type >  results, OnTile &&  on_tile, device_span< uint32_t >  result_match_counts, search_options  options, cuda::stream_ref  stream, all_to_all_shard  shard = {}  ) const

inline

Indexed search of every unique database-row pair using bounded storage.

on_tile receives each tile's batch_result_tile, holding exactly the pairs with at least options.minimum_matches index matches. Before returning, it must consume the tile synchronously or enqueue dependent work on stream because the supplied storage is reused by the next tile. shard restricts the search to one part of the pairs.

Definition at line 1741 of file reference_database.cuh.

                                 {}
    ) const {
        auto tiles = CUDDL_TRY(all_to_all_tile_queue::make(metadata_.reference_count, shard));
        return search_all_to_all_indexed_async(
            workspace,
            results,
            std::forward<OnTile>(on_tile),
            result_match_counts,
            options,
            stream,
            tiles
        );
    }

search_all_to_all_indexed_async() [2/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

template<typename OnTile >

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_all_to_all_indexed_async ( device_span< uint8_t >  workspace, device_span< batch_result_type >  results, OnTile &&  on_tile, device_span< uint32_t >  result_match_counts, search_options  options, cuda::stream_ref  stream, all_to_all_tile_queue &  tiles  ) const

inline

Indexed all-to-all search of the tiles claimed from tiles, which other searches may claim from concurrently.

Definition at line 1765 of file reference_database.cuh.

            {
        device_scope const scope{stream};
        CUDDL_TRY(validate_stored_rows());
        CUDDL_TRY(validate_tile_queue(tiles));
        auto const requirements = CUDDL_TRY(indexed_all_to_all_search_requirements());
        CUDDL_TRY(validate_indexed_batch_inputs(
            requirements, workspace, results, result_match_counts, options
        ));
        // Every tile reuses the same storage, so the next tile is claimed once the previous
        // tile's consumers are done; tiles claimed from a shared queue then go to whichever
        // device frees up first.
        std::optional<cuda::event> consumed;
        while (true) {
            if (consumed) CUDDL_CUDA_TRY(consumed->sync());
            auto const tile = tiles.claim();
            if (!tile) break;
            on_tile(CUDDL_TRY(
                launch_batch_indexed<true>(
                    nullptr,
                    tile->first_query_id,
                    tile->query_count,
                    tile->first_query_id,
                    workspace,
                    results,
                    result_match_counts,
                    options,
                    stream
                )
            ));
            if (!consumed) consumed.emplace(CUDDL_CUDA_TRY(cuda::event(stream.device())));
            CUDDL_CUDA_TRY(consumed->record(stream));
        }
        return Ok();
    }

Here is the call graph for this function:

search_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_async ( device_span< score_type const >  query, score_compatibility const &  query_compatibility, device_span< uint8_t >  workspace, device_span< result_type >  results, cuda::stream_ref  stream  ) const

inline

Compares one compact query row with every reference row on stream.

Definition at line 1330 of file reference_database.cuh.

            {
        device_scope const scope{stream};
        auto const expected_scores = static_cast<size_t>(metadata_.reference_count) * BucketCount;
        if (!rows_match_metadata(expected_scores)) {
            return Err(Error::invalid_argument("database extent does not match its metadata"));
        }
        if (query.size() != BucketCount || query.data() == nullptr) {
            return Err(Error::invalid_argument("query must contain one complete score row"));
        }
        CUDDL_TRY(
            (detail::validate_score_compatibility<K, BucketCount, Layout>(query_compatibility))
        );
        if (query_compatibility != metadata_.compatibility) {
            return Err(Error::invalid_argument("query construction metadata is incompatible"));
        }
        if (workspace.size_bytes() < single_query_workspace_bytes()) {
            return Err(Error::resource("single-query workspace is too small"));
        }
        if (results.size() < metadata_.reference_count) {
            return Err(Error::resource("single-query result capacity is too small"));
        }
        if (metadata_.reference_count == 0U) {
            return Ok();
        }
        if (results.data() == nullptr) {
            return Err(Error::invalid_argument("result storage must be device accessible"));
        }
 
        if (workspace.data() == nullptr) {
            return Err(Error::invalid_argument("single-query workspace must be device accessible"));
        }
        auto* query_planes = CUDDL_TRY(stage_query_planes(query, 0U, 1U, workspace, stream));
        detail::single_query_planes_kernel<BucketCount>
            <<<detail::warp_grid_blocks(metadata_.reference_count),
               detail::block_size,
               0,
               stream.get()>>>(
                query_planes,
                planes_.data(),
                metadata_.reference_count,
                nullptr,
                nullptr,
                results.data()
            );
        return cuda_try(cudaGetLastError());
    }

Here is the call graph for this function:

Here is the caller graph for this function:

search_batch_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

template<typename OnTile >

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_batch_async ( device_span< score_type const >  queries, score_compatibility const &  query_compatibility, uint32_t  query_id_offset, device_span< uint8_t >  workspace, device_span< batch_result_type >  results, OnTile &&  on_tile, device_span< uint32_t >  result_match_counts, cuda::stream_ref  stream  ) const

inline

Exhaustively searches compact queries using bounded reusable storage.

on_tile receives each tile's batch_result_tile. Before returning, it must consume the tile synchronously or enqueue dependent work on stream because the supplied storage is reused by the next tile.

Definition at line 1563 of file reference_database.cuh.

            {
        device_scope const scope{stream};
        auto const query_count =
            CUDDL_TRY(validate_batch_queries(queries, query_compatibility, query_id_offset));
        auto const requirements = CUDDL_TRY(batch_search_requirements(query_count, stream));
        CUDDL_TRY(validate_batch_outputs(requirements, results, result_match_counts));
        CUDDL_TRY(validate_workspace(requirements, workspace));
        // Result storage holds one full-size tile; two half-size tiles alternate in it.
        auto const storage_queries = external_batch_tile(query_count, stream);
        auto const tile_size = storage_queries < 2U ? storage_queries : storage_queries / 2U;
        auto const tiles = tile_size == 0U ? 0U : (query_count + tile_size - 1U) / tile_size;
        auto const slot = static_cast<size_t>(tile_size) * metadata_.reference_count;
        uint32_t next_tile = 0U;
        return pipeline_tiles(
            [&]() -> std::optional<uint32_t> {
                if (next_tile >= tiles) return std::nullopt;
                return next_tile++;
            },
            slot,
            results,
            result_match_counts,
            [&](uint32_t tile,
                device_span<batch_result_type> slot_results,
                device_span<uint32_t> slot_matches,
                cuda::stream_ref work) {
                auto const first_query_id = tile * tile_size;
                return launch_batch_exhaustive<false>(
                    queries,
                    first_query_id,
                    std::min(tile_size, query_count - first_query_id),
                    query_id_offset + first_query_id,
                    workspace,
                    slot_results,
                    slot_matches,
                    work
                );
            },
            on_tile,
            stream
        );
    }

Here is the caller graph for this function:

search_batch_indexed_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

template<typename OnTile >

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_batch_indexed_async ( device_span< score_type const >  queries, score_compatibility const &  query_compatibility, uint32_t  query_id_offset, device_span< uint8_t >  workspace, device_span< batch_result_type >  results, OnTile &&  on_tile, device_span< uint32_t >  result_match_counts, search_options  options, cuda::stream_ref  stream  ) const

inline

Indexed search for compact queries using bounded reusable storage.

on_tile receives each tile's batch_result_tile, holding exactly the pairs with at least options.minimum_matches index matches. Before returning, it must consume the tile synchronously or enqueue dependent work on stream because the supplied storage is reused by the next tile.

Definition at line 1623 of file reference_database.cuh.

            {
        device_scope const scope{stream};
        auto const query_count =
            CUDDL_TRY(validate_batch_queries(queries, query_compatibility, query_id_offset));
        auto const requirements = CUDDL_TRY(indexed_batch_search_requirements(query_count, stream));
        CUDDL_TRY(validate_indexed_batch_inputs(
            requirements, workspace, results, result_match_counts, options
        ));
        auto const tile_size = external_batch_tile(query_count, stream);
        for (uint32_t first_query_id = 0U; tile_size != 0U && first_query_id < query_count;
             first_query_id += tile_size) {
            auto const tile_query_count = std::min(tile_size, query_count - first_query_id);
            on_tile(CUDDL_TRY(
                launch_batch_indexed<false>(
                    queries.data() + static_cast<size_t>(first_query_id) * BucketCount,
                    first_query_id,
                    tile_query_count,
                    query_id_offset + first_query_id,
                    workspace,
                    results,
                    result_match_counts,
                    options,
                    stream
                )
            ));
        }
        return Ok();
    }

Here is the call graph for this function:

search_indexed_async()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

Result< void > cuddl::detail::reference_database_view< K, BucketCount, Layout >::search_indexed_async ( device_span< score_type const >  query, score_compatibility const &  query_compatibility, device_span< uint8_t >  workspace, device_span< result_type >  results, device_span< uint32_t >  result_count, search_options  options, cuda::stream_ref  stream  ) const

inline

Finds and exactly refines references meeting options on stream.

Definition at line 1384 of file reference_database.cuh.

            {
        device_scope const scope{stream};
        auto const expected_scores = static_cast<size_t>(metadata_.reference_count) * BucketCount;
        if (!rows_match_metadata(expected_scores)) {
            return Err(Error::invalid_argument("database extent does not match its metadata"));
        }
        CUDDL_TRY(validate_index_storage());
        if (query.size() != BucketCount || query.data() == nullptr) {
            return Err(Error::invalid_argument("query must contain one complete score row"));
        }
        CUDDL_TRY(
            (detail::validate_score_compatibility<K, BucketCount, Layout>(query_compatibility))
        );
        if (query_compatibility != metadata_.compatibility) {
            return Err(Error::invalid_argument("query construction metadata is incompatible"));
        }
        if (options.minimum_matches > metadata_.compatibility.indexed_bucket_count) {
            return Err(Error::invalid_argument("minimum matches exceeds indexed bucket count"));
        }
        if (results.size() < metadata_.reference_count) {
            return Err(Error::resource("indexed result capacity is too small"));
        }
        if (result_count.empty()) {
            return Err(Error::resource("indexed result count capacity is too small"));
        }
        if (result_count.data() == nullptr) {
            return Err(Error::invalid_argument("result count storage must be device accessible"));
        }
        if (metadata_.reference_count != 0U && results.data() == nullptr) {
            return Err(Error::invalid_argument("result storage must be device accessible"));
        }
 
        if (options.minimum_matches == 0U) {
            CUDDL_TRY(search_async(query, query_compatibility, workspace, results, stream));
            return write_batch_result_count(metadata_.reference_count, result_count, stream);
        }
        if (metadata_.reference_count == 0U) {
            return write_batch_result_count(0U, result_count, stream);
        }
 
        auto const required_workspace = CUDDL_TRY(indexed_single_query_workspace_bytes(stream));
        if (workspace.size_bytes() < required_workspace) {
            return Err(Error::resource("indexed single-query workspace is too small"));
        }
        if (workspace.data() == nullptr) {
            return Err(Error::invalid_argument("indexed workspace must be device accessible"));
        }
 
        auto const workspace_begin = reinterpret_cast<uintptr_t>(workspace.data());
        auto const workspace_end = workspace_begin + workspace.size_bytes();
        if (workspace_end < workspace_begin) {
            return Err(Error::resource("indexed workspace address range overflows"));
        }
        auto* query_planes = CUDDL_TRY(stage_query_planes(query, 0U, 1U, workspace, stream));
        auto address = detail::align_up(
            reinterpret_cast<uintptr_t>(query_planes) + query_plane_bytes(1U), alignof(uint32_t)
        );
        auto* match_counts = reinterpret_cast<uint32_t*>(address);
        address += static_cast<size_t>(metadata_.reference_count) * sizeof(uint32_t);
        address = detail::align_up(address, alignof(uint32_t));
        auto* candidate_ids = reinterpret_cast<uint32_t*>(address);
        address += static_cast<size_t>(metadata_.reference_count) * sizeof(uint32_t);
        address = detail::align_up(address, 256U);
        if (address > workspace_end) {
            return Err(Error::resource("indexed workspace layout exceeds its capacity"));
        }
        auto* selection_workspace = reinterpret_cast<void*>(address);
        auto selection_bytes = static_cast<size_t>(workspace_end - address);
 
        if (!indexed_) {
            detail::count_plane_matches_kernel<BucketCount>
                <<<detail::warp_grid_blocks(metadata_.reference_count),
                   detail::block_size,
                   0,
                   stream.get()>>>(
                    query_planes,
                    1U,
                    planes_.data(),
                    metadata_.reference_count,
                    metadata_.compatibility.indexed_bucket_count,
                    metadata_.compatibility.key_mask,
                    match_counts
                );
            CUDDL_CUDA_TRY(cudaGetLastError());
        } else {
            CUDDL_CUDA_TRY(
                cuda::fill_bytes(
                    stream,
                    cuda::std::span{match_counts, static_cast<size_t>(metadata_.reference_count)},
                    0
                )
            );
            // The posting ranges live in the selection workspace, which is free until the
            // selection below runs on the same stream.
            auto const indexed_bucket_count = metadata_.compatibility.indexed_bucket_count;
            auto* range_begin = static_cast<uint32_t*>(selection_workspace);
            auto* range_end_scan = range_begin + indexed_bucket_count;
            detail::index_posting_ranges_kernel<<<
                1,
                detail::posting_range_block_size,
                0,
                stream.get()>>>(
                query.data(),
                index_offsets_.data(),
                index_keys_.empty() ? nullptr : index_keys_.data(),
                metadata_.reference_count,
                indexed_bucket_count,
                metadata_.compatibility.key_mask,
                range_begin,
                range_end_scan
            );
            CUDDL_CUDA_TRY(cudaGetLastError());
            auto const counting_blocks =
                static_cast<uint32_t>(
                    stream.device().attribute(cuda::device_attributes::multiprocessor_count)
                ) *
                static_cast<uint32_t>(stream.device().attribute(
                    cuda::device_attributes::max_threads_per_multiprocessor
                )) /
                detail::block_size;
            detail::count_balanced_index_matches_kernel<<<
                counting_blocks,
                detail::block_size,
                0,
                stream.get()>>>(
                range_begin,
                range_end_scan,
                indexed_bucket_count,
                index_postings_.data(),
                match_counts
            );
            CUDDL_CUDA_TRY(cudaGetLastError());
        }
 
        auto const ids = cuda::make_counting_iterator(uint32_t{0});
        CUDDL_CUDA_TRY(
            cub::DeviceSelect::If(
                selection_workspace,
                selection_bytes,
                ids,
                candidate_ids,
                result_count.data(),
                static_cast<int64_t>(metadata_.reference_count),
                detail::minimum_match_predicate{match_counts, options.minimum_matches},
                stream.get()
            )
        );
 
        detail::single_query_planes_kernel<BucketCount>
            <<<detail::warp_grid_blocks(metadata_.reference_count),
               detail::block_size,
               0,
               stream.get()>>>(
                query_planes,
                planes_.data(),
                metadata_.reference_count,
                candidate_ids,
                result_count.data(),
                results.data()
            );
        return cuda_try(cudaGetLastError());
    }

Here is the call graph for this function:

single_query_result_count() [1/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr uint32_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::single_query_result_count ( ) const

inlineconstexprnoexcept

Results one exhaustive query writes: one per reference.

Definition at line 1221 of file reference_database.cuh.

                                                                                {
        return metadata_.reference_count;
    }

Here is the caller graph for this function:

single_query_result_count() [2/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

static constexpr uint32_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::single_query_result_count ( uint32_t  reference_count )

inlinestaticconstexprnoexcept

Definition at line 1214 of file reference_database.cuh.

               {
        return reference_count;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

single_query_workspace_bytes() [1/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

constexpr size_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::single_query_workspace_bytes ( ) const

inlineconstexprnoexcept

Caller-owned bytes for one exhaustive query on this view.

Definition at line 1210 of file reference_database.cuh.

                                                                                 {
        return single_query_workspace_bytes(metadata_.reference_count);
    }

Here is the call graph for this function:

Here is the caller graph for this function:

single_query_workspace_bytes() [2/2]

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>

static constexpr size_t cuddl::detail::reference_database_view< K, BucketCount, Layout >::single_query_workspace_bytes ( uint32_t  reference_count )

inlinestaticconstexprnoexcept

Caller-owned bytes for one exhaustive query: the query's bit-planes.

Definition at line 1203 of file reference_database.cuh.

               {
        return reference_count == 0U ? 0U : query_plane_bytes(1U) + plane_alignment - 1U;
    }

Here is the call graph for this function:

Here is the caller graph for this function:

---

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

• reference_database.cuh