cuDDL
Loading...
Searching...
No Matches
replicated_database.cuh
Go to the documentation of this file.
1#pragma once
2
3#include <cuda/buffer>
4#include <cuda/memory_pool>
5#include <cuda/stream>
6
7#include <condition_variable>
8#include <exception>
9#include <filesystem>
10#include <memory>
11#include <mutex>
12#include <optional>
13#include <span>
14#include <thread>
15#include <vector>
16
18
19namespace cuddl {
20
26 cuda::stream_ref stream;
28 uint32_t replica{};
30
35 return static_cast<host_batch_results const*>(host);
36 }
37};
38
46 bool in_query_order = false;
50 std::function<uint64_t()> next_tile{};
51};
52
71template <uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
73 public:
76
78 [[nodiscard]] static Result<replicated_database>
79 upload(reference_database_file const& file, std::span<cuda::device_ref const> devices) {
80 auto group = CUDDL_TRY(make(devices));
81 CUDDL_TRY(group.on_each([&](replica& r, uint32_t) -> Result<void> {
82 r.database.emplace(CUDDL_TRY((file.upload<K, BucketCount, Layout>(r.stream))));
83 return Ok();
84 }));
85 return Result<replicated_database>::ok(std::move(group));
86 }
87
91 [[nodiscard]] static Result<replicated_database> adopt(std::vector<database_type> databases) {
92 std::vector<cuda::device_ref> devices;
93 for (auto const& database : databases) {
94 if (database.metadata() != databases.front().metadata()) {
95 return Err(Error::invalid_argument("replicas must share their database metadata"));
96 }
97 devices.push_back(database.device());
98 }
99 auto group = CUDDL_TRY(make(devices));
100 for (size_t i = 0; i < databases.size(); ++i) {
101 group.replicas_[i]->database.emplace(std::move(databases[i]));
102 }
103 return Result<replicated_database>::ok(std::move(group));
104 }
105
108 return on_each([&](replica& r, uint32_t) -> Result<void> {
109 r.index.emplace(CUDDL_TRY(index_type::build_async(*r.database, r.stream, storage)));
110 r.stream.sync();
111 return Ok();
112 });
113 }
114
116 [[nodiscard]] Result<void> load_index(std::filesystem::path const& path) {
117 return on_each([&](replica& r, uint32_t) -> Result<void> {
118 r.index.emplace(CUDDL_TRY(reference_index_file::load(path, *r.database, r.stream)));
119 return Ok();
120 });
121 }
122
134 template <typename OnTile>
135 [[nodiscard]] Result<void> search_all_to_all(OnTile&& on_tile, replica_search split = {}) {
136 return search(on_tile, split, std::nullopt);
137 }
138
141 template <typename OnTile>
142 [[nodiscard]] Result<void>
143 search_all_to_all(OnTile&& on_tile, search_options options, replica_search split = {}) {
144 return search(on_tile, split, options);
145 }
146
148 [[nodiscard]] size_t size() const noexcept {
149 return replicas_.size();
150 }
151
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());
156 return devices;
157 }
158
160 [[nodiscard]] database_type const& database() const noexcept {
161 return *replicas_.front()->database;
162 }
163
164 private:
165 struct replica {
166 explicit replica(cuda::device_ref device)
167 : stream(device),
168 host(
169 stream,
170 detail::device_pinned_pool(device, detail::pinned_placement::device_node)
171 ),
172 workspace(stream, cuda::device_default_memory_pool(device)),
173 results(stream, cuda::device_default_memory_pool(device)) {}
174
175 // Declaration order is construction order: the stream outlives its users.
176 cuda::stream stream;
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;
182
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
189 );
190 }
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
194 );
195 }
196 }
197 };
198
199 replicated_database() = default;
200
201 static Result<replicated_database> make(std::span<cuda::device_ref const> devices) {
202 if (devices.empty()) {
203 return Err(Error::invalid_argument("a replicated database needs at least one device"));
204 }
205 replicated_database group;
206 try {
207 for (auto const device : devices) {
208 group.replicas_.push_back(std::make_unique<replica>(device));
209 }
210 } catch (std::exception const& error) {
211 return Err(Error::resource(error.what()));
212 }
213 return Result<replicated_database>::ok(std::move(group));
214 }
215
225 template <typename Body, typename Stop = void (*)()>
226 Result<void> on_each(Body&& body, Stop&& stop = [] {}) {
227 std::mutex mutex;
228 std::optional<Error> error;
229 std::exception_ptr exception;
230 auto const run = [&](size_t i) {
231 // An empty outcome means the replica stopped without an error to report.
232 auto const outcome = [&]() -> std::optional<Result<void>> {
233 try {
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&) {
238 // Another replica already recorded why the search stopped.
239 return std::nullopt;
240 } catch (...) {
241 std::lock_guard lock(mutex);
242 if (!exception) exception = std::current_exception();
243 stop();
244 return std::nullopt;
245 }
246 }();
247 if (outcome && !*outcome) {
248 std::lock_guard lock(mutex);
249 if (!error) error = outcome->error();
250 stop();
251 }
252 };
253 {
254 std::vector<std::jthread> threads;
255 for (size_t i = 1; i < replicas_.size(); ++i) threads.emplace_back(run, i);
256 run(0);
257 }
258 if (exception) std::rethrow_exception(exception);
259 if (error) return Err(*error);
260 return Ok();
261 }
262
264 struct search_stopped {};
265
266 template <typename OnTile>
267 Result<void>
268 search(OnTile& on_tile, replica_search split, std::optional<search_options> options) {
269 if (split.next_tile && split.in_query_order) {
270 return Err(
272 "tiles claimed across processes cannot be delivered in query order"
273 )
274 );
275 }
276 auto tiles = CUDDL_TRY(
279 split.shard,
280 static_cast<uint32_t>(replicas_.size()),
281 std::move(split.next_tile)
282 )
283 );
284 std::mutex order;
285 std::condition_variable turn;
286 size_t delivered = 0U;
287 bool stopped = false;
288 return on_each(
289 [&](replica& r, uint32_t replica_index) -> Result<void> {
290 auto const* index = r.index ? &*r.index : nullptr;
291 // Without an index or a threshold every pair is reported, so the dedicated
292 // exhaustive kernels run instead of counting matches only to discard nothing.
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);
300 return;
301 }
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);
307 ++delivered;
308 turn.notify_all();
309 };
310 auto const workspace = device_span<uint8_t>{r.workspace.data(), r.workspace.size()};
311 auto const results =
312 device_span<packed_pairwise_counts>{r.results.data(), r.results.size()};
313 if (exhaustive) {
314 CUDDL_TRY(r.database->search_all_to_all_async(
315 workspace, results, deliver, {}, r.stream, tiles
316 ));
317 } else {
318 CUDDL_TRY(r.database->search_all_to_all_async(
319 workspace, results, deliver, {}, *options, r.stream, index, tiles
320 ));
321 }
322 CUDDL_CUDA_TRY(r.stream.sync());
323 return Ok();
324 },
325 [&] {
326 tiles.cancel();
327 {
328 std::lock_guard lock(order);
329 stopped = true;
330 }
331 turn.notify_all();
332 }
333 );
334 }
335
336 std::vector<std::unique_ptr<replica>> replicas_;
337};
338
339} // namespace cuddl
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>.
Definition error.hpp:248
#define CUDDL_TRY(expr)
Propagates a cuddl::Result failure from the enclosing function (GNU statement expression).
Definition error.hpp:225
__global__ uint32_t reference_count
Definition kernels.cuh:437
__global__ size_t uint32_t const uint32_t const uint32_t uint32_t uint16_t uint32_t tile
Definition kernels.cuh:922
__global__ uint32_t const uint32_t uint32_t const uint32_t const SearchResult * results
Definition kernels.cuh:1544
constexpr uint32_t group
Definition kernels.cuh:1858
auto const tiles
Definition kernels.cuh:937
Batch operations over contiguous device rows.
Definition a48.hpp:14
Result< void > Ok() noexcept
Success return for Result<void>; same as return {}.
Definition error.hpp:181
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.
Definition error.hpp:176
static Error cuda(cudaError_t code, std::source_location location=std::source_location::current())
Definition error.hpp:90
static Error invalid_argument(std::string message)
Definition error.hpp:99
static Error resource(std::string message)
Definition error.hpp:103
Fallible API result: cuda::std::expected<T, Error> with cuDDL factories.
Definition error.hpp:118
static Result ok(T value)
Definition error.hpp:121
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.