nvCOMP

https://docs.nvidia.com/cuda/nvcomp/
Admin
nvCOMP is a CUDA-accelerated library that provides high-throughput, lossless data compression and...
Tokens:183,459
Snippets:1,214
Trust Score:10
Update:1 month ago
Show doc for...
Tokens:
Raw
### ANS Compression Example

Source: https://docs.nvidia.com/cuda/nvcomp/device_api.html

Example of using the nvcomp_device_execution API with ANS compression, showing type setup and compression call.

```cpp
using namespace nvcomp::device;
constexpr size_t max_uncomp_chunk_size = 1<<15; constexpr size_t num_warps_per_block = 1;
using ans_base_type = decltype(Grouptype<nvcomp_grouptype::warp>() + Algo<nvcomp_algo::ans>() + Datatype<nvcomp_datatype::float16>());
using ans_compressor = nvcomp_device_execution<ans_base_type>;

```

```cpp
auto block = cg::this_thread_block(); auto warp = cg::tiled_partition<32>(block);
// one warp per CTA in this example auto wid = blockIdx.x; auto wid_in_block = threadIdx.x/32;
compressor.compress( uncomp_chunks[wid], comp_chunks[wid], max_uncomp_chunk_size, &comp_chunk_sizes[wid], shared_mem, tmp_buf + wid*num_warps_per_block*shmem_size_group(), warp);
```

--------------------------------

### High-Level C++ Quickstart Code Snippets

Source: https://docs.nvidia.com/cuda/nvcomp/samples/highlevel_cpp_quickstart.html

Code snippets from the high-level C++ quickstart guide, including includes, namespace usage, and manager construction.

```cpp
#include "nvcomp/lz4.hpp"
#include "nvcomp.hpp"
```

```cpp
using namespace nvcomp;
```

```cpp
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));

std::shared_ptr<nvcompManagerBase> decomp_nvcomp_manager = create_manager(comp_buffer, stream);
```

--------------------------------

### Download example text files

Source: https://docs.nvidia.com/cuda/nvcomp/samples/nvcomp.html

Downloads two text files used in later examples: 'locke-essay-113.txt' and 'mobydick.txt'.

```python
import urllib.request
urllib.request.urlretrieve("http://textfiles.com/etext/NONFICTION/locke-essay-113.txt", "locke-essay-113.txt")
urllib.request.urlretrieve("http://textfiles.com/etext/FICTION/mobydick.txt", "mobydick.txt")
```

--------------------------------

### Custom Allocator Example

Source: https://docs.nvidia.com/cuda/nvcomp/py_api.html

Example of a custom allocator function signature for set_device_allocator.

```python
def my_allocator(nbytes: int, stream: nvcomp.Stream) -> PtrProtocol:
    return MyBuffer(nbytes, stream)
```

--------------------------------

### LZ4 CPU compression and NVRTC decompression example

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/examples/example_nvrtc.html

This example shows how to use nvCOMPDx with NVRTC and nvJitLink for LZ4 decompression. The device code is defined in a string and compiled at runtime. The host code uses the lz4 CPU library for compression. The example is located in the `example/nvcompdx/05_lz4_cpu_and_nvrtc` folder.

```cpp
// Copyright (c) 2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
//
// NVIDIA CORPORATION and its licensors retain all intellectual property
// and proprietary rights in and to this software, related documentation
// and any modifications thereto.  Any use, reproduction, disclosure or
// distribution of this software and related documentation without an express
// license agreement from NVIDIA CORPORATION is strictly prohibited.

#include <lz4.h>
#include <lz4hc.h>

#include <nvrtc.h>
#include <nvJitLink.h>

#include <nvcompdx.hpp>
#include "../common/batch_data.hpp"
#include "../common/util_nvrtc.hpp"

using namespace nvcompdx;

// This sample demonstrates the usage of the warp-level device API for
// LZ4 GPU decompression. The decompression kernel is compiled and linked
// during runtime. The compression happens through the host-side
// lz4 CPU library.

// LZ4 decompression kernel, using the preconfigured decompressor
// 1 warp per chunk, but multiple chunks per thread block

const char* decomp_kernel = R"kernel(
#include <nvcompdx.hpp>

using namespace nvcompdx;

extern "C" __global__ void decomp_warp_kernel(
    size_t batch_size,
    const void * const * comp_chunks,
    void * const * uncomp_chunks,
    const size_t * comp_chunk_sizes,
    size_t * decomp_chunk_sizes) {

  const unsigned int global_chunk_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
  const unsigned int local_chunk_id = threadIdx.x / 32;
  if(global_chunk_id >= batch_size) {
    return;
  }

  using decompressor_type =
    decltype(Algorithm<algorithm::ALGORITHM>() +
             DataType<datatype::DATATYPE>() +
             MaxUncompChunkSize<MAX_UNCOMP_CHUNK_SIZE>() +
             Direction<direction::decompress>() +
             Warp() +
             SM<ARCH>());

  auto decompressor = decompressor_type();
  constexpr auto shmem_size_warp = decompressor.shmem_size_group();
  extern __shared__ __align__(decompressor.shmem_alignment()) uint8_t shared_scratch_decomp_buffer[];

  decompressor.execute(
    comp_chunks[global_chunk_id],
    uncomp_chunks[global_chunk_id],
    comp_chunk_sizes[global_chunk_id],
    decomp_chunk_sizes + global_chunk_id,
    shared_scratch_decomp_buffer + shmem_size_warp * local_chunk_id,
    nullptr);
}
)kernel";

static std::string get_device_architecture_option(CUdevice& device)
{
  // Note:
  // -arch=compute_...       will generate PTX
  // -arch=sm_...            will generate SASS
  // -arch=sm_... with -dlto will generate LTO IR
  std::string option = "-arch=sm_" + std::to_string(get_device_architecture(device));
  return option;
}

```

--------------------------------

### Install nvCOMP C++ for CUDA 13.x

Source: https://docs.nvidia.com/cuda/nvcomp/installation.html

Install the nvCOMP C++ library for CUDA 13.x using pip. This installs the nvidia-libnvcomp-cu13 package.

```bash
pip install nvidia-libnvcomp-cu13
```

--------------------------------

### Install nvCOMP C++ for CUDA 12.x

Source: https://docs.nvidia.com/cuda/nvcomp/installation.html

Install the nvCOMP C++ library for CUDA 12.x using pip. This installs the nvidia-libnvcomp-cu12 package.

```bash
pip install nvidia-libnvcomp-cu12
```

--------------------------------

### NVRTC Example Code

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/examples/example_nvrtc.html

Two C++ code snippets from the NVRTC example: one for collecting compiler include directories, and one for running the NVRTC benchmark with LZ4 compression.

```cpp
static std::vector<std::string> get_comp_include_dirs()
{
#ifndef NVCOMPDX_INCLUDE_DIRS
  return std::vector<std::string>();
#else
  std::vector<std::string> comp_include_dirs_array;
  {
    std::string comp_include_dirs = NVCOMPDX_INCLUDE_DIRS;
    std::string delim             = ",";
    size_t      start             = 0U;
    size_t      end               = comp_include_dirs.find(delim);
    while (end != std::string::npos) {
      comp_include_dirs_array.push_back("--include-path=" +
        comp_include_dirs.substr(start, end - start));
      start = end + delim.length();
      end   = comp_include_dirs.find(delim);
    }
    if (end != std::string::npos) {
      comp_include_dirs_array.push_back("--include-path=" +
        comp_include_dirs.substr(start, end - start));
    }
  }
  return comp_include_dirs_array;
#endif
}
```

```cpp
template<unsigned int Arch>
static int run_nvrtc_example(const std::vector<std::vector<char>>& data)
{
  assert(!data.empty());

  size_t total_bytes = 0;
  for (const std::vector<char>& part : data) {
    total_bytes += part.size();
  }
  std::cout << "----------" << std::endl;
  std::cout << "files: " << data.size() << std::endl;
  std::cout << "uncompressed (B): " << total_bytes << std::endl;

  // Compile-time (de)compression parameters
  constexpr size_t num_warps_per_chunk = 1;
  constexpr size_t num_chunks_per_block = 4;
  constexpr size_t num_warps_per_block = num_warps_per_chunk * num_chunks_per_block;
  constexpr unsigned int block_size = static_cast<unsigned int>(num_warps_per_block * 32);
  constexpr size_t chunk_size = 1 << 16; // [bytes]

  // Build up input batch on CPU
  BatchDataCPU input_data_cpu(data, chunk_size);
  size_t batch_size = input_data_cpu.batch_size();
  std::cout << "chunks: " << batch_size << std::endl;

  // Allocate and prepare output/compressed batch
  BatchDataCPU compressed_data_cpu(
      LZ4_compressBound(chunk_size), batch_size);

  // Compressing on the CPU
  // loop over chunks on the CPU, compressing each one one by one
  for (size_t i = 0; i < batch_size; ++i) {
    // could use LZ4_compress_default or LZ4_compress_fast instead
    const int size = LZ4_compress_HC(
        static_cast<const char*>(input_data_cpu.chunk_ptrs()[i]),
        static_cast<char*>(compressed_data_cpu.chunk_ptrs()[i]),
        12 /* compression level */);
    if (size == 0) {
      throw std::runtime_error(
          "LZ4 CPU failed to compress chunk " + std::to_string(i) + ".");
    }

    // Set the actual compressed size
    compressed_data_cpu.chunk_sizes()[i] = static_cast<size_t>(size);
  }

  // Compute compression ratio
  size_t* compressed_sizes_host = compressed_data_cpu.chunk_sizes();
  size_t comp_bytes =
    std::accumulate(compressed_sizes_host, compressed_sizes_host + batch_size, size_t(0));

  std::cout << "comp_size: " << comp_bytes
            << ", compressed ratio: " << std::fixed << std::setprecision(2)
            << (double)total_bytes / comp_bytes << std::endl;

  // Configure the GPU decompressor
  using lz4_decompressor_type =
    decltype(Algorithm<algorithm::lz4>() +
             DataType<datatype::uint8>() +
             MaxUncompChunkSize<chunk_size>() +
             Direction<direction::decompress>() +
             Warp() +
             SM<Arch>());

  // Runtime decompression parameters
  const auto block_count =
    static_cast<unsigned int>((batch_size + num_chunks_per_block - 1) / num_chunks_per_block);
}
```

--------------------------------

### Main Function and Argument Parsing

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/examples/intro_example.html

This snippet shows the main function of the example. It parses command-line arguments for input and output file names, reads the input file, validates its size, runs the compression using the appropriate architecture runner, and writes the compressed output to a file.

```cpp
int main(int argc, char* argv[])
{
  std::string file_name_in;
  std::string file_name_out;

  size_t warmup_iteration_count = 2;
  size_t total_iteration_count = 5;

  do {
    if (argc < 3) {
      break;
    }

    int i = 1;
    while (i < argc) {
      const char* current_argv = argv[i++];
      if (strcmp(current_argv, "-f") == 0) {
        if(i < argc) {
          file_name_in = argv[i++];
        }
      }
      else if (strcmp(current_argv, "-o") == 0) {
        if(i < argc) {
          file_name_out = argv[i++];
        }
      }
      else {
        std::cerr << "Unknown option: " << current_argv << std::endl;
        print_usage();
        return 1;
      }
    }
  } while (0);

  if (file_name_in.empty()) {
    std::cerr << "Must specify one input file via '-f <file>'." << std::endl;
    print_usage();
    return 1;
  }

  if (file_name_out.empty()) {
    file_name_out = file_name_in + ".lz4";
  }

  auto input = read_file(file_name_in);

  if (input.size() > chunk_size) {
    std::cerr << "The input file is too large for this example. " << std::endl
              << "Select an input file up to " << chunk_size << " bytes." << std::endl;
    return 1;
  }

  std::vector<char> output;
  int ret = run_with_current_arch<Runner>(input,
                                          output,
                                          warmup_iteration_count,
                                          total_iteration_count);

  write_file(file_name_out, output);

  return 0;
}
```

--------------------------------

### Install nvCOMP Python for CUDA 13.x

Source: https://docs.nvidia.com/cuda/nvcomp/installation.html

Install the nvCOMP Python module for CUDA 13.x using pip. This installs the nvidia-nvcomp-cu13 package.

```bash
pip install nvidia-nvcomp-cu13
```

--------------------------------

### NVRTC and nvJitLink Example

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/examples/example_nvrtc.html

This code snippet demonstrates the process of setting up decompression buffers, compiling a CUDA kernel with NVRTC, and linking it with nvCOMPDx using nvJitLink. It includes scratch buffer allocation, NVRTC program creation, compilation options, LTO IR extraction, and linking with the nvCOMPDx library or fatbinary.

```cpp
  // Global scratch buffer
  // Note: lz4 decompression requires no global scratch buffer
  static_assert(lz4_decompressor_type().tmp_size_group() == 0);

  // Shared scratch buffer
  const auto decomp_shared_memory =
    static_cast<unsigned int>(lz4_decompressor_type().shmem_size_group() * num_chunks_per_block);

  // Copy compressed data to GPU
  BatchData compressed_data(compressed_data_cpu, true, lz4_decompressor_type().input_alignment());

  // Allocate and build up decompression batch on GPU
  BatchData decomp_data(input_data);
```

```cpp
  nvrtcProgram program;
  NVRTC_CHECK(nvrtcCreateProgram(&program,
                                 decomp_kernel,
                                 NULL /* CUDA program name */,
                                 0 /* numHeaders */,
                                 NULL /* headers */,
                                 NULL /* includeNames */));

  // Prepare compilation options
  CUdevice cuDevice;
  CU_CHECK(cuInit(0 /* flags */));
  CU_CHECK(cuDeviceGet(&cuDevice, 0 /* by default using the first device */));
  const auto gpu_architecture_option = get_device_architecture_option(cuDevice);
  std::vector<const char*> opts = {
    "--std=c++17",
    "--device-as-default-execution-space",
    "--include-path=" CUDAToolkit_INCLUDE_DIR, // Path to the CUDA include directory
    "--include-path=" CUDAToolkit_INCLUDE_DIR "/cuda/std", // Path to standard headers
    "--include-path=" CUDAToolkit_INCLUDE_DIR "/cuda/std" // Path to standard headers
  };
```

```cpp
  for (auto& config : comp_config_values) {
    opts.push_back(config.c_str());
  }

  nvrtcResult compileResult = nvrtcCompileProgram(program,
                                                  static_cast<int>(opts.size()),
                                                  opts.data());
  if (compileResult != NVRTC_SUCCESS) {
      // Obtain compilation log from the program if unsuccessful
      for (auto option : opts) {
          // (log handling omitted)
      }
  }

  // Obtain generated LTO IR from the program
  size_t lto_size;
  NVRTC_CHECK(nvrtcGetLTOIRSize(program, &lto_size));
  auto ltoir = std::make_unique<char[]>(lto_size);
  NVRTC_CHECK(nvrtcGetLTOIR(program, ltoir.get()));
```

```cpp
  // Load the generated Cubin and get a handle to our kernel
  CUcontext context;
  CU_CHECK(cuCtxCreate(&context, 0 /* flags */, cuDevice));

  // Load the generated LTO IR and the static nvCOMPDx LTO library
  nvJitLinkHandle linker;
  std::vector<const char*> lopts;
  lopts.emplace_back("-lto");
  lopts.emplace_back(gpu_architecture_option.c_str());
  NVJITLINK_CHECK(linker, nvJitLinkCreate(&linker,
                                          static_cast<uint32_t>(lopts.size()),
                                          lopts.data()));

  // Add the runtime-compiled kernel LTO IR
  NVJITLINK_CHECK(linker,
    nvJitLinkAddData(linker, NVJITLINK_INPUT_LTOIR, ltoir.get(), lto_size, "lto_online"));

  // Add nvCOMPDx LTO library or the nvCOMPDx fatbinary
  const char* fatbin_env_ptr = std::getenv("NVCOMPDX_EXAMPLE_NVCOMPDX_FATBIN");
  const char* library_env_ptr = std::getenv("NVCOMPDX_EXAMPLE_NVCOMPDX_LIBRARY");
  if(fatbin_env_ptr) {
    NVJITLINK_CHECK(linker, nvJitLinkAddFile(linker, NVJITLINK_INPUT_FATBIN, fatbin_env_ptr));
  } else if(library_env_ptr) {
    NVJITLINK_CHECK(linker, nvJitLinkAddFile(linker, NVJITLINK_INPUT_LIBRARY, library_env_ptr));
  } else {
#if defined(NVCOMPDX_FATBIN)
    NVJITLINK_CHECK(linker, nvJitLinkAddFile(linker, NVJITLINK_INPUT_FATBIN, NVCOMPDX_FATBIN));
#elif defined(NVCOMPDX_LIBRARY)
    NVJITLINK_CHECK(linker, nvJitLinkAddFile(linker, NVJITLINK_INPUT_LIBRARY, NVCOMPDX_LIBRARY));
#else
    std::cerr << "Please set one of the environment variables: "
              << "NVCOMPDX_EXAMPLE_NVCOMPDX_LIBRARY, "
              << "NVCOMPDX_EXAMPLE_NVCOMPDX_FATBIN, " << std::endl
              << "or during compilation define NVCOMPDX_LIBRARY or "
              << "NVCOMPDX_FATBIN." << std::endl;
    return 1;
#endif
  }
```

--------------------------------

### Install nvCOMP Python for CUDA 12.x

Source: https://docs.nvidia.com/cuda/nvcomp/installation.html

Install the nvCOMP Python module for CUDA 12.x using pip. This installs the nvidia-nvcomp-cu12 package.

```bash
pip install nvidia-nvcomp-cu12
```

--------------------------------

### Decompressing with nvcompdx

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/introduction.html

This example demonstrates a full decompression procedure using a warp-level kernel. It assumes the maximum uncompressed chunk size is known in advance.

```cpp
#include <nvcompdx.hpp>

using namespace nvcompdx;

// For simplicity, let's assume the maximum uncompressed
// chunk size is known prior to decompression.
constexpr size_t max_uncomp_chunk_size = 65536;

using DECOMP = decltype(Algorithm<algorithm::lz4>() +
                        DataType<datatype::uint8>() +
                        Direction<direction::decompress>() +
                        Warp() +
                        SM<800>());

__global__ void decomp_warp_kernel(const void * const comp_chunk,
                                   const size_t comp_chunk_size,
                                   void * uncomp_chunk,
                                   size_t * uncomp_chunk_size,
                                   uint8_t * tmp_buffer) {
   auto decompressor = DECOMP();
   constexpr size_t shmem_alignment = decompressor.shmem_alignment();
   extern __shared__ __align__(shmem_alignment) uint8_t shared_comp_scratch_buffer[];

   decompressor.execute(
      comp_chunk,
      uncomp_chunk,
      comp_chunk_size,
      uncomp_chunk_size,
      shared_comp_scratch_buffer,
      tmp_buffer);

   // At this point, the uncompressed chunk resides in the buffer pointed by
   // `uncomp_chunk`. The compressed input data `comp_chunk` is
   // left untouched.
}

int lz4_gpu_decomp_introduction(const std::vector<char>& compressed,
                                std::vector<char>& uncompressed)
{
  size_t total_bytes = data.size();

  // The input data consists of a single compressed chunk
  size_t num_chunks = 1;

  // Allocate buffer for the input (compressed) data
  // Note: with cudaMalloc() the input alignment is implicitly met
  void* d_input_data;
  cudaMalloc(&d_input_data, total_bytes);
  cudaMemcpy(d_input_data, data.data(), total_bytes, cudaMemcpyHostToDevice);

  // Allocate buffer for the input/output sizes
  size_t* d_output_size;
  cudaMalloc(&d_output_size, sizeof(size_t));

  // Allocate temporary scratch space
  // Note: with cudaMalloc() the global temporary memory alignment is implicitly met
  uint8_t* d_decomp_temp;
  cudaMalloc(&d_decomp_temp, DECOMP().tmp_size_total(num_chunks));

  // Allocate buffer for the output (uncompressed) data
  // The size of the output buffer is known in advance in this example
  // Note: with cudaMalloc() the output alignment is implicitly met
  void* d_output_data;
  cudaMalloc(&d_output_data, max_uncomp_chunk_size);

  // Decompression parameters
  // We are decompressing 1 chunk per thread block
  const unsigned int block_size = 32; // 1 warp
  const unsigned int block_count = static_cast<unsigned int>(num_chunks); // 1 chunk
  const auto decomp_shared_memory = DECOMP().shmem_size_group();

  decomp_warp_kernel<<<block_count, block_size, decomp_shared_memory>>>(
      d_input_data,
      total_bytes,
      d_output_data,
      d_output_size,
      d_decomp_temp
    );

  size_t uncomp_chunk_size;
  cudaMemcpy(&uncomp_chunk_size,
             d_output_size,
             sizeof(size_t),
             cudaMemcpyDeviceToHost);

  // Copy data back to host for write out
  uncompressed.resize(uncomp_chunk_size);
  cudaMemcpy(uncompressed.data(),
             d_output_data,
             uncomp_chunk_size,
             cudaMemcpyDeviceToHost);

  cudaFree(d_input_data);
  cudaFree(d_output_size);
  cudaFree(d_decomp_temp);
  cudaFree(d_output_data);

  return 0;
}
```

--------------------------------

### LZ4 GPU Compression Introduction

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/examples/intro_example.html

This code snippet is the main introductory example for nvCOMPDx, demonstrating GPU LZ4 compression using the warp-level device API. It includes a templated compression kernel, configuration of the compressor with the LZ4 algorithm, uint8 data type, compression direction, maximum uncompressed chunk size, warp execution, and SM architecture, followed by CUDA memory allocation and copying of input data.

```cpp
#include <cassert>
#include <iostream>
#include <algorithm>
#include <string>
#include <functional>

#include <nvcompdx.hpp>
#include "../common/util.hpp"

using namespace nvcompdx;


// This introductory sample demonstrates the usage of the warp-level device API
// for LZ4 GPU compression. The compressed buffer afterwards is written to disk.
// Note, however, that for optimal performance, one should perform multiple chunk
// compressions simultaneously. The example here only intends to showcase usage.

// LZ4 compression kernel, using the preconfigured compressor
// 1 warp per chunk
template<typename compressor_type>
__global__ void comp_warp_kernel(
  const void * const uncomp_chunk,
  const size_t uncomp_chunk_size,
  void * comp_chunk,
  size_t * comp_chunk_size,
  uint8_t * tmp_buffer) {
  // Note:
  // Given the (de)compressor expression has an SM<> operator,
  // it makes the fully-typed kernel only applicable on one targeted device architecture.
  // We need to signal to the compiler not to continue compiling this kernel whenever
  // the current compilation architecture is different from the one specified in
  // the SM<> operator.
  NVCOMPDX_SKIP_IF_NOT_APPLICABLE(compressor_type);

  auto compressor = compressor_type();
  constexpr size_t shmem_alignment = compressor.shmem_alignment();
  extern __shared__ __align__(shmem_alignment) uint8_t shared_comp_scratch_buffer[];
  assert(reinterpret_cast<uintptr_t>(shared_comp_scratch_buffer) % compressor.shmem_alignment() == 0);

  compressor.execute(
    uncomp_chunk,
    comp_chunk,
    uncomp_chunk_size,
    comp_chunk_size,
    shared_comp_scratch_buffer,
    tmp_buffer);
}

static constexpr size_t chunk_size = 1 << 16;

// Benchmark performance from the binary data file
template<unsigned int Arch>
int lz4_gpu_comp_introduction(const std::vector<char>& data,
                              std::vector<char>& compressed,
                              size_t warmup_iteration_count,
                              size_t total_iteration_count)
{
  assert(!data.empty());

  size_t total_bytes = data.size();
  size_t batch_size = 1; // a single chunk
  std::cout << "----------" << std::endl;
  std::cout << "uncompressed (B): " << total_bytes << std::endl;
  std::cout << "chunks: " << 1 << std::endl;

  // Configure the GPU compressor
  using lz4_compressor_type =
    decltype(Algorithm<algorithm::lz4>() +
             DataType<datatype::uint8>() +
             Direction<direction::compress>() +
             MaxUncompChunkSize<chunk_size>() +
             Warp() +
             SM<Arch>());

  // Allocate buffer for the input (uncompressed) data
  // Note: with cudaMalloc() the input alignment is implicitly met
  void* d_input_data;
  CUDA_CHECK(cudaMalloc(&d_input_data, total_bytes));
  CUDA_CHECK(cudaMemcpy(d_input_data, data.data(), total_bytes, cudaMemcpyHostToDevice));

  // Allocate buffer for the input/output sizes
  size_t* d_output_size;
  CUDA_CHECK(cudaMalloc(&d_output_size, sizeof(size_t)));

```

--------------------------------

### Compile and link with libnvcompdx.fatbin library

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/installation.html

Use this command for host platform agnostic builds, such as AArch64. Requires NVCC 12.8 or newer.

```bash
nvcc -dlto -std=c++17 -arch sm_XY (...) -I<mathdx_include_dir> <your_source_file>.cu -o <your_binary> libnvcompdx.fatbin
```

--------------------------------

### Compile and link with libnvcompdx.a static library

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/installation.html

Use this command for x86_64 Linux builds when linking against the static library.

```bash
nvcc -dlto -std=c++17 -arch sm_XY (...) -I<mathdx_include_dir> <your_source_file>.cu -o <your_binary> -lnvcompdx
```

--------------------------------

### nvcompBatchedGzipDecompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Gets the minimum buffer alignment requirements for decompression.

```APIDOC
## nvcompBatchedGzipDecompressGetRequiredAlignments

### Description
Gets the minimum buffer alignment requirements for decompression. Providing buffers with alignments above the minimum requirements (e.g., 16- or 32-byte alignment) may help improve performance.

### Method
Function call

### Endpoint
nvcompBatchedGzipDecompressGetRequiredAlignments

### Parameters
#### Input Parameters
- **decompress_opts** (nvcompBatchedGzipDecompressOpts_t) - [in] Decompression options.

#### Output Parameters
- **alignment_requirements** (nvcompAlignmentRequirements_t *) - [out] The minimum buffer alignment requirements for decompression.

### Returns
nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### nvcompBatchedCascadedDecompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Gets the minimum buffer alignment requirements for decompression.

```APIDOC
## nvcompBatchedCascadedDecompressGetRequiredAlignments

### Description
Gets the minimum buffer alignment requirements for decompression.

### Method
Function

### Signature
```c
nvcompStatus_t nvcompBatchedCascadedDecompressGetRequiredAlignments(
    nvcompBatchedCascadedDecompressOpts_t decompress_opts,
    nvcompAlignmentRequirements_t *alignment_requirements
);
```

### Parameters
- **decompress_opts** (nvcompBatchedCascadedDecompressOpts_t) - [in] Decompression options.
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - [out] The minimum buffer alignment requirements for decompression.

### Returns
- **nvcompStatus_t** - nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### Example: Query description traits to check completeness and retrieve values

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/api/traits_description.html

Demonstrates using `nvcompdx.hpp` and the description traits API to define a complete compression description with the ANS algorithm, uint8 datatype, compression direction, max uncompressed chunk size of 32768, and SM 800. Then it uses `is_complete_comp<COMP>` to check completeness and retrieves `max_uncomp_chunk_size_of<COMP>` and `sm_of<COMP>` for output.

```cpp
#include <nvcompdx.hpp>

using namespace nvcompdx;

using COMP = decltype(Algorithm<algorithm::ans>() +
                      DataType<datatype::uint8>() +
                      Direction<direction::compress>() +
                      MaxUncompChunkSize<32768>() +
                      SM<800>());

if constexpr (is_complete_comp<COMP>::value) {
  std::cout << "Compressor (maximum uncompressed chunk size, SM): "
            << max_uncomp_chunk_size_of<COMP>::value << ", "
            << sm_of<COMP>::value << std::endl;
}
```

--------------------------------

### Zero-copy import host array

Source: https://docs.nvidia.com/cuda/nvcomp/samples/nvcomp.html

Creates a numpy array and imports it as an nvCOMP array without copying, then prints its properties.

```python
ascending = np.arange(0, 4096, dtype=np.int32)
nvarr_h = nvcomp.as_array(ascending)
```

```python
print(ascending. __array_interface__ )
print(nvarr_h. __array_interface__ )
print(nvarr_h. __cuda_array_interface__ )
print(nvarr_h.buffer_size)
print(nvarr_h.buffer_kind)
print(nvarr_h.ndim)
print(nvarr_h.dtype)
print(nvarr_h.shape)
print(nvarr_h.strides)
print(nvarr_h.item_size)
print(nvarr_h.size)
```

--------------------------------

### nvcompBatchedANSDecompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Gets the minimum buffer alignment requirements for decompression.

```APIDOC
## nvcompBatchedANSDecompressGetRequiredAlignments

### Description
Gets the minimum buffer alignment requirements for decompression. Note: Providing buffers with alignments above the minimum requirements (e.g., 16- or 32-byte alignment) may help improve performance.

### Method
Function

### Endpoint
nvcompStatus_t nvcompBatchedANSDecompressGetRequiredAlignments( nvcompBatchedANSDecompressOpts_t decompress_opts, nvcompAlignmentRequirements_t *alignment_requirements )

### Parameters
#### Input
- **decompress_opts** (nvcompBatchedANSDecompressOpts_t) - Required - Decompression options.

#### Output
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - Required - The minimum buffer alignment requirements for decompression.

### Returns
nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### bitcompCreateBatchPlan

Source: https://docs.nvidia.com/cuda/nvcomp/native_api.html

Creates a bitcomp plan for batched compression and decompression, supporting lossless and lossy modes.

```APIDOC
## bitcompCreateBatchPlan

### Description
Creates a bitcomp plan for compression and decompression of batched inputs, lossy or lossless. The lossless compression can be used on any data type, viewed as integral type. Lossy compression is only available for floating point data types and is based on quantization.

### Method
C function

### Endpoint
bitcompCreateBatchPlan(bitcompHandle_t* handle, size_t nbatch, bitcompDataType_t dataType, bitcompMode_t mode, bitcompAlgorithm_t algo)

### Parameters
- **handle** (bitcompHandle_t*) - [out] Handle created.
- **nbatch** (size_t) - [in] Number of batches to process.
- **dataType** (bitcompDataType_t) - [in] Datatype of the uncompressed data.
- **mode** (bitcompMode_t) - [in] Compression mode, lossless or lossy to signed / lossy to unsigned.
- **algo** (bitcompAlgorithm_t) - [in] Which compression algorithm to use.

### Returns
Returns BITCOMP_SUCCESS if successful, or an error.
```

--------------------------------

### nvcompBatchedANSCompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Get the minimum buffer alignment requirements for compression.

```APIDOC
## nvcompBatchedANSCompressGetRequiredAlignments

### Description
Get the minimum buffer alignment requirements for compression. Providing buffers with alignments above the minimum requirements (e.g., 16- or 32-byte alignment) may help improve performance.

### Method
Function

### Signature
nvcompStatus_t nvcompBatchedANSCompressGetRequiredAlignments(
    nvcompBatchedANSCompressOpts_t compress_opts,
    nvcompAlignmentRequirements_t *alignment_requirements
)

### Parameters
- **compress_opts** (nvcompBatchedANSCompressOpts_t) - [in] Compression options.
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - [out] The minimum buffer alignment requirements for compression.

### Returns
- nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### nvcompBatchedDeflateDecompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Function to get the minimum buffer alignment requirements for decompression.

```APIDOC
## nvcompBatchedDeflateDecompressGetRequiredAlignments

### Description
Gets the minimum buffer alignment requirements for decompression.

### Method
Function call

### Endpoint
nvcompBatchedDeflateDecompressGetRequiredAlignments

### Parameters
#### Input Parameters
- **decompress_opts** (nvcompBatchedDeflateDecompressOpts_t) - [in] Decompression options.

#### Output Parameters
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - [out] The minimum buffer alignment requirements for decompression.

### Returns
nvcompStatus_t - nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### bitcompCreatePlanFromCompressedData

Source: https://docs.nvidia.com/cuda/nvcomp/native_api.html

Creates a handle from existing compressed data, extracting all handle parameters from the provided compressed data pointer.

```APIDOC
## bitcompCreatePlanFromCompressedData

### Description
Creates a handle from existing compressed data.

### Signature
bitcompResult_t bitcompCreatePlanFromCompressedData(bitcompHandle_t* handle, const void* data)

### Parameters
- **handle** ([out] bitcompHandle_t*) - Handle created.
- **data** ([in] const void*) - Pointer to the compressed data, from which all the handle parameters will be extracted.

### Returns
Returns BITCOMP_SUCCESS if successful, or an error.
```

--------------------------------

### nvcompBatchedLZ4DecompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/genindex.html

Gets the required alignment requirements for batched LZ4 decompression.

```APIDOC
## nvcompBatchedLZ4DecompressGetRequiredAlignments

### Description
Gets the required alignment requirements for batched LZ4 decompression.

### Method
C++ function

### Signature
```cpp
nvcompBatchedLZ4DecompressGetRequiredAlignments(nvcompBatchedLZ4DecompressOpts_t format_opts, nvcompAlignmentRequirements_t* alignment_requirements)
```

### Parameters
- **format_opts** (nvcompBatchedLZ4DecompressOpts_t) - Options for LZ4 decompression.
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - Output: alignment requirements.

### Returns
- **void** - This function does not return a value; results are written to output parameters.

```

--------------------------------

### bitcompCreatePlan

Source: https://docs.nvidia.com/cuda/nvcomp/native_api.html

Creates a bitcomp plan for compression and decompression. Supports lossless compression on integral types and lossy compression on floating point types with quantization to signed or unsigned integers.

```APIDOC
## bitcompCreatePlan

### Description
Creates a bitcomp plan for compression and decompression, lossy or lossless. The lossless compression can be used on any data type, viewed as integral type. Lossy compression is only available for floating point data types and is based on quantization.

### Signature
bitcompResult_t bitcompCreatePlan(bitcompHandle_t* handle, size_t n, bitcompDataType_t dataType, bitcompMode_t mode, bitcompAlgorithm_t algo)

### Parameters
- **handle** ([out] bitcompHandle_t*) - Output handle created.
- **n** ([in] size_t) - Size of the uncompressed data in bytes.
- **dataType** ([in] bitcompDataType_t) - Datatype of the uncompressed data.
- **mode** ([in] bitcompMode_t) - Compression mode: lossless or lossy to signed / lossy to unsigned.
- **algo** ([in] bitcompAlgorithm_t) - Which compression algorithm to use.

### Returns
Returns BITCOMP_SUCCESS if successful, or an error.
```

--------------------------------

### Zero-copy import device array

Source: https://docs.nvidia.com/cuda/nvcomp/samples/nvcomp.html

Creates a CuPy array and imports it as an nvCOMP array without copying, then prints its properties.

```python
data_gpu = cp.array(ascending)
nvarr_d = nvcomp.as_array(data_gpu)
print(data_gpu. __cuda_array_interface__ )
print(nvarr_d. __cuda_array_interface__ )
print(nvarr_d.buffer_kind)
print(nvarr_d.ndim)
print(nvarr_d.dtype)
print(nvarr_d.shape)
print(nvarr_d.strides)
print(nvarr_d.item_size)
print(nvarr_d.size)
```

--------------------------------

### nvcompBatchedDeflateCompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Function to get the minimum buffer alignment requirements for compression.

```APIDOC
## nvcompBatchedDeflateCompressGetRequiredAlignments

### Description
Gets the minimum buffer alignment requirements for compression.

### Method
Function call

### Endpoint
nvcompBatchedDeflateCompressGetRequiredAlignments

### Parameters
#### Input Parameters
- **compress_opts** (nvcompBatchedDeflateCompressOpts_t) - [in] The Deflate compression options to use.

#### Output Parameters
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - [out] The minimum buffer alignment requirements for compression.

### Returns
nvcompStatus_t - nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### nvcompBatchedLZ4CompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/genindex.html

Gets the required alignment requirements for batched LZ4 compression.

```APIDOC
## nvcompBatchedLZ4CompressGetRequiredAlignments

### Description
Gets the required alignment requirements for batched LZ4 compression.

### Method
C++ function

### Signature
```cpp
nvcompBatchedLZ4CompressGetRequiredAlignments(nvcompBatchedLZ4CompressOpts_t format_opts, nvcompAlignmentRequirements_t* alignment_requirements)
```

### Parameters
- **format_opts** (nvcompBatchedLZ4CompressOpts_t) - Options for LZ4 compression.
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - Output: alignment requirements.

### Returns
- **void** - This function does not return a value; results are written to output parameters.

```

--------------------------------

### bitcompCreateBatchPlanFromCompressedData

Source: https://docs.nvidia.com/cuda/nvcomp/native_api.html

Creates a batch handle from existing compressed data, allowing subsequent compression or decompression operations.

```APIDOC
## bitcompCreateBatchPlanFromCompressedData

### Description
Creates a batch handle from batch-compressed data. The data must be device-visible. Returns an error if the compressed data is invalid or if batches have not been compressed with the same parameters. This call triggers synchronous activity in the default stream of the GPU.

### Method
C function

### Endpoint
bitcompCreateBatchPlanFromCompressedData(bitcompHandle_t* handle, const void* const* data, size_t batches, void* const device_temp_ptr = nullptr, cudaStream_t stream = 0)

### Parameters
- **handle** (bitcompHandle_t*) - [out] Output handle, which can be used for batch compression or decompression.
- **data** (const void* const*) - [in] Device-visible pointers to the device-visible data of each batch.
- **batches** (size_t) - [in] Number of batches.
- **device_temp_ptr** (void* const) - [in] Temporary scratch memory used to store compression information. Optional, defaults to nullptr.
- **stream** (cudaStream_t) - [in] CUDA stream to use. Optional, defaults to 0.

### Returns
Returns BITCOMP_SUCCESS if successful, or an error.
```

--------------------------------

### nvcompBatchedCascadedCompressGetTempSize

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Gets the amount of temporary memory required on the GPU for compression.

```APIDOC
## nvcompBatchedCascadedCompressGetTempSize

### Description
Gets the amount of temporary memory required on the GPU for compression.

### Method
Function

### Signature
```c
nvcompStatus_t nvcompBatchedCascadedCompressGetTempSize(
    size_t num_chunks,
    size_t max_uncompressed_chunk_bytes,
    nvcompBatchedCascadedCompressOpts_t compress_opts,
    size_t *temp_bytes
);
```

### Parameters
- **num_chunks** (size_t) - [in] Number of chunks of data to compress.
- **max_uncompressed_chunk_bytes** (size_t) - [in] The size of the largest uncompressed chunk. This parameter is currently unused. Set it to either the actual value or zero.
- **compress_opts** (nvcompBatchedCascadedCompressOpts_t) - [in] The cascaded format options. The format must be valid.
- **temp_bytes** (size_t*) - [out] The amount of GPU memory that will be temporarily required during compression. The value is returned on the host side.

### Returns
- **nvcompStatus_t** - nvcompSuccess if successful, and an error code otherwise.
```

--------------------------------

### Create LZ4 codec

Source: https://docs.nvidia.com/cuda/nvcomp/samples/nvcomp.html

Initializes a Codec instance for the LZ4 algorithm.

```python
lz4_codec = nvcomp.Codec(algorithm="LZ4")
```

--------------------------------

### nvcompBatchedCascadedCompressGetMaxOutputChunkSize

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Function to get the maximum size that a chunk could compress to.

```APIDOC
## nvcompBatchedCascadedCompressGetMaxOutputChunkSize

### Description
Gets the maximum size that a chunk of size at most max_uncompressed_chunk_bytes could compress to. That is, the minimum amount of output memory required to be given nvcompBatchedCascadedCompressAsync for each chunk.

### Method
Function

### Endpoint
nvcompBatchedCascadedCompressGetMaxOutputChunkSize

### Parameters
#### Input Parameters
- **max_uncompressed_chunk_bytes** (size_t) - [in] The maximum size of a chunk before compression.
- **compress_opts** (nvcompBatchedCascadedCompressOpts_t) - [in] The Cascaded compression options to use.

#### Output Parameters
- **max_compressed_chunk_bytes** (size_t*) - [out] The maximum possible compressed size of the chunk.

### Returns
nvcompSuccess if successful, and an error code otherwise.

```

--------------------------------

### Execute COMP in a CUDA Kernel

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/introduction.html

Demonstrates instantiating the COMP class and executing it within a global CUDA kernel. It includes querying the shared memory alignment and using a scratch buffer for the compression process.

```cpp
__global__ void comp_warp_kernel(const void * const uncomp_chunk,
                                 const size_t uncomp_chunk_size,
                                 void * comp_chunk,
                                 size_t * comp_chunk_size,
                                 uint8_t * tmp_buffer) {
   auto compressor = COMP();
   constexpr size_t shmem_alignment = compressor.shmem_alignment();
   extern __shared__ __align__(shmem_alignment) uint8_t shared_comp_scratch_buffer[];

   compressor.execute(
      uncomp_chunk,
      comp_chunk,
      uncomp_chunk_size,
      comp_chunk_size,
      shared_comp_scratch_buffer,
      tmp_buffer);

   // At this point, the compressed chunk resides in the buffer pointed by
   // `comp_chunk`. The uncompressed input data `uncomp_chunk` is
   // left untouched.
}
```

--------------------------------

### Verify completeness with execution traits

Source: https://docs.nvidia.com/cuda/nvcompdx/0.1.0/api/traits_execution.html

Example of using execution traits to verify completeness before invoking execution.

```cpp
#include <nvcompdx.hpp>

using namespace nvcompdx;

using COMP = decltype(Algorithm<algorithm::ans>() +
                      DataType<datatype::uint8>() +
                      Direction<direction::compress>() +
                      MaxUncompChunkSize<32768>() +
                      SM<800>() +
                      Block() +
                      BlockWarp<8, true>());

if constexpr (is_complete_comp_execution<COMP>::value) {
  // Perform the described compression operation
  COMP::execute(...);
}

```

--------------------------------

### nvcompBatchedCascadedCompressGetRequiredAlignments

Source: https://docs.nvidia.com/cuda/nvcomp/c_api.html

Function to get the minimum buffer alignment requirements for compression.

```APIDOC
## nvcompBatchedCascadedCompressGetRequiredAlignments

### Description
Gets the minimum buffer alignment requirements for compression. Providing buffers with alignments above the minimum requirements (e.g., 16- or 32-byte alignment) may help improve performance.

### Method
Function

### Endpoint
nvcompBatchedCascadedCompressGetRequiredAlignments

### Parameters
#### Input Parameters
- **compress_opts** (nvcompBatchedCascadedCompressOpts_t) - [in] Compression options.

#### Output Parameters
- **alignment_requirements** (nvcompAlignmentRequirements_t*) - [out] The minimum buffer alignment requirements for compression.

### Returns
nvcompSuccess if successful, and an error code otherwise.

```
Tokens:
Raw