cuSPARSELt開發(fā)NVIDIA Ampere結(jié)構(gòu)化稀疏性
深度神經(jīng)網(wǎng)絡(luò)在各種領(lǐng)域（例如計(jì)算機(jī)視覺，語音識別和自然語言處理）中均具有出色的性能。處理這些神經(jīng)網(wǎng)絡(luò)所需的計(jì)算能力正在迅速提高，因此有效的模型和計(jì)算至關(guān)重要。神經(jīng)網(wǎng)絡(luò)剪枝（刪除不必要的模型參數(shù)以生成稀疏網(wǎng)絡(luò)）是一種在保持準(zhǔn)確性的同時(shí)降低模型復(fù)雜性的有用方法。
為了利用細(xì)粒度的網(wǎng)絡(luò)剪枝，NVIDIA Ampere GPU架構(gòu)引入了細(xì)粒度的結(jié)構(gòu)稀疏性的概念。在NVIDIA A100 GPU上，結(jié)構(gòu)顯示為2：4模式：每四個元素中至少有兩個必須為零。通過使用新的NVIDIA Sparse Tensor Core跳過零值的計(jì)算，這可以將一個矩陣乘法（也稱為GEMM）操作數(shù)的數(shù)據(jù)占用空間和帶寬減少2倍，并使吞吐量翻倍。
cuSPARSELt：用于稀疏矩陣-密集矩陣乘法的高性能CUDA庫
為了簡化NVIDIA Ampere架構(gòu)稀疏功能的使用，NVIDIA引入了cuSPARSELt ，這是一種高性能CUDA庫，專用于常規(guī)矩陣操作，其中至少一個操作數(shù)是稀疏矩陣。cuSPARSELt庫可以使用NVIDIA第三代Tensor Core稀疏矩陣乘累加（SpMMA）操作，而無需進(jìn)行底層編程。該庫還提供用于剪枝和壓縮矩陣的輔助函數(shù)。
cuSPARSELt的主要功能包括：
· NVIDIA Sparse Tensor Core支持
· 混合精度支持：
o FP16輸入/輸出，FP32張量核心累積
o BFLOAT16輸入/輸出，FP32張量核心累積
o INT8輸入/輸出，INT32張量核心累積
· Row-major and column-major memory layouts的內(nèi)存布局
· 矩陣剪枝和壓縮實(shí)用程序
· 自動調(diào)整功能
NVIDIA Sparse Tensor Core support
Mixed-precision support:
FP16 inputs/output, FP32 Tensor Core accumulation
BFLOAT16 inputs/output, FP32 Tensor Core accumulation
INT8 inputs/output, INT32 Tensor Core accumulation
Row-major and column-major memory layouts
Matrix pruning and compression utilities
Auto-tuning functionality
定制工作流程
cuSPARSELt庫遵循等效方法，并采用與cuBLASLt和cuTENSOR類似的概念。庫編程模型要求以某種方式組織計(jì)算，以使相同的設(shè)置可以重復(fù)用于不同的輸入。
該模型尤其依賴于以下高層階段：
· 問題定義：指定矩陣形狀，數(shù)據(jù)類型，操作等。
· 用戶偏好和約束：提供算法選擇或限制可行實(shí)現(xiàn)（候選）的搜索空間。
· 計(jì)劃：收集執(zhí)行的描述符，并在需要時(shí)“找到”最佳實(shí)施。
· 執(zhí)行：執(zhí)行實(shí)際計(jì)算。
通用工作流程包括以下步驟：

1. 初始化庫句柄：cusparseLtInit。
2. 指定輸入/輸出矩陣特征：cusparseLtDenseDescriptorInit， cusparseLtStructuredDescriptorInit。
3. 初始化矩陣乘法描述符和它的屬性（例如操作，計(jì)算類型等）： cusparseLtMatmulDescriptorInit。
4. 初始化算法選擇描述符：cusparseLtMatmulAlgSelectionInit。
5. 初始化矩陣乘法計(jì)劃：cusparseLtMatmulPlanInit。
6. 剪枝A矩陣：cusparseLtSpMMAPrune。如果用戶提供已經(jīng)滿足2：4結(jié)構(gòu)化稀疏性約束的矩陣，例如由ASP庫生成的權(quán)重矩陣，則不需要此步驟。
7. 壓縮剪枝后的矩陣：cusparseLtSpMMACompress。
8. 執(zhí)行矩陣乘法：cusparseLtMatmul。可以使用不同的輸入多次重復(fù)此步驟。
9. 銷毀矩陣乘法計(jì)劃和庫句柄：cusparseLtMatmulPlanDestroy，cusparseLtDestroy。
   稀疏的GEMM性能
   與密集矩陣乘法一樣，稀疏矩陣乘法的性能隨GEMM尺寸，布局和數(shù)據(jù)類型而變化。這是當(dāng)前軟件與稀疏GEMM相對性能的快照。
   下表顯示了cuSPARSELt和cuBLAS在以下操作中的性能：
   D = alpha * op（A）* op（B）+ beta * C
   在該操作中，A，B，和 D=C分別是尺寸的密集矩陣MXK，KXN，和M×N個。矩陣的布局A和B與?為列主順序（OP是非轉(zhuǎn)置）和?為行優(yōu)先順序（OP調(diào)換）。
   為了展示使用cuSPARSELt可以針對實(shí)際工作負(fù)載實(shí)現(xiàn)的性能，下表顯示了帶有主要列TN FP16內(nèi)核的剪枝后的BERT-Large模型（seqlen = 128，BS = 128）使用的一些常見GEMM大小。通常，工作量越大，稀疏性可以提供的幫助越多。
   

表1. BERT-Large模型和不同層的cuSPARSELt性能。
結(jié)構(gòu)化稀疏矩陣-矩陣乘法代碼示例
已經(jīng)看到了可用的性能，下面是一個示例，該示例使用NVIDIA A100或GA100 GPU中的稀疏Tensor內(nèi)核在cuSPARSELt庫中執(zhí)行具有結(jié)構(gòu)稀疏性的矩陣乘法。有關(guān)更多信息，請參見NVIDIA / CUDALibrarySamples / tree / master / cuSPARSELt / spmma GitHub存儲庫。
首先，包括cuSPARSELt標(biāo)頭，設(shè)置一些設(shè)備指針和數(shù)據(jù)結(jié)構(gòu)，并初始化cuSPARSELt句柄。
#include <cusparseLt.h> // cusparseLt header // Device pointers and coefficient definitions float alpha = 1.0f; float beta = 0.0f; __half* dA = … __half* dB = … __half* dC = … // cusparseLt data structures and handle initialization cusparseLtHandle_t handle; cusparseLtMatDescriptor_t matA, matB, matC; cusparseLtMatmulDescriptor_t matmul; cusparseLtMatmulAlgSelection_t alg_sel; cusparseLtMatmulPlan_t plan; cudaStream_t stream = nullptr; cusparseLtInit(&handle);
接下來，初始化結(jié)構(gòu)化的稀疏輸入矩陣（matrix A），密集輸入矩陣（matrix B）和密集輸出矩陣（matrix C）描述符。
cusparseLtStructuredDescriptorInit(&handle, &matA, num_A_rows, num_A_cols, lda, alignment, type, order, CUSPARSELT_SPARSITY_50_PERCENT); cusparseLtDenseDescriptorInit(&handle, &matB, num_B_rows, num_B_cols, ldb, alignment, type, order); cusparseLtDenseDescriptorInit(&handle, &matC, num_C_rows, num_C_cols, ldc, alignment, type, order);
準(zhǔn)備好描述符后，可以準(zhǔn)備矩陣乘法運(yùn)算的描述符，選擇用于執(zhí)行matmul運(yùn)算的算法，并初始化matmul計(jì)劃。
cusparseLtMatmulDescriptorInit(&handle, &matmul, opA, opB, &matA, &matB, &matC, &matC, compute_type); cusparseLtMatmulAlgSelectionInit(&handle, &alg_sel, &matmul, CUSPARSELT_MATMUL_ALG_DEFAULT); int alg = 0; // set algorithm ID cusparseLtMatmulAlgSetAttribute(&handle, &alg_sel, CUSPARSELT_MATMUL_ALG_CONFIG_ID, &alg, sizeof(alg)); size_t workspace_size, compressed_size; cusparseLtMatmulGetWorkspace(&handle, &alg_sel, &workspace_size); cusparseLtMatmulPlanInit(&handle, &plan, &matmul, &alg_sel, workspace_size);
如果稀疏矩陣尚未被其他進(jìn)程剪枝，則可以在此時(shí)進(jìn)行。不要忘記檢查稀疏模式的有效性，以確保可以使用稀疏張量核心來加速它。
cusparseLtSpMMAPrune(&handle, &matmul, dA, dA, CUSPARSELT_PRUNE_SPMMA_TILE, stream); // checking the correctness int is_valid = 0; cusparseLtSpMMAPruneCheck(&handle, &matmul, dA, &is_valid, stream); if (is_valid != 0) { std::printf("!!! The matrix does not conform to the SpMMA sparsity pattern. " “cusparseLtMatmul does not provide correct results\n”); return EXIT_FAILURE; }
現(xiàn)在已將矩陣A剪枝為2：4稀疏度，可以將其壓縮到大約原始大小的一半。與實(shí)際的矩陣乘法（小于5％）相比，該步驟的執(zhí)行時(shí)間可以忽略不計(jì)。
cusparseLtSpMMACompressedSize(&handle, &plan, &compressed_size); cudaMalloc((void**) &dA_compressed, compressed_size); cusparseLtSpMMACompress(&handle, &plan, dA, dA_compressed, stream);
設(shè)置完成后，執(zhí)行matmul操作。cusparseLtMatmul使用不同的B矩陣可以多次重復(fù)調(diào)用 。只需設(shè)置一次稀疏矩陣。對于A矩陣值更改的用例，cusparseLtSpMMACompress必須再次調(diào)用該例程以設(shè)置稀疏矩陣的數(shù)據(jù)結(jié)構(gòu)。
void* d_workspace = nullptr; int num_streams = 0; cudaStream_t* streams = nullptr; cusparseLtMatmul(&handle, &plan, &alpha, dA_compressed, dB, &beta, dC, dD, d_workspace, streams, num_streams) )
最后，通過破壞matmul計(jì)劃和cuSPARSELt句柄來清理已使用的內(nèi)存。
cusparseLtMatmulPlanDestroy（＆plan）;
cusparseLtDestroy（＆handle）;

cuSPARSELt
通過cuSPARSELt庫，可以輕松利用NVIDIA Sparse Tensor Core運(yùn)算，從而在不降低網(wǎng)絡(luò)準(zhǔn)確性的情況下，顯著提高了用于深度學(xué)習(xí)應(yīng)用程序的矩陣矩陣乘法的性能。該庫還提供了用于矩陣壓縮，剪枝和性能自動調(diào)整的實(shí)用程序。簡而言之，與普通的密集數(shù)學(xué)方法相比，cuSPARSELt減少了計(jì)算，功耗，執(zhí)行時(shí)間和內(nèi)存存儲。

總結(jié)

以上是生活随笔為你收集整理的cuSPARSELt开发NVIDIA Ampere结构化稀疏性的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。