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()) {
216 auto const bucket = hash & (buckets - 1);
219 }
220
221 for (size_t b = 1; b <= buckets; ++b) {
223 }
224
226 std::vector<uint64_t>
keys(source_.
keys().size());
227 for (
auto key : source_.
keys()) {
229 }
230
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) {
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.
static __global__ uint32_t const * offsets
__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.
__host__ static __device__ uint32_t presence_slot(uint64_t key, uint32_t bits) noexcept