cuDDL
Loading...
Searching...
No Matches
Classes | Public Member Functions | Static Public Member Functions | List of all members
cuddl::all_to_all_tile_queue Class Reference

Thread-safe source of the all-to-all tiles one process searches. More...

Classes

struct  tile
 

Public Member Functions

 all_to_all_tile_queue (all_to_all_tile_queue &&other) noexcept
 
std::optional< tile > claim () noexcept
 The next unclaimed tile, or nothing once all are claimed or after cancel.
 
void cancel () noexcept
 Makes every later claim of this queue return nothing, for example after a device failed.
 
bool shared () const noexcept
 Whether processes share the tiles through a next_tile counter.
 
uint32_t tile_queries () const noexcept
 Query rows per tile (the last tile may hold fewer).
 
size_t size () const noexcept
 Tiles this process owns, or every tile of the search when they are shared.
 
uint32_t reference_count () const noexcept
 Reference count the tiles cover.
 
size_t position (uint32_t first_query_id) const noexcept
 Position, in claim order, of the tile starting at first_query_id.
 

Static Public Member Functions

static Result< all_to_all_tile_queue > make (uint32_t reference_count, all_to_all_shard shard={}, uint32_t searchers=1U, std::function< uint64_t()> next_tile={})
 

Detailed Description

Thread-safe source of the all-to-all tiles one process searches.

Tiles hold detail::all_to_all_tile_queries(reference_count, shard.count * searchers) query rows, so each of the shard.count * searchers searches gets about four. The process keeps the tiles shard owns, and concurrent searches, typically one per device, claim them one at a time, largest first. A device that finishes early simply claims more, so devices of different speeds share the work without a fixed split.

With next_tile, processes share the tiles instead of keeping those shard assigns: next_tile returns the next unclaimed tile number of the whole search, counting from 0 across every process, and must not throw. A faster process then claims more tiles, as a faster device does within one process. shard.count still sizes the tiles; pass the number of processes. See shared_tile_counter in distributed.cuh for a counter shared over MPI.

Definition at line 1023 of file reference_database.cuh.

Constructor & Destructor Documentation

◆ all_to_all_tile_queue()

cuddl::all_to_all_tile_queue::all_to_all_tile_queue ( all_to_all_tile_queue &&  other)
inlinenoexcept

Definition at line 1058 of file reference_database.cuh.

1059 : tiles_(std::move(other.tiles_)),
1060 next_(other.next_.load()),
1061 next_tile_(std::move(other.next_tile_)),
1062 cancelled_(other.cancelled_.load()),
1063 tile_queries_(other.tile_queries_),
1064 reference_count_(other.reference_count_) {}

Member Function Documentation

◆ cancel()

void cuddl::all_to_all_tile_queue::cancel ( )
inlinenoexcept

Makes every later claim of this queue return nothing, for example after a device failed.

Other processes sharing the tiles keep claiming.

Definition at line 1077 of file reference_database.cuh.

1077 {
1078 cancelled_.store(true);
1079 }

◆ claim()

std::optional< tile > cuddl::all_to_all_tile_queue::claim ( )
inlinenoexcept

The next unclaimed tile, or nothing once all are claimed or after cancel.

Definition at line 1067 of file reference_database.cuh.

1067 {
1068 if (cancelled_.load(std::memory_order_relaxed)) return std::nullopt;
1069 auto const i = next_tile_ ? next_tile_() : next_.fetch_add(1U, std::memory_order_relaxed);
1070 if (i >= tiles_.size()) return std::nullopt;
1071 auto const first = tiles_[i] * tile_queries_;
1072 return tile{first, std::min(tile_queries_, reference_count_ - first)};
1073 }
__global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t tile
Definition kernels.cuh:922

◆ make()

static Result< all_to_all_tile_queue > cuddl::all_to_all_tile_queue::make ( uint32_t  reference_count,
all_to_all_shard  shard = {},
uint32_t  searchers = 1U,
std::function< uint64_t()>  next_tile = {} 
)
inlinestatic

Definition at line 1030 of file reference_database.cuh.

1032 {},
1033 uint32_t searchers = 1U,
1034 std::function<uint64_t()> next_tile = {}
1035 ) {
1036 if (shard.count == 0U || shard.index >= shard.count) {
1037 return Err(Error::invalid_argument("all-to-all shard index must be below its count"));
1038 }
1039 if (searchers == 0U || shard.count > std::numeric_limits<uint32_t>::max() / searchers) {
1040 return Err(Error::invalid_argument("all-to-all searcher count is out of range"));
1041 }
1043 queue.reference_count_ = reference_count;
1044 queue.tile_queries_ =
1045 detail::all_to_all_tile_queries(reference_count, shard.count * searchers);
1046 auto const count = queue.tile_queries_ == 0U
1047 ? 0U
1048 : (reference_count + queue.tile_queries_ - 1U) / queue.tile_queries_;
1049 for (uint32_t t = 0U; t < count; ++t) {
1050 auto const lane = t % shard.count;
1051 auto const owner = (t / shard.count) % 2U == 0U ? lane : shard.count - 1U - lane;
1052 if (next_tile || owner == shard.index) queue.tiles_.push_back(t);
1053 }
1054 queue.next_tile_ = std::move(next_tile);
1055 return Result<all_to_all_tile_queue>::ok(std::move(queue));
1056 }
all_to_all_tile_queue(all_to_all_tile_queue &&other) noexcept
uint32_t reference_count() const noexcept
Reference count the tiles cover.
auto const lane
Definition kernels.cuh:966
constexpr uint32_t all_to_all_tile_queries(uint32_t reference_count, uint32_t parts=1U) noexcept
Query rows one all-to-all tile compares: about 2^25 pairs, and at least batch_query_tile_count.
auto const count
Definition kernels.cuh:1031
cuda::std::unexpected< Error > Err(Error error)
Failure return; converts to any Result<T> via cuda::std::unexpected.
Definition error.hpp:176
static Error invalid_argument(std::string message)
Definition error.hpp:99
static Result ok(T value)
Definition error.hpp:121

◆ position()

size_t cuddl::all_to_all_tile_queue::position ( uint32_t  first_query_id) const
inlinenoexcept

Position, in claim order, of the tile starting at first_query_id.

Definition at line 1102 of file reference_database.cuh.

1102 {
1103 return static_cast<size_t>(
1104 std::lower_bound(tiles_.begin(), tiles_.end(), first_query_id / tile_queries_) -
1105 tiles_.begin()
1106 );
1107 }

◆ reference_count()

uint32_t cuddl::all_to_all_tile_queue::reference_count ( ) const
inlinenoexcept

Reference count the tiles cover.

Definition at line 1097 of file reference_database.cuh.

1097 {
1098 return reference_count_;
1099 }

◆ shared()

bool cuddl::all_to_all_tile_queue::shared ( ) const
inlinenoexcept

Whether processes share the tiles through a next_tile counter.

Definition at line 1082 of file reference_database.cuh.

1082 {
1083 return static_cast<bool>(next_tile_);
1084 }

◆ size()

size_t cuddl::all_to_all_tile_queue::size ( ) const
inlinenoexcept

Tiles this process owns, or every tile of the search when they are shared.

Definition at line 1092 of file reference_database.cuh.

1092 {
1093 return tiles_.size();
1094 }

◆ tile_queries()

uint32_t cuddl::all_to_all_tile_queue::tile_queries ( ) const
inlinenoexcept

Query rows per tile (the last tile may hold fewer).

Definition at line 1087 of file reference_database.cuh.

1087 {
1088 return tile_queries_;
1089 }

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