NVIDIA nvCOMP中的兩個(gè)統(tǒng)一接口和使用

壓縮可以在各種用例中提高性能，例如 DL 工作負(fù)載、數(shù)據(jù)庫(kù)和通用 HPC 。在 GPU 上，壓縮可以加速協(xié)作工作流的 GPU 間通信。它可以通過(guò)在數(shù)據(jù)存儲(chǔ)到全局內(nèi)存之前壓縮數(shù)據(jù)來(lái)增加單個(gè) GPU 可以處理的數(shù)據(jù)集的大小。它還可以加速 CPU 和 GPU 之間的數(shù)據(jù)鏈路。

為了讓這些工作流中的任何一個(gè)都變得有用，壓縮和解壓縮必須快速，并且在給定數(shù)據(jù)集上以足夠高的壓縮比運(yùn)行，才能發(fā)揮作用。然而，不同算法的壓縮比和吞吐量因數(shù)據(jù)集而異。如果沒(méi)有大量關(guān)于算法和數(shù)據(jù)統(tǒng)計(jì)的專(zhuān)業(yè)知識(shí)，可能很難選擇最佳算法。

NVIDIA NVCOMP 庫(kù)使您可以在應(yīng)用程序中結(jié)合高性能 GPU 壓縮和解壓縮。該庫(kù)提供了一組統(tǒng)一的 API ，允許您快速交換壓縮格式，以在數(shù)據(jù)集上實(shí)現(xiàn)最佳性能，而對(duì)代碼的更改最少。

使用 nvCOMP ，您可以快速、輕松地使用不同的算法進(jìn)行實(shí)驗(yàn)，以找到最適合您的用例的算法。在最近的版本中，我們更新了 nvCOMP 以進(jìn)一步改進(jìn)和統(tǒng)一接口。在新發(fā)布的 2.2 版本中，我們提供了一個(gè)易于使用的、高級(jí)的 C ++ API 和一個(gè)通用的低級(jí)別批量 C API 。在本文中，我們將詳細(xì)介紹這兩個(gè)接口。你還可以學(xué)習(xí)如何有效地使用它們，以及何時(shí)應(yīng)該選擇其中一個(gè)。

高級(jí) API

高級(jí) API 更易于使用，并抽象了向 GPU 公開(kāi)并行性的工作。當(dāng)您必須將連續(xù)緩沖區(qū)壓縮為連續(xù)的壓縮緩沖區(qū)時(shí)，它最有用。例如，在通過(guò)網(wǎng)絡(luò)發(fā)送緩沖區(qū)或?qū)⑵浔４娴酱疟P(pán)之前壓縮緩沖區(qū)時(shí)，這種方法效果很好。

以下示例使用高通量GDeflate壓縮格式。GDeflate類(lèi)似于 deflate ，可以有效地映射到數(shù)據(jù)并行架構(gòu)，如 GPU 。如果您對(duì)使用的壓縮格式?jīng)]有限制，那么這是一個(gè)很好的起點(diǎn)。

高級(jí)接口是基于nvcompManagerBase類(lèi)層次結(jié)構(gòu)的C++ API。每個(gè)派生的Manager類(lèi)都在nvcomp/include中與其關(guān)聯(lián)的頭中聲明。例如，本文中使用的GDeflateManager在nvcomp/include/gdeflate.hpp中聲明。

首先，構(gòu)建所需的Manager類(lèi)。每個(gè)Manager構(gòu)造函數(shù)都有一組唯一的參數(shù)；然而，有一些觀點(diǎn)通常是一致的。所有子類(lèi)都允許使用指定的流 ID 構(gòu)造用于所有內(nèi)核和內(nèi)存?zhèn)鬏?。您還可以指定要使用的設(shè)備 ID 。如果不為這兩個(gè)參數(shù)指定值，則使用默認(rèn)的流和設(shè)備。

另一個(gè)常見(jiàn)的輸入是未壓縮的塊大小。這在壓縮過(guò)程中用于將緩沖區(qū)拆分為獨(dú)立的塊進(jìn)行處理。較大的塊大小通常會(huì)導(dǎo)致更高的壓縮比，但代價(jià)是暴露在 GPU 中的并行性會(huì)降低。一個(gè)好的起始數(shù)據(jù)塊大小是 64KB ，但是可以自由地使用這些值來(lái)探索數(shù)據(jù)集的相關(guān)權(quán)衡。

Manager類(lèi)也使用特定于格式的參數(shù)構(gòu)造。您可以查看nvcomp/include中的相關(guān)標(biāo)題，了解Manager類(lèi)構(gòu)造函數(shù)參數(shù)的描述，并了解如何為所選格式構(gòu)造Manager對(duì)象。

const size_t uncomp_chunk_size = 64 * 1024; cudaStream_t stream;
cudaStreamCreate(&stream));
const int gdeflate_algorithm = 0; // Use standard GDeflate
const int device_id = 0; // Use the default device GdeflateManager gdeflate_manager{chunk_size, gdeflate_algorithm, stream, device_id};

nvcompManager需要一個(gè)臨時(shí)的 scratch 工作區(qū)來(lái)進(jìn)行壓縮和解壓縮。根據(jù)特定的壓縮格式參數(shù)以及壓縮和解壓縮內(nèi)核的最大占用率，所需的暫存空間大小是固定的。如果對(duì)您的用例有意義，您可以在構(gòu)造后使用set_scratch_buffer為nvcompManager對(duì)象提供一個(gè)臨時(shí)緩沖區(qū)。

size_t scratch_buffer_size = gdeflate_manager.get_required_scratch_buffer_size();
uint8_t* scratch_buffer;
cudaMalloc(&scratch_buffer, scratch_buffer_size);
gdeflate_manager.set_scratch_buffer(scratch_buffer);

手動(dòng)設(shè)置暫存緩沖區(qū)可能有助于控制用于此分配的內(nèi)存分配方案。如果您同意默認(rèn)設(shè)置，我們建議跳過(guò)此步驟并啟用nvcompManager對(duì)象來(lái)處理分配。

此緩沖區(qū)可用于nvcompManager執(zhí)行的所有壓縮和解壓縮操作。如果nvcompManager對(duì)象分配了暫存緩沖區(qū)，則在銷(xiāo)毀該對(duì)象時(shí)會(huì)釋放該緩沖區(qū)。

壓縮

現(xiàn)在可以壓縮緩沖區(qū)了。首先，使用configure_compression API 配置壓縮。此異步操作返回CompressionConfig對(duì)象。

配置步驟只需要input-uncompressed緩沖區(qū)的大小。您必須分配一個(gè) GPU 可訪問(wèn)的內(nèi)存緩沖區(qū)，其大小至少為該大小，以用作壓縮例程的結(jié)果緩沖區(qū)。有了這些信息，可以執(zhí)行壓縮，如下面的代碼示例所示：

CompressionConfig comp_config = gdeflate_manager.configure_compression(input_buffer_len); uint8_t* comp_buffer;
cudaMallocAsync(&comp_buffer, comp_config.max_compressed_buffer_size, stream); gdeflate_manager.compress(uncomp_buffer, comp_buffer, comp_config);

您還可以在 GPU 上排隊(duì)進(jìn)行其他壓縮。

uint8_t* comp_buffer1, comp_buffer2;
CompressionConfig comp_config1 = gdeflate_manager.configure_compression(input_buffer_len1); cudaMallocAsync(&comp_buffer1, comp_config1.max_compressed_buffer_size, stream); gdeflate_manager.compress(uncomp_buffer1, comp_buffer1, comp_config1); CompressionConfig comp_config2 = gdeflate_manager.configure_compression(input_buffer_len2); cudaMallocAsync(&comp_buffer2, comp_config2.max_compressed_buffer_size, stream); gdeflate_manager.compress(uncomp_buffer2, comp_buffer2, comp_config2); cudaStreamSynchronize(stream);

減壓

高級(jí)接口壓縮產(chǎn)生的緩沖區(qū)在壓縮數(shù)據(jù)之前包含一個(gè)頭（圖 1 ）。此標(biāo)題包含有關(guān)緩沖區(qū)如何壓縮的信息，因此您可以從壓縮的緩沖區(qū)構(gòu)造nvcompManager對(duì)象，而不知道它是如何壓縮的。這使您可以在不知道緩沖區(qū)是如何壓縮的情況下對(duì)其進(jìn)行解壓縮。

圖 1 HLIF 壓縮數(shù)據(jù)格式

為此，請(qǐng)使用nvcompManagerFactory.hpp中聲明的create_manager API 。這個(gè)同步 API 將壓縮的緩沖區(qū)以及可選的流和設(shè)備 ID 作為輸入。

auto decomp_nvcomp_manager = create_manager(comp_buffer, stream);

如果您已經(jīng)掌握了有關(guān)緩沖區(qū)壓縮方式的信息，那么可以使用前面描述的配置構(gòu)造一個(gè)新的管理器。您還可以重用用于壓縮的同一nvcompManager對(duì)象來(lái)執(zhí)行解壓縮。這些方法的優(yōu)點(diǎn)是不需要同步流。

給定一個(gè)nvcompManager對(duì)象和一個(gè)壓縮的緩沖區(qū)，解壓的執(zhí)行與壓縮類(lèi)似，但有幾個(gè)細(xì)微的區(qū)別。首先，有兩種可能的方式來(lái)進(jìn)行解壓縮配置。如果壓縮使用CompressionConfig對(duì)象，則可以完全異步配置解壓縮。

DecompressionConfig decomp_config = gdeflate_manager->configure_decompression(comp_config);

該 API 的一個(gè)示例用例是大型神經(jīng)網(wǎng)絡(luò)的訓(xùn)練?？梢允褂玫纳窠?jīng)網(wǎng)絡(luò)或訓(xùn)練集的大小取決于 GPU 的內(nèi)存容量。使用壓縮，您可以有效地增加此容量，而無(wú)需將數(shù)據(jù)卸載到 CPU 或使用多個(gè) GPU 。

具體來(lái)說(shuō)，基于反向傳播的訓(xùn)練包括在向前傳球時(shí)計(jì)算激活圖，然后在向后傳球的計(jì)算中重用它們。這些激活映射比較大且相對(duì)稀疏，因此非常適合壓縮。使用gdeflate_manager壓縮地圖，并在內(nèi)存中保存網(wǎng)絡(luò)各層的壓縮緩沖區(qū)和CompressionConfig對(duì)象。這可以實(shí)現(xiàn)完全異步的反向傳播，包括解壓縮。

如果沒(méi)有使用的CompressionConfig對(duì)象，也可以使用壓縮緩沖區(qū)配置解壓縮。這是一個(gè)同步操作，必須從設(shè)備執(zhí)行cudaMemcpyAsync操作。所有同步都在nvcompManager構(gòu)造函數(shù)中指定的流上，并且不是設(shè)備范圍的。

DecompressionConfig decomp_config = gdeflate_manager->configure_decompression(comp_buffer);

與壓縮一樣，在同步流之前，您可以一次將多個(gè)解壓縮項(xiàng)目排隊(duì)。

uint8_t* res_decomp_buffer1, res_decomp_buffer2;
DecompressionConfig decomp_config1 = gdeflate_manager->configure_decompression(comp_config1);
DecompressionConfig decomp_config2 = gdeflate_manager->configure_decompression(comp_config2); cudaMallocAsync(&res_decomp_buffer1, decomp_config1.decomp_data_size, stream);
cudaMallocAsync(&res_decomp_buffer2, decomp_config2.decomp_data_size, stream); gdeflate_manager->decompress(res_decomp_buffer1, comp_buffer1, decomp_config1);
gdeflate_manager->decompress(res_decomp_buffer2, comp_buffer2, decomp_config2); cudaStreamSynchronize(stream));

最后，在高級(jí) API 中有兩種類(lèi)型的錯(cuò)誤檢查：std：：runtime_error異常和檢查nvcompStatus_t值。

如果任何 CUDA API 失敗，就會(huì)引發(fā)std：：runtime_error異常。您可以在應(yīng)用程序中捕獲這些錯(cuò)誤，也可以不處理它們，在這種情況下，您的應(yīng)用程序會(huì)失敗，并顯示一條描述錯(cuò)誤的消息。例如，如果您提供的輸出緩沖區(qū)大小不足或無(wú)法在 GPU 上訪問(wèn)，就會(huì)發(fā)生這種情況。

錯(cuò)誤檢查的第二種形式是檢查CompressionConfig或DecompressionConfig對(duì)象中的nvcompStatus_t值。此狀態(tài)在相關(guān)的內(nèi)核調(diào)用期間設(shè)置。損壞的輸入緩沖區(qū)和其他錯(cuò)誤會(huì)觸發(fā)它。

低級(jí) API

低級(jí) API 為更高級(jí)的工作流提供了 C API 。低級(jí) API 同時(shí)壓縮和解壓縮您提供的一批獨(dú)立塊。這取決于您對(duì)數(shù)據(jù)進(jìn)行分塊，并提供足夠數(shù)量的分塊來(lái)利用 GPU 的并行處理能力。

如果有許多獨(dú)立的、不連續(xù)的緩沖區(qū)，這是處理數(shù)據(jù)最有效的方法。低級(jí) API 避免了將生成的壓縮塊打包到單個(gè)連續(xù)的壓縮緩沖區(qū)的工作量。它還避免了與在高級(jí) API 中保存有關(guān)緩沖區(qū)如何壓縮的信息相關(guān)的壓縮比開(kāi)銷(xiāo)。

該工作流非常適合數(shù)據(jù)庫(kù)應(yīng)用程序，例如，在這些應(yīng)用程序中，往往需要壓縮或解壓縮許多獨(dú)立的列。這個(gè) API 用于 RAPIDS 和 NVIDIA Spark 實(shí)現(xiàn)。

壓縮

對(duì)于低級(jí) API 中的壓縮，必須分配一個(gè)臨時(shí)暫存緩沖區(qū)。臨時(shí)緩沖區(qū)與高級(jí) API 中描述的類(lèi)似。然而，緩沖區(qū)大小取決于輸入緩沖區(qū)的大小，因此必須重新定義它，并可能與每一組新的用戶(hù)輸入一起重新分配。

size_t temp_bytes;
nvcompBatchedGdeflateCompressGetTempSize(batch_size, chunk_size, nvcompBatchedGdeflateDefaultOpts, &temp_bytes); void* device_temp_ptr;
cudaMalloc(&device_temp_ptr, temp_bytes);

接下來(lái)，應(yīng)該計(jì)算批處理中壓縮塊的最大大小。這允許您分配一組結(jié)果緩沖區(qū)。在下面的示例中，batch_size是要處理的塊數(shù)。結(jié)果指針的設(shè)備數(shù)組在復(fù)制到設(shè)備之前在固定的主機(jī)內(nèi)存中構(gòu)造。

size_t max_out_bytes;
nvcompBatchedGdeflateCompressGetMaxOutputChunkSize(chunk_size, nvcompBatchedGdeflateDefaultOpts, &max_out_bytes); // Allocate output space on the device
void ** host_compressed_ptrs;
cudaMallocHost((void**)&host_compressed_ptrs, sizeof(size_t) * batch_size);
for(size_t ix_chunk = 0; ix_chunk < batch_size; ++ix_chunk) { cudaMalloc(&host_compressed_ptrs[ix_chunk], max_out_bytes);
} void** device_compressed_ptrs;
cudaMalloc(&device_compressed_ptrs, sizeof(size_t) * batch_size);
cudaMemcpy( device_compressed_ptrs, host_compressed_ptrs, sizeof(size_t) * batch_size,cudaMemcpyHostToDevice);

通過(guò)計(jì)算所有這些輸入，您現(xiàn)在可以異步進(jìn)行壓縮，如圖所示。

nvcompStatus_t comp_res = nvcompBatchedGdeflateCompressAsync( device_uncompressed_ptrs, device_uncompressed_bytes, chunk_size, batch_size, device_temp_ptr, temp_bytes, device_compressed_ptrs, device_compressed_bytes, nvcompBatchedGdeflateDefaultOpts,

減壓

要開(kāi)始解壓，請(qǐng)根據(jù)壓縮的緩沖區(qū)預(yù)計(jì)算解壓的大小。如果您已經(jīng)有此信息，請(qǐng)?zhí)^(guò)此步驟。

nvcompBatchedGdeflateGetDecompressSizeAsync( device_compressed_ptrs, device_compressed_bytes, device_uncompressed_bytes, batch_size, stream);

與壓縮類(lèi)似，您還必須計(jì)算所需的臨時(shí)大小，并分配臨時(shí)暫存緩沖區(qū)。

size_t decomp_temp_bytes;
nvcompBatchedGdeflateDecompressGetTempSize(batch_size, chunk_size, &decomp_temp_bytes);
void * device_decomp_temp;
cudaMalloc(&device_decomp_temp, decomp_temp_bytes);

最后，可以進(jìn)行異步解壓縮。

nvcompStatus_t decomp_res = nvcompBatchedGdeflateDecompressAsync( device_compressed_ptrs, device_compressed_bytes, device_uncompressed_bytes, device_actual_uncompressed_bytes, batch_size, device_decomp_temp, decomp_temp_bytes, device_uncompressed_ptrs, device_statuses, stream);

標(biāo)桿管理

nvCOMP 為低級(jí)和高級(jí)格式的每種格式提供了一組基準(zhǔn)。圖 2 比較了在幾個(gè)不同的數(shù)據(jù)集上使用大型連續(xù)緩沖區(qū)時(shí)高級(jí)和低級(jí)的性能。使用 A100 GPU 收集結(jié)果。

圖 2a 各種數(shù)據(jù)集的解壓縮吞吐量。

圖 2b 各種數(shù)據(jù)集的壓縮比。

圖 2c 各種數(shù)據(jù)集的壓縮吞吐量

從結(jié)果中可以看出，在使用大型連續(xù)緩沖區(qū)時(shí)，低級(jí)和高級(jí) API 之間的性能差異可以忽略不計(jì)。使用哪一個(gè)取決于您的用例。如果有許多小緩沖區(qū)，請(qǐng)使用低級(jí) API ，或者避免與高級(jí) API 相關(guān)的內(nèi)存占用。

圖 3 顯示了日志規(guī)模下不同緩沖區(qū)大小的性能。為了產(chǎn)生這些結(jié)果，圖 2 中顯示的 mortgage int 數(shù)據(jù)集被分成許多批batchSize，如圖所示。該文件超過(guò) 314 MB 。對(duì)于 1 MB 的批量大小，執(zhí)行 315 次壓縮和解壓縮操作。批量大小為 400 MB 時(shí)，執(zhí)行單個(gè)壓縮和解壓縮操作。

以這種方式批處理數(shù)據(jù)不會(huì)影響低級(jí)批處理 API 。

圖 3a ：在 314 MB 文件上運(yùn)行的各種批量大小的壓縮吞吐量。

在 314 MB 文件上操作的各種批量大小的解壓縮吞吐量。

正如所證明的，對(duì)于小批量，高級(jí)接口的性能會(huì)嚴(yán)重下降。這顯示了在壓縮或解壓縮許多較小的緩沖區(qū)時(shí)使用低級(jí)批處理 API 的實(shí)用性。低級(jí)批處理 API 可以使用更少、占用率更高的內(nèi)核來(lái)完成操作，而高級(jí) API 需要許多具有相關(guān)尾部效應(yīng)和占用率問(wèn)題的小型內(nèi)核啟動(dòng)。

我們?cè)趲?kù)中加入了基準(zhǔn)測(cè)試應(yīng)用程序，以便您可以嘗試不同的壓縮格式，并查看哪種格式對(duì)您的數(shù)據(jù)最有效。提供的基準(zhǔn)是benchmark_hlif和benchmark_《format》_chunked。有關(guān)更多信息，請(qǐng)參閱 nvCOMP README 。

總結(jié)

現(xiàn)在，您已經(jīng)了解了如何使用高級(jí) nvCOMP API 來(lái)輕松壓縮和解壓縮。您已經(jīng)了解了何時(shí)使用低級(jí) API 更好，以及如何使用它。

關(guān)于作者

Eric Schmidt 是 NVIDIA 的高級(jí)開(kāi)發(fā)技術(shù)工程師。 Eric 目前正在 GPU 上加速壓縮例程。在 2021 加入 NVIDIA 之前，埃里克在航天工業(yè)中花了 11 年的時(shí)間在應(yīng)用數(shù)學(xué)中開(kāi)發(fā)軟件和研究算法。

審核編輯：郭婷