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

Reusable, move-only device lookup. More...

Public Member Functions

 device_blacklist (kmer_blacklist source, size_t buckets, cuda::stream_ref stream)
 
 device_blacklist (device_blacklist const &)=delete
 
device_blacklist & operator= (device_blacklist const &)=delete
 
 device_blacklist (device_blacklist &&) noexcept=default
 
device_blacklist & operator= (device_blacklist &&) noexcept=default
 
kmer_blacklist const & source () const noexcept
 
cuda::std::optional< detail::blacklist_view > view () const noexcept
 
Result< void > validate (uint32_t k, size_t buckets, cuda::stream_ref stream) const
 

Detailed Description

Reusable, move-only device lookup.

Construction synchronizes the upload stream. Keep this owner and its allocation stream alive until every consuming kernel completes.

Definition at line 201 of file blacklist.cuh.

Constructor & Destructor Documentation

◆ device_blacklist() [1/3]

cuddl::device_blacklist::device_blacklist ( kmer_blacklist  source,
size_t  buckets,
cuda::stream_ref  stream 
)
inline

Definition at line 203 of file blacklist.cuh.

204 : source_(std::move(source)), buckets_(buckets), device_(stream.device().get()) {
205 if (buckets < 2048 || buckets > 8192 || !std::has_single_bit(buckets)) {
206 throw std::invalid_argument("invalid blacklist bucket count");
207 }
208
209 if (source_.keys().empty()) return;
210
211 std::vector<uint32_t> offsets(buckets + 1, 0);
212 std::vector<uint64_t> fingerprints(buckets, 0);
213
214 for (auto key : source_.keys()) {
215 auto const hash = detail::hash_kmer(key);
216 auto const bucket = hash & (buckets - 1);
217 ++offsets[bucket + 1];
218 fingerprints[bucket] |= uint64_t{1} << detail::blacklist_view::fingerprint(hash);
219 }
220
221 for (size_t b = 1; b <= buckets; ++b) {
222 offsets[b] += offsets[b - 1];
223 }
224
225 auto cursor = offsets;
226 std::vector<uint64_t> keys(source_.keys().size());
227 for (auto key : source_.keys()) {
228 keys[cursor[detail::hash_kmer(key) & (buckets - 1)]++] = key;
229 }
230 // 32 to 64 map bits per key hold false positives to 1.5-3% at any list size.
231 auto const presence_bits =
232 std::max(10U, static_cast<uint32_t>(std::bit_width(keys.size())) + 5U);
233 std::vector<uint32_t> presence((size_t{1} << presence_bits) / 32);
234 for (auto key : keys) {
235 auto const slot = detail::blacklist_view::presence_slot(key, presence_bits);
236 presence[slot / 32] |= uint32_t{1} << (slot % 32);
237 }
238 presence_bits_ = presence_bits;
239
240 keys_.emplace(cuda::make_device_buffer<uint64_t>(stream, stream.device(), keys));
241 offsets_.emplace(cuda::make_device_buffer<uint32_t>(stream, stream.device(), offsets));
242 presence_.emplace(cuda::make_device_buffer<uint32_t>(stream, stream.device(), presence));
243 fingerprints_.emplace(
244 cuda::make_device_buffer<uint64_t>(stream, stream.device(), fingerprints)
245 );
246 stream.sync();
247 }
kmer_blacklist const & source() const noexcept
std::vector< uint64_t > const & keys() const noexcept
__global__ uint32_t uint32_t uint64_t * keys
__host__ __device__ constexpr uint64_t hash_kmer(uint64_t packed_kmer) noexcept
Hashes a canonical packed k-mer into the value space used by DDL registers.
Definition hash.cuh:21
static __global__ uint32_t const * offsets
Definition kernels.cuh:536
__host__ static __device__ uint32_t fingerprint(uint64_t hash) noexcept
Bits 13-18 sit above every supported bucket index and below the scored high word.
Definition blacklist.cuh:41
__host__ static __device__ uint32_t presence_slot(uint64_t key, uint32_t bits) noexcept
Definition blacklist.cuh:45

◆ device_blacklist() [2/3]

cuddl::device_blacklist::device_blacklist ( device_blacklist const &  )
delete

◆ device_blacklist() [3/3]

cuddl::device_blacklist::device_blacklist ( device_blacklist &&  )
defaultnoexcept

Member Function Documentation

◆ operator=() [1/2]

device_blacklist & cuddl::device_blacklist::operator= ( device_blacklist &&  )
defaultnoexcept

◆ operator=() [2/2]

device_blacklist & cuddl::device_blacklist::operator= ( device_blacklist const &  )
delete

◆ source()

kmer_blacklist const & cuddl::device_blacklist::source ( ) const
inlinenoexcept

Definition at line 253 of file blacklist.cuh.

253 {
254 return source_;
255 }

◆ validate()

Result< void > cuddl::device_blacklist::validate ( uint32_t  k,
size_t  buckets,
cuda::stream_ref  stream 
) const
inline

Definition at line 266 of file blacklist.cuh.

266 {
267 if ((!source_.keys().empty() && source_.kmer_length() != k) || buckets != buckets_ ||
268 stream.device().get() != device_) {
269 return Err(
270 Error::invalid_argument("blacklist construction configuration or device mismatch")
271 );
272 }
273 return Ok();
274 }
uint32_t kmer_length() const noexcept
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 invalid_argument(std::string message)
Definition error.hpp:99
Here is the call graph for this function:

◆ view()

cuda::std::optional< detail::blacklist_view > cuddl::device_blacklist::view ( ) const
inlinenoexcept

Definition at line 256 of file blacklist.cuh.

256 {
257 if (!keys_) return cuda::std::nullopt;
258 return detail::blacklist_view{
259 keys_->data(),
260 offsets_->data(),
261 presence_->data(),
262 fingerprints_->data(),
263 presence_bits_
264 };
265 }

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