CUDA 8混合精度編程
Mixed-Precision Programming with CUDA 8
論文地址：https://devblogs.nvidia.com/mixed-precision-programming-cuda-8/
更新，2019年3月25日：最新的Volta和Turing GPU現在加入了張量核，加速了某些類型的FP16矩陣數學。這使得在流行的人工智能框架中進行更快、更容易的混合精度計算成為可能。使用張量磁芯需要使用CUDA9或更高版本。NVIDIA還為TensorFlow、PyTorch和MXNet添加了自動混合精度功能。想多學點還是自己試試？在這里獲取流行的人工智能框架的tensor核心優化示例。
在軟件開發的實踐中，程序員很早就學會了使用正確的工具來完成工作的重要性。當涉及到數值計算時，這一點尤其重要，因為在精度、精度和性能之間的權衡使得選擇數據的最佳表示非常重要。隨着Pascal GPU體系結構和CUDA 8的引入，NVIDIA正在擴展可用於混合精度計算的工具集，包括新的16位浮點和8/16位整數計算功能。
“隨着在不同精度下計算的相對成本和易用性的發展，由於體系結構和軟件的變化，以及GPU等加速器的破壞性影響，將看到混合精度算法的開發和使用越來越多。”—Nick Higham，Richardson應用數學教授，曼徹斯特大學。
許多技術和高性能計算機應用需要32位（單浮點數，或FP32）或64位（雙浮點數，或FP64）浮點的高精度計算，甚至還有依賴更高精度（128位或256位浮點）的GPU加速應用。但是有許多應用需要低精度的算法。例如，在快速增長的深度學習領域的研究人員發現，由於訓練深層神經網絡時使用的反向傳播算法，深層神經網絡結構對錯誤具有自然的彈性，一些人認為16位浮點（半精度，或FP16）足以訓練神經網絡。
與精度更高的FP32或FP64相比，存儲FP16（半精度）數據減少了神經網絡的內存使用，允許訓練和部署更大的網絡，並且FP16數據傳輸比FP32或FP64傳輸花費的時間更少。此外，對於許多網絡，可以使用8位整數計算來執行深度學習推斷，而不會對精度產生顯著影響。
除了深度學習之外，使用攝像機或其真實傳感器數據的應用程序通常不需要高精度浮點計算，因為傳感器生成低精度或低動態范圍數據。射電望遠鏡處理的數據就是一個很好的例子。正如將在本文后面看到的，使用8位整數計算可以大大加快用於處理射電望遠鏡數據的互相關算法。
在計算方法中結合使用不同的數值精度稱為混合精度。NVIDIA Pascal體系結構通過在32位數據路徑中添加向量指令（將多個操作打包到一個數據路徑中），為能夠利用較低精度計算的應用程序提供了旨在提供更高性能的功能。具體地說，這些指令操作16位浮點數據（“半”或FP16）和8位和16位整數數據（INT8和INT16）。
新的NVIDIA Tesla P100由GP100 GPU供電，可以以FP32的兩倍吞吐量執行FP16算法。GP102（Tesla P40和NVIDIA Titan X）、GP104（Tesla P4）和GP106 gpu都支持指令，這些指令可以對2和4元素8位向量執行整數點積，並累加為32位整數。這些指令對於實現高效的深度學習推理以及射電天文學等其應用具有重要價值。
在這篇文章中，將提供一些有關半精度浮點的詳細信息，並提供使用FP16和INT8矢量計算在Pascal gpu上可實現的性能的詳細信息。還將討論各種CUDA平台庫和api提供的混合精度計算能力。
A Bit (or 16) about Floating Point Precision
每一位計算機科學家都應該知道，浮點數提供了一種表示法，允許在計算機上對實數進行近似，同時在范圍和精度之間進行權衡。浮點數將實值近似為一組有效數字（稱為尾數或有效位），然后按固定基數的指數縮放（當前大多數計算機上使用的IEEE標准浮點數的基數為2）。
常見的浮點格式包括32位，稱為“單精度”（“float”在C派生的編程語言中）和64位，稱為“雙精度”（“double”）。根據IEEE 754標准的定義，32位浮點值包括符號位、8個指數位和23個尾數位。64位雙精度包含一個符號位、11個指數位和52個尾數位。在本文中，對（較新的）IEEE754標准16位浮點半類型感興趣，包含一個符號位、5個指數位和10個尾數位，如圖1所示。
Figure 1: 16-bit half-precision floating point (FP16) representation: 1 sign bit, 5 exponent bits, and 10 mantissa bits.
為了了解精度16位之間的差異，FP16可以表示2-14和215（其指數范圍）之間2的每個冪的1024個值。這是30720個值。與之形成對比的是FP32，在2-126和2127之間，每2次冪的值約為800萬。這大約是20億的價值，差別很大。那么為什么要使用像FP16這樣的小浮點格式呢？一句話，表演。
NVIDIA Tesla P100（基於GP100 GPU）支持雙向矢量半精度融合乘法加法（FMA）指令（操作碼HFMA2），可以以與32位FMA指令相同的速率發出該指令。這意味着半精度算法在P100上的吞吐量是單精度算法的兩倍，是雙精度算法的四倍。具體來說，啟用NVLink的P100（SXM2模塊）能夠達到21.2teraflop/s的半精度。有了這么大的性能優勢，應該看看如何使用。
在使用降低精度時要記住的一點是，由於FP16的標准化范圍較小，生成次標准化數（也稱為非標准化數）的概率增加。因此，NVIDIA的gpu必須在低標准數上實現FMA操作，並具有完整的性能。有些處理器沒有，性能會受到影響。（注意：啟用“flush to zero”仍有好處）。請參閱文章“CUDA Pro Tip:Flush Denormals with Confidence”。）
High Performance with Low-Precision Integers
浮點數結合了高動態范圍和高精度，但也有不需要動態范圍的情況，因此整數可以完成這項工作。甚至有些應用程序處理的數據精度很低，因此可以使用非常低的精度存儲（如C short或char/byte類型）。
Figure 2: New DP4A and DP2A instructions in Tesla P4 and P40 GPUs provide fast 2- and 4-way 8-bit/16-bit integer vector dot products with 32-bit integer accumulation.
對於此類應用，最新的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上實現的改進的功率效率。
Figure 3: Using INT8 computation on the Tesla P4 for deep learning inference provides a very large improvement in power efficiency for image recognition using AlexNet and other deep neural networks, when compared to FP32 on previous generation Tesla M4 GPUs. Efficiency of this computation on Tesla P4 is up to 8x more efficient than an Arria10 FPGA, and up to 40x more efficient than an Intel Xeon CPU. (AlexNet, batch size = 128, CPU: Intel E5-2690v4 using Intel MKL 2017, FPGA is Arria10-115. 1x M4/P4 in node, P4 board power at 56W, P4 GPU power at 36W, M4 board power at 57W, M4 GPU power at 39W, Perf/W chart using GPU power.)
DP4A計算總共8個整數操作的等效值，DP2A計算4個。這使Tesla P40（基於GP102）的峰值整數吞吐量達到47 TOP/s（Tera操作/秒）。
DP4A的一個應用實例是在射電望遠鏡數據處理管道中常用的互相關算法。與光學望遠鏡一樣，較大的射電望遠鏡可以分辨宇宙中較暗和較遠的物體；但是，建造越來越大的單片單天線射電望遠鏡是不實際的。取而代之的是，射電天文學家在大面積上建造了許多天線陣列。要使用這些望遠鏡，來自所有天線的信號必須是互相關的，這是一種高度並行的計算，其成本與天線數量成正比。由於射電望遠鏡元件通常捕獲非常低精度的數據，所以信號的互相關不需要浮點運算。gpu已經被用於射電天文學互相關的制作，但通常使用FP32計算。DP4A的引入為這種計算提供了更高的功率效率。
圖4顯示了修改互相關代碼以使用DP4A的結果，從而在具有默認時鍾的Tesla P40 GPU上提高了4.5倍的效率（與P40上的FP32計算相比），並在設置GPU時鍾以降低溫度（從而降低泄漏電流）的情況下提高了6.4倍。總的來說，新代碼比上一代Tesla M40 GPU上的FP32交叉相關效率高出近12倍（圖片來源：Kate Clark）。
Figure 4: INT8 vector dot products (DP4A) improve the efficiency of radio astronomy cross-correlation by a large factor compared to FP32 computation.
Mixed Precision Performance on Pascal GPUs
半精度（FP16）格式對gpu來說並不新鮮。事實上，FP16作為一種存儲格式在NVIDIA GPUs上已經支持了很多年，主要用於降低精度的浮點紋理存儲和過濾等特殊用途的操作。Pascal GPU架構實現了通用的IEEE 754 FP16算法。高性能FP16在Tesla P100（GP100）上以全速支持，在其Pascal gpu（GP102、GP104和GP106）上以較低的吞吐量（類似於雙精度）支持，如下表所示。
GP102-GP106支持8位和16位DP4A和DP2A點產品指令，但GP100不支持。表1顯示了基於Pascal的Tesla gpu上不同數值指令的算術吞吐量。
Table 1: Pascal-based Tesla GPU peak arithmetic throughput for half-, single-, and double-precision fused multiply-add instructions, and for 8- and 16-bit vector dot product instructions. (Boost clock rates are used in calculating peak throughputs. TFLOP/s: Tera Floating-point Operations per Second. TIOP/s: Tera Integer Operations per Second.
Mixed-Precision Programming with NVIDIA Libraries
從應用程序的混合精度中獲益的最簡單方法是利用NVIDIA GPU庫中對FP16和INT8計算的支持。NVIDIA SDK的密鑰庫現在支持計算和存儲的各種精度。
表2顯示了當前對FC16和It8在關鍵CUDA庫以及PTX組件和CUDA C/C++內部的支持。
Table 2: CUDA 8 FP16 and INT8 API and library support.
cuDNN
cuDNN是一個原始程序庫，用於訓練和部署深層神經網絡。cuDNN 5.0包括對前向卷積的FP16支持，以及對FP16后向卷積的5.1附加支持。庫中的所有其例程都是內存綁定的，因此FP16計算不利於性能。因此，這些例程使用FP32計算，但支持FP16數據輸入和輸出。cuDNN 6將增加對INT8推理卷積的支持。
TensorRT
TensorRT是一個高性能的深度學習推理機，用於深度學習應用程序的生產部署，自動優化訓練神經網絡的運行時性能。TensorRT v1支持FP16進行推理卷積，v2支持INT8進行推理卷積。
cuBLAS
cuBLAS是一個用於密集線性代數的GPU庫，是基本線性代數子程序BLAS的一個實現。cuBLAS在幾個矩陣乘法例程中支持混合精度。cubrashgemm是一個FP16密集矩陣乘法例程，使用FP16進行計算以及輸入和輸出。cubassgemex（）在FP32中計算，但輸入數據可以是FP32、FP16或INT8，輸出可以是FP32或FP16。cublasgem（）是CUDA 8中的一個新例程，允許指定計算精度，包括INT8計算（使用DP4A）。
將根據需要添加對具有FP16計算和/或存儲的更多BLAS級別3例程的支持，因此如果需要，請與聯系。級別1和級別2的BLAS例程是內存限制的，因此減少精度計算是不利的。
cuFFT
cuft是CUDA中一種流行的快速傅立葉變換庫。從CUDA 7.5開始，cuft支持單GPU fft的FP16計算和存儲。FP16 FFT比FP32快2倍。FP16計算需要一個計算能力為5.3或更高的GPU（Maxwell架構）。當前大小限制為2的冪，並且不支持R2C或C2R轉換的實際部分上的跨步。
cuSPARSE
cuSPARSE是一個用於稀疏矩陣的GPU加速線性代數例程庫。cuSPARSE支持幾個例程的FP16存儲（`cusparseXtcsrmv（）`、`cusparseCsrsv_analysisEx（）`、`cusparseCsrsv_solvex（）`、`cusparseScsr2cscEx（）`和`cusparseCsrilu0Ex（）`）。正在研究cuSPARSE的FP16計算。
Using Mixed Precision in your own CUDA Code
對於定制的CUDA C++內核和推力並行算法庫的用戶，CUDA提供了需要從FP16和It8計算、存儲和I/O.中充分利用的類型定義和API。
FP16 types and intrinsics
對於FP16，CUDA定義了CUDA include路徑中包含的頭文件“CUDA_FP16.h”中的“half”和“half 2”類型。此頭還定義了一組完整的內部函數，用於對“半”數據進行操作。例如，下面顯示標量FP16加法函數“hadd（）”和雙向向量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中的峰值吞吐量數字采用“半2”矢量計算。如果使用標量“half”指令，則可以達到峰值吞吐量的50%。同樣，要在從FP16陣列加載和存儲到FP16陣列時獲得最大帶寬，需要對“半2”數據進行矢量訪問。理想情況下，可以通過加載和存儲“float2”或“float4”類型並強制轉換到“half2”或從“half2”轉換到“half2”，進一步將加載矢量化以獲得更高的帶寬。有關相關示例，請參閱所有Pro-Tip博客文章的上一篇平行文章。
下面的示例代碼演示如何使用CUDA的uu hfma（）（半精度融合乘法加法）和其內部函數計算半精度AXPY（a*X+Y）。該示例的完整代碼在Github上提供，展示了如何在主機上初始化半精度數組。重要的是，當開始使用半類型時，可能需要在主機端代碼中的半值和浮點值之間進行轉換。這篇來自FabianGiesen的博客文章包含了一些快速CPU類型轉換例程（請參閱相關的要點以獲得完整的源代碼）。在這個例子中使用了一些Giesen的代碼。
__global__
void haxpy(int n, half a, const half *x, half *y)
{
int start = 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 = start; i < n2; i+= stride)
y2[i] = __hfma2(__halves2half2(a, a), x2[i], y2[i]);
// first thread handles singleton for odd arrays
if (start == 0 && (n%2))
y[n-1] = __hfma(a, x[n-1], y[n-1]);
#else
for (int i = start; i < n; i+= stride) {
y[i] = __float2half(__half2float(a) * __half2float(x[i])
+ __half2float(y[i]));
}
#endif
}
Integer Dot Product Intrinsics
CUDA在頭文件“smɤu intrinsics.h”（smɤ61是對應於GP102、GP104和GP106的sm體系結構）中定義8位和16位點產品（前面描述的DP4A和DP2A指令）的內部函數。也稱為計算能力6.1。為了方便起見，DP4A內部函數有“int”和“char4”兩種版本，有符號和無符號兩種：
__device__ int __dp4a(int srcA, int srcB, int c);
__device__ int __dp4a(char4 srcA, char4 srcB, int c);
__device__ unsigned int __dp4a(unsigned int srcA, unsigned int srcB, unsigned int c);
__device__ unsigned int __dp4a(uchar4 srcA, uchar4 srcB, unsigned int c);
兩個版本都假設A和B的四個向量元素被壓縮到32位字的四個相應字節中。char4`/`uchar4`版本使用帶有顯式字段的CUDA結構類型，而包裝在'int'版本中是隱式的。
如前所述，DP2A具有用於分別選擇輸入B的高或低兩個字節的“高”和“低”版本。
// Generic [_lo]
__device__ int __dp2a_lo(int srcA, int srcB, int c);
__device__ unsigned int __dp2a_lo(unsigned int srcA, unsigned int srcB, unsigned int c);
// Vector-style [_lo]
__device__ int __dp2a_lo(short2 srcA, char4 srcB, int c);
__device__ unsigned int __dp2a_lo(ushort2 srcA, uchar4 srcB, unsigned int c);
// Generic [_hi]
__device__ int __dp2a_hi(int srcA, int srcB, int c);
__device__ unsigned int __dp2a_hi(unsigned int srcA, unsigned int srcB, unsigned int c);
// Vector-style [_hi]
__device__ int __dp2a_hi(short2 srcA, char4 srcB, int c);
__device__ unsigned int __dp2a_hi(ushort2 srcA, uchar4 srcB, unsigned int c);
請記住，基於GP102、GP104和GP106 GPU的Tesla、GeForce和Quadro加速器上提供了DP2A和DP4A，而不是Tesla P100（基於GP100 GPU）。
Download CUDA 8
要充分利用GPU上的混合精度計算，請下載免費的NVIDIA CUDA工具包版本8。要了解CUDA 8的所有強大功能，請查看后cuda8顯示的功能。