在cuDNN中簡(jiǎn)化Tensor Ops
在Tesla V100 GPU中引入神經(jīng)網(wǎng)絡(luò)模型以來，神經(jīng)網(wǎng)絡(luò)模型已迅速利用NVIDIA Tensor Cores進(jìn)行深度學(xué)習(xí)。例如，基于Tensor Core的解決方案宣布了ResNet50訓(xùn)練的性能記錄。
NVIDIA的cuDNN庫 使CUDA程序員能夠優(yōu)化循環(huán)神經(jīng)網(wǎng)絡(luò)和卷積神經(jīng)網(wǎng)絡(luò)，以實(shí)現(xiàn)GPU加速。概述了cuDNN用戶使用Tensor Core 進(jìn)行卷積的簡(jiǎn)便方法，并附有說明和示例代碼。該文章為cuDNN應(yīng)用提供了一些簡(jiǎn)單的規(guī)則：FP16數(shù)據(jù)規(guī)則，張量維數(shù)規(guī)則，ALGO_1的使用等。
cuDNN版本解除了大多數(shù)限制。cuDNN 7.2版本取消了FP16數(shù)據(jù)約束，而cuDNN 7.3刪除了張量尺寸約束（對(duì)于打包的NCHW張量數(shù)據(jù)），直接進(jìn)行改進(jìn)。
將FP32數(shù)據(jù)用于Tensor Ops
關(guān)于在CUDA中使用Tensor Core的帖子討論了將FP16輸入用于張量操作，如圖1所示。雖然張量操作仍然使用FP16數(shù)據(jù)，但卷積cuDNN API允許用戶選擇將FP32輸入數(shù)據(jù)轉(zhuǎn)換為FP16。如果需要，卷積的輸出數(shù)據(jù)也將轉(zhuǎn)換為FP32。

圖1. FP32數(shù)據(jù)現(xiàn)在可以用作輸入
該CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION枚舉值，在cuDNN 7.2，使cuDNN應(yīng)用程序員選擇轉(zhuǎn)換FP32數(shù)據(jù)運(yùn)算使用。該枚舉值與枚舉值一樣傳遞給cudnnSetConvolutionMathType()調(diào)用CUDNN_TENSOR_OP_MATH。此代碼段顯示了如何執(zhí)行此操作：
//設(shè)置數(shù)學(xué)類型以允許cuDNN使用Tensor Core：
checkCudnnErr（cudnnSetConvolutionMathType（cudnnConvDesc，CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION））;
將在后面的部分中看到使用代碼片段的上下文。
FP32數(shù)據(jù)也用于RNN
現(xiàn)在還為RNN啟用了類似的FP32數(shù)據(jù)轉(zhuǎn)換。只需將CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION枚舉值傳遞給cudnnSetRNNMatrixMathType()調(diào)用，即可將FP32數(shù)據(jù)轉(zhuǎn)換為在RNN中使用。如下使用：
//設(shè)置數(shù)學(xué)類型以允許cuDNN使用Tensor Core：
checkCudnnErr（cudnnSetRNNMatrixMathType（cudnnRnnDesc，CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION））;
消除了NCHW張量尺寸約束
早期版本的cuDNN要求所有張量的通道維數(shù)必須為8的倍數(shù)。cuDNN可以根據(jù)需要自動(dòng)填充張量。
在CUDNN_TENSOR_OP_MATH和CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION情況下，對(duì)于填充的NCHW數(shù)據(jù)，此填充都是自動(dòng)的。發(fā)生填充時(shí)，性能損失可忽略不計(jì)。
//設(shè)置NCHW張量尺寸，不必設(shè)置為8的倍數(shù)（此處僅顯示輸入張量）：
int dimA [] = {1，7，32，32};
int strideA [] = {7168，1024，32，1};
下一節(jié)中的示例代碼演示了如何使用。
樣例代碼
將張量運(yùn)算用于FP32數(shù)據(jù)和任何通道尺寸的邏輯類似于為cuDNN的早期版本編寫時(shí)使用的邏輯。只有維度和數(shù)據(jù)類型發(fā)生了變化（以及使用CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION)：
//創(chuàng)建一個(gè)cuDNN句柄：
checkCudnnErr（cudnnCreate（＆handle_））;

//創(chuàng)建張量描述符：
checkCudnnErr（cudnnCreateTensorDescriptor（＆cudnnIdesc））;
checkCudnnErr（cudnnCreateFilterDescriptor（＆cudnnFdesc））;
checkCudnnErr（cudnnCreateTensorDescriptor（＆cudnnOdesc））;
checkCudnnErr（cudnnCreateConvolutionDescriptor（＆cudnnConvDesc））;

//設(shè)置NCHW張量尺寸，不必設(shè)置為8的倍數(shù)（此處僅顯示輸入張量）：
int dimA [] = {1，7，32，32};
int strideA [] = {7168，1024，32，1};

checkCudnnErr（cudnnSetTensorNdDescriptor（cudnnIdesc，CUDNN_DATA_FLOAT，
convDim + 2，dimA，strideA））;

//分配和初始化張量（同樣，僅顯示輸入張量）：
checkCudaErr（cudaMalloc（（void *）＆（devPtrI），（insize） sizeof（devPtrI [0]））））;;
hostI =（T_ELEM *）calloc（insize，sizeof（hostI [0]））;

initImage（hostI，insize）;

checkCudaErr（cudaMemcpy（devPtrI，hostI，sizeof（hostI [0]）* insize，cudaMemcpyHostToDevice））;

//設(shè)置計(jì)算數(shù)據(jù)類型（以下為CUDNN_DATA_FLOAT）：
checkCudnnErr（cudnnSetConvolutionNdDescriptor（cudnnConvDesc，convDim，padA，convstrideA，dilationA，CUDNN_CONVOLUTION，CUDNN_DATA_FLOAT））;;

//設(shè)置數(shù)學(xué)類型以允許cuDNN使用Tensor Core：
checkCudnnErr（cudnnSetConvolutionMathType（cudnnConvDesc，CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION））;

//選擇支持的算法：
cudnnConvolutionFwdAlgo_t算法= CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM;

//分配的工作空間：
checkCudnnErr（cudnnGetConvolutionForwardWorkspaceSize（handle_，cudnnIdesc，
cudnnFdesc，cudnnConvDesc，
cudnnOdesc，algo，＆workSpaceSize））;

如果（workSpaceSize> 0）{
cudaMalloc（＆workSpace，workSpaceSize）;
}

//調(diào)用卷積：
checkCudnnErr（cudnnConvolutionForward（handle_，（void *）（＆alpha），cudnnIdesc，devPtrI，
cudnnFdesc，devPtrF，cudnnConvDesc，algo，
workSpace，workSpaceSize，（void *）（＆beta），
cudnnOdesc，devPtrO））;
FP32性能

圖2顯示了將Tensor Core用于FP32張量數(shù)據(jù)時(shí)卷積的比較性能。該圖表將V100張量運(yùn)算與V100 FMA運(yùn)算進(jìn)行了比較，因此，其增益并不像以前的將V100性能與P100 FMA進(jìn)行比較的圖表那樣明顯。但是，與使用FMA ops相比，與FP32輸入一起使用的Tensor ops仍然代表了可觀的收益。
圖2.具有Tensor Core的Tesla V100（Volta）與Tesla V100（Volta）的卷積性能比較。比較是在每個(gè)神經(jīng)網(wǎng)絡(luò)的卷積層運(yùn)行時(shí)間的幾何方法之間進(jìn)行的。兩種情況都使用FP32輸入/輸出數(shù)據(jù)和FP32計(jì)算。一種使用Tensor Core，另一種使用FP32融合乘加（FMA）。
剩余約束
盡管解除了在cuDNN中使用張量運(yùn)算的主要限制，但仍然存在一些次要限制。一個(gè)限制是使用ALGO_1（IMPLICIT_PRECOMP_GEMM用于轉(zhuǎn)發(fā)）。cuDNN中還沒有其他卷積算法使用張量運(yùn)算。
另一個(gè)較小的限制是卷積濾波器的大小，特別是空間尺寸（r和s）。但是，用于卷積的FFT算法非常適合于濾波器尺寸較大的用例。只需在超出張量運(yùn)算濾波器限制以達(dá)到最佳性能之前就將卷積切換為使用FFT算法即可。

總結(jié)

以上是生活随笔為你收集整理的在cuDNN中简化Tensor Ops的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問題。

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