CUDA 8混合精度編程

Mixed-Precision Programming with CUDA 8

論文地址：https://devblogs.nvidia.com/mixed-precision-programming-cuda-8/

更新，2019年3月25日：最新的Volta和Turing GPU現(xiàn)在加入了張量核，加速了某些類(lèi)型的FP16矩陣數(shù)學(xué)。這使得在流行的人工智能框架中進(jìn)行更快、更容易的混合精度計(jì)算成為可能。使用張量磁芯需要使用CUDA9或更高版本。NVIDIA還為T(mén)ensorFlow、PyTorch和MXNet添加了自動(dòng)混合精度功能。想多學(xué)點(diǎn)還是自己試試？在這里獲取流行的人工智能框架的tensor核心優(yōu)化示例。

在軟件開(kāi)發(fā)的實(shí)踐中，程序員很早就學(xué)會(huì)了使用正確的工具來(lái)完成工作的重要性。當(dāng)涉及到數(shù)值計(jì)算時(shí)，這一點(diǎn)尤其重要，因?yàn)樵诰取⒕群托阅苤g的權(quán)衡使得選擇數(shù)據(jù)的最佳表示非常重要。隨著Pascal GPU體系結(jié)構(gòu)和CUDA 8的引入，NVIDIA正在擴(kuò)展可用于混合精度計(jì)算的工具集，包括新的16位浮點(diǎn)和8/16位整數(shù)計(jì)算功能。

“隨著在不同精度下計(jì)算的相對(duì)成本和易用性的發(fā)展，由于體系結(jié)構(gòu)和軟件的變化，以及GPU等加速器的破壞性影響，將看到混合精度算法的開(kāi)發(fā)和使用越來(lái)越多。”—Nick Higham，Richardson應(yīng)用數(shù)學(xué)教授，曼徹斯特大學(xué)。

許多技術(shù)和高性能計(jì)算機(jī)應(yīng)用需要32位（單浮點(diǎn)數(shù)，或FP32）或64位（雙浮點(diǎn)數(shù)，或FP64）浮點(diǎn)的高精度計(jì)算，甚至還有依賴(lài)更高精度（128位或256位浮點(diǎn)）的GPU加速應(yīng)用。但是有許多應(yīng)用需要低精度的算法。例如，在快速增長(zhǎng)的深度學(xué)習(xí)領(lǐng)域的研究人員發(fā)現(xiàn)，由于訓(xùn)練深層神經(jīng)網(wǎng)絡(luò)時(shí)使用的反向傳播算法，深層神經(jīng)網(wǎng)絡(luò)結(jié)構(gòu)對(duì)錯(cuò)誤具有自然的彈性，一些人認(rèn)為16位浮點(diǎn)（半精度，或FP16）足以訓(xùn)練神經(jīng)網(wǎng)絡(luò)。

與精度更高的FP32或FP64相比，存儲(chǔ)FP16（半精度）數(shù)據(jù)減少了神經(jīng)網(wǎng)絡(luò)的內(nèi)存使用，允許訓(xùn)練和部署更大的網(wǎng)絡(luò)，并且FP16數(shù)據(jù)傳輸比FP32或FP64傳輸花費(fèi)的時(shí)間更少。此外，對(duì)于許多網(wǎng)絡(luò)，可以使用8位整數(shù)計(jì)算來(lái)執(zhí)行深度學(xué)習(xí)推斷，而不會(huì)對(duì)精度產(chǎn)生顯著影響。

除了深度學(xué)習(xí)之外，使用攝像機(jī)或其真實(shí)傳感器數(shù)據(jù)的應(yīng)用程序通常不需要高精度浮點(diǎn)計(jì)算，因?yàn)閭鞲衅魃傻途然虻蛣?dòng)態(tài)范圍數(shù)據(jù)。射電望遠(yuǎn)鏡處理的數(shù)據(jù)就是一個(gè)很好的例子。正如將在本文后面看到的，使用8位整數(shù)計(jì)算可以大大加快用于處理射電望遠(yuǎn)鏡數(shù)據(jù)的互相關(guān)算法。

在計(jì)算方法中結(jié)合使用不同的數(shù)值精度稱(chēng)為混合精度。NVIDIA Pascal體系結(jié)構(gòu)通過(guò)在32位數(shù)據(jù)路徑中添加向量指令（將多個(gè)操作打包到一個(gè)數(shù)據(jù)路徑中），為能夠利用較低精度計(jì)算的應(yīng)用程序提供了旨在提供更高性能的功能。具體地說(shuō)，這些指令操作16位浮點(diǎn)數(shù)據(jù)（“半”或FP16）和8位和16位整數(shù)數(shù)據(jù)（INT8和INT16）。

新的NVIDIA Tesla P100由GP100 GPU供電，可以以FP32的兩倍吞吐量執(zhí)行FP16算法。GP102（Tesla P40和NVIDIA Titan X）、GP104（Tesla P4）和GP106 gpu都支持指令，這些指令可以對(duì)2和4元素8位向量執(zhí)行整數(shù)點(diǎn)積，并累加為32位整數(shù)。這些指令對(duì)于實(shí)現(xiàn)高效的深度學(xué)習(xí)推理以及射電天文學(xué)等其應(yīng)用具有重要價(jià)值。

在這篇文章中，將提供一些有關(guān)半精度浮點(diǎn)的詳細(xì)信息，并提供使用FP16和INT8矢量計(jì)算在Pascal gpu上可實(shí)現(xiàn)的性能的詳細(xì)信息。還將討論各種CUDA平臺(tái)庫(kù)和api提供的混合精度計(jì)算能力。

A Bit (or 16) about Floating Point Precision

每一位計(jì)算機(jī)科學(xué)家都應(yīng)該知道，浮點(diǎn)數(shù)提供了一種表示法，允許在計(jì)算機(jī)上對(duì)實(shí)數(shù)進(jìn)行近似，同時(shí)在范圍和精度之間進(jìn)行權(quán)衡。浮點(diǎn)數(shù)將實(shí)值近似為一組有效數(shù)字（稱(chēng)為尾數(shù)或有效位），然后按固定基數(shù)的指數(shù)縮放（當(dāng)前大多數(shù)計(jì)算機(jī)上使用的IEEE標(biāo)準(zhǔn)浮點(diǎn)數(shù)的基數(shù)為2）。
常見(jiàn)的浮點(diǎn)格式包括32位，稱(chēng)為“單精度”（“float”在C派生的編程語(yǔ)言中）和64位，稱(chēng)為“雙精度”（“double”）。根據(jù)IEEE 754標(biāo)準(zhǔn)的定義，32位浮點(diǎn)值包括符號(hào)位、8個(gè)指數(shù)位和23個(gè)尾數(shù)位。64位雙精度包含一個(gè)符號(hào)位、11個(gè)指數(shù)位和52個(gè)尾數(shù)位。在本文中，對(duì)（較新的）IEEE754標(biāo)準(zhǔn)16位浮點(diǎn)半類(lèi)型感興趣，包含一個(gè)符號(hào)位、5個(gè)指數(shù)位和10個(gè)尾數(shù)位，如圖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（其指數(shù)范圍）之間2的每個(gè)冪的1024個(gè)值。這是30720個(gè)值。與之形成對(duì)比的是FP32，在2-126和2127之間，每2次冪的值約為800萬(wàn)。這大約是20億的價(jià)值，差別很大。那么為什么要使用像FP16這樣的小浮點(diǎn)格式呢？一句話，表演。

NVIDIA Tesla P100（基于GP100 GPU）支持雙向矢量半精度融合乘法加法（FMA）指令（操作碼HFMA2），可以以與32位FMA指令相同的速率發(fā)出該指令。這意味著半精度算法在P100上的吞吐量是單精度算法的兩倍，是雙精度算法的四倍。具體來(lái)說(shuō)，啟用NVLink的P100（SXM2模塊）能夠達(dá)到21.2teraflop/s的半精度。有了這么大的性能優(yōu)勢(shì)，應(yīng)該看看如何使用。

在使用降低精度時(shí)要記住的一點(diǎn)是，由于FP16的標(biāo)準(zhǔn)化范圍較小，生成次標(biāo)準(zhǔn)化數(shù)（也稱(chēng)為非標(biāo)準(zhǔn)化數(shù)）的概率增加。因此，NVIDIA的gpu必須在低標(biāo)準(zhǔn)數(shù)上實(shí)現(xiàn)FMA操作，并具有完整的性能。有些處理器沒(méi)有，性能會(huì)受到影響。（注意：啟用“flush to zero”仍有好處）。請(qǐng)參閱文章“CUDA Pro Tip:Flush Denormals with Confidence”。）

High Performance with Low-Precision Integers

浮點(diǎn)數(shù)結(jié)合了高動(dòng)態(tài)范圍和高精度，但也有不需要?jiǎng)討B(tài)范圍的情況，因此整數(shù)可以完成這項(xiàng)工作。甚至有些應(yīng)用程序處理的數(shù)據(jù)精度很低，因此可以使用非常低的精度存儲(chǔ)（如C short或char/byte類(lèi)型）。

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.

對(duì)于此類(lèi)應(yīng)用，最新的Pascal gpu（GP102、GP104和GP106）引入了新的8位整數(shù)4元向量點(diǎn)積（DP4A）和16位2元向量點(diǎn)積（DP2A）指令。DP4A執(zhí)行兩個(gè)4元素向量A和B（每個(gè)向量包含存儲(chǔ)在32位字中的4個(gè)單字節(jié)值）之間的向量點(diǎn)積，將結(jié)果存儲(chǔ)為32位整數(shù)，并將其添加到第三個(gè)參數(shù)C（也是32位整數(shù)）中。見(jiàn)圖2。DP2A是類(lèi)似的指令，其中a是16位值的2元向量，B是8位值的4元向量，不同類(lèi)型的DP2A為2路點(diǎn)積選擇高字節(jié)對(duì)或低字節(jié)對(duì)。這些靈活的指令對(duì)于線性代數(shù)計(jì)算（如矩陣乘法和卷積）非常有用。對(duì)于實(shí)現(xiàn)用于深度學(xué)習(xí)推理的8位整數(shù)卷積特別強(qiáng)大，通常用于部署用于圖像分類(lèi)和對(duì)象檢測(cè)的深度神經(jīng)網(wǎng)絡(luò)。圖3顯示了在AlexNet上使用INT8卷積在Tesla P4 GPU上實(shí)現(xiàn)的改進(jìn)的功率效率。

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計(jì)算總共8個(gè)整數(shù)操作的等效值，DP2A計(jì)算4個(gè)。這使Tesla P40（基于GP102）的峰值整數(shù)吞吐量達(dá)到47 TOP/s（Tera操作/秒）。

DP4A的一個(gè)應(yīng)用實(shí)例是在射電望遠(yuǎn)鏡數(shù)據(jù)處理管道中常用的互相關(guān)算法。與光學(xué)望遠(yuǎn)鏡一樣，較大的射電望遠(yuǎn)鏡可以分辨宇宙中較暗和較遠(yuǎn)的物體；但是，建造越來(lái)越大的單片單天線射電望遠(yuǎn)鏡是不實(shí)際的。取而代之的是，射電天文學(xué)家在大面積上建造了許多天線陣列。要使用這些望遠(yuǎn)鏡，來(lái)自所有天線的信號(hào)必須是互相關(guān)的，這是一種高度并行的計(jì)算，其成本與天線數(shù)量成正比。由于射電望遠(yuǎn)鏡元件通常捕獲非常低精度的數(shù)據(jù)，所以信號(hào)的互相關(guān)不需要浮點(diǎn)運(yùn)算。gpu已經(jīng)被用于射電天文學(xué)互相關(guān)的制作，但通常使用FP32計(jì)算。DP4A的引入為這種計(jì)算提供了更高的功率效率。

圖4顯示了修改互相關(guān)代碼以使用DP4A的結(jié)果，從而在具有默認(rèn)時(shí)鐘的Tesla P40 GPU上提高了4.5倍的效率（與P40上的FP32計(jì)算相比），并在設(shè)置GPU時(shí)鐘以降低溫度（從而降低泄漏電流）的情況下提高了6.4倍。總的來(lái)說(shuō)，新代碼比上一代Tesla M40 GPU上的FP32交叉相關(guān)效率高出近12倍（圖片來(lái)源：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）格式對(duì)gpu來(lái)說(shuō)并不新鮮。事實(shí)上，FP16作為一種存儲(chǔ)格式在NVIDIA GPUs上已經(jīng)支持了很多年，主要用于降低精度的浮點(diǎn)紋理存儲(chǔ)和過(guò)濾等特殊用途的操作。Pascal GPU架構(gòu)實(shí)現(xiàn)了通用的IEEE 754 FP16算法。高性能FP16在Tesla P100（GP100）上以全速支持，在其Pascal gpu（GP102、GP104和GP106）上以較低的吞吐量（類(lèi)似于雙精度）支持，如下表所示。

GP102-GP106支持8位和16位DP4A和DP2A點(diǎn)產(chǎn)品指令，但GP100不支持。表1顯示了基于Pascal的Tesla gpu上不同數(shù)值指令的算術(shù)吞吐量。

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

從應(yīng)用程序的混合精度中獲益的最簡(jiǎn)單方法是利用NVIDIA GPU庫(kù)中對(duì)FP16和INT8計(jì)算的支持。NVIDIA SDK的密鑰庫(kù)現(xiàn)在支持計(jì)算和存儲(chǔ)的各種精度。

表2顯示了當(dāng)前對(duì)FC16和It8在關(guān)鍵CUDA庫(kù)以及PTX組件和CUDA C/C++內(nèi)部的支持。

Table 2: CUDA 8 FP16 and INT8 API and library support.

cuDNN

cuDNN是一個(gè)原始程序庫(kù)，用于訓(xùn)練和部署深層神經(jīng)網(wǎng)絡(luò)。cuDNN 5.0包括對(duì)前向卷積的FP16支持，以及對(duì)FP16后向卷積的5.1附加支持。庫(kù)中的所有其例程都是內(nèi)存綁定的，因此FP16計(jì)算不利于性能。因此，這些例程使用FP32計(jì)算，但支持FP16數(shù)據(jù)輸入和輸出。cuDNN 6將增加對(duì)INT8推理卷積的支持。

TensorRT

TensorRT是一個(gè)高性能的深度學(xué)習(xí)推理機(jī)，用于深度學(xué)習(xí)應(yīng)用程序的生產(chǎn)部署，自動(dòng)優(yōu)化訓(xùn)練神經(jīng)網(wǎng)絡(luò)的運(yùn)行時(shí)性能。TensorRT v1支持FP16進(jìn)行推理卷積，v2支持INT8進(jìn)行推理卷積。

cuBLAS

cuBLAS是一個(gè)用于密集線性代數(shù)的GPU庫(kù)，是基本線性代數(shù)子程序BLAS的一個(gè)實(shí)現(xiàn)。cuBLAS在幾個(gè)矩陣乘法例程中支持混合精度。cubrashgemm是一個(gè)FP16密集矩陣乘法例程，使用FP16進(jìn)行計(jì)算以及輸入和輸出。cubassgemex（）在FP32中計(jì)算，但輸入數(shù)據(jù)可以是FP32、FP16或INT8，輸出可以是FP32或FP16。cublasgem（）是CUDA 8中的一個(gè)新例程，允許指定計(jì)算精度，包括INT8計(jì)算（使用DP4A）。

將根據(jù)需要添加對(duì)具有FP16計(jì)算和/或存儲(chǔ)的更多BLAS級(jí)別3例程的支持，因此如果需要，請(qǐng)與聯(lián)系。級(jí)別1和級(jí)別2的BLAS例程是內(nèi)存限制的，因此減少精度計(jì)算是不利的。

cuFFT

cuft是CUDA中一種流行的快速傅立葉變換庫(kù)。從CUDA 7.5開(kāi)始，cuft支持單GPU fft的FP16計(jì)算和存儲(chǔ)。FP16 FFT比FP32快2倍。FP16計(jì)算需要一個(gè)計(jì)算能力為5.3或更高的GPU（Maxwell架構(gòu)）。當(dāng)前大小限制為2的冪，并且不支持R2C或C2R轉(zhuǎn)換的實(shí)際部分上的跨步。

cuSPARSE

cuSPARSE是一個(gè)用于稀疏矩陣的GPU加速線性代數(shù)例程庫(kù)。cuSPARSE支持幾個(gè)例程的FP16存儲(chǔ)（cusparseXtcsrmv（）、cusparseCsrsv_analysisEx（）、cusparseCsrsv_solvex（）、cusparseScsr2cscEx（）和cusparseCsrilu0Ex（））。正在研究cuSPARSE的FP16計(jì)算。

Using Mixed Precision in your own CUDA Code

對(duì)于定制的CUDA C++內(nèi)核和推力并行算法庫(kù)的用戶(hù)，CUDA提供了需要從FP16和It8計(jì)算、存儲(chǔ)和I/O.中充分利用的類(lèi)型定義和API。

FP16 types and intrinsics

對(duì)于FP16，CUDA定義了CUDA include路徑中包含的頭文件“CUDA_FP16.h”中的“half”和“half 2”類(lèi)型。此頭還定義了一組完整的內(nèi)部函數(shù)，用于對(duì)“半”數(shù)據(jù)進(jìn)行操作。例如，下面顯示標(biāo)量FP16加法函數(shù)“hadd（）”和雙向向量FP16加法函數(shù)“hadd2（）”的聲明。

device __half
__hadd ( const __half a, const __half b );

device __half2
__hadd2 ( const __half2 a, const __half2 b );

cuda_fp16.h為算術(shù)、比較、轉(zhuǎn)換和數(shù)據(jù)移動(dòng)以及其數(shù)學(xué)函數(shù)定義了一整套半精度的內(nèi)部函數(shù)。所有這些都在CUDA Math API文檔中描述。

盡可能使用“half2”向量類(lèi)型和內(nèi)部函數(shù)以獲得最高的吞吐量。GPU硬件算術(shù)指令一次對(duì)2個(gè)FP16值進(jìn)行操作，并打包在32位寄存器中。表1中的峰值吞吐量數(shù)字采用“半2”矢量計(jì)算。如果使用標(biāo)量“half”指令，則可以達(dá)到峰值吞吐量的50%。同樣，要在從FP16陣列加載和存儲(chǔ)到FP16陣列時(shí)獲得最大帶寬，需要對(duì)“半2”數(shù)據(jù)進(jìn)行矢量訪問(wèn)。理想情況下，可以通過(guò)加載和存儲(chǔ)“float2”或“float4”類(lèi)型并強(qiáng)制轉(zhuǎn)換到“half2”或從“half2”轉(zhuǎn)換到“half2”，進(jìn)一步將加載矢量化以獲得更高的帶寬。有關(guān)相關(guān)示例，請(qǐng)參閱所有Pro-Tip博客文章的上一篇平行文章。

下面的示例代碼演示如何使用CUDA的uu hfma（）（半精度融合乘法加法）和其內(nèi)部函數(shù)計(jì)算半精度AXPY（a*X+Y）。該示例的完整代碼在Github上提供，展示了如何在主機(jī)上初始化半精度數(shù)組。重要的是，當(dāng)開(kāi)始使用半類(lèi)型時(shí)，可能需要在主機(jī)端代碼中的半值和浮點(diǎn)值之間進(jìn)行轉(zhuǎn)換。這篇來(lái)自FabianGiesen的博客文章包含了一些快速CPU類(lèi)型轉(zhuǎn)換例程（請(qǐng)參閱相關(guān)的要點(diǎn)以獲得完整的源代碼）。在這個(gè)例子中使用了一些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是對(duì)應(yīng)于GP102、GP104和GP106的sm體系結(jié)構(gòu)）中定義8位和16位點(diǎn)產(chǎn)品（前面描述的DP4A和DP2A指令）的內(nèi)部函數(shù)。也稱(chēng)為計(jì)算能力6.1。為了方便起見(jiàn)，DP4A內(nèi)部函數(shù)有“int”和“char4”兩種版本，有符號(hào)和無(wú)符號(hào)兩種：

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);

兩個(gè)版本都假設(shè)A和B的四個(gè)向量元素被壓縮到32位字的四個(gè)相應(yīng)字節(jié)中。char4/uchar4`版本使用帶有顯式字段的CUDA結(jié)構(gòu)類(lèi)型，而包裝在’int’版本中是隱式的。

如前所述，DP2A具有用于分別選擇輸入B的高或低兩個(gè)字節(jié)的“高”和“低”版本。

// 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);

請(qǐng)記住，基于GP102、GP104和GP106 GPU的Tesla、GeForce和Quadro加速器上提供了DP2A和DP4A，而不是Tesla P100（基于GP100 GPU）。

Download CUDA 8

要充分利用GPU上的混合精度計(jì)算，請(qǐng)下載免費(fèi)的NVIDIA CUDA工具包版本8。要了解CUDA 8的所有強(qiáng)大功能，請(qǐng)查看后cuda8顯示的功能。

總結(jié)

以上是生活随笔為你收集整理的CUDA 8混合精度编程的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問(wèn)題。

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