Device API#

注意

免责声明：Device API 是实验性的，可能会在 nvCOMP 的后续版本中更改。

Operators#

namespace device#

Enums

enum nvcomp_direction#

压缩或解压缩的选择。

Values

enumerator compress#
enumerator decompress#
enum nvcomp_algo#

要选择的压缩算法。

Values

enumerator ans#
enumerator zstd#
enumerator bitcomp#
enumerator lz4#
enumerator deflate#
enumerator gdeflate#
enum nvcomp_datatype#

压缩算法将解释输入数据的方式。

Values

enumerator uint8#

要解释为连续字节的数据。如果输入数据类型未包含在以下选项中，则应选择 uint8。

enumerator float16#

要解释为连续 IEEE 半精度浮点数的数据。要求每个块的输入字节总数可被 2 整除。

enumerator bfloat16#

要解释为连续 bfloat16 值的数据。要求每个块的输入字节总数可被 2 整除。

enum nvcomp_grouptype#

要处理的线程组类型。

Values

enumerator warp#

提供给 API 的组，预计为单 warp 大小。

namespace detail#

Enums

enum nvcomp_operator#

Values

enumerator direction#

选择压缩/解压缩。

enumerator algo#

选择压缩算法。

enumerator grouptype#

选择要处理的线程组类型。

enumerator datatype#

输入数据的格式。

enumerator max_uncomp_chunk_size#

最大未压缩块大小。（仅用于压缩）。

User API#

namespace device
namespace detail
template<class ...Operators>
class nvcomp_device_execution#
#include <user_api.hpp>

为了使用设备端 API，需要构造 `nvcomp_device_execution` 类型的对象。为此，首先使用以下 operators 声明用户定义的类型

`Direction()` `Algo()` `Grouptype()` `Datatype()` `MaxUncompChunkSize()`

请参阅 `operators.hpp` 中这些 operators 的描述

自定义类型是使用可用 operators 之间的 `+` 创建的。然后，生成的类型用于初始化 `nvcomp_device_execution` 类型的对象，并且可以从该对象进行设备 API 调用。

使用 ANS 压缩的示例： ``` 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_type = decltype(ans_base_type() + Direction<nvcomp_direction::compress>() + MaxUncompChunkSize<max_uncomp_chunk_size>());

auto compressor = ans_compressor_type(); size_t shmem_size_warp = compressor.shmem_size_group(); size_t tmp_size_warp = compressor.tmp_size_group(); constexpr size_t scs = compressor.shmem_size_block(num_warps_per_block); constexpr size_t shmem_alignment = compressor.shmem_alignment(); __shared__ uint8_t __align__(shmem_alignment) shared_buffer[scs];

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], uncomp_chunk_sizes[wid], comp_chunk_sizes + wid, shared_buffer + shmem_size_warp*wid_in_block, tmp_buf + tmp_size_warp*wid, warp); ```

Public Functions

template<typename CG> inline void __device__ compress (const void *input_chunk, void *output_chunk, const size_t input_size, size_t *output_size, uint8_t *shared_mem_buf, uint8_t *tmp_buf, CG &group)

压缩连续的数据缓冲区。

Parameters:

  • input_chunk – [in] 要压缩的块。应与 16 字节边界对齐。

  • output_chunk – [in] 生成的压缩块。应与 8 字节边界对齐。

  • input_size – [in] 要压缩的块的大小（以字节为单位）

  • output_size – [in] 生成的压缩块的大小（以字节为单位）

  • shared_mem_buf – [in] API 内部使用的共享内存缓冲区

  • tmp_buf – [in] API 内部使用的全局暂存缓冲区

  • group – [in] 压缩输入的协同组

template<typename CG> inline void __device__ decompress (const void *input_chunk, void *output_chunk, uint8_t *shared_mem_buf, uint8_t *tmp_buf, CG &group)

解压缩连续的数据缓冲区。

Parameters:

  • input_chunk – [in] 要解压缩的块。应与 8 字节边界对齐。

  • output_chunk – [in] 生成的解压缩块

  • shared_mem_buf – [in] API 内部使用的共享内存缓冲区

  • tmp_buf – [in] API 内部使用的全局暂存缓冲区

  • group – [in] 解压缩输入的协同组

template<typename CG> inline void __device__ decompress (const void *input_chunk, void *output_chunk, const size_t comp_chunk_size, size_t *const decomp_chunk_size, uint8_t *shared_mem_buf, uint8_t *tmp_buf, CG &group)

解压缩连续的数据缓冲区。

Parameters:

  • input_chunk – [in] 要解压缩的块。应与 8 字节边界对齐。

  • output_chunk – [in] 生成的解压缩块

  • comp_chunk_size – [in] 压缩块的大小

  • decomp_chunk_size – [in] 生成的解压缩块的大小

  • shared_mem_buf – [in] API 内部使用的共享内存缓冲区

  • tmp_buf – [in] API 内部使用的全局暂存缓冲区

  • group – [in] 解压缩输入的协同组

Public Static Functions

static inline __device__ __host__ constexpr size_t shmem_size_block (size_t num_warps_per_block)

返回给定 CTA 所需的共享内存量。

Parameters:

num_warps_per_block – [in] 每个块的 warp 数量

Returns:

共享内存量（以字节为单位）

static inline __device__ __host__ constexpr size_t shmem_size_group ()

返回每个协同组所需的共享内存量。与 `shmem_size_block` 不同，因为可能存在多个 API 调用，每个调用都有不同的协同组，它们都属于同一个 CTA。

Returns:

共享内存量（以字节为单位）

static inline __device__ __host__ constexpr size_t shmem_alignment ()

返回 CG 的共享内存分配所需的对齐方式。

Returns:

共享内存对齐大小

static inline constexpr size_t max_comp_chunk_size()#

返回最大压缩块大小。

Returns:

最大压缩块大小（以字节为单位）

static inline constexpr size_t tmp_size_total(size_t num_warps)#

返回整个内核所需的暂存空间大小。

Parameters:

num_warps – [in] 所有内核块中的 warp 总数。

Returns:

内存暂存空间大小（以字节为单位）

static inline __device__ __host__ constexpr size_t tmp_size_group ()

返回每个协同组所需的全局内存暂存空间。与 `tmp_size_total` 不同，因为每个内核可能存在多个 API 调用，每个调用都需要全局暂存内存总量的一部分。

Returns:

全局内存暂存空间大小（以字节为单位）