CUDA 8的混合精度編程
Volta和Turing GPU包含 Tensor Cores，可加速某些類型的FP16矩陣數(shù)學(xué)運(yùn)算。這樣可以在流行的AI框架內(nèi)更快，更輕松地進(jìn)行混合精度計(jì)算。要使用Tensor Core，需要使用 CUDA 9 或更高版本。NVIDIA還 為T(mén)ensorFlow，PyTorch和MXNet添加了 自動(dòng)混合精度功能。
流行AI框架的張量核心優(yōu)化示例 。
在軟件開(kāi)發(fā)的實(shí)踐中，程序員通常會(huì)及早學(xué)習(xí)到使用正確的工具完成工作的重要性。當(dāng)涉及數(shù)值計(jì)算時(shí)，這一點(diǎn)尤其重要，因?yàn)樵跀?shù)值計(jì)算中，精度，準(zhǔn)確性和性能之間的折衷使得必須選擇最佳的數(shù)據(jù)表示形式。隨著Pascal GPU架構(gòu)和CUDA 8的推出，NVIDIA正在利用新的16位浮點(diǎn)和8/16位整數(shù)計(jì)算功能擴(kuò)展可用于混合精度計(jì)算的工具集。
“隨著架構(gòu)和軟件的不斷變化以及GPU等加速器的破壞性影響，隨著不同精度的相對(duì)成本和易用性的發(fā)展，將看到越來(lái)越多的混合精度算法得到開(kāi)發(fā)和使用。” —曼徹斯特大學(xué)理查森應(yīng)用數(shù)學(xué)教授尼克·海姆（Nick Higham）
許多技術(shù)和HPC應(yīng)用程序都要求使用32位（單浮點(diǎn)或FP32）或64位（雙浮點(diǎn)或FP64）浮點(diǎn)進(jìn)行高精度計(jì)算，甚至GPU加速的應(yīng)用都依賴于更高的精度（128） -或256位浮點(diǎn)數(shù)！）。在許多應(yīng)用中，低精度的算術(shù)就足夠了。例如，在快速發(fā)展的深度學(xué)習(xí)領(lǐng)域中的研究人員發(fā)現(xiàn)，由于深度神經(jīng)網(wǎng)絡(luò)體系結(jié)構(gòu)用于訓(xùn)練反向傳播算法，因此對(duì)誤差具有自然的抵抗力，并且有人認(rèn)為16位浮點(diǎn)數(shù)（半精度或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ì)算執(zhí)行深度學(xué)習(xí)推理，而不會(huì)對(duì)準(zhǔn)確性產(chǎn)生重大影響。
除了深度學(xué)習(xí)之外，使用來(lái)自相機(jī)或其它實(shí)際傳感器的數(shù)據(jù)的應(yīng)用程序通常不需要高精度的浮點(diǎn)計(jì)算，因?yàn)閭鞲衅鲿?huì)生成低精度或低動(dòng)態(tài)范圍的數(shù)據(jù)。射電望遠(yuǎn)鏡處理的數(shù)據(jù)就是一個(gè)很好的例子。正如將在本文后面看到的那樣，通過(guò)使用8位整數(shù)計(jì)算，可以大大加速用于處理射電望遠(yuǎn)鏡數(shù)據(jù)的互相關(guān)算法。
在計(jì)算方法中不同數(shù)值精度的組合使用稱為混合精度。NVIDIA Pascal架構(gòu)通過(guò)添加將多個(gè)操作打包到32位數(shù)據(jù)路徑中的矢量指令，旨在為可以利用較低精度計(jì)算的應(yīng)用程序提供更高的性能。具體來(lái)說(shuō)，這些指令對(duì)16位浮點(diǎn)數(shù)據(jù)（“ half”或FP16）以及8位和16位整數(shù)數(shù)據(jù)（INT8和INT16）進(jìn)行操作。
由GP100 GPU驅(qū)動(dòng)的新型NVIDIA Tesla P100可以以FP32兩倍的吞吐量執(zhí)行FP16算術(shù)運(yùn)算。GP102（Tesla P40和NVIDIA Titan X），GP104（Tesla P4）和GP106 GPU均支持可在2和4元素8位向量上執(zhí)行整數(shù)點(diǎn)積的指令，并累加為32位整數(shù)。這些指令對(duì)于實(shí)現(xiàn)高效的深度學(xué)習(xí)推理以及射電天文學(xué)等其它應(yīng)用程序非常有價(jià)值。
在本文中，將提供有關(guān)半精度浮點(diǎn)的一些詳細(xì)信息，并提供有關(guān)使用FP16和INT8矢量計(jì)算的Pascal GPU可獲得的性能的詳細(xì)信息。還將討論各種CUDA平臺(tái)庫(kù)和API提供的混合精度計(jì)算功能。
浮點(diǎn)精度（或16）
正如每位計(jì)算機(jī)科學(xué)家都應(yīng)該知道的那樣，浮點(diǎn)數(shù)提供了一種表示形式，可以在范圍和精度之間進(jìn)行權(quán)衡的情況下，在計(jì)算機(jī)上近似實(shí)數(shù)。浮點(diǎn)數(shù)將實(shí)際值近似為一組有效數(shù)字（稱為尾數(shù)或有效位數(shù)），然后以固定基數(shù)（今天大多數(shù)計(jì)算機(jī)上使用的IEEE標(biāo)準(zhǔn)浮點(diǎn)數(shù)的基數(shù)2）進(jìn)行縮放。
常見(jiàn)的浮點(diǎn)格式包括32位（稱為“單精度”）（在C派生的編程語(yǔ)言中為“ float”）和64位（稱為“雙精度”（double））。如IEEE 754標(biāo)準(zhǔn)所定義，一個(gè)32位浮點(diǎn)值包括一個(gè)符號(hào)位，8個(gè)指數(shù)位和23個(gè)尾數(shù)位。64位雙精度數(shù)包括一個(gè)符號(hào)位，11個(gè)指數(shù)位和52個(gè)尾數(shù)位。在本文中，對(duì)（較新的）IEEE 754標(biāo)準(zhǔn)16位浮點(diǎn)半類型感興趣，該類型包括一個(gè)符號(hào)位，5個(gè)指數(shù)位和10個(gè)尾數(shù)位，如圖1所示。

圖1：16位半精度浮點(diǎn)（FP16）表示形式：1個(gè)符號(hào)位，5個(gè)指數(shù)位和10個(gè)尾數(shù)位。
要了解16位精度會(huì)有什么不同，FP16可以表示2 -14和2 15（指數(shù)范圍）之間的2的冪的1024個(gè)值。那是30,720個(gè)值。將此與FP32相比，FP32可以表示2 -126與2 127之間的2的冪的大約800萬(wàn)個(gè)值。大約有20億個(gè)值，相差很大。那么，為什么要使用像FP16這樣的小浮點(diǎn)格式呢？因?yàn)樾阅堋?br /> NVIDIA Tesla P100（基于GP100 GPU）支持2路矢量半精度融合乘加（FMA）指令（操作碼HFMA2），該指令的發(fā)布速度與32位FMA指令相同。這意味著半精度算法的吞吐量是P100上單精度算法的兩倍，是雙精度算法的四倍。具體而言，支持NVLink的P100（SXM2模塊）的半精度精度為21.2 Teraflop / s。憑借如此巨大的性能優(yōu)勢(shì)，值得研究如何使用它。
使用降低的精度時(shí)要記住的一件事是，由于FP16的規(guī)格化范圍較小，因此生成次正規(guī)數(shù)（也稱為非正規(guī)數(shù)）的可能性增加。因此，重要的是，NVIDIA GPU必須以低于正常水平的性能實(shí)現(xiàn)FMA操作。某些處理器不會(huì)這樣做，并且性能可能會(huì)受到影響。（注意：啟用“刷新到零”可能仍然會(huì)帶來(lái)好處。請(qǐng)參閱“ CUDA Pro提示：放心刷新異常”。）
高性能與低精度整數(shù)
浮點(diǎn)數(shù)將高動(dòng)態(tài)范圍與高精度結(jié)合在一起，但是在某些情況下，不需要?jiǎng)討B(tài)范圍，因此整數(shù)可以勝任。甚至在某些應(yīng)用中，正在處理的數(shù)據(jù)的精度也很低，因此可以使用非常低精度的存儲(chǔ)（例如C short或char / byte類型）。

圖2：Tesla P4和P40 GPU中的新DP4A和DP2A指令提供具有32位整數(shù)累加的快速2和4路8位/ 16位整數(shù)矢量點(diǎn)積。
對(duì)于此類應(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是類似的指令，其中A是16位值的2元素向量，而B(niǎo)是8位值的4元素向量，并且DP2A的不同形式為2選擇高字節(jié)或低字節(jié)對(duì)。雙向點(diǎn)積。這些靈活的指令可用于線性代數(shù)計(jì)算，例如矩陣乘法和卷積。對(duì)于實(shí)現(xiàn)用于深度學(xué)習(xí)推理的8位整數(shù)卷積特別強(qiáng)大，這在部署用于圖像分類和目標(biāo)檢測(cè)的深度神經(jīng)網(wǎng)絡(luò)中很常見(jiàn)。圖3顯示了在AlexNet上使用INT8卷積在Tesla P4 GPU上實(shí)現(xiàn)的提高的電源效率。

圖3：與上一代Tesla M4 GPU上的FP32相比，在Tesla P4上使用INT8計(jì)算進(jìn)行深度學(xué)習(xí)推理可大大提高使用AlexNet和其它深度神經(jīng)網(wǎng)絡(luò)進(jìn)行圖像識(shí)別的電源效率。Tesla P4的計(jì)算效率比Arria10 FPGA高出8倍，比Intel Xeon CPU高40倍。（AlexNet，批處理大小= 128，CPU：使用Intel MKL 2017的Intel E5-2690v4，FPGA為Arria10-115.1x M4 / P4節(jié)點(diǎn)，P4板功率為56W，P4 GPU功率為36W，M4板功率為57W， M4 GPU功率為39W，Perf / W圖表使用GPU功率。）
DP4A計(jì)算總計(jì)八個(gè)整數(shù)運(yùn)算的等效項(xiàng)，DP2A計(jì)算四個(gè)整數(shù)運(yùn)算。這樣，Tesla P40（基于GP102）的峰值整數(shù)吞吐量為47 TOP / s（每秒Tera操作）。
DP4A的一個(gè)示例應(yīng)用是通常在射電望遠(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)計(jì)算。GPU已用于生產(chǎn)射電天文學(xué)互相關(guān)，但他們通常使用FP32計(jì)算。DP4A的引入保證了該計(jì)算的更高功率效率。圖4顯示了修改a的結(jié)果互相關(guān)代碼以使用DP4A，從而在具有默認(rèn)時(shí)鐘的Tesla P40 GPU上效率提高了4.5倍（與P40上的FP32計(jì)算相比）在GPU時(shí)鐘上設(shè)置了6.4倍的提高，從而降低了溫度（從而降低了泄漏電流） ）。總體而言，新代碼比上一代Tesla M40 GPU上的FP32互相關(guān)效率高近12倍（來(lái)源：Kate Clark）。

圖4：與FP32計(jì)算相比，INT8矢量點(diǎn)積（DP4A）在很大程度上提高了射電天文互相關(guān)的效率。
Pascal GPU上的混合精度性能
半精度（FP16）格式對(duì)于GPU來(lái)說(shuō)并不是新事物。實(shí)際上，FP16作為存儲(chǔ)格式已經(jīng)在NVIDIA GPU上得到了多年的支持，主要用于降低精度的浮點(diǎn)紋理存儲(chǔ)和過(guò)濾以及其它特殊用途。Pascal GPU體系結(jié)構(gòu)實(shí)現(xiàn)了通用的IEEE 754 FP16算法。如下表所示，Tesla P100（GP100）上全速支持高性能FP16，而其它Pascal GPU（GP102，GP104和GP106）則以較低的吞吐量（類似于雙精度）支持。
GP102-GP106支持8位和16位DP4A和DP2A點(diǎn)產(chǎn)品指令，但GP100不支持。表1顯示了基于Pascal的Tesla GPU上不同數(shù)字指令的算術(shù)吞吐量。

表1：基于Pascal的Tesla GPU的半，單精度和雙精度融合乘法加法指令以及8位和16位矢量點(diǎn)乘積指令的峰值算術(shù)吞吐量。（Boost時(shí)鐘速率用于計(jì)算峰值吞吐量。TFLOP / s：每秒Tera浮點(diǎn)運(yùn)算。TIOP / s：每秒Tera整數(shù)運(yùn)算。）
NVIDIA庫(kù)的混合精度編程
從應(yīng)用程序的混合精度中受益的最簡(jiǎn)單方法是利用NVIDIA GPU庫(kù)中對(duì)FP16和INT8計(jì)算的支持。NVIDIA SDK的密鑰庫(kù)支持計(jì)算和存儲(chǔ)的多種精度。
表2顯示了關(guān)鍵CUDA庫(kù)以及PTX匯編和CUDA C / C ++內(nèi)部函數(shù)中對(duì)FP16和INT8的當(dāng)前支持。

表2：CUDA 8 FP16和INT8 API和庫(kù)支持。
神經(jīng)網(wǎng)絡(luò)
cuDNN是用于訓(xùn)練和部署深度神經(jīng)網(wǎng)絡(luò)的原始例程庫(kù)。cuDNN 5.0包括對(duì)前向卷積的FP16支持，并增加了對(duì)FP16后向卷積的支持。庫(kù)中的所有其它例程均受內(nèi)存限制，因此FP16計(jì)算對(duì)性能無(wú)益。因此，這些例程使用FP32計(jì)算，但支持FP16數(shù)據(jù)輸入和輸出。cuDNN 6將增加對(duì)INT8推理卷積的支持。
TensorRT
TensorRT是用于深度學(xué)習(xí)應(yīng)用程序生產(chǎn)部署的高性能深度學(xué)習(xí)推理引擎，該引擎自動(dòng)優(yōu)化訓(xùn)練有素的神經(jīng)網(wǎng)絡(luò)以實(shí)現(xiàn)運(yùn)行時(shí)性能。TensorRT v1支持FP16進(jìn)行推理卷積，而v2支持INT8進(jìn)行推理卷積。
cuBlas
cuBLAS是用于密集線性代數(shù)的GPU庫(kù)，它是BLAS（基本線性代數(shù)子例程）的實(shí)現(xiàn)。cuBLAS支持幾種矩陣矩陣乘法例程中的混合精度。cublasHgemm是FP16密集矩陣矩陣乘法例程，使用FP16進(jìn)行計(jì)算以及輸入和輸出。cublasSgemmEx（）在FP32中計(jì)算，但是輸入數(shù)據(jù)可以是FP32，FP16或INT8，輸出可以是FP32或FP16。cublasGemm（）是CUDA 8中的新例程，它允許指定計(jì)算精度，包括INT8計(jì)算（使用DP4A）。
將根據(jù)需求增加對(duì)更多具有FP16計(jì)算和/或存儲(chǔ)功能的BLAS 3級(jí)例程的支持。1級(jí)和2級(jí)BLAS例程受內(nèi)存限制，因此降低精度的計(jì)算是無(wú)益的。
傅立葉變換
cuFFT是在CUDA中實(shí)現(xiàn)的流行的快速傅立葉變換庫(kù)。從CUDA 7.5開(kāi)始，cuFFT支持FP16的單GPU FFT計(jì)算和存儲(chǔ)。FP16 FFT的速度比FP32快2倍。FP16計(jì)算需要具有Compute Capability 5.3或更高版本（Maxwell架構(gòu)）的GPU。大小目前限制為2的冪，并且不支持R2C或C2R轉(zhuǎn)換的實(shí)部上的跨步。
cuSPARSE
cuSPARSE是用于稀疏矩陣的GPU加速線性代數(shù)例程庫(kù)。cuSPARSE支持FP16的多個(gè)例程存儲(chǔ)（cusparseXtcsrmv（），cusparseCsrsv_analysisEx（），cusparseCsrsv_solveEx（），cusparseScsr2cscEx（）和cusparseCsrilu0Ex（））。正在研究cuSPARSE的FP16計(jì)算。
在CUDA代碼中使用混合精度
對(duì)于自定義CUDA C ++內(nèi)核的開(kāi)發(fā)人員和Thrust并行算法庫(kù)的用戶，CUDA提供了從FP16和INT8計(jì)算，存儲(chǔ)和I / O中獲得最大收益所需的類型定義和API。
FP16類型和內(nèi)在函數(shù)
對(duì)于FP16，CUDA在CUDA包含路徑中包含的標(biāo)頭“ cuda_fp16.h”中定義了“ half”和“ half2”類型。該頭文件還定義了一套完整的內(nèi)部函數(shù)，用于對(duì)“半”數(shù)據(jù)進(jìn)行操作。例如，下面顯示了標(biāo)量FP16加法函數(shù)“ hadd（）”和2路矢量FP16加法函數(shù)“ hadd2（）”的聲明。
device __half __hadd（const __half a，const __half b）;
device __half2 __hadd2（const __half2 a，const __half2 b）;
cuda_fp16.h定義了一套完整的半精度內(nèi)在函數(shù)，用于算術(shù)，比較，轉(zhuǎn)換和數(shù)據(jù)移動(dòng)以及其它數(shù)學(xué)函數(shù)。所有這些都在CUDA Math API文檔中進(jìn)行了描述。
在可能的情況下使用“ half2”向量類型和內(nèi)在函數(shù)來(lái)實(shí)現(xiàn)最高吞吐量。GPU硬件算術(shù)指令一次對(duì)2個(gè)FP16值進(jìn)行運(yùn)算，并打包在32位寄存器中。表1中的峰值吞吐率假設(shè)為“ half2”矢量計(jì)算。如果使用標(biāo)量“半”指令，則可以達(dá)到峰值吞吐量的50％。同樣，在從FP16陣列加載和存儲(chǔ)到FP16陣列時(shí)要實(shí)現(xiàn)最大帶寬，需要向量訪問(wèn)“ half2”數(shù)據(jù)。理想情況下，可以通過(guò)加載和存儲(chǔ)“ float2”或“ float4”類型并強(qiáng)制轉(zhuǎn)換為“ half2”或從“ half2”進(jìn)行轉(zhuǎn)換，來(lái)進(jìn)一步矢量化負(fù)載以實(shí)現(xiàn)更高的帶寬。
以下示例代碼演示了如何使用CUDA __hfma() （半精度融合乘加）和其它內(nèi)在函數(shù)來(lái)計(jì)算半精度AXPY（A * X + Y）。該示例的完整代碼在Github上可用，并且顯示了如何在主機(jī)上初始化半精度數(shù)組。重要的是，當(dāng)開(kāi)始使用half類型時(shí)，可能需要 在主機(jī)端代碼中的half 和float值之間進(jìn)行轉(zhuǎn)換。包括一些快速的CPU類型轉(zhuǎn)換例程（有關(guān)完整源代碼，請(qǐng)參見(jiàn)相關(guān)的Gist）。在此示例中，使用了Giesen的一些代碼。
全球
void haxpy（int n，half a，const half * x，half * y）
{
整數(shù)開(kāi)始= 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 =開(kāi)始； i <n2； i + =步幅）
y2 [i] = hfma2（ halves2half2（a，a），x2 [i]，y2 [i]）;

//第一個(gè)線程處理奇數(shù)數(shù)組的單例

如果（開(kā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]））;
}
＃萬(wàn)一
}
整數(shù)點(diǎn)乘本征
CUDA在標(biāo)頭“ sm_61_intrinsics.h”（sm_61是與GP102，GP104和GP106對(duì)應(yīng)的SM架構(gòu)）中為8位和16位點(diǎn)乘積（先前描述的DP4A和DP2A指令）定義了內(nèi)部函數(shù)。）。為方便起見(jiàn)，DP4A內(nèi)部函數(shù)有int和char4版本，有符號(hào)和無(wú)符號(hào)兩種形式：
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的四個(gè)向量元素被打包到32位字的四個(gè)相應(yīng)字節(jié)中。char4 /uchar4版本使用帶有顯式字段的CUDA的struct類型，而打包在int版本中是隱式的。
如前所述，DP2A具有“高”和“低”版本，分別用于選擇輸入B的高或低兩個(gè)字節(jié)。
//通用[_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 ）;
請(qǐng)記住，DP2A和DP4A在基于GP102，GP104和GP106 GPU的Tesla，GeForce和Quadro加速器上可用，但在基于Tesla P100（基于GP100 GPU）上不可用。

總結(jié)

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

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