4#include <cuda/memory_pool>
7#include <condition_variable>
71template <u
int32_t K,
size_t BucketCount,
typename Layout = default_register_layout>
82 r.database.emplace(CUDDL_TRY((file.upload<K, BucketCount, Layout>(r.stream))));
92 std::vector<cuda::device_ref>
devices;
93 for (
auto const&
database : databases) {
100 for (
size_t i = 0; i < databases.size(); ++i) {
101 group.replicas_[i]->database.emplace(std::move(databases[i]));
108 return on_each([&](replica& r, uint32_t) ->
Result<void> {
117 return on_each([&](replica& r, uint32_t) ->
Result<void> {
134 template <
typename OnTile>
136 return search(on_tile, split, std::nullopt);
141 template <
typename OnTile>
142 [[nodiscard]] Result<void>
144 return search(on_tile, split, options);
148 [[nodiscard]]
size_t size() const noexcept {
149 return replicas_.size();
153 [[nodiscard]] std::vector<cuda::device_ref>
devices()
const {
154 std::vector<cuda::device_ref>
devices;
155 for (
auto const& r : replicas_)
devices.push_back(r->stream.device());
161 return *replicas_.front()->database;
166 explicit replica(cuda::device_ref device)
170 detail::device_pinned_pool(device, detail::pinned_placement::device_node)
172 workspace(stream,
cuda::device_default_memory_pool(device)),
173 results(stream,
cuda::device_default_memory_pool(device)) {}
177 host_batch_results host;
178 std::optional<database_type> database;
179 std::optional<index_type> index;
180 cuda::device_buffer<uint8_t> workspace;
181 cuda::device_buffer<packed_pairwise_counts> results;
184 void reserve(batch_search_requirements
const& requirements) {
185 auto const device = stream.device();
186 if (workspace.size() < requirements.workspace_bytes) {
187 workspace = cuda::make_device_buffer<uint8_t>(
188 stream, device, requirements.workspace_bytes, cuda::no_init
191 if (
results.size() < requirements.maximum_pair_count) {
192 results = cuda::make_device_buffer<packed_pairwise_counts>(
193 stream, device, requirements.maximum_pair_count, cuda::no_init
199 replicated_database() =
default;
201 static Result<replicated_database> make(std::span<cuda::device_ref const>
devices) {
205 replicated_database
group;
207 for (
auto const device :
devices) {
208 group.replicas_.push_back(std::make_unique<replica>(device));
210 }
catch (std::exception
const& error) {
225 template <
typename Body,
typename Stop =
void (*)()>
226 Result<void> on_each(Body&& body, Stop&& stop = [] {}) {
228 std::optional<Error> error;
229 std::exception_ptr exception;
230 auto const run = [&](
size_t i) {
232 auto const outcome = [&]() -> std::optional<Result<void>> {
234 return body(*replicas_[i],
static_cast<uint32_t
>(i));
235 }
catch (cuda::cuda_error
const& failure) {
236 return Err(
Error::cuda(
static_cast<cudaError_t
>(failure.status())));
237 }
catch (search_stopped
const&) {
241 std::lock_guard lock(mutex);
242 if (!exception) exception = std::current_exception();
247 if (outcome && !*outcome) {
248 std::lock_guard lock(mutex);
249 if (!error) error = outcome->error();
254 std::vector<std::jthread> threads;
255 for (
size_t i = 1; i < replicas_.size(); ++i) threads.emplace_back(run, i);
258 if (exception) std::rethrow_exception(exception);
259 if (error)
return Err(*error);
264 struct search_stopped {};
266 template <
typename OnTile>
268 search(OnTile& on_tile, replica_search split, std::optional<search_options> options) {
269 if (split.next_tile && split.in_query_order) {
272 "tiles claimed across processes cannot be delivered in query order"
280 static_cast<uint32_t
>(replicas_.size()),
281 std::move(split.next_tile)
285 std::condition_variable turn;
286 size_t delivered = 0U;
287 bool stopped =
false;
289 [&](replica& r, uint32_t replica_index) -> Result<void> {
290 auto const* index = r.index ? &*r.index :
nullptr;
293 bool const exhaustive =
294 !options || (index ==
nullptr && options->minimum_matches == 0U);
295 r.reserve(
CUDDL_TRY(r.database->all_to_all_search_requirements(index)));
296 auto const deliver = [&](batch_result_tile
const&
tile) {
297 replica_tile
const delivered_tile{
tile, r.stream, replica_index, &r.host};
298 if (!split.in_query_order) {
299 on_tile(delivered_tile);
302 auto const position =
tiles.position(
tile.first_query_id);
303 std::unique_lock lock(order);
304 turn.wait(lock, [&] {
return stopped || delivered == position; });
305 if (stopped)
throw search_stopped{};
306 on_tile(delivered_tile);
310 auto const workspace = device_span<uint8_t>{r.workspace.data(), r.workspace.size()};
312 device_span<packed_pairwise_counts>{r.results.data(), r.results.size()};
314 CUDDL_TRY(r.database->search_all_to_all_async(
315 workspace, results, deliver, {}, r.stream, tiles
318 CUDDL_TRY(r.database->search_all_to_all_async(
319 workspace, results, deliver, {}, *options, r.stream, index, tiles
328 std::lock_guard lock(order);
336 std::vector<std::unique_ptr<replica>> replicas_;
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={})
Host copy of a batch_result_tile holding only the passing results, packed in query-major,...
Host-owned reference sketches and labels, ready for binary storage or GPU upload.
Move-only owner of one immutable contiguous reference database.
reference_database_metadata metadata() const noexcept
Compatibility metadata and reference count.
cuda::device_ref device() const noexcept
Device holding the rows: the device of the stream that built them.
static Result< reference_index< K, BucketCount, Layout > > load(std::filesystem::path const &path, reference_database< K, BucketCount, Layout > const &database, cuda::stream_ref stream)
Validate the file and its database binding once, then restore a ready-to-search index.
Move-only acceleration data for an immutable reference database.
static Result< reference_index > build_async(database_type const &database, cuda::stream_ref stream, index_storage storage=index_storage::automatic)
Build acceleration data without taking ownership of or copying reference rows.
One reference database replicated on explicitly chosen devices, searched all-to-all by all of them at...
Result< void > search_all_to_all(OnTile &&on_tile, replica_search split={})
Exhaustively searches every unique database-row pair in split.shard on all replicas.
static Result< replicated_database > upload(reference_database_file const &file, std::span< cuda::device_ref const > devices)
Uploads file to each of devices, concurrently.
Result< void > load_index(std::filesystem::path const &path)
Loads the index file at path onto every replica.
database_type const & database() const noexcept
The database every replica holds, as uploaded to replica 0.
size_t size() const noexcept
Number of replicas, one per device.
std::vector< cuda::device_ref > devices() const
The devices, in replica order.
static Result< replicated_database > adopt(std::vector< database_type > databases)
Adopts databases that already hold identical rows, one per device, each searched on the device that b...
Result< void > build_index(index_storage storage=index_storage::automatic)
Builds an index on every replica.
Result< void > search_all_to_all(OnTile &&on_tile, search_options options, replica_search split={})
Threshold search of every unique pair in split.shard: pairs with at least options....
#define CUDDL_CUDA_TRY(...)
Propagates a CUDA error wrapped in cuddl::Result<void>.
#define CUDDL_TRY(expr)
Propagates a cuddl::Result failure from the enclosing function (GNU statement expression).
__global__ uint32_t reference_count
__global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t tile
__global__ uint32_t const uint32_t uint32_t const uint32_t const SearchResult * results
Batch operations over contiguous device rows.
Result< void > Ok() noexcept
Success return for Result<void>; same as return {}.
Result< void > download(batch_result_tile const &tile, cuda::stream_ref stream, host_batch_results &host)
Copies a device tile's passing results into host's pinned buffers, growing them only when needed,...
index_storage
Dense offsets favor query latency; sparse sorted keys reduce index memory and build time.
cuda::std::unexpected< Error > Err(Error error)
Failure return; converts to any Result<T> via cuda::std::unexpected.
static Error cuda(cudaError_t code, std::source_location location=std::source_location::current())
static Error invalid_argument(std::string message)
static Error resource(std::string message)
Fallible API result: cuda::std::expected<T, Error> with cuDDL factories.
static Result ok(T value)
Part cuDDL of count disjoint parts of an all-to-all search.
Non-owning view of one batch-search tile's results.
How a replicated_database search splits and delivers its tiles.
all_to_all_shard shard
This process's part of the pairs, when several processes share the search.
bool in_query_order
Deliver tiles one at a time in increasing query order, for example to write output that does not depe...
std::function< uint64_t()> next_tile
Claim tiles from a counter that several processes share instead of keeping the tiles shard assigns,...
One all-to-all tile as a replicated_database search delivers it.
host_batch_results * host
Result< host_batch_results const * > download() const
Downloads the tile into the replica's reusable pinned buffer, which lives on its device's host NUMA n...
batch_result_tile tile
Results in the replica's device memory; valid only during the callback.
uint32_t replica
Position of the replica's device in the devices the group was made for.
cuda::stream_ref stream
Stream ordering work on the tile's storage; work enqueued here finishes before reuse.
Matching-bucket threshold, independent of whether an index accelerates the search.