cuddl::host_batch_results Class Reference

Host copy of a batch_result_tile holding only the passing results, packed in query-major, reference order, in pinned memory; tile views it. More...

Public Member Functions
	host_batch_results (cuda::stream_ref stream)
	Empty results to download into, with pinned buffers for stream's device.
	host_batch_results (cuda::stream_ref stream, cuda::pinned_memory_pool_ref pool)
	Empty results to download into, with pinned buffers taken from pool.
	host_batch_results (host_batch_results const &)=delete
host_batch_results &	operator= (host_batch_results const &)=delete
__host__	host_batch_results (host_batch_results &&other) noexcept
__host__	~host_batch_results ()
std::vector< batch_search_result >	passing () const
	Every passing result, query-major and in reference order, as full values.
std::span< uint32_t const >	passing_match_counts () const &noexcept
	The passing results' index match counts, in passing order; empty unless match counts were requested.
void	passing_match_counts () const &&=delete
batch_result_tile	tile () const &noexcept
	View of the packed host results; valid while this object lives.
void	tile () const &&=delete

Friends
Result< void >	download (batch_result_tile const &, cuda::stream_ref, host_batch_results &)
	Copies a device tile's passing results into host's pinned buffers, growing them only when needed, and synchronizes stream.

Detailed Description

Host copy of a batch_result_tile holding only the passing results, packed in query-major, reference order, in pinned memory; tile views it.

Downloading tile after tile into one object reuses its pinned buffers, which only grow when a tile needs more. That keeps pinned allocation, and the pool growth behind it, out of every tile after the first.

Definition at line 334 of file reference_database.cuh.

Constructor & Destructor Documentation

host_batch_results() [1/4]

cuddl::host_batch_results::host_batch_results ( cuda::stream_ref  stream )

inlineexplicit

Empty results to download into, with pinned buffers for stream's device.

Definition at line 337 of file reference_database.cuh.

        : host_batch_results(
              stream,
              detail::device_pinned_pool(stream.device(), detail::pinned_placement::device_node)
          ) {}

host_batch_results() [2/4]

cuddl::host_batch_results::host_batch_results ( cuda::stream_ref  stream, cuda::pinned_memory_pool_ref  pool  )

inline

Empty results to download into, with pinned buffers taken from pool.

Definition at line 344 of file reference_database.cuh.

        : pool_(pool),
          results_(stream, pool, 0, cuda::no_init),
          match_counts_(stream, pool, 0, cuda::no_init),
          pass_bits_(stream, pool, 0, cuda::no_init),
          word_offsets_(stream, pool, 0, cuda::no_init) {}

host_batch_results() [3/4]

cuddl::host_batch_results::host_batch_results ( host_batch_results const &  )

delete

host_batch_results() [4/4]

__host__ cuddl::host_batch_results::host_batch_results ( host_batch_results &&  other )

inlinenoexcept

Definition at line 354 of file reference_database.cuh.

        : pool_(other.pool_),
          results_(std::move(other.results_)),
          match_counts_(std::move(other.match_counts_)),
          pass_bits_(std::move(other.pass_bits_)),
          word_offsets_(std::move(other.word_offsets_)),
          count_(std::exchange(other.count_, 0U)),
          has_match_counts_(other.has_match_counts_),
          layout_(other.layout_) {}

~host_batch_results()

__host__ cuddl::host_batch_results::~host_batch_results ( )

inline

Definition at line 364 of file reference_database.cuh.

{}

Member Function Documentation

operator=()

host_batch_results & cuddl::host_batch_results::operator= ( host_batch_results const &  )

delete

passing()

std::vector< batch_search_result > cuddl::host_batch_results::passing ( ) const

inline

Every passing result, query-major and in reference order, as full values.

This builds a 24-byte value per result; tile and for_each_passing read the 8-byte packed results without the copy.

Definition at line 370 of file reference_database.cuh.

                                                                 {
        std::vector<batch_search_result> kept;
        kept.reserve(static_cast<size_t>(count_));
        tile().for_each_passing([&](batch_search_result const& result) { kept.push_back(result); });
        return kept;
    }

Here is the call graph for this function:

passing_match_counts() [1/2]

void cuddl::host_batch_results::passing_match_counts ( ) const &&

delete

passing_match_counts() [2/2]

std::span< uint32_t const > cuddl::host_batch_results::passing_match_counts ( ) const &

inlinenoexcept

The passing results' index match counts, in passing order; empty unless match counts were requested.

Definition at line 380 of file reference_database.cuh.

                                                                                 {
        return {match_counts_.data(), has_match_counts_ ? static_cast<size_t>(count_) : 0U};
    }

tile() [1/2]

void cuddl::host_batch_results::tile ( ) const &&

delete

tile() [2/2]

batch_result_tile cuddl::host_batch_results::tile ( ) const &

inlinenoexcept

View of the packed host results; valid while this object lives.

Definition at line 386 of file reference_database.cuh.

                                                           {
        auto view = layout_;
        view.results = results_.data();
        view.match_counts = has_match_counts_ ? match_counts_.data() : nullptr;
        view.pass_bits = layout_.pass_bits == nullptr ? nullptr : pass_bits_.data();
        view.word_offsets = layout_.pass_bits == nullptr ? nullptr : word_offsets_.data();
        return view;
    }

Here is the caller graph for this function:

Friends And Related Symbol Documentation

download

Result< void > download ( batch_result_tile const &  tile, cuda::stream_ref  stream, host_batch_results &  host  )

friend

Copies a device tile's passing results into host's pinned buffers, growing them only when needed, and synchronizes stream.

Exhaustive tiles are copied as they are, since every slot passes in order. Threshold tiles are packed on the device: a scan of the pass-bit words gives each word's first packed position, and a gather writes the passing records straight into the host buffer.

Definition at line 427 of file reference_database.cuh.

                                                                                         {
    detail::device_scope const scope{stream};
    host.layout_ = tile;
    host.layout_.word_offsets = nullptr;
    host.has_match_counts_ = tile.match_counts != nullptr;
    if (tile.pass_bits == nullptr) {
        auto const slots = static_cast<size_t>(tile.slot_count());
        host.count_ = slots;
        host.reserve(host.results_, slots, stream);
        if (slots != 0U) {
            CUDDL_CUDA_TRY(
                cuda::copy_bytes(
                    stream,
                    cuda::std::span{tile.results, slots},
                    cuda::std::span{host.results_.data(), slots}
                )
            );
        }
        if (host.has_match_counts_) {
            host.reserve(host.match_counts_, slots, stream);
            if (slots != 0U) {
                CUDDL_CUDA_TRY(
                    cuda::copy_bytes(
                        stream,
                        cuda::std::span{tile.match_counts, slots},
                        cuda::std::span{host.match_counts_.data(), slots}
                    )
                );
            }
        }
        CUDDL_CUDA_TRY(stream.sync());
        return Ok();
    }
 
    auto const words = static_cast<size_t>(tile.query_count) * tile.words_per_query();
    host.reserve(host.pass_bits_, words, stream);
    host.reserve(host.word_offsets_, words, stream);
    if (words == 0U) {
        host.count_ = 0U;
        CUDDL_CUDA_TRY(stream.sync());
        return Ok();
    }
    auto const popcounts = cuda::make_transform_iterator(tile.pass_bits, detail::word_popcount{});
    size_t scan_bytes = 0;
    CUDDL_CUDA_TRY(
        cub::DeviceScan::ExclusiveSum(
            nullptr,
            scan_bytes,
            popcounts,
            host.word_offsets_.data(),
            static_cast<int64_t>(words),
            stream.get()
        )
    );
    auto scan_storage = CUDDL_CUDA_TRY(
        cuda::make_device_buffer<uint8_t>(stream, stream.device(), scan_bytes, cuda::no_init)
    );
    CUDDL_CUDA_TRY(
        cub::DeviceScan::ExclusiveSum(
            scan_storage.data(),
            scan_bytes,
            popcounts,
            host.word_offsets_.data(),
            static_cast<int64_t>(words),
            stream.get()
        )
    );
    CUDDL_CUDA_TRY(
        cuda::copy_bytes(
            stream,
            cuda::std::span{tile.pass_bits, words},
            cuda::std::span{host.pass_bits_.data(), words}
        )
    );
    CUDDL_CUDA_TRY(stream.sync());
    host.count_ = static_cast<uint64_t>(host.word_offsets_.data()[words - 1U]) +
                  static_cast<uint64_t>(cuda::std::popcount(host.pass_bits_.data()[words - 1U]));
    auto const count = static_cast<size_t>(host.count_);
    host.reserve(host.results_, count, stream);
    if (host.has_match_counts_) host.reserve(host.match_counts_, count, stream);
    if (count == 0U) {
        CUDDL_CUDA_TRY(stream.sync());
        return Ok();
    }
    // Passing records are gathered into bounded device staging, so each chunk crosses the bus as
    // one bulk copy rather than as scattered small writes.
    constexpr size_t staging_records = size_t{1} << 23U;
    auto const staged = std::min(count, staging_records);
    auto staging = CUDDL_CUDA_TRY(
        cuda::make_device_buffer<packed_pairwise_counts>(
            stream, stream.device(), staged, cuda::no_init
        )
    );
    auto staging_matches = CUDDL_CUDA_TRY(
        cuda::make_device_buffer<uint32_t>(
            stream, stream.device(), tile.match_counts == nullptr ? 0U : staged, cuda::no_init
        )
    );
    auto const kernel = tile.upper_triangle
                            ? detail::gather_passing_kernel<true, packed_pairwise_counts>
                            : detail::gather_passing_kernel<false, packed_pairwise_counts>;
    auto const* offsets = host.word_offsets_.data();
    for (size_t first_word = 0U; first_word < words;) {
        auto const base = offsets[first_word];
        // The last word whose results all fit the staging buffer ends the chunk.
        auto const last_word = static_cast<size_t>(
            std::upper_bound(offsets + first_word + 1U, offsets + words, base + staged - 32U) -
            offsets
        );
        auto const end = last_word == words ? count : static_cast<size_t>(offsets[last_word]);
        kernel<<<
            detail::warp_grid_blocks(last_word - first_word),
            detail::block_size,
            0,
            stream.get()>>>(
            tile.results,
            tile.match_counts,
            tile.pass_bits,
            host.word_offsets_.data(),
            tile.first_query_id,
            tile.reference_count,
            first_word,
            last_word,
            base,
            staging.data(),
            tile.match_counts == nullptr ? nullptr : staging_matches.data()
        );
        CUDDL_CUDA_TRY(cudaGetLastError());
        auto const records = end - base;
        if (records != 0U) {
            CUDDL_CUDA_TRY(
                cuda::copy_bytes(
                    stream,
                    cuda::std::span{staging.data(), records},
                    cuda::std::span{host.results_.data() + base, records}
                )
            );
            if (tile.match_counts != nullptr) {
                CUDDL_CUDA_TRY(
                    cuda::copy_bytes(
                        stream,
                        cuda::std::span{staging_matches.data(), records},
                        cuda::std::span{host.match_counts_.data() + base, records}
                    )
                );
            }
        }
        first_word = last_word;
    }
    CUDDL_CUDA_TRY(stream.sync());
    return Ok();
}

---

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

• reference_database.cuh