Is this page helpful?

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`	16,777,216 bytes (16 MiB)

Note

ANS payloads use the nvCOMP 6.0 / nvCOMPDx 0.2.0 compressed chunk format, which is incompatible with earlier nvCOMP/nvCOMPDx releases. Compression and decompression must use compatible versions; existing ANS data from the previous format requires its original decoder.

See the nvCOMPDx release notes.

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. Code generation requires exactly one of NVCOMPDX_OPERATOR_BLOCK_DIM (block dimensions) or NVCOMPDX_OPERATOR_BLOCK_WARP (number of warps).

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, except for MAX_COMP_CHUNK_SIZE queries): a commondxValueType (e.g. COMMONDX_R_8UI).
- NVCOMPDX_OPERATOR_SM (required for code generation and SYMBOL_NAME): target SM architecture (e.g. 800 for SM 8.0). Numeric traits do not require it.
- NVCOMPDX_OPERATOR_EXECUTION (required, except for MAX_COMP_CHUNK_SIZE queries): COMMONDX_EXECUTION_WARP or COMMONDX_EXECUTION_BLOCK.
- NVCOMPDX_OPERATOR_MAX_UNCOMP_CHUNK_SIZE: maximum uncompressed chunk size in bytes. Required for compression code generation and its TMP_SIZE_GROUP and MAX_COMP_CHUNK_SIZE traits. Other traits and decompression do not require it.
- 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. Block execution requires a layout for code generation, shared-memory size, layout traits, and the symbol; scratch size and alignments do not require it.
- NVCOMPDX_OPERATOR_BLOCK_WARP (block only): {num_warps, complete} — number of warps (> 0) and completeness enum (NVCOMPDX_BLOCK_WARP_INCOMPLETE or NVCOMPDX_BLOCK_WARP_COMPLETE). Mutually exclusive with NVCOMPDX_OPERATOR_BLOCK_DIM and subject to the same layout requirements.
- NVCOMPDX_OPERATOR_FUNCTION_ABI (optional): COMMONDX_FUNCTION_ABI_NATIVE or COMMONDX_FUNCTION_ABI_TILE. The native function ABI is used when this operator is omitted.
Query traits (locks the descriptor):

Set the chunk size before querying traits if the compression descriptor will also be used for size queries or code generation.

A standalone MAX_COMP_CHUNK_SIZE query needs only the algorithm, direction (compress), and maximum uncompressed chunk size. Datatype, execution, layout, SM, and function ABI are unnecessary.

Scratch-size and alignment queries need algorithm, direction, datatype, and execution; compression scratch size also needs the maximum uncompressed chunk size. They do not require SM, block layout, or an explicit function ABI. Shared-memory size and layout queries additionally require a valid block layout for block execution, but still do not require SM. Set all operators needed for later queries or code generation before the first query 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__. Compression and decompression use the same signature for a selected function ABI:

Native function ABI	Tile function ABI
`void symbol(const void* input, void* output, const unsigned long long* input_size, unsigned long long* output_size, unsigned char* shmem, unsigned char* tmp)`	`void symbol(const void* input, const unsigned long long input_size, unsigned char* tmp, void* output, unsigned long long* output_size, void* shmem)`

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: input chunk size in bytes, passed by pointer for Native and by value for Tile.
output_size: pointer written with the output size in bytes.
shmem: Native and Tile shared memory scratch buffer, at least NVCOMPDX_TRAIT_SHMEM_SIZE_GROUP bytes, aligned to NVCOMPDX_TRAIT_SHMEM_ALIGNMENT. The Native caller provides it; for Tile, the compiler supplies it through the last parameter.
tmp: Native 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.

For both function ABIs, input is read-only and output is a destination; the global scratch buffer is an inout operand.

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(const void* input, void* output, const unsigned long long* input_size, unsigned long long* output_size, unsigned char* shmem, unsigned char* 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(const void* input, void* output, const unsigned long long* input_size, unsigned long long* output_size, unsigned char* shmem, unsigned char* 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.

With COMMONDX_FUNCTION_ABI_TILE, both directions use: void symbol(const void* input, const unsigned long long input_size, unsigned char* tmp, void* output, unsigned long long* output_size, void* shmem) where input_size and the global-memory scratch pointer tmp are passed by value.

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 nvcompdxBlockWarpCompleteness_t#

Whether every warp in a block is complete.

Used as the second element of NVCOMPDX_OPERATOR_BLOCK_WARP and NVCOMPDX_TRAIT_BLOCK_WARP. The first element is the warp count.

Values:

enumerator NVCOMPDX_BLOCK_WARP_INCOMPLETE#

enumerator NVCOMPDX_BLOCK_WARP_COMPLETE#

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, except for MAX_COMP_CHUNK_SIZE queries.

enumerator NVCOMPDX_OPERATOR_MAX_UNCOMP_CHUNK_SIZE#: Operator data type: long long int. Expected content: maximum uncompressed chunk size in bytes. Required for compression code generation and compression size traits (TMP_SIZE_GROUP and MAX_COMP_CHUNK_SIZE). Optional for alignment, shared-memory and layout traits, and for decompression.

enumerator NVCOMPDX_OPERATOR_SM#: Operator data type: long long int. Expected content: 750 (Turing), 800 (Ampere), …, Operator definition: required for code generation and SYMBOL_NAME; optional for numeric traits.

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, except for MAX_COMP_CHUNK_SIZE queries.

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; block execution requires one layout operator for code generation, SHMEM_SIZE_GROUP, layout traits, and SYMBOL_NAME. Scratch size and alignments do not require it.

enumerator NVCOMPDX_OPERATOR_BLOCK_WARP#: Operator data type: long long int followed by nvcompdxBlockWarpCompleteness_t. Expected content: <num_warps, complete>. num_warps: number of warps (> 0). complete: nvcompdxBlockWarpCompleteness_t. Mutually exclusive with NVCOMPDX_OPERATOR_BLOCK_DIM. Operator definition: optional; block execution requires one layout operator for code generation, SHMEM_SIZE_GROUP, layout traits, and SYMBOL_NAME. Scratch size and alignments do not require it.

enumerator NVCOMPDX_OPERATOR_FUNCTION_ABI#: Operator data type: commondxFunctionAbi_t. Operator definition: optional. Defaults to COMMONDX_FUNCTION_ABI_NATIVE. Does not affect numeric traits.

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. Requires ALGORITHM, DIRECTION=COMPRESS, and MAX_UNCOMP_CHUNK_SIZE only. Execution, layout, SM, and function ABI do not affect this query.

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; complete is a nvcompdxBlockWarpCompleteness_t. 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: