CUDA 8的混合精度程式設計

CUDA 8的混合精度程式設計
Volta和Turing GPU包含 Tensor Cores，可加速某些型別的FP16矩陣數學運算。這樣可以在流行的AI框架內更快，更輕鬆地進行混合精度計算。要使用Tensor Core，需要使用 CUDA 9 或更高版本。NVIDIA還 為TensorFlow，PyTorch和MXNet新增了 自動混合精度功能。
流行AI框架的張量核心優化示例 。
在軟體開發的實踐中，程式設計師通常會及早學習到使用正確的工具完成工作的重要性。當涉及數值計算時，這一點尤其重要，因為在數值計算中，精度，準確性和效能之間的折衷使得必須選擇最佳的資料表示形式。隨著Pascal GPU架構和CUDA 8的推出，NVIDIA正在利用新的16位浮點和8/16位整數計算功能擴充套件可用於混合精度計算的工具集。
“隨著架構和軟體的不斷變化以及GPU等加速器的破壞性影響，隨著不同精度的相對成本和易用性的發展，將看到越來越多的混合精度演算法得到開發和使用。” —曼徹斯特大學理查森應用數學教授尼克·海姆（Nick Higham）
許多技術和HPC應用程式都要求使用32位（單浮點或FP32）或64位（雙浮點或FP64）浮點進行高精度計算，甚至GPU加速的應用都依賴於更高的精度（128） -或256位浮點數！）。在許多應用中，低精度的算術就足夠了。例如，在快速發展的深度學習領域中的研究人員發現，由於深度神經網路體系結構用於訓練反向傳播演算法，因此對誤差具有自然的抵抗力，並且有人認為16位浮點數（半精度或FP16）足以訓練神經網路。
與更高精度的FP32或FP64相比，儲存FP16（半精度）資料可減少神經網路的記憶體使用量，從而可以訓練和部署更大的網路，並且FP16資料傳輸比FP32或FP64傳輸花費的時間更少。此外，對於許多網路而言，可以使用8位整數計算執行深度學習推理，而不會對準確性產生重大影響。
除了深度學習之外，使用來自相機或其它實際感測器的資料的應用程式通常不需要高精度的浮點計算，因為感測器會生成低精度或低動態範圍的資料。射電望遠鏡處理的資料就是一個很好的例子。正如將在本文後面看到的那樣，通過使用8位整數計算，可以大大加速用於處理射電望遠鏡資料的互相關演算法。
在計算方法中不同數值精度的組合使用稱為混合精度。NVIDIA Pascal架構通過新增將多個操作打包到32位資料路徑中的向量指令，旨在為可以利用較低精度計算的應用程式提供更高的效能。具體來說，這些指令對16位浮點資料（“ half”或FP16）以及8位和16位整數資料（INT8和INT16）進行操作。
由GP100 GPU驅動的新型NVIDIA Tesla P100可以以FP32兩倍的吞吐量執行FP16算術運算。GP102（Tesla P40和NVIDIA Titan X），GP104（Tesla P4）和GP106 GPU均支援可在2和4元素8位向量上執行整數點積的指令，並累加為32位整數。這些指令對於實現高效的深度學習推理以及射電天文學等其它應用程式非常有價值。
在本文中，將提供有關半精度浮點的一些詳細資訊，並提供有關使用FP16和INT8向量計算的Pascal GPU可獲得的效能的詳細資訊。還將討論各種CUDA平臺庫和API提供的混合精度計算功能。
浮點精度（或16）
正如每位電腦科學家都應該知道的那樣，浮點數提供了一種表示形式，可以在範圍和精度之間進行權衡的情況下，在計算機上近似實數。浮點數將實際值近似為一組有效數字（稱為尾數或有效位數），然後以固定基數（今天大多數計算機上使用的IEEE標準浮點數的基數2）進行縮放。
常見的浮點格式包括32位（稱為“單精度”）（在C派生的程式語言中為“ float”）和64位（稱為“雙精度”（double））。如IEEE 754標準所定義，一個32位浮點值包括一個符號位，8個指數位和23個尾數位。64位雙精度數包括一個符號位，11個指數位和52個尾數位。在本文中，對（較新的）IEEE 754標準16位浮點半型別感興趣，該型別包括一個符號位，5個指數位和10個尾數位，如圖1所示。

圖1：16位半精度浮點（FP16）表示形式：1個符號位，5個指數位和10個尾數位。
要了解16位精度會有什麼不同，FP16可以表示2 -14和2 15（指數範圍）之間的2的冪的1024個值。那是30,720個值。將此與FP32相比，FP32可以表示2 -126與2 127之間的2的冪的大約800萬個值。大約有20億個值，相差很大。那麼，為什麼要使用像FP16這樣的小浮點格式呢？因為效能。
NVIDIA Tesla P100（基於GP100 GPU）支援2路向量半精度融合乘加（FMA）指令（操作碼HFMA2），該指令的釋出速度與32位FMA指令相同。這意味著半精度演算法的吞吐量是P100上單精度演算法的兩倍，是雙精度演算法的四倍。具體而言，支援NVLink的P100（SXM2模組）的半精度精度為21.2 Teraflop / s。憑藉如此巨大的效能優勢，值得研究如何使用它。
使用降低的精度時要記住的一件事是，由於FP16的規格化範圍較小，因此生成次正規數（也稱為非正規數）的可能性增加。因此，重要的是，NVIDIA GPU必須以低於正常水平的效能實現FMA操作。某些處理器不會這樣做，並且效能可能會受到影響。（注意：啟用“重新整理到零”可能仍然會帶來好處。請參閱“ CUDA Pro提示：放心重新整理異常”。）
高效能與低精度整數
浮點數將高動態範圍與高精度結合在一起，但是在某些情況下，不需要動態範圍，因此整數可以勝任。甚至在某些應用中，正在處理的資料的精度也很低，因此可以使用非常低精度的儲存（例如C short或char / byte型別）。

圖2：Tesla P4和P40 GPU中的新DP4A和DP2A指令提供具有32位整數累加的快速2和4路8位/ 16位整數向量點積。
對於此類應用，最新的Pascal GPU（GP102，GP104和GP106）引入了新的8位整數4元素向量點積（DP4A）和16位2元素向量點積（DP2A）指令。DP4A執行兩個4元素向量A和B（每個向量都包含儲存在32位字中的4個單位元組值）之間的向量點積，將結果儲存在32位整數中，並將其新增到第三個引數C中，也是32位整數。參見圖2。DP2A是類似的指令，其中A是16位值的2元素向量，而B是8位值的4元素向量，並且DP2A的不同形式為2選擇高位元組或低位元組對。雙向點積。這些靈活的指令可用於線性代數計算，例如矩陣乘法和卷積。對於實現用於深度學習推理的8位整數卷積特別強大，這在部署用於影像分類和目標檢測的深度神經網路中很常見。圖3顯示了在AlexNet上使用INT8卷積在Tesla P4 GPU上實現的提高的電源效率。

圖3：與上一代Tesla M4 GPU上的FP32相比，在Tesla P4上使用INT8計算進行深度學習推理可大大提高使用AlexNet和其它深度神經網路進行影像識別的電源效率。Tesla P4的計算效率比Arria10 FPGA高出8倍，比Intel Xeon CPU高40倍。（AlexNet，批處理大小= 128，CPU：使用Intel MKL 2017的Intel E5-2690v4，FPGA為Arria10-115.1x M4 / P4節點，P4板功率為56W，P4 GPU功率為36W，M4板功率為57W， M4 GPU功率為39W，Perf / W圖表使用GPU功率。）
DP4A計算總計八個整數運算的等效項，DP2A計算四個整數運算。這樣，Tesla P40（基於GP102）的峰值整數吞吐量為47 TOP / s（每秒Tera操作）。
DP4A的一個示例應用是通常在射電望遠鏡資料處理管道中使用的互相關演算法。與光學望遠鏡一樣，大型射電望遠鏡可以分辨宇宙中微弱的物體和更遠的物體。但是建造越來越大的單片單天線射電射電望遠鏡是不切實際的。取而代之的是，射電天文學家建立了分佈在大面積上的許多天線陣列。要使用這些望遠鏡，來自所有天線的訊號必須是互相關的-高度並行的計算，其成本隨天線數量成倍增加。由於射電望遠鏡元件通常捕獲非常低的精度資料，因此訊號的互相關不需要浮點計算。GPU已用於生產射電天文學互相關，但他們通常使用FP32計算。DP4A的引入保證了該計算的更高功率效率。圖4顯示了修改a的結果互相關程式碼以使用DP4A，從而在具有預設時鐘的Tesla P40 GPU上效率提高了4.5倍（與P40上的FP32計算相比）在GPU時鐘上設定了6.4倍的提高，從而降低了溫度（從而降低了洩漏電流） ）。總體而言，新程式碼比上一代Tesla M40 GPU上的FP32互相關效率高近12倍（來源：Kate Clark）。

圖4：與FP32計算相比，INT8向量點積（DP4A）在很大程度上提高了射電天文互相關的效率。
Pascal GPU上的混合精度效能
半精度（FP16）格式對於GPU來說並不是新事物。實際上，FP16作為儲存格式已經在NVIDIA GPU上得到了多年的支援，主要用於降低精度的浮點紋理儲存和過濾以及其它特殊用途。Pascal GPU體系結構實現了通用的IEEE 754 FP16演算法。如下表所示，Tesla P100（GP100）上全速支援高效能FP16，而其它Pascal GPU（GP102，GP104和GP106）則以較低的吞吐量（類似於雙精度）支援。
GP102-GP106支援8位和16位DP4A和DP2A點產品指令，但GP100不支援。表1顯示了基於Pascal的Tesla GPU上不同數字指令的算術吞吐量。

表1：基於Pascal的Tesla GPU的半，單精度和雙精度融合乘法加法指令以及8位和16位向量點乘積指令的峰值算術吞吐量。（Boost時鐘速率用於計算峰值吞吐量。TFLOP / s：每秒Tera浮點運算。TIOP / s：每秒Tera整數運算。）
NVIDIA庫的混合精度程式設計
從應用程式的混合精度中受益的最簡單方法是利用NVIDIA GPU庫中對FP16和INT8計算的支援。NVIDIA SDK的金鑰庫支援計算和儲存的多種精度。
表2顯示了關鍵CUDA庫以及PTX彙編和CUDA C / C ++內部函式中對FP16和INT8的當前支援。

表2：CUDA 8 FP16和INT8 API和庫支援。
神經網路
cuDNN是用於訓練和部署深度神經網路的原始例程庫。cuDNN 5.0包括對前向卷積的FP16支援，並增加了對FP16後向卷積的支援。庫中的所有其它例程均受記憶體限制，因此FP16計算對效能無益。因此，這些例程使用FP32計算，但支援FP16資料輸入和輸出。cuDNN 6將增加對INT8推理卷積的支援。
TensorRT
TensorRT是用於深度學習應用程式生產部署的高效能深度學習推理引擎，該引擎自動優化訓練有素的神經網路以實現執行時效能。TensorRT v1支援FP16進行推理卷積，而v2支援INT8進行推理卷積。
cuBlas
cuBLAS是用於密集線性代數的GPU庫，它是BLAS（基本線性代數子例程）的實現。cuBLAS支援幾種矩陣矩陣乘法例程中的混合精度。cublasHgemm是FP16密集矩陣矩陣乘法例程，使用FP16進行計算以及輸入和輸出。cublasSgemmEx（）在FP32中計算，但是輸入資料可以是FP32，FP16或INT8，輸出可以是FP32或FP16。cublasGemm（）是CUDA 8中的新例程，它允許指定計算精度，包括INT8計算（使用DP4A）。
將根據需求增加對更多具有FP16計算和/或儲存功能的BLAS 3級例程的支援。1級和2級BLAS例程受記憶體限制，因此降低精度的計算是無益的。
傅立葉變換
cuFFT是在CUDA中實現的流行的快速傅立葉變換庫。從CUDA 7.5開始，cuFFT支援FP16的單GPU FFT計算和儲存。FP16 FFT的速度比FP32快2倍。FP16計算需要具有Compute Capability 5.3或更高版本（Maxwell架構）的GPU。大小目前限制為2的冪，並且不支援R2C或C2R轉換的實部上的跨步。
cuSPARSE
cuSPARSE是用於稀疏矩陣的GPU加速線性代數例程庫。cuSPARSE支援FP16的多個例程儲存（cusparseXtcsrmv（），cusparseCsrsv_analysisEx（），cusparseCsrsv_solveEx（），cusparseScsr2cscEx（）和cusparseCsrilu0Ex（））。正在研究cuSPARSE的FP16計算。
在CUDA程式碼中使用混合精度
對於自定義CUDA C ++核心的開發人員和Thrust並行演算法庫的使用者，CUDA提供了從FP16和INT8計算，儲存和I / O中獲得最大收益所需的型別定義和API。
FP16型別和內在函式
對於FP16，CUDA在CUDA包含路徑中包含的標頭“ cuda_fp16.h”中定義了“ half”和“ half2”型別。該標頭檔案還定義了一套完整的內部函式，用於對“半”資料進行操作。例如，下面顯示了標量FP16加法函式“ hadd（）”和2路向量FP16加法函式“ hadd2（）”的宣告。
device __half __hadd（const __half a，const __half b）;
device __half2 __hadd2（const __half2 a，const __half2 b）;
cuda_fp16.h定義了一套完整的半精度內在函式，用於算術，比較，轉換和資料移動以及其它數學函式。所有這些都在CUDA Math API文件中進行了描述。
在可能的情況下使用“ half2”向量型別和內在函式來實現最高吞吐量。GPU硬體算術指令一次對2個FP16值進行運算，並打包在32位暫存器中。表1中的峰值吞吐率假設為“ half2”向量計算。如果使用標量“半”指令，則可以達到峰值吞吐量的50％。同樣，在從FP16陣列載入和儲存到FP16陣列時要實現最大頻寬，需要向量訪問“ half2”資料。理想情況下，可以通過載入和儲存“ float2”或“ float4”型別並強制轉換為“ half2”或從“ half2”進行轉換，來進一步向量化負載以實現更高的頻寬。
以下示例程式碼演示瞭如何使用CUDA __hfma() （半精度融合乘加）和其它內在函式來計算半精度AXPY（A * X + Y）。該示例的完整程式碼在Github上可用，並且顯示瞭如何在主機上初始化半精度陣列。重要的是，當開始使用half型別時，可能需要 在主機端程式碼中的half 和float值之間進行轉換。包括一些快速的CPU型別轉換例程（有關完整原始碼，請參見相關的Gist）。在此示例中，使用了Giesen的一些程式碼。
全球
void haxpy（int n，half a，const half * x，half * y）
{
整數開始= threadIdx.x + blockDim.x * blockIdx.x;
int stride = blockDim.x * gridDim.x;

#if CUDA_ARCH> = 530
int n2 = n / 2;
half2 * x2 =（half2 ）x， y2 =（half2 *）y;

for（int i =開始； i <n2； i + =步幅）
y2 [i] = hfma2（ halves2half2（a，a），x2 [i]，y2 [i]）;

//第一個執行緒處理奇數陣列的單例

如果（開始== 0 &&（n％2））
y [n-1] = __hfma（a，x [n-1]，y [n-1]）;

＃其它
for（int i = start; i <n; i + = stride）{
y [i] = float2half（ half2float（a）* __half2float（x [i]）
+ __half2float（y [i]））;
}
＃萬一
}
整數點乘本徵
CUDA在標頭“ sm_61_intrinsics.h”（sm_61是與GP102，GP104和GP106對應的SM架構）中為8位和16位點乘積（先前描述的DP4A和DP2A指令）定義了內部函式。）。為方便起見，DP4A內部函式有int和char4版本，有符號和無符號兩種形式：
device int __dp4a（int srcA，int srcB，int c）;int __dp4a （int srcA ，int srcB ，int c ）;
device int __dp4a（char4 srcA，char4 srcB，int c）;int __dp4a （char4 srcA ，char4 srcB ，int c ）;
device unsigned int __dp4a（unsigned int srcA，unsigned int srcB，unsigned int c）;unsigned int __dp4a （unsigned int srcA ，unsigned int srcB ，unsigned int c ）;
device unsigned int __dp4a（uchar4 srcA，uchar4 srcB，unsigned int c）;unsigned int __dp4a （uchar4 srcA ，uchar4 srcB ，unsigned int c ）;
兩種版本均假定A和B的四個向量元素被打包到32位字的四個相應位元組中。char4 /uchar4版本使用帶有顯式欄位的CUDA的struct型別，而打包在int版本中是隱式的。
如前所述，DP2A具有“高”和“低”版本，分別用於選擇輸入B的高或低兩個位元組。
//通用[_lo]
device int __dp2a_lo（int srcA，int srcB，int c）;int __dp2a_lo （int srcA ，int srcB ，int c ）;
device unsigned int __dp2a_lo（unsigned int srcA，unsigned int srcB，unsigned int c）;unsigned int __dp2a_lo （unsigned int srcA ，unsigned int srcB ，unsigned int c ）;

//向量樣式[_lo]//向量樣式[_lo]
device int __dp2a_lo（short2 srcA，char4 srcB，int c）;int __dp2a_lo （short2 srcA ，char4 srcB ，int c ）;
device unsigned int __dp2a_lo（ushort2 srcA，uchar4 srcB，unsigned int c）;unsigned int __dp2a_lo （ushort2 srcA ，uchar4 srcB ，unsigned int c ）;

//通用[_hi]//通用[_hi]
device int __dp2a_hi（int srcA，int srcB，int c）;int __dp2a_hi （int srcA ，int srcB ，int c ）;
device unsigned int __dp2a_hi（unsigned int srcA，unsigned int srcB，unsigned int c）;unsigned int __dp2a_hi （unsigned int srcA ，unsigned int srcB ，unsigned int c ）;

//向量樣式[_hi]//向量樣式[_hi]
device int __dp2a_hi（short2 srcA，char4 srcB，int c）;int __dp2a_hi （short2 srcA ，char4 srcB ，int c ）;
device unsigned int __dp2a_hi（ushort2 srcA，uchar4 srcB，unsigned int c）;unsigned int __dp2a_hi （ushort2 srcA ，uchar4 srcB ，unsigned int c ）;
請記住，DP2A和DP4A在基於GP102，GP104和GP106 GPU的Tesla，GeForce和Quadro加速器上可用，但在基於Tesla P100（基於GP100 GPU）上不可用。