cuddl::detail::blacklist_view Struct Reference

Exact membership within the candidate's destination bucket. More...

Public Member Functions
__device__ bool	contains (uint64_t key, uint64_t hash, size_t bucket) const noexcept
	hash must be hash_kmer(key); construction already holds it.
__device__ __noinline__ bool	search (uint64_t key, size_t bucket) const noexcept

Static Public Member Functions
__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

Public Attributes
uint64_t const *	keys = nullptr
uint32_t const *	offsets = nullptr
uint32_t const *	presence = nullptr
uint64_t const *	fingerprints = nullptr
uint32_t	presence_bits = 0

Detailed Description

Exact membership within the candidate's destination bucket.

Three filters run from cheapest to most selective. Each bucket has a 64-bit mask with one bit per 6-bit hash fingerprint of its keys, 16 KiB for 2,048 buckets; one load rejects keys whose bit is clear. A one-bit-per-slot presence map sized to the key count then rejects most of the rest before the bucket's sorted keys are searched.

Definition at line 33 of file blacklist.cuh.

Member Function Documentation

contains()

__device__ bool cuddl::detail::blacklist_view::contains ( uint64_t  key, uint64_t  hash, size_t  bucket  ) const

inlinenoexcept

hash must be hash_kmer(key); construction already holds it.

Definition at line 50 of file blacklist.cuh.

                                                                                        {
        if ((__ldg(fingerprints + bucket) >> fingerprint(hash) & 1U) == 0) return false;
        auto const slot = presence_slot(key, presence_bits);
        if ((__ldg(presence + slot / 32) >> (slot % 32) & 1U) == 0) return false;
        return search(key, bucket);
    }

Here is the call graph for this function:

fingerprint()

__host__ static __device__ uint32_t cuddl::detail::blacklist_view::fingerprint ( uint64_t  hash )

inlinestaticnoexcept

Bits 13-18 sit above every supported bucket index and below the scored high word.

Definition at line 41 of file blacklist.cuh.

                                                                            {
        return static_cast<uint32_t>(hash >> 13) & 63U;
    }

Here is the caller graph for this function:

presence_slot()

__host__ static __device__ uint32_t cuddl::detail::blacklist_view::presence_slot ( uint64_t  key, uint32_t  bits  )

inlinestaticnoexcept

Definition at line 45 of file blacklist.cuh.

                                                                                            {
        return static_cast<uint32_t>((key * 0x9E3779B97F4A7C15ULL) >> (64 - bits));
    }

Here is the caller graph for this function:

search()

__device__ __noinline__ bool cuddl::detail::blacklist_view::search ( uint64_t  key, size_t  bucket  ) const

inlinenoexcept

Definition at line 58 of file blacklist.cuh.

                                                                                    {
        auto first = offsets[bucket];
        auto last = offsets[bucket + 1];
        while (first < last) {
            auto const middle = first + (last - first) / 2;
            auto const value = __ldg(keys + middle);
            if (value == key) return true;
            if (value < key) {
                first = middle + 1;
            } else {
                last = middle;
            }
        }
        return false;
    }

Here is the caller graph for this function:

Member Data Documentation

fingerprints

uint64_t const* cuddl::detail::blacklist_view::fingerprints = nullptr

Definition at line 37 of file blacklist.cuh.

keys

uint64_t const* cuddl::detail::blacklist_view::keys = nullptr

Definition at line 34 of file blacklist.cuh.

offsets

uint32_t const* cuddl::detail::blacklist_view::offsets = nullptr

Definition at line 35 of file blacklist.cuh.

presence

uint32_t const* cuddl::detail::blacklist_view::presence = nullptr

Definition at line 36 of file blacklist.cuh.

presence_bits

uint32_t cuddl::detail::blacklist_view::presence_bits = 0

Definition at line 38 of file blacklist.cuh.

---

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

• blacklist.cuh