cuDDL
Loading...
Searching...
No Matches
reference_index_file.cuh
Go to the documentation of this file.
1#pragma once
2
6
7#include <algorithm>
8#include <atomic>
9#include <thread>
10#include <vector>
11
12namespace cuddl::detail {
13
15 uint64_t digest = 0;
16
17 Result<void> bytes(void const* data, size_t size) {
18 digest = xxhash64(static_cast<uint8_t const*>(data), size, digest);
19 return Ok();
20 }
21
22 template <typename T>
24 item = database_file_little_endian(item);
25 return bytes(&item, sizeof(item));
26 }
27
28 template <typename T>
29 Result<void> words(std::span<T const> items) {
30 if constexpr (std::endian::native == std::endian::little) {
31 return bytes(items.data(), items.size_bytes());
32 } else {
33 std::vector<T> little_endian(items.begin(), items.end());
34 for (auto& item : little_endian) {
35 item = database_file_little_endian(item);
36 }
37 return bytes(little_endian.data(), items.size_bytes());
38 }
39 }
40};
41
44 CUDDL_TRY(hash.bytes("CUDDLDB-XXH64-v1", sizeof("CUDDLDB-XXH64-v1") - 1));
45 auto const metadata = db.metadata();
46 auto const& c = metadata.compatibility;
47 CUDDL_TRY(hash.value(c.kmer_length));
48 CUDDL_TRY(hash.value(c.bucket_count));
49 CUDDL_TRY(hash.value(c.indexed_bucket_count));
50 CUDDL_TRY(hash.value(c.score_encoder_identity));
51 CUDDL_TRY(hash.value(c.exponent_bits));
52 CUDDL_TRY(hash.value(c.mantissa_bits));
53 CUDDL_TRY(hash.value(c.hash_identity));
54 CUDDL_TRY(hash.value(c.hash_seed));
55 CUDDL_TRY(hash.value(c.canonicalisation_policy));
56 CUDDL_TRY(hash.value(c.blacklist_identity));
57 CUDDL_TRY(hash.value(c.blacklist_version));
58 CUDDL_TRY(hash.value(c.key_mask));
59 CUDDL_TRY(hash.value(metadata.reference_count));
60 CUDDL_TRY(hash.value(static_cast<uint64_t>(db.names().size())));
61 for (auto const& name : db.names()) {
62 CUDDL_TRY(hash.value(static_cast<uint64_t>(name.size())));
63 CUDDL_TRY(hash.bytes(name.data(), name.size()));
64 }
65 CUDDL_TRY(hash.words(db.rows()));
66 return hash.digest;
67}
68
69} // namespace cuddl::detail
70
71namespace cuddl {
72
75 public:
77 template <uint32_t K, size_t BucketCount, typename Layout>
78 [[nodiscard]] static Result<void> save(
81 std::filesystem::path const& path,
82 cuda::stream_ref stream
83 ) try {
84 return [&]() -> Result<void> {
85 detail::device_scope const scope{stream};
86 CUDDL_TRY(source.view(&index));
87 auto database = CUDDL_TRY(snapshot(source, stream));
88 auto storage = index.storage();
89 auto offsets = CUDDL_TRY(
90 download(index.index_offsets_.data(), index.index_offsets_.size(), stream)
91 );
92 auto posting_count =
93 storage == index_storage::dense ? offsets.back() : index.index_posting_capacity_;
94 auto postings =
95 CUDDL_TRY(download(index.index_postings_.data(), posting_count, stream));
96 auto keys =
97 CUDDL_TRY(download(index.index_keys_.data(), index.index_keys_.size(), stream));
98 auto digest = CUDDL_TRY(detail::reference_database_digest(database));
100 writer.temporary = path.string() + ".tmp.XXXXXX";
101 auto fd = ::mkstemp(writer.temporary.data());
102 if (fd < 0) {
103 return Err(Error::resource("cannot create index file beside: " + path.string()));
104 }
105 writer.output = ::fdopen(fd, "wb");
106 if (!writer.output) {
107 ::close(fd);
108 return Err(Error::resource("cannot open temporary index file"));
109 }
110 CUDDL_TRY(writer.bytes("CUDDLIX\0", 8));
111 CUDDL_TRY(writer.value(uint32_t{1}));
112 CUDDL_TRY(writer.value(storage == index_storage::dense ? uint32_t{0} : uint32_t{1}));
113 CUDDL_TRY(writer.value(digest));
114 CUDDL_TRY(writer.value(static_cast<uint64_t>(postings.size())));
115 CUDDL_TRY(writer.value(index.pair_work_));
116 CUDDL_TRY(writer.words(offsets));
117 CUDDL_TRY(writer.words(postings));
118 if constexpr (std::endian::native == std::endian::little) {
119 CUDDL_TRY(writer.bytes(keys.data(), keys.size() * sizeof(uint16_t)));
120 } else {
121 for (auto key : keys) {
122 CUDDL_TRY(writer.value(key));
123 }
124 }
125 CUDDL_TRY(writer.value(static_cast<uint32_t>(writer.checksum)));
126 if (std::fclose(std::exchange(writer.output, nullptr)) != 0) {
127 return Err(Error::resource("cannot close index file"));
128 }
129 std::filesystem::rename(writer.temporary, path);
130 writer.temporary.clear();
131 return Ok();
132 }();
133 } catch (cuda::cuda_error const& error) {
134 return Err(Error::cuda(static_cast<cudaError_t>(error.status())));
135 } catch (std::system_error const& error) {
136 return Err(Error::resource(error.what()));
137 } catch (std::bad_alloc const& error) {
138 return Err(Error::resource(error.what()));
139 }
140
143 template <uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
145 std::filesystem::path const& path,
147 cuda::stream_ref stream
148 ) {
149 detail::device_scope const scope{stream};
150 auto source = CUDDL_TRY(snapshot(database, stream));
151 auto decoded = CUDDL_TRY(read(path, std::move(source)));
152 auto index = CUDDL_CUDA_TRY((reference_index<K, BucketCount, Layout>(database, stream)));
153 index.index_offsets_ = CUDDL_CUDA_TRY(
154 cuda::make_device_buffer<uint32_t>(stream, stream.device(), decoded.offsets_)
155 );
156 index.index_keys_ = CUDDL_CUDA_TRY(
157 cuda::make_device_buffer<uint16_t>(stream, stream.device(), decoded.keys_)
158 );
159 index.index_posting_capacity_ = static_cast<size_t>(detail::indexed_posting_count(
160 database.reference_count(), database.metadata().compatibility
161 ));
162 index.index_postings_ = CUDDL_CUDA_TRY(
163 cuda::make_device_buffer<uint32_t>(
164 stream, stream.device(), index.index_posting_capacity_, cuda::no_init
165 )
166 );
167 if (!decoded.postings_.empty()) {
169 cuda::copy_bytes(
170 stream,
171 cuda::std::span{decoded.postings_.data(), decoded.postings_.size()},
172 cuda::std::span{index.index_postings_.data(), decoded.postings_.size()}
173 )
174 );
175 }
176 // The driver may still read the pageable host sources after the copies return.
177 CUDDL_CUDA_TRY(stream.sync());
178 CUDDL_TRY(index.build_key_directory(
179 database.reference_count(),
181 stream
182 ));
183 index.indexed_ = true;
184 index.pair_work_ = decoded.pair_work_;
185 {
186 auto const reference_count = database.reference_count();
187 auto const indexed_buckets = database.metadata().compatibility.indexed_bucket_count;
188 index.pair_fraction_ = reference_count == 0U || indexed_buckets == 0U
189 ? 0.0
190 : static_cast<double>(index.pair_work_) /
191 (static_cast<double>(reference_count) *
192 reference_count * indexed_buckets);
193 }
194 CUDDL_TRY(index.build_posting_bitmaps(database, stream));
195 return index;
196 }
197
198 private:
199 template <uint32_t K, size_t BucketCount, typename Layout>
201 snapshot(reference_database<K, BucketCount, Layout> const& source, cuda::stream_ref stream) {
202 CUDDL_TRY((detail::validate_non_indexed_score_compatibility<K, BucketCount, Layout>(
203 source.metadata().compatibility
204 )));
206 database.metadata_ = source.metadata();
207 database.names_.assign(source.names().begin(), source.names().end());
208 if (database.names_.size() != source.reference_count()) {
209 return Err(Error::invalid_argument("index files require named reference rows"));
210 }
211 auto scores = CUDDL_CUDA_TRY(
212 cuda::make_device_buffer<uint16_t>(
213 stream,
214 stream.device(),
215 static_cast<size_t>(source.reference_count()) * BucketCount,
216 cuda::no_init
217 )
218 );
219 CUDDL_TRY(source.copy_scores_async({scores.data(), scores.size()}, stream));
220 database.rows_ = CUDDL_TRY(download(scores.data(), scores.size(), stream));
221 return database;
222 }
223
224 reference_index_file() = delete;
225
226 struct decoded_index {
227 reference_database_file database_;
228 index_storage storage_;
229 unsigned long long pair_work_{};
230 std::vector<uint32_t> offsets_{};
231 std::vector<uint32_t> postings_{};
232 std::vector<uint16_t> keys_{};
233 };
234
235 [[nodiscard]] static Result<decoded_index>
236 read(std::filesystem::path const& path, reference_database_file database) try {
237 return [&]() -> Result<decoded_index> {
238 detail::database_file_reader reader{
239 std::ifstream(path, std::ios::binary | std::ios::ate)
240 };
241 auto size = reader.input.tellg();
242 if (size < 0) {
243 return Err(Error::resource("cannot open index file: " + path.string()));
244 }
245 reader.remaining = static_cast<uint64_t>(size);
246 reader.input.seekg(0);
247 char magic[8];
248 CUDDL_TRY(reader.bytes(magic, sizeof(magic)));
249 if (std::string_view(magic, 8) != std::string_view("CUDDLIX\0", 8)) {
250 return Err(Error::invalid_argument("not a cuDDL reference index file"));
251 }
252 uint32_t version{}, kind{};
253 CUDDL_TRY(reader.value(version));
254 CUDDL_TRY(reader.value(kind));
255 if (version != 1 || kind > 1) {
256 return Err(Error::invalid_argument("unsupported reference index format"));
257 }
258 uint64_t digest{};
259 CUDDL_TRY(reader.value(digest));
260 auto const source = database.metadata();
261 if (database.rows().size() !=
262 static_cast<uint64_t>(source.reference_count) * source.compatibility.bucket_count) {
263 return Err(
264 Error::invalid_argument("reference database has no complete stored rows")
265 );
266 }
267 if (digest != CUDDL_TRY(detail::reference_database_digest(database))) {
268 return Err(
269 Error::invalid_argument("reference index belongs to a different database")
270 );
271 }
272 uint64_t count{};
273 CUDDL_TRY(reader.value(count));
274 unsigned long long pair_work{};
275 CUDDL_TRY(reader.value(pair_work));
276 auto const metadata = database.metadata();
277 auto const capacity =
278 detail::indexed_posting_count(metadata.reference_count, metadata.compatibility);
279 auto const offset_count =
280 kind == 0 ? detail::indexed_cell_count(metadata.compatibility) + 1 : 0;
281 auto const key_count = kind == 1 ? capacity : 0;
282 if (capacity > std::numeric_limits<uint32_t>::max() || count > capacity ||
283 (kind == 1 && count != capacity) ||
284 offset_count > std::numeric_limits<size_t>::max() / sizeof(uint32_t) ||
285 capacity > std::numeric_limits<size_t>::max() / sizeof(uint32_t) ||
286 reader.remaining != (offset_count + count) * sizeof(uint32_t) +
287 key_count * sizeof(uint16_t) + sizeof(uint32_t)) {
288 return Err(Error::invalid_argument("invalid reference index extents"));
289 }
290 decoded_index result{
291 std::move(database), kind == 0 ? index_storage::dense : index_storage::sparse
292 };
293 result.pair_work_ = pair_work;
294 result.offsets_.resize(static_cast<size_t>(offset_count));
295 result.postings_.resize(static_cast<size_t>(count));
296 result.keys_.resize(static_cast<size_t>(key_count));
297 CUDDL_TRY(reader.words(result.offsets_));
298 CUDDL_TRY(reader.words(result.postings_));
299 CUDDL_TRY(reader.bytes(result.keys_.data(), result.keys_.size() * sizeof(uint16_t)));
300 if constexpr (std::endian::native != std::endian::little) {
301 for (auto& key : result.keys_) {
303 }
304 }
305 auto expected = static_cast<uint32_t>(reader.checksum);
306 uint32_t checksum{};
307 CUDDL_TRY(reader.value(checksum));
308 if (checksum != expected) {
309 return Err(Error::invalid_argument("reference index checksum mismatch"));
310 }
311 CUDDL_TRY(validate_payload(result));
312 return result;
313 }();
314 } catch (std::bad_alloc const& error) {
315 return Err(Error::resource(error.what()));
316 }
317
318 template <typename T>
319 static Result<std::vector<T>> download(T const* data, size_t size, cuda::stream_ref stream) {
320 std::vector<T> host(size);
321 if (size != 0) {
323 cuda::copy_bytes(
324 stream, cuda::std::span{data, size}, cuda::std::span{host.data(), host.size()}
325 )
326 );
327 CUDDL_CUDA_TRY(stream.sync());
328 }
329 return host;
330 }
331
334 static Result<void> validate_payload(decoded_index const& index) {
335 auto const& offsets_ = index.offsets_;
336 auto const& postings_ = index.postings_;
337 auto const& keys_ = index.keys_;
338 auto const m = index.database_.metadata();
339 auto const& c = m.compatibility;
340 auto const rows = index.database_.rows();
341 auto const sparse = index.storage_ == index_storage::sparse;
342 auto const key_count = static_cast<size_t>(c.key_mask) + 1;
343 if (!sparse) {
344 if (offsets_.front() != 0 || offsets_.back() != postings_.size()) {
345 return Err(Error::invalid_argument("invalid dense index endpoints"));
346 }
347 }
348 // Each bucket either passes or names why it fails; seen[id] == bucket marks an ID
349 // already listed in that bucket.
350 auto const check_bucket =
351 [&](uint32_t bucket, std::vector<uint32_t>& seen, size_t& indexed) -> char const* {
352 auto const score_of = [&](uint32_t id) {
353 return rows[static_cast<size_t>(id) * c.bucket_count + bucket];
354 };
355 if (sparse) {
356 auto start = static_cast<size_t>(bucket) * m.reference_count;
357 for (uint32_t i = 0; i < m.reference_count; ++i) {
358 auto pos = start + i;
359 auto id = postings_[pos];
360 if (id >= m.reference_count || seen[id] == bucket ||
361 (i != 0 && keys_[pos] < keys_[pos - 1]) ||
362 (i != 0 && keys_[pos] == keys_[pos - 1] && id <= postings_[pos - 1])) {
363 return "invalid sparse index postings";
364 }
365 seen[id] = bucket;
366 auto score = score_of(id);
367 auto key =
368 score == 0 || c.key_mask == 0xffffU ? score : (score & c.key_mask) + 1U;
369 if (keys_[pos] != key) return "sparse index key does not match reference";
370 }
371 return nullptr;
372 }
373 for (uint32_t id = 0; id < m.reference_count; ++id) indexed += score_of(id) != 0;
374 for (size_t key = 0; key < key_count; ++key) {
375 auto cell = static_cast<size_t>(bucket) * key_count + key;
376 auto begin = offsets_[cell], end = offsets_[cell + 1];
377 if (begin > end || end > postings_.size()) return "invalid dense index offsets";
378 for (auto pos = begin; pos < end; ++pos) {
379 auto id = postings_[pos];
380 if (id >= m.reference_count || seen[id] == bucket ||
381 (pos != begin && id <= postings_[pos - 1])) {
382 return "invalid dense index postings";
383 }
384 seen[id] = bucket;
385 auto score = score_of(id);
386 if (score == 0 || (score & c.key_mask) != key) {
387 return "dense index key does not match reference";
388 }
389 }
390 }
391 return nullptr;
392 };
393 auto const workers = std::max(
394 1U, std::min<unsigned>(std::thread::hardware_concurrency(), c.indexed_bucket_count)
395 );
396 std::atomic<uint32_t> next{0};
397 std::atomic<size_t> indexed{0};
398 std::atomic<char const*> failure{nullptr};
399 {
400 std::vector<std::jthread> pool;
401 for (unsigned worker = 0; worker < workers; ++worker) {
402 pool.emplace_back([&] {
403 std::vector<uint32_t> seen(
404 m.reference_count, std::numeric_limits<uint32_t>::max()
405 );
406 size_t local = 0;
407 for (uint32_t bucket; failure.load(std::memory_order_relaxed) == nullptr &&
408 (bucket = next++) < c.indexed_bucket_count;) {
409 if (auto const error = check_bucket(bucket, seen, local)) {
410 char const* none = nullptr;
411 failure.compare_exchange_strong(none, error);
412 }
413 }
414 indexed += local;
415 });
416 }
417 }
418 if (auto const error = failure.load()) return Err(Error::invalid_argument(error));
419 if (!sparse && indexed != postings_.size()) {
420 return Err(Error::invalid_argument("dense index omits reference scores"));
421 }
422 return Ok();
423 }
424};
425
426} // namespace cuddl
Host-owned reference sketches and labels, ready for binary storage or GPU upload.
std::span< std::string const > names() const noexcept
Reference labels in reference-ID order.
std::span< uint16_t const > rows() const noexcept
Winner scores in reference-ID order, one bucket row per genome.
reference_database_metadata metadata() const noexcept
Compatibility metadata and reference count.
Move-only owner of one immutable contiguous reference database.
uint32_t reference_count() const noexcept
Reference count.
Result< void > copy_scores_async(device_span< score_type > scores, cuda::stream_ref stream) const
Writes the reference scores, row-major in reference-ID order, to scores.
std::span< std::string const > names() const noexcept
Reference labels in reference-ID order.
reference_database_metadata metadata() const noexcept
Compatibility metadata and reference count.
Index-file persistence and the validation boundary for reference_index.
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.
static Result< void > save(reference_index< K, BucketCount, Layout > const &index, reference_database< K, BucketCount, Layout > const &source, std::filesystem::path const &path, cuda::stream_ref stream)
Atomically save an already built or loaded index.
Move-only acceleration data for an immutable reference database.
index_storage storage() const noexcept
Dense index when offsets cover every cell, sparse when only keys are stored.
#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
T database_file_little_endian(T value)
auto const count
Definition kernels.cuh:1031
__global__ uint16_t const uint32_t uint32_t uint32_t key_count
Definition kernels.cuh:472
@ host
Any node: for buffers host threads fill.
__host__ __device__ uint64_t xxhash64(const uint8_t *bytes, uint64_t size, uint64_t seed)
Computes xxHash-64 over a byte range without allocating or copying it.
Definition hash.cuh:63
Result< uint64_t > reference_database_digest(reference_database_file const &db)
unsigned long long local
Definition kernels.cuh:478
constexpr uint64_t indexed_posting_count(uint32_t reference_count, score_compatibility const &compatibility) noexcept
::cuda::__ensure_current_context device_scope
Makes stream's device current for the enclosing scope and restores the previous context on exit.
constexpr uint64_t indexed_cell_count(score_compatibility const &compatibility) noexcept
auto const items
Definition kernels.cuh:1549
__host__ __device__ uint16_t score(uint64_t hash) noexcept
Computes the 16-bit DDL rarity score for hash.
Definition register.cuh:43
__global__ size_t uint16_t * rows
Definition kernels.cuh:1499
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
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
Result< void > words(std::vector< T > const &items)
Result< void > bytes(void const *source, size_t size)
Result< void > words(std::span< T const > items)
Result< void > bytes(void const *data, size_t size)