cuddl::detail::device_gzip Namespace Reference

Classes
struct	kept_byte
	The bytes a FASTA record keeps: everything but sequence whitespace. More...

Functions
__device__ __forceinline__ bool	is_whitespace (char c)
__device__ uint32_t	crc_multiply (uint32_t a, uint32_t b)
	a * b modulo the CRC-32 polynomial, bit-reflected (zlib's multmodp).
__device__ uint32_t	crc_combine (uint32_t first, uint32_t second, uint64_t second_bytes, uint32_t const *x2n)
	crc32(A || B) from crc32(A), crc32(B) and |B| (zlib's crc32_combine).
__device__ __forceinline__ uint32_t	slot_containing (uint64_t const *slots, uint32_t files, uint64_t position)
	The last of files slots that starts at or before position.
__device__ __forceinline__ uint32_t	kept_in_word (uint64_t word, uint32_t first, uint32_t last)
	How many of word's bytes [first, last) kept_byte keeps.
size_t	hardware_inflate_limit (int device)
	Largest output, in bytes, that device's hardware decompression engine (Blackwell B200, B300, GB200, GB300) inflates in one operation, or 0 for a device without one.
cuda::device_memory_pool_ref	hardware_inflate_pool (cuda::device_ref device)
	Device memory device's decompression engine can read and write.
PFN_cuMemBatchDecompressAsync_v12060	batch_decompress_entry ()
	The driver's cuMemBatchDecompressAsync, or nullptr where the driver lacks it.
size_t	deflate_offset (char const *bytes, size_t size)
	Where the DEFLATE stream starts in the size bytes of a gzip candidate (one member, no FEXTRA), or 0 if its header runs into the trailer.

Variables
constexpr uint32_t	crc_polynomial = 0xEDB88320U
constexpr uint32_t	block_size = 256
constexpr uint64_t	kept_tile = uint64_t{block_size} * 16 * 8
	Bytes a tile of kept_count_kernel covers: 8 aligned 16-byte loads per thread.

Function Documentation

batch_decompress_entry()

PFN_cuMemBatchDecompressAsync_v12060 cuddl::detail::device_gzip::batch_decompress_entry ( )

inline

The driver's cuMemBatchDecompressAsync, or nullptr where the driver lacks it.

Definition at line 654 of file device_gzip.cuh.

                                                                     {
    void* address = nullptr;
    cudaDriverEntryPointQueryResult found{};
    if (cudaGetDriverEntryPointByVersion(
            "cuMemBatchDecompressAsync", &address, 12060U, cudaEnableDefault, &found
        ) != cudaSuccess ||
        found != cudaDriverEntryPointSuccess) {
        return nullptr;
    }
    return reinterpret_cast<PFN_cuMemBatchDecompressAsync_v12060>(address);
}

Here is the call graph for this function:

Here is the caller graph for this function:

crc_combine()

__device__ uint32_t cuddl::detail::device_gzip::crc_combine ( uint32_t  first, uint32_t  second, uint64_t  second_bytes, uint32_t const *  x2n  )

inline

crc32(A || B) from crc32(A), crc32(B) and |B| (zlib's crc32_combine).

Definition at line 351 of file device_gzip.cuh.

                                                                                         {
    uint32_t shift = 1U << 31U;
    for (uint32_t k = 3; second_bytes != 0; second_bytes >>= 1U, ++k) {
        if ((second_bytes & 1U) != 0) shift = crc_multiply(x2n[k & 31U], shift);
    }
    return crc_multiply(shift, first) ^ second;
}

Here is the call graph for this function:

crc_multiply()

__device__ uint32_t cuddl::detail::device_gzip::crc_multiply ( uint32_t  a, uint32_t  b  )

inline

a * b modulo the CRC-32 polynomial, bit-reflected (zlib's multmodp).

Definition at line 336 of file device_gzip.cuh.

                                                                {
    uint32_t m = 1U << 31U, p = 0;
    while (true) {
        if ((a & m) != 0) {
            p ^= b;
            if ((a & (m - 1)) == 0) break;
        }
        m >>= 1U;
        b = (b & 1U) != 0 ? (b >> 1U) ^ crc_polynomial : b >> 1U;
    }
    return p;
}

Here is the caller graph for this function:

deflate_offset()

size_t cuddl::detail::device_gzip::deflate_offset ( char const *  bytes, size_t  size  )

inline

Where the DEFLATE stream starts in the size bytes of a gzip candidate (one member, no FEXTRA), or 0 if its header runs into the trailer.

Definition at line 668 of file device_gzip.cuh.

                                                             {
    auto const flags = static_cast<unsigned char>(bytes[3]);
    size_t at = 10;
    // FNAME, then FCOMMENT: zero-terminated strings.
    for (auto const flag : {8U, 16U}) {
        if ((flags & flag) == 0) continue;
        auto const* const end =
            static_cast<char const*>(std::memchr(bytes + at, '\0', size - std::min(at, size)));
        if (end == nullptr) return 0;
        at = static_cast<size_t>(end - bytes) + 1;
    }
    if ((flags & 2) != 0) at += 2;  // FHCRC
    return at + 8 < size ? at : 0;
}

Here is the call graph for this function:

Here is the caller graph for this function:

hardware_inflate_limit()

size_t cuddl::detail::device_gzip::hardware_inflate_limit ( int  device )

inline

Largest output, in bytes, that device's hardware decompression engine (Blackwell B200, B300, GB200, GB300) inflates in one operation, or 0 for a device without one.

nvCOMP sends a whole batch to the SMs once any output buffer in it exceeds this.

Definition at line 618 of file device_gzip.cuh.

                                                 {
    void* address = nullptr;
    cudaDriverEntryPointQueryResult found{};
    if (cudaGetDriverEntryPointByVersion(
            "cuDeviceGetAttribute", &address, 12000U, cudaEnableDefault, &found
        ) != cudaSuccess ||
        found != cudaDriverEntryPointSuccess) {
        return 0;
    }
    int limit = 0;
    auto const attribute = reinterpret_cast<decltype(&cuDeviceGetAttribute)>(address);
    auto const status =
        attribute(&limit, CU_DEVICE_ATTRIBUTE_MEM_DECOMPRESS_MAXIMUM_LENGTH, device);
    return status == CUDA_SUCCESS && limit > 0 ? static_cast<size_t>(limit) : 0;
}

Here is the call graph for this function:

Here is the caller graph for this function:

hardware_inflate_pool()

cuda::device_memory_pool_ref cuddl::detail::device_gzip::hardware_inflate_pool ( cuda::device_ref  device )

inline

Device memory device's decompression engine can read and write.

CUDA's default pool is not: nvCOMP then falls back to the SMs, which on a B300 inflated 2,000 genomes at 7.3 GB/s of gzip against 57.6 GB/s from this pool. One pool per device for the life of the process, like device_pinned_pool.

Definition at line 638 of file device_gzip.cuh.

                                                                                             {
    static std::mutex mutex;
    static std::vector<cudaMemPool_t> pools(cuda::devices.size());
    std::lock_guard lock(mutex);
    auto& pool = pools.at(static_cast<size_t>(device.get()));
    if (pool == nullptr) {
        cudaMemPoolProps properties{};
        properties.allocType = cudaMemAllocationTypePinned;
        properties.location = {cudaMemLocationTypeDevice, device.get()};
        properties.usage = cudaMemPoolCreateUsageHwDecompress;
        CUDDL_CUDA_CALL(cudaMemPoolCreate(&pool, &properties));
    }
    return cuda::device_memory_pool_ref(pool);
}

Here is the call graph for this function:

Here is the caller graph for this function:

is_whitespace()

__device__ __forceinline__ bool cuddl::detail::device_gzip::is_whitespace

(

char

c

)

Definition at line 324 of file device_gzip.cuh.

                                                      {
    return c == '\n' || c == '\r' || c == ' ' || c == '\t';
}

Here is the caller graph for this function:

kept_in_word()

__device__ __forceinline__ uint32_t cuddl::detail::device_gzip::kept_in_word	(	uint64_t	word,
		uint32_t	first,
		uint32_t	last
	)

How many of word's bytes [first, last) kept_byte keeps.

Definition at line 513 of file device_gzip.cuh.

                                                                                               {
    constexpr uint64_t low7 = 0x7F7F7F7F7F7F7F7FULL, ones = 0x0101010101010101ULL;
    // The high bit of every zero byte of x, exactly: the addition carries within bytes only.
    auto const zero_bytes = [](uint64_t x) {
        return ~(((x & low7) + low7) | x | low7);
    };
    auto const blank = zero_bytes(word ^ (ones * '\n')) | zero_bytes(word ^ (ones * '\r')) |
                       zero_bytes(word ^ (ones * ' ')) | zero_bytes(word ^ (ones * '\t'));
    auto const width = last - first;
    auto const span = (width == 8 ? ~uint64_t{0} : (uint64_t{1} << (8 * width)) - 1) << (8 * first);
    return static_cast<uint32_t>(__popcll(~blank & span & ~low7));
}

Here is the call graph for this function:

Here is the caller graph for this function:

slot_containing()

__device__ __forceinline__ uint32_t cuddl::detail::device_gzip::slot_containing	(	uint64_t const *	slots,
		uint32_t	files,
		uint64_t	position
	)

The last of files slots that starts at or before position.

Definition at line 468 of file device_gzip.cuh.

                                                                          {
    uint32_t low = 0, high = files;
    while (high - low > 1) {
        auto const middle = (low + high) / 2;
        if (slots[middle] <= position) {
            low = middle;
        } else {
            high = middle;
        }
    }
    return low;
}

Here is the call graph for this function:

Here is the caller graph for this function:

Variable Documentation

block_size

constexpr uint32_t cuddl::detail::device_gzip::block_size = 256

inlineconstexpr

Definition at line 322 of file device_gzip.cuh.

crc_polynomial

constexpr uint32_t cuddl::detail::device_gzip::crc_polynomial = 0xEDB88320U

inlineconstexpr

Definition at line 321 of file device_gzip.cuh.

kept_tile

constexpr uint64_t cuddl::detail::device_gzip::kept_tile = uint64_t{block_size} * 16 * 8

inlineconstexpr

Bytes a tile of kept_count_kernel covers: 8 aligned 16-byte loads per thread.

Definition at line 527 of file device_gzip.cuh.

{block_size} * 16 * 8;