cuddl::detail::device_fasta_pipeline Class Reference

Normalises FASTA batches supplied compressed, uncompressed, or already inflated. More...

Public Member Functions
	device_fasta_pipeline (cuda::stream_ref stream, size_t slot_capacity, size_t compressed_capacity, size_t workers, bool hardware_inflate)
size_t	slot_capacity () const noexcept
size_t	compressed_capacity () const noexcept
size_t	hardware_limit () const noexcept
	Largest inflated file the device's decompression engine takes; 0 without one.
size_t	batch_files (std::span< size_t const > ids, std::span< gzip_file_probe const > probes) const noexcept
	How many ids from the front of ids one lane holds.
Result< size_t >	submit (size_t lane_index, std::span< size_t const > ids, std::span< std::filesystem::path const > paths, std::span< gzip_file_probe const > probes, bool direct_reads)
	Starts inflating files from the front of ids on lane.
Result< void >	submit_decoded (size_t lane_index, std::span< decoded_fasta const > sources)
	Copies CPU-inflated FASTA into a lane and enqueues GPU normalisation.
Result< void >	wait_lanes (cuda::stream_ref consumer)
	Completes lane work before callers release outstanding input buffers.
Result< void >	finish (size_t lane_index, std::span< gzip_file_probe const > probes, std::vector< inflated_genome > &genomes, std::vector< size_t > &fallback)
	Waits for lane's batch and splits it into genomes the device holds and ids the host loader must take.
Result< void >	release (size_t lane_index, cuda::stream_ref consumer)
	Lets lane reuse its slots once consumer reaches this point.

Static Public Member Functions
static size_t	lane_overhead (size_t slot_capacity) noexcept
	Device bytes a lane needs besides its slots and compressed bytes.

Static Public Attributes
static constexpr size_t	lane_count = 6
static constexpr size_t	max_files = 4096
static constexpr size_t	max_members = 65536
	Gzip members one batch inflates: one per file, or one per non-empty BGZF block.

Detailed Description

Normalises FASTA batches supplied compressed, uncompressed, or already inflated.

Each lane receives raw bytes through a batched copy, or reads files itself and inflates the gzip ones through nvCOMP. Kernels blank header text and leading bytes, then compaction drops whitespace. GPU-inflated data reuses compressed-input storage. CPU-inflated data compacts in place. One genome comes out as one run of bases with its records joined by the '>' that opened each one, a byte no k-mer window accepts, so it sketches exactly as its separate records would.

A file is handed back to the host loader when anything is off: an inflate failure, an output whose size or CRC32 disagrees with the trailer (a multi-member stream inflates only its first member here), FASTQ content, or no header. The host loader then gives it its usual result or error.

Definition at line 711 of file device_gzip.cuh.

Constructor & Destructor Documentation

device_fasta_pipeline()

cuddl::detail::device_fasta_pipeline::device_fasta_pipeline ( cuda::stream_ref  stream, size_t  slot_capacity, size_t  compressed_capacity, size_t  workers, bool  hardware_inflate  )

inline

Parameters

slot_capacity	Inflated bytes one lane holds, one guard byte per file included.
compressed_capacity	Compressed bytes per lane; zero selects decoded input.
hardware_inflate	Whether nvCOMP may use the device's decompression engine; without it, nvCOMP inflates on the SMs from ordinary device memory.

Definition at line 725 of file device_gzip.cuh.

        : workers_(std::max<size_t>(1, workers)),
          compressed_capacity_(compressed_capacity),
          slot_capacity_(slot_capacity),
          header_capacity_(
              static_cast<uint32_t>(std::clamp<size_t>(
                  slot_capacity / 4096,
                  1 << 16,
                  std::numeric_limits<uint32_t>::max()
              ))
          ) {
        auto const options = nvcompBatchedDeflateDecompressDefaultOpts;
        nvcomp_temp_bytes_ = 0;
        if (compressed_capacity_ != 0 &&
            nvcompBatchedDeflateDecompressGetTempSizeAsync(
                max_members,
                std::min<size_t>(slot_capacity_, std::numeric_limits<uint32_t>::max()),
                options,
                &nvcomp_temp_bytes_,
                slot_capacity_
            ) != nvcompSuccess) {
            throw std::bad_alloc();
        }
        select_temp_bytes_ = 0;
        auto const selected = compressed_capacity_ == 0 ? cub::DeviceSelect::If(
                                                              nullptr,
                                                              select_temp_bytes_,
                                                              static_cast<char*>(nullptr),
                                                              static_cast<int64_t*>(nullptr),
                                                              static_cast<int64_t>(slot_capacity_),
                                                              device_gzip::kept_byte{},
                                                              stream.get()
                                                          )
                                                        : cub::DeviceSelect::If(
                                                              nullptr,
                                                              select_temp_bytes_,
                                                              static_cast<char*>(nullptr),
                                                              static_cast<char*>(nullptr),
                                                              static_cast<int64_t*>(nullptr),
                                                              static_cast<int64_t>(slot_capacity_),
                                                              device_gzip::kept_byte{},
                                                              stream.get()
                                                          );
        if (selected != cudaSuccess) throw std::runtime_error(cudaGetErrorString(selected));
        auto const device = stream.device();
        if (hardware_inflate) {
            batch_decompress_ = device_gzip::batch_decompress_entry();
            if (batch_decompress_ != nullptr) {
                hardware_limit_ = device_gzip::hardware_inflate_limit(device.get());
            }
        }
        streaming_window_ =
            compressed_capacity_ == 0
                ? std::min(
                      slot_capacity_,
                      static_cast<size_t>(
                          device.attribute(cuda::device_attributes::max_access_policy_window_size)
                      )
                  )
                : size_t{0};
        for (auto& lane : lanes_) {
            lane.stream.emplace(device);
            auto const s = cuda::stream_ref{*lane.stream};
            lane.uploaded.emplace(s);
            lane.classified.emplace(s);
            lane.released.emplace(stream);
        }
        // The other lanes allocate on first use, so the device starts on the first batch while
        // they page-lock their compressed buffers.
        allocate(lanes_[0]);
    }

Member Function Documentation

batch_files()

size_t cuddl::detail::device_fasta_pipeline::batch_files ( std::span< size_t const >  ids, std::span< gzip_file_probe const >  probes  ) const

inlinenoexcept

How many ids from the front of ids one lane holds.

Definition at line 822 of file device_gzip.cuh.

                     {
        size_t files = 0, compressed = 0, slots = 0, members = 0;
        while (files < ids.size() && files < max_files) {
            auto const& probe = probes[ids[files]];
            if (compressed + probe.host_bytes() > compressed_capacity_ ||
                slots + probe.isize + 1 > slot_capacity_ || members + probe.members > max_members) {
                break;
            }
            compressed += probe.host_bytes();
            slots += probe.isize + 1;
            members += probe.members;
            ++files;
        }
        return files;
    }

compressed_capacity()

size_t cuddl::detail::device_fasta_pipeline::compressed_capacity ( ) const

inlinenoexcept

Definition at line 807 of file device_gzip.cuh.

                                                              {
        return compressed_capacity_;
    }

finish()

Result< void > cuddl::detail::device_fasta_pipeline::finish ( size_t  lane_index, std::span< gzip_file_probe const >  probes, std::vector< inflated_genome > &  genomes, std::vector< size_t > &  fallback  )

inline

Waits for lane's batch and splits it into genomes the device holds and ids the host loader must take.

The genomes stay valid until release, which must follow once the kernels that read them are enqueued.

Definition at line 1259 of file device_gzip.cuh.

      {
        auto& lane = lanes_[lane_index];
        fallback.insert(fallback.end(), lane.fallback.begin(), lane.fallback.end());
        if (lane.files.empty()) return Ok();
        CUDDL_CUDA_TRY(lane.classified->sync());
        auto const table = host_table(lane);
        auto const overflow = *table.header_count > header_capacity_;
        auto const* bases = lane.decoded ? lane.slots->data() : lane.payload->data();
        uint64_t offset = 0;
        for (size_t file = 0; file < lane.files.size(); ++file) {
            auto const id = lane.files[file];
            auto const size = table.kept[file];
            bool intact = lane.decoded || probes[id].plain;
            if (!intact) {
                uint64_t inflated = 0;
                bool succeeded = true;
                for (auto m = lane.member_first[file]; m < lane.member_first[file + 1]; ++m) {
                    succeeded = succeeded && table.member_statuses[m] == nvcompSuccess;
                    inflated += table.member_inflated[m];
                }
                intact =
                    succeeded && inflated == probes[id].isize && table.crcs[file] == probes[id].crc;
            }
            auto const fasta =
                table.first_char[file] != '@' && table.first_header[file] < table.slots[file + 1];
            // Each header contributes one separator, including the first header.
            // ponytail: CPU-validate shorter outputs; count per-file headers if this dominates.
            auto const has_sequence = size > *table.header_count;
            if (!overflow && intact && fasta && has_sequence) {
                genomes.push_back({id, bases + offset, size});
            } else {
                fallback.push_back(id);
            }
            offset += size;
        }
        if (static_cast<int64_t>(offset) != *table.selected) {
            return Err(Error::resource("device FASTA compaction lost track of its bytes"));
        }
        return Ok();
    }

Here is the call graph for this function:

hardware_limit()

size_t cuddl::detail::device_fasta_pipeline::hardware_limit ( ) const

inlinenoexcept

Largest inflated file the device's decompression engine takes; 0 without one.

Definition at line 812 of file device_gzip.cuh.

                                                         {
        return hardware_limit_;
    }

lane_overhead()

static size_t cuddl::detail::device_fasta_pipeline::lane_overhead ( size_t  slot_capacity )

inlinestaticnoexcept

Device bytes a lane needs besides its slots and compressed bytes.

Definition at line 817 of file device_gzip.cuh.

                                                                             {
        return slot_capacity / 4096 * sizeof(uint64_t) + (size_t{64} << 20);
    }

release()

Result< void > cuddl::detail::device_fasta_pipeline::release ( size_t  lane_index, cuda::stream_ref  consumer  )

inline

Lets lane reuse its slots once consumer reaches this point.

Definition at line 1306 of file device_gzip.cuh.

                                                                                   {
        CUDDL_CUDA_TRY(lanes_[lane_index].released->record(consumer));
        return Ok();
    }

Here is the call graph for this function:

slot_capacity()

size_t cuddl::detail::device_fasta_pipeline::slot_capacity ( ) const

inlinenoexcept

Definition at line 803 of file device_gzip.cuh.

                                                        {
        return slot_capacity_;
    }

submit()

Result< size_t > cuddl::detail::device_fasta_pipeline::submit ( size_t  lane_index, std::span< size_t const >  ids, std::span< std::filesystem::path const >  paths, std::span< gzip_file_probe const >  probes, bool  direct_reads  )

inline

Starts inflating files from the front of ids on lane.

Every id must be a device candidate that fits one lane on its own. Plain files skip nvCOMP: they lead the lane's slots and reach them as read. Files with extra member signatures reach finish without device inflation.

Parameters

direct_reads

reads the files with O_DIRECT, as read_whole_file does.

Returns

How many ids the batch took: batch_files of them.

Definition at line 848 of file device_gzip.cuh.

      {
        auto& lane = lanes_[lane_index];
        CUDDL_TRY(cuda_try([&] { allocate(lane); }));
        lane.decoded = false;
        if (ids.empty()) return Err(Error::invalid_argument("an inflate batch needs a file"));
        auto const files = batch_files(ids, probes);
        // The pinned bytes are rewritten below, so the previous batch's upload must be done.
        CUDDL_CUDA_TRY(lane.uploaded->sync());
        auto table = host_table(lane);
        // Plain files lie in the page-locked buffer exactly as they will in the slots, which
        // they lead, so one copy places them all. Compressed bytes follow them.
        uint64_t plain_at = 0;
        for (size_t file = 0; file < files; ++file) {
            auto const& probe = probes[ids[file]];
            if (probe.plain) plain_at += probe.host_bytes();
        }
        uint64_t compressed_at = plain_at, slot_at = 0;
        for (size_t file = 0; file < files; ++file) {
            auto const& probe = probes[ids[file]];
            table.compressed_bytes[file] = probe.compressed;
            if (probe.plain) {
                table.compressed_offsets[file] = slot_at + 1;
                slot_at += probe.host_bytes();
            } else {
                table.compressed_offsets[file] = compressed_at;
                compressed_at += probe.compressed;
            }
        }
        std::atomic<bool> failed{false};
        readers_.for_each(files, [&](size_t file) {
            auto const* const path = paths[ids[file]].c_str();
            auto* const out = lane.pinned->data() + table.compressed_offsets[file];
            auto const bytes = table.compressed_bytes[file];
            int const fd = ::open(path, O_RDONLY);
            // Files lie back to back in the lane, so a direct read goes through a bounce buffer.
            bool const read = fd != -1 && read_whole_file(fd, out, bytes, direct_reads);
            auto const done = read ? bytes : 0;
            if (!read) failed = true;
            // A matching trailer cannot distinguish repeated first and last members.
            // A signature inside compressed data is ambiguous too, so use the host parser.
            // BGZF names every member in its block headers.
            auto const& probe = probes[ids[file]];
            lane.possible_members[file] =
                !probe.plain && !probe.bgzf && done > 10 && has_gzip_signature(out + 10, done - 10);
        });
        if (failed) return Err(Error::invalid_argument("cannot read reference inputs"));
        lane.files.clear();
        lane.fallback.clear();
        lane.member_first.clear();
        engine_params_.clear();
        slot_at = 0;
        uint32_t members = 0;
        auto const accept = [&](size_t file) {
            table.slots[lane.files.size()] = slot_at;
            lane.member_first.push_back(members);
            slot_at += probes[ids[file]].isize + 1;
            lane.files.push_back(ids[file]);
        };
        for (size_t file = 0; file < files; ++file) {
            if (probes[ids[file]].plain) accept(file);
        }
        auto const plain = lane.files.size();
        for (size_t file = 0; file < files; ++file) {
            if (!probes[ids[file]].plain && lane.possible_members[file]) {
                lane.fallback.push_back(ids[file]);
            }
        }
        // One member: raw DEFLATE at @p deflate in the page-locked copy, which the payload buffer
        // mirrors, inflating to @p size bytes at @p out.
        auto const add_member = [&](unsigned char const* deflate,
                                    uint64_t deflate_bytes,
                                    uint32_t size,
                                    char* out) {
            auto* const source = lane.payload->data() +
                                 (reinterpret_cast<char const*>(deflate) - lane.pinned->data());
            table.member_sources[members] = source;
            table.member_source_bytes[members] = deflate_bytes;
            table.member_outputs[members] = out;
            table.member_output_bytes[members] = size;
            if (hardware_limit_ != 0 && size <= hardware_limit_) {
                CUmemDecompressParams params{};
                params.src = source;
                params.srcNumBytes = deflate_bytes;
                params.dst = out;
                params.dstNumBytes = size;
                // The low half of the member's 64-bit inflated size, zeroed before the batch.
                params.dstActBytes =
                    reinterpret_cast<cuuint32_t*>(device_tables(lane).member_inflated + members);
                params.algo = CU_MEM_DECOMPRESS_ALGORITHM_DEFLATE;
                engine_params_.push_back(params);
            }
            ++members;
        };
        // Gzip files whose members all fit the decompression engine come first, then those
        // nvCOMP inflates on the SMs, so each group's members form one range. Both take raw
        // DEFLATE: the host steps over each member's header in its copy. Without an engine, every
        // gzip file is in the second group.
        auto const gzip_pass = [&](bool engine) {
            for (size_t file = 0; file < files; ++file) {
                auto const& probe = probes[ids[file]];
                if (probe.plain || lane.possible_members[file]) continue;
                bool const fits =
                    hardware_limit_ != 0 && (probe.bgzf || probe.isize <= hardware_limit_);
                if (fits != engine) continue;
                auto const* const bytes = reinterpret_cast<unsigned char const*>(
                    lane.pinned->data() + table.compressed_offsets[file]
                );
                auto* out = lane.slots->data() + slot_at + 1;
                if (probe.bgzf) {
                    // The blocks must still be the ones probed, or they would overrun the slot.
                    auto const blocks = walk_bgzf(bytes, probe.compressed, [](auto...) {});
                    if (!blocks || blocks->first != probe.members ||
                        blocks->second != probe.isize) {
                        lane.fallback.push_back(ids[file]);
                        continue;
                    }
                    accept(file);
                    (void)walk_bgzf(
                        bytes, probe.compressed, [&](uint64_t at, uint64_t length, uint32_t size) {
                            add_member(bytes + at, length, size, out);
                            out += size;
                        }
                    );
                    continue;
                }
                auto const deflate_at = device_gzip::deflate_offset(
                    reinterpret_cast<char const*>(bytes), probe.compressed
                );
                if (deflate_at == 0) {
                    lane.fallback.push_back(ids[file]);
                    continue;
                }
                accept(file);
                add_member(bytes + deflate_at, probe.compressed - deflate_at - 8, probe.isize, out);
            }
        };
        gzip_pass(true);
        auto const engine_members = members;
        gzip_pass(false);
        lane.member_first.push_back(members);
        table.slots[lane.files.size()] = slot_at;
        lane.slot_bytes = slot_at;
        if (lane.files.empty()) return files;
 
        auto const s = cuda::stream_ref{*lane.stream};
        auto* const stream = s.get();
        auto const device_table = device_tables(lane);
        // The slots are rewritten below, so the sketch kernel reading them must be done.
        CUDDL_CUDA_TRY(s.wait(*lane.released));
        // Plain bytes land in their slots; slot_init_kernel then writes the guard bytes between.
        if (plain_at != 0) {
            CUDDL_CUDA_TRY(
                cuda::copy_bytes(
                    s,
                    cuda::std::span{lane.pinned->data(), plain_at},
                    device_span<char>{lane.slots->data(), plain_at}
                )
            );
        }
        if (compressed_at != plain_at) {
            CUDDL_CUDA_TRY(
                cuda::copy_bytes(
                    s,
                    cuda::std::span{lane.pinned->data() + plain_at, compressed_at - plain_at},
                    device_span<char>{lane.payload->data() + plain_at, compressed_at - plain_at}
                )
            );
        }
        CUDDL_CUDA_TRY(
            cuda::copy_bytes(
                s,
                cuda::std::span{lane.host->data(), inputs_bytes()},
                device_span<char>{lane.tables->data(), inputs_bytes()}
            )
        );
        CUDDL_CUDA_TRY(lane.uploaded->record(s));
        auto const count = static_cast<uint32_t>(lane.files.size());
        device_gzip::slot_init_kernel<<<(count + 255) / 256, 256, 0, stream>>>(
            lane.slots->data(),
            device_table.slots,
            count,
            device_table.first_header,
            device_table.header_count
        );
        CUDDL_CUDA_TRY(cudaGetLastError());
        if (engine_members != 0) {
            // The driver takes the parameters from host memory. nvCOMP's own engine path builds
            // them on the device and synchronises the stream to read them back, which stalls this
            // thread and every lane it drives. The engine reports only inflated sizes, so the
            // statuses read as success; finish still checks each size and CRC32.
            static_assert(nvcompSuccess == 0);
            static_assert(std::endian::native == std::endian::little);
            CUDDL_CUDA_TRY(cudaMemsetAsync(
                device_table.member_inflated, 0, engine_members * sizeof(uint64_t), stream
            ));
            CUDDL_CUDA_TRY(cudaMemsetAsync(
                device_table.member_statuses, 0, engine_members * sizeof(nvcompStatus_t), stream
            ));
            size_t error_index = 0;
            auto const submitted = batch_decompress_(
                engine_params_.data(), engine_params_.size(), 0, &error_index, stream
            );
            if (submitted != CUDA_SUCCESS) {
                return Err(
                    Error::resource(
                        "hardware decompression launch failed: " + std::to_string(submitted)
                    )
                );
            }
        }
        if (engine_members != members) {
            auto options = nvcompBatchedDeflateDecompressDefaultOpts;
            options.backend = NVCOMP_DECOMPRESS_BACKEND_CUDA;
            auto const first = engine_members;
            auto const inflated = nvcompBatchedDeflateDecompressAsync(
                reinterpret_cast<void const* const*>(device_table.member_sources + first),
                device_table.member_source_bytes + first,
                device_table.member_output_bytes + first,
                device_table.member_inflated + first,
                members - first,
                lane.temp->data(),
                nvcomp_temp_bytes_,
                reinterpret_cast<void* const*>(device_table.member_outputs + first),
                options,
                device_table.member_statuses + first,
                stream
            );
            if (inflated != nvcompSuccess) {
                return Err(
                    Error::resource("nvCOMP DEFLATE launch failed: " + std::to_string(inflated))
                );
            }
        }
        if (auto const gzip = count - static_cast<uint32_t>(plain); gzip != 0) {
            device_gzip::slot_crc_kernel<<<gzip, device_gzip::block_size, 0, stream>>>(
                lane.slots->data(), device_table.slots + plain, device_table.crcs + plain
            );
            CUDDL_CUDA_TRY(cudaGetLastError());
        }
        CUDDL_TRY(normalise(lane, slot_at));
        return files;
    }

Here is the call graph for this function:

submit_decoded()

Result< void > cuddl::detail::device_fasta_pipeline::submit_decoded ( size_t  lane_index, std::span< decoded_fasta const >  sources  )

inline

Copies CPU-inflated FASTA into a lane and enqueues GPU normalisation.

Retain the source bytes until this lane finishes, or until wait_lanes returns.

Definition at line 1181 of file device_gzip.cuh.

                                                                            {
        if (sources.empty() || sources.size() > max_files || compressed_capacity_ != 0) {
            return Err(Error::invalid_argument("invalid decoded FASTA batch"));
        }
        auto& lane = lanes_[lane_index];
        CUDDL_TRY(cuda_try([&] { allocate(lane); }));
        lane.uploaded->sync();
        lane.decoded = true;
        lane.files.clear();
        lane.fallback.clear();
        auto const table = host_table(lane);
        uint64_t bytes = 0;
        for (auto const& source : sources) {
            if (source.bytes.size() >= slot_capacity_ - bytes) {
                return Err(Error::invalid_argument("decoded FASTA batch exceeds its lane"));
            }
            auto const file = lane.files.size();
            table.slots[file] = bytes;
            table.compressed_ptrs[file] = const_cast<char*>(source.bytes.data());
            table.compressed_bytes[file] = source.bytes.size();
            table.output_ptrs[file] = lane.slots->data() + bytes + 1;
            bytes += source.bytes.size() + 1;
            lane.files.push_back(source.id);
        }
        table.slots[sources.size()] = bytes;
        auto const s = cuda::stream_ref{*lane.stream};
        auto const device_table = device_tables(lane);
        CUDDL_CUDA_TRY(s.wait(*lane.released));
        CUDDL_CUDA_TRY(
            cuda::copy_bytes(
                s,
                cuda::std::span{lane.host->data(), inputs_bytes()},
                device_span<char>{lane.tables->data(), inputs_bytes()}
            )
        );
        auto const count = static_cast<uint32_t>(sources.size());
        device_gzip::slot_init_kernel<<<(count + 255) / 256, 256, 0, s.get()>>>(
            lane.slots->data(),
            device_table.slots,
            count,
            device_table.first_header,
            device_table.header_count
        );
        CUDDL_CUDA_TRY(cudaGetLastError());
        static_assert(CUDART_VERSION >= 13000, "cuDDL needs CUDA 13.0 or newer");
        cudaMemcpyAttributes attributes{};
        attributes.srcAccessOrder = cudaMemcpySrcAccessOrderAny;
        attributes.srcLocHint = {cudaMemLocationTypeHost, 0};
        attributes.flags = cudaMemcpyFlagPreferOverlapWithCompute;
        size_t attributes_index = 0;
        CUDDL_CUDA_TRY(cudaMemcpyBatchAsync(
            reinterpret_cast<void* const*>(table.output_ptrs),
            reinterpret_cast<void const* const*>(table.compressed_ptrs),
            table.compressed_bytes,
            count,
            &attributes,
            &attributes_index,
            1,
            s.get()
        ));
        CUDDL_CUDA_TRY(lane.uploaded->record(s));
        return normalise(lane, bytes);
    }

Here is the call graph for this function:

wait_lanes()

Result< void > cuddl::detail::device_fasta_pipeline::wait_lanes ( cuda::stream_ref  consumer )

inline

Completes lane work before callers release outstanding input buffers.

Definition at line 1246 of file device_gzip.cuh.

                                                                   {
        for (auto& lane : lanes_) {
            CUDDL_CUDA_TRY(lane.stream->sync());
        }
        CUDDL_CUDA_TRY(consumer.sync());
        return Ok();
    }

Here is the call graph for this function:

Member Data Documentation

lane_count

constexpr size_t cuddl::detail::device_fasta_pipeline::lane_count = 6

staticconstexpr

Definition at line 715 of file device_gzip.cuh.

max_files

constexpr size_t cuddl::detail::device_fasta_pipeline::max_files = 4096

staticconstexpr

Definition at line 716 of file device_gzip.cuh.

max_members

constexpr size_t cuddl::detail::device_fasta_pipeline::max_members = 65536

staticconstexpr

Gzip members one batch inflates: one per file, or one per non-empty BGZF block.

With BGZF's 64 KiB blocks, about 4 GiB of output.

Definition at line 719 of file device_gzip.cuh.

---

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

• device_gzip.cuh