nvCOMPDx#

nvCOMPDx provides runtime code-generation for GPU compression and decompression using the nvCOMP library’s device extensions. It supports LZ4 and ANS compression algorithms with warp-level and block-level cooperative execution. nvCompDx can be found here.

Overview#

nvCOMPDx generates extern "C" __device__ functions that perform compression or decompression of a single data chunk. The generated code is distributed as LTO-IR and must be linked with the nvCOMPDx universal fatbin (via nvJitLink) before loading onto the GPU.

Supported configurations:

Algorithm support matrix#
Algorithm	Supported data types	Max uncompressed chunk size
LZ4	`COMMONDX_R_8UI`, `COMMONDX_R_16UI`, `COMMONDX_R_32UI`	16,777,216 bytes (16 MiB)
ANS	`COMMONDX_R_8UI`, `COMMONDX_R_16F`	262,144 bytes (256 KiB)

Execution modes:

Warp (COMMONDX_EXECUTION_WARP): A single warp (32 threads) processes one chunk. Participating threads must start at %laneid = 0.
Block (COMMONDX_EXECUTION_BLOCK): An entire thread block cooperates on one chunk. Exactly one of NVCOMPDX_OPERATOR_BLOCK_DIM (block dimensions) or NVCOMPDX_OPERATOR_BLOCK_WARP (number of warps) must be set.

Workflow#

Create a descriptor using nvcompdxCreateDescriptor().
Set operators:
- NVCOMPDX_OPERATOR_ALGORITHM (required): LZ4 or ANS.
- NVCOMPDX_OPERATOR_DIRECTION (required): compress or decompress.
- NVCOMPDX_OPERATOR_DATATYPE (required): a commondxValueType (e.g. COMMONDX_R_8UI).
- NVCOMPDX_OPERATOR_SM (required): target SM architecture (e.g. 800 for SM 8.0).
- NVCOMPDX_OPERATOR_EXECUTION (required): COMMONDX_EXECUTION_WARP or COMMONDX_EXECUTION_BLOCK.
- NVCOMPDX_OPERATOR_MAX_UNCOMP_CHUNK_SIZE (required for compress): maximum uncompressed chunk size in bytes. Not needed for decompress.
- NVCOMPDX_OPERATOR_BLOCK_DIM (block only): {x, y, z} thread block dimensions; total threads must be at least 32. Mutually exclusive with NVCOMPDX_OPERATOR_BLOCK_WARP. Exactly one of BLOCK_DIM or BLOCK_WARP must be set for COMMONDX_EXECUTION_BLOCK.
- NVCOMPDX_OPERATOR_BLOCK_WARP (block only): {num_warps, complete} — number of warps (> 0) and completeness flag (0 or 1). Mutually exclusive with NVCOMPDX_OPERATOR_BLOCK_DIM. Exactly one of BLOCK_DIM or BLOCK_WARP must be set for COMMONDX_EXECUTION_BLOCK.
Query traits (locks the descriptor):
- NVCOMPDX_TRAIT_SHMEM_SIZE_GROUP: shared memory bytes needed per warp/block.
- NVCOMPDX_TRAIT_TMP_SIZE_GROUP: global temporary memory bytes per warp/block.
- NVCOMPDX_TRAIT_MAX_COMP_CHUNK_SIZE (compress only): worst-case compressed size.
- NVCOMPDX_TRAIT_SHMEM_ALIGNMENT, NVCOMPDX_TRAIT_TMP_ALIGNMENT, NVCOMPDX_TRAIT_INPUT_ALIGNMENT, NVCOMPDX_TRAIT_OUTPUT_ALIGNMENT: alignment requirements (power of 2).
- NVCOMPDX_TRAIT_BLOCK_DIM: actual block dimensions. Only queryable when NVCOMPDX_OPERATOR_BLOCK_DIM was set explicitly; not derivable from NVCOMPDX_OPERATOR_BLOCK_WARP.
- NVCOMPDX_TRAIT_BLOCK_WARP: {num_warps, complete} block warp configuration. Queryable when NVCOMPDX_OPERATOR_BLOCK_WARP was set explicitly (echoed) or when NVCOMPDX_OPERATOR_BLOCK_DIM was set with COMMONDX_EXECUTION_BLOCK (derived).
- NVCOMPDX_TRAIT_SYMBOL_NAME: the extern "C" device function symbol.
Generate code:
1. Create a code handle with commondxCreateCode().
2. Set the target SM with commondxSetCodeOptionInt64().
3. Generate code with nvcompdxFinalizeCode() and extract with commondxGetCodeLTOIR().
4. Retrieve the universal fatbin with nvcompdxGetUniversalFATBINSize() / nvcompdxGetUniversalFATBIN().
5. Link the generated LTO-IR together with the universal fatbin using nvJitLink.
Destroy the descriptor with nvcompdxDestroyDescriptor().

Device function ABI#

The generated device function is extern "C" __device__ with the following signature for both compression and decompression:

void symbol(void*     input,
            void*     output,
            size_t*   input_size,
            size_t*   output_size,
            uint8_t*  shmem,
            uint8_t*  tmp);

Parameters:

input: pointer to the input chunk (uncompressed data for compress, compressed data for decompress).
output: pointer to the output buffer (compressed output for compress, decompressed output for decompress). Pass nullptr during decompression to query only the output size.
input_size: pointer to the input chunk size in bytes.
output_size: pointer written with the output size in bytes.
shmem: shared memory scratch buffer, at least NVCOMPDX_TRAIT_SHMEM_SIZE_GROUP bytes, aligned to NVCOMPDX_TRAIT_SHMEM_ALIGNMENT.
tmp: global memory scratch buffer, at least NVCOMPDX_TRAIT_TMP_SIZE_GROUP bytes, aligned to NVCOMPDX_TRAIT_TMP_ALIGNMENT.

For compression, the MaxUncompChunkSize is a compile-time parameter baked into the generated code and does not appear in the runtime signature.

Example: compress and decompress roundtrip#

// 1. Create compress descriptor
nvcompdxDescriptor comp;
nvcompdxCreateDescriptor(&comp);
nvcompdxSetOperatorInt64(comp, NVCOMPDX_OPERATOR_ALGORITHM,             NVCOMPDX_ALGORITHM_LZ4);
nvcompdxSetOperatorInt64(comp, NVCOMPDX_OPERATOR_DIRECTION,             NVCOMPDX_DIRECTION_COMPRESS);
nvcompdxSetOperatorInt64(comp, NVCOMPDX_OPERATOR_DATATYPE,              COMMONDX_R_8UI);
nvcompdxSetOperatorInt64(comp, NVCOMPDX_OPERATOR_SM,                    800);
nvcompdxSetOperatorInt64(comp, NVCOMPDX_OPERATOR_EXECUTION,             COMMONDX_EXECUTION_WARP);
nvcompdxSetOperatorInt64(comp, NVCOMPDX_OPERATOR_MAX_UNCOMP_CHUNK_SIZE, 65536);

// 2. Query traits
long long int shmem_size, max_comp_size, tmp_size;
nvcompdxGetTraitInt64(comp, NVCOMPDX_TRAIT_SHMEM_SIZE_GROUP,    &shmem_size);
nvcompdxGetTraitInt64(comp, NVCOMPDX_TRAIT_MAX_COMP_CHUNK_SIZE, &max_comp_size);
nvcompdxGetTraitInt64(comp, NVCOMPDX_TRAIT_TMP_SIZE_GROUP,      &tmp_size);

// 3. Generate code and link with fatbin
commondxCode code;
commondxCreateCode(&code);
commondxSetCodeOptionInt64(code, COMMONDX_OPTION_TARGET_SM, 800);
nvcompdxFinalizeCode(code, comp);
// ... extract LTO-IR with commondxGetCodeLTOIR, get fatbin with nvcompdxGetUniversalFATBIN,
// link with nvJitLink, load cubin with cuModuleLoadData, launch kernel ...

// 4. Cleanup
commondxDestroyCode(code);
nvcompdxDestroyDescriptor(comp);

API reference#

typedef long long int nvcompdxDescriptor#

An nvCOMPDx descriptor.

Equivalent to using COMP = decltype(...) in nvCOMPDx C++.

enum nvcompdxAlgorithm_t#

Compression algorithm.

Values:

enumerator NVCOMPDX_ALGORITHM_LZ4#: LZ4 - general-purpose byte-level compressor

enumerator NVCOMPDX_ALGORITHM_ANS#: ANS - proprietary entropy encoder

enum nvcompdxDirection_t#

Compression direction.

For NVCOMPDX_DIRECTION_COMPRESS the device function signature is: void compress(void* input, void* output, size_t* input_size, size_t* output_size, uint8_t* shmem, uint8_t* tmp) where input points to uncompressed data, output receives compressed data, *input_size is the uncompressed chunk size in bytes, *output_size is written with the compressed size, shmem is a shared-memory scratch buffer of at least NVCOMPDX_TRAIT_SHMEM_SIZE_GROUP bytes, and tmp is a global-memory scratch buffer of at least NVCOMPDX_TRAIT_TMP_SIZE_GROUP bytes.

For NVCOMPDX_DIRECTION_DECOMPRESS the device function signature is: void decompress(void* input, void* output, size_t* input_size, size_t* output_size, uint8_t* shmem, uint8_t* tmp) where input points to compressed data, output receives decompressed data, *input_size is the compressed chunk size, and *output_size is written with the decompressed size. Pass output = nullptr to query only the decompressed size without decompressing.

The actual symbol name can be retrieved using NVCOMPDX_TRAIT_SYMBOL_NAME .

Values:

enumerator NVCOMPDX_DIRECTION_COMPRESS#: Compress raw data into compressed format

enumerator NVCOMPDX_DIRECTION_DECOMPRESS#: Decompress compressed data back to original format

enum nvcompdxOperatorType_t#

Operators.

The set of supported nvCOMPDx operators.

Values:

enumerator NVCOMPDX_OPERATOR_ALGORITHM#: Operator data type: nvcompdxAlgorithm_t. Operator definition: required

enumerator NVCOMPDX_OPERATOR_DIRECTION#: Operator data type: nvcompdxDirection_t. Operator definition: required

enumerator NVCOMPDX_OPERATOR_DATATYPE#: Operator data type: commondxValueType_t. Supported: COMMONDX_R_8UI, COMMONDX_R_16UI, COMMONDX_R_32UI (LZ4), COMMONDX_R_8UI, COMMONDX_R_16F (ANS). Operator definition: required

enumerator NVCOMPDX_OPERATOR_MAX_UNCOMP_CHUNK_SIZE#: Operator data type: long long int. Expected content: maximum uncompressed chunk size in bytes. Operator definition: required for compress; not needed for decompress

enumerator NVCOMPDX_OPERATOR_SM#: Operator data type: long long int. Expected content: 750 (Turing), 800 (Ampere), …, Operator definition: required

enumerator NVCOMPDX_OPERATOR_EXECUTION#: Operator data type: commondxExecution_t. COMMONDX_EXECUTION_WARP maps to Warp execution (32 threads), COMMONDX_EXECUTION_BLOCK maps to Block execution. Operator definition: required

enumerator NVCOMPDX_OPERATOR_BLOCK_DIM#: Operator data type: long long int * 3. Expected content: <x, y, z> block dimensions. Total threads must be >= 32 (at least one warp). Only valid with COMMONDX_EXECUTION_BLOCK. Mutually exclusive with NVCOMPDX_OPERATOR_BLOCK_WARP. Operator definition: optional, but exactly one of BLOCK_DIM or BLOCK_WARP must be set for COMMONDX_EXECUTION_BLOCK.

enumerator NVCOMPDX_OPERATOR_BLOCK_WARP#: Operator data type: long long int * 2. Expected content: <num_warps, complete>. num_warps: number of warps (> 0). complete: 0 or 1. Mutually exclusive with NVCOMPDX_OPERATOR_BLOCK_DIM. Operator definition: optional, but exactly one of BLOCK_DIM or BLOCK_WARP must be set for COMMONDX_EXECUTION_BLOCK.

enum nvcompdxTraitType_t#

Traits.

The set of supported types of traits that can be accessed from finalized sources that use nvCOMPDx.

Values:

enumerator NVCOMPDX_TRAIT_SHMEM_SIZE_GROUP#: Trait data type: long long int. Value: shared memory size per warp/block, in bytes.

enumerator NVCOMPDX_TRAIT_TMP_SIZE_GROUP#: Trait data type: long long int. Value: global temporary memory size per warp/block, in bytes.

enumerator NVCOMPDX_TRAIT_MAX_COMP_CHUNK_SIZE#: Trait data type: long long int. Value: maximum compressed chunk size, in bytes. Only valid for compressors.

enumerator NVCOMPDX_TRAIT_SHMEM_ALIGNMENT#: Trait data type: long long int. Value: shared memory alignment requirement (power of 2).

enumerator NVCOMPDX_TRAIT_TMP_ALIGNMENT#: Trait data type: long long int. Value: global temporary memory alignment requirement.

enumerator NVCOMPDX_TRAIT_INPUT_ALIGNMENT#: Trait data type: long long int. Value: input buffer alignment requirement.

enumerator NVCOMPDX_TRAIT_OUTPUT_ALIGNMENT#: Trait data type: long long int. Value: output buffer alignment requirement.

enumerator NVCOMPDX_TRAIT_BLOCK_DIM#: Trait data type: long long int * 3. Expected content: <x, y, z> block dimensions. Only queryable when NVCOMPDX_OPERATOR_BLOCK_DIM was set explicitly; the trait is not derivable from NVCOMPDX_OPERATOR_BLOCK_WARP.

enumerator NVCOMPDX_TRAIT_SYMBOL_NAME#: Trait data type: C-string Value: symbol (device function) name.

enumerator NVCOMPDX_TRAIT_BLOCK_WARP#: Trait data type: long long int * 2. Value: <num_warps, complete> block warp configuration. Queryable when NVCOMPDX_OPERATOR_BLOCK_WARP was set explicitly (echoed) or when NVCOMPDX_OPERATOR_BLOCK_DIM was set with COMMONDX_EXECUTION_BLOCK (derived as <total_threads / 32, total_threads % 32 == 0>).

commondxStatusType nvcompdxGetVersion( int *major, int *minor, int *patch )#

Returns the major.minor.patch version of nvCOMPDx.

Parameters:

major – [out] The major version
minor – [out] The minor version
patch – [out] The patch version

Returns:

COMMONDX_SUCCESS

commondxStatusType nvcompdxCreateDescriptor( nvcompdxDescriptor *handle )#

Creates an nvCOMPDx descriptor.

Parameters:: handle – [inout] A pointer to a descriptor handle. As output, an initialized nvCOMPDx descriptor
Returns:: COMMONDX_SUCCESS on success, or an error code.

commondxStatusType nvcompdxSetOperatorInt64( nvcompdxDescriptor handle, nvcompdxOperatorType op, long long int value )#

Sets an integer operator on an nvCOMPDx descriptor.

Parameters:

handle – [in] An nvCOMPDx descriptor, output of nvcompdxCreateDescriptor
op – [in] The operator to set.
value – [in] A value for the operator

Returns: