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

Index-file persistence and the validation boundary for reference_index. More...

Static Public Member Functions

template<uint32_t K, size_t BucketCount, typename Layout >
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.
 
template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
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.
 

Detailed Description

Index-file persistence and the validation boundary for reference_index.

Definition at line 74 of file reference_index_file.cuh.

Member Function Documentation

◆ load()

template<uint32_t K, size_t BucketCount, typename Layout = default_register_layout>
static Result< reference_index< K, BucketCount, Layout > > cuddl::reference_index_file::load ( std::filesystem::path const &  path,
reference_database< K, BucketCount, Layout > const &  database,
cuda::stream_ref  stream 
)
inlinestatic

Validate the file and its database binding once, then restore a ready-to-search index.

Definition at line 144 of file reference_index_file.cuh.

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(),
180 database.metadata().compatibility.indexed_bucket_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 }
#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
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.
Here is the call graph for this function:
Here is the caller graph for this function:

◆ save()

template<uint32_t K, size_t BucketCount, typename Layout >
static Result< void > cuddl::reference_index_file::save ( reference_index< K, BucketCount, Layout > const &  index,
reference_database< K, BucketCount, Layout > const &  source,
std::filesystem::path const &  path,
cuda::stream_ref  stream 
)
inlinestatic

Atomically save an already built or loaded index.

Definition at line 78 of file reference_index_file.cuh.

83 {
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));
99 detail::database_file_writer writer;
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 }
__global__ uint32_t uint32_t uint64_t * keys
Result< uint64_t > reference_database_digest(reference_database_file const &db)
__shared__ block_scan::TempStorage storage
Definition kernels.cuh:545
static __global__ uint32_t const * offsets
Definition kernels.cuh:536
static __global__ uint32_t const uint32_t uint32_t const * postings
Definition kernels.cuh:574
Result< void > Ok() noexcept
Success return for Result<void>; same as return {}.
Definition error.hpp:181
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 resource(std::string message)
Definition error.hpp:103
Here is the call graph for this function:

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