硬件解压引擎工作原理
某机构Blackwell架构中的新型解压引擎（DE）是一个固定功能硬件模块，旨在加速Snappy、LZ4和基于Deflate的流解压。通过硬件处理解压，DE释放了宝贵的流式多处理器（SM）资源用于计算，而不是将周期消耗在数据移动上。
作为复制引擎的一部分集成，DE消除了顺序主机到设备复制后跟软件解压的需求。相反，压缩数据可以直接通过PCIe或C2C传输，并在传输过程中解压，从而减少主要的I/O瓶颈。
除了原始吞吐量之外，DE还实现了数据移动和计算的真正并发。多流工作负载可以并行发出解压操作与SM内核，保持GPU完全利用。实际上，这意味着数据密集型应用程序（如训练LLM、分析大规模基因组数据集或运行HPC模拟）可以跟上下一代Blackwell GPU的带宽，而不会因I/O而停滞。
nvCOMP的GPU加速压缩优势
某机构nvCOMP库提供GPU加速的压缩和解压例程。它支持各种标准格式，以及某机构为最佳GPU性能优化的格式。
对于标准格式，由于可用并行性有限，CPU和固定功能硬件通常比GPU具有架构优势。解压引擎是针对一系列工作负载的解决方案。以下部分将讨论如何利用nvCOMP使用DE。
如何使用DE和nvCOMP
开发者最好通过nvCOMP API利用DE。由于DE仅在选定的GPU上可用（截至目前，B200、B300、GB200和GB300），使用nvCOMP使开发者能够编写可移植的代码，随着DE覆盖范围的演变，跨GPU扩展和工作。当DE可用时，nvCOMP将使用它而无需更改用户代码。如果不可用，nvCOMP将回退到其基于SM的加速实现。
在支持DE的GPU上，需要做一些事情来确保此行为。nvCOMP通常允许任何设备可访问的输入和输出缓冲区类型。DE有更严格的要求。如果缓冲区不满足这些要求，nvCOMP将在SM上执行解压。参见表1以了解允许的分配类型及其预期用途。
表1. 允许的分配类型及其预期用途
| 分配类型 | 描述 | 可访问性 |
|---|---|---|
| cudaMalloc | 标准仅设备分配 | 设备 |
| cudaMallocFromPoolAsync | 易于使用的基于池的分配，具有更多功能 | 主机/设备 |
| cuMemCreate | 主机/设备分配的低级控制 | 主机/设备 |
cudaMalloc分配可以正常分配用于设备到设备解压。如果使用cudaMallocFromPoolAsync或cuMemCreate，主机到设备甚至主机到主机解压也是可能的，但必须注意正确设置分配器。
以下部分将提供如何使用这些不同分配器的工作示例。注意，在这两种情况下，标准使用这些API的唯一区别是添加了cudaMemPoolCreateUsageHwDecompress和CU_MEM_CREATE_USAGE_HW_DECOMPRESS标志。在两个示例中，这些分配都放置在第一个CPU NUMA节点上。
使用cudaMallocFromPoolAsync
以下代码示例显示如何创建具有cudaMemPoolCreateUsageHwDecompress标志的固定主机内存池，启用与DE兼容的分配。
cudaMemPoolProps props = {};
props.location.type = cudaMemLocationTypeHostNuma;
props.location.id = 0;
props.allocType = cudaMemAllocationTypePinned;
props.usage = cudaMemPoolCreateUsageHwDecompress;
cudaMemPool_t mem_pool;
CUDA_CHECK(cudaMemPoolCreate(&mem_pool, &props));
char* mem_pool_ptr;
CUDA_CHECK(cudaMallocFromPoolAsync(&mem_pool_ptr, 1024, mem_pool, stream));
使用cuMemCreate
此示例演示如何使用低级CUDA驱动程序API（cuMemCreate）分配具有CU_MEM_CREATE_USAGE_HW_DECOMPRESS标志的固定主机内存。它确保缓冲区与DE兼容。
CUdeviceptr mem_create_ptr;
CUmemGenericAllocationHandle allocHandle;
CUmemAllocationProp props = {};
props.location.type = CU_MEM_LOCATION_TYPE_HOST_NUMA;
props.location.id = 0;
props.type = CU_MEM_ALLOCATION_TYPE_PINNED;
props.allocFlags.usage = CU_MEM_CREATE_USAGE_HW_DECOMPRESS;
size_t granularity;
CU_CHECK(cuMemGetAllocationGranularity(&granularity, &props, CU_MEM_ALLOC_GRANULARITY_MINIMUM));
// 创建分配句柄
CU_CHECK(cuMemCreate(&allocHandle, granularity, &props, 0));
// 保留虚拟地址空间
CU_CHECK(cuMemAddressReserve(&mem_create_ptr, granularity, 0, 0, 0));
// 将物理内存映射到虚拟地址
CU_CHECK(cuMemMap(mem_create_ptr, granularity, 0, allocHandle, 0));
缓冲区批处理最佳实践
为了获得最佳性能，用于解压的缓冲区批次（输入/输出/大小）应是指向相同分配中的偏移的指针。如果提供来自不同分配的缓冲区批次，主机驱动程序启动开销可能很大。
uint8_t* d_decompressed_buffer;
CUDA_CHECK(cudaMalloc(&d_decompressed_buffer, total_decompressed_size));
// 为设备解压指针创建固定主机数组
uint8_t** h_d_decompressed_ptrs;
CUDA_CHECK(cudaHostAlloc(&h_d_decompressed_ptrs, actual_num_buffers * sizeof(uint8_t*), cudaHostAllocDefault));
// 使用偏移填充设备解压的固定主机指针数组
size_t decompressed_offset = 0;
for (int i = 0; i < actual_num_buffers; ++i) {
h_d_decompressed_ptrs[i] = d_decompressed_buffer + decompressed_offset;
decompressed_offset += input_sizes[i];
}
注意，由于与DE相关的同步要求，nvCOMP的异步API将与调用流同步。通常，nvCOMP仍将在API完成之前返回，因此如果解压到主机，在使用解压结果之前需要再次同步调用流。对于设备端访问，解压结果在正常流排序中可用。
在B200上，如果任何缓冲区大于4 MB，nvCOMP将回退到基于SM的实现。此限制将来可能会更改，并可以通过以下代码查询：
int max_supported_size = 0;
res = CudaDriver::cuDeviceGetAttribute(&max_supported_size,
CU_DEVICE_ATTRIBUTE_MEM_DECOMPRESS_MAXIMUM_LENGTH,
device_id);
SM性能与DE比较
DE提供更快的解压，同时释放SM用于其他工作。DE提供数十个执行单元，而SM上有数千个warp可用。每个DE执行单元在执行解压时比SM快得多，但在某些工作负载中，当完全饱和SM资源时，SM速度将接近DE。SM或DE都可以使用主机固定数据作为输入执行，启用零拷贝解压。
下图将显示在Silesia基准测试中LZ4、Deflate和Snappy算法的SM与DE性能。注意，Snappy在nvCOMP 5.0中进行了新优化，并且Deflate和LZ4有进一步的软件优化机会。
性能测量使用"小"和"大"数据集对64 KiB和512 KiB块大小进行。大数据集是完整的Silesia数据集，而小数据集是Silesia.tar的前约50 MB（可在此处获得）。
图1. 在六个示例中比较流式多处理器与解压引擎的性能。
开始使用
Blackwell中的解压引擎使处理数据繁重工作负载中最大挑战之一变得更加容易：快速、高效的解压。通过将这项工作转移到专用硬件，应用程序不仅看到更快的结果，而且释放GPU计算用于其他任务。
通过nvCOMP自动处理集成，开发者可以在不更改代码的情况下利用这些改进，从而实现更流畅的管道和更好的性能。
要开始使用这些新功能，请探索以下资源：
- 了解有关nvCOMP和硬件解压引擎的更多信息，以及如何轻松将它们集成到现有工作流程中。
- 了解有关最新nvCOMP API示例和基准测试的更多信息。
- 下载最新版本的nvCOMP以开始使用。
