cuDDL
Loading...
Searching...
No Matches
Classes | Namespaces | Functions | Variables
device_gzip.cuh File Reference
#include <cuda.h>
#include <cudaTypedefs.h>
#include <nvcomp/deflate.h>
#include <nvcomp/gzip.h>
#include <algorithm>
#include <array>
#include <atomic>
#include <bit>
#include <condition_variable>
#include <cstdint>
#include <cstring>
#include <filesystem>
#include <functional>
#include <limits>
#include <mutex>
#include <optional>
#include <span>
#include <stdexcept>
#include <string>
#include <string_view>
#include <thread>
#include <vector>
#include <fcntl.h>
#include <sys/stat.h>
#include <unistd.h>
#include <cub/block/block_reduce.cuh>
#include <cub/device/device_copy.cuh>
#include <cub/device/device_select.cuh>
#include <cuda/algorithm>
#include <cuda/buffer>
#include <cuda/iterator>
#include <cuda/memory_pool>
#include <cuda/stream>
#include <cuddl/detail/fastx_sequence_file.hpp>
#include <cuddl/detail/pinned_pool.cuh>
#include <cuddl/detail/worker_pool.hpp>
#include <cuddl/device_span.cuh>
#include <cuddl/error.hpp>

Go to the source code of this file.

Classes

struct  cuddl::detail::gzip_file_probe
 What a gzip file's header and trailer say, or that a file is plain FASTA, read without inflating it. More...
 
struct  cuddl::detail::device_gzip::kept_byte
 The bytes a FASTA record keeps: everything but sequence whitespace. More...
 
struct  cuddl::detail::inflated_genome
 One genome the device inflated: its compacted bases, resident on the device. More...
 
struct  cuddl::detail::decoded_fasta
 
class  cuddl::detail::device_fasta_pipeline
 Normalises FASTA batches supplied compressed, uncompressed, or already inflated. More...
 

Namespaces

namespace  cuddl
 Batch operations over contiguous device rows.
 
namespace  cuddl::detail
 
namespace  cuddl::detail::device_gzip
 

Functions

uint32_t cuddl::detail::bgzf_block_size (unsigned char const *header) noexcept
 A BGZF block's total size from its 18-byte header, or 0 if it is not one: a gzip header with only FEXTRA set, holding just the BC subfield.
 
uint32_t cuddl::detail::crc32_combine (uint32_t a, uint32_t b, uint64_t length) noexcept
 CRC32 of A followed by B, from CRC32(A), CRC32(B) and B's length, as zlib's crc32_combine: CRC32(A) times x^(8 * length) modulo the reflected polynomial, plus CRC32(B).
 
template<typename Block >
std::optional< std::pair< uint32_t, uint64_t > > cuddl::detail::walk_bgzf (unsigned char const *bytes, uint64_t size, Block &&block)
 Calls block(deflate_at, deflate_bytes, inflated) for each non-empty block of the BGZF file in bytes.
 
gzip_file_probe cuddl::detail::probe_gzip_file (std::filesystem::path const &path)
 Reads the header and trailer of path, and of every block of a BGZF file.
 
template<typename Work >
void cuddl::detail::parallel_for (size_t count, size_t workers, Work &&work)
 Runs work(index) for every index below count on workers threads.
 
bool cuddl::detail::has_gzip_signature_portable (char const *data, size_t size)
 True when data holds what could start another gzip member: 1f 8b 08, then a flag byte with its reserved bits clear.
 
auto cuddl::detail::resolve_gzip_signature_scanner () -> bool(*)(char const *, size_t)
 Selects the widest signature scanner the host supports.
 
__device__ __forceinline__ bool cuddl::detail::device_gzip::is_whitespace (char c)
 
__device__ uint32_t cuddl::detail::device_gzip::crc_multiply (uint32_t a, uint32_t b)
 a * b modulo the CRC-32 polynomial, bit-reflected (zlib's multmodp).
 
__device__ uint32_t cuddl::detail::device_gzip::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 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.
 
__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.
 
size_t cuddl::detail::device_gzip::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 cuddl::detail::device_gzip::hardware_inflate_pool (cuda::device_ref device)
 Device memory device's decompression engine can read and write.
 
PFN_cuMemBatchDecompressAsync_v12060 cuddl::detail::device_gzip::batch_decompress_entry ()
 The driver's cuMemBatchDecompressAsync, or nullptr where the driver lacks it.
 
size_t cuddl::detail::device_gzip::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 cuddl::detail::device_gzip::crc_polynomial = 0xEDB88320U
 
constexpr uint32_t cuddl::detail::device_gzip::block_size = 256
 
constexpr uint64_t cuddl::detail::device_gzip::kept_tile = uint64_t{block_size} * 16 * 8
 Bytes a tile of kept_count_kernel covers: 8 aligned 16-byte loads per thread.