CUDA 混合精度编程 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可以以FP吞吐量的32倍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位浮点值包括一个符号位数位和23个尾数位。64位双精度数包括符号位、11个指数位和52个尾数位。在本文中，对(更新)IEEE 754标准16位浮点半类型感兴趣，包括符号位、5个指数位和10个尾数位，如图1所示。 
图1:16半精度浮点（FP16)表示形式:1个符号位，5个指数位，10个尾数位。 了解16位精度的差异，FP16可以表示2 -14和2 15米(指数范围)之间的1024个值。30720个值。将此与FP32相比，FP32可以表示2 -126与2 127之间的200万米左右。大约有20亿个值，差别很大。那么，为什么要使用像FP16这浮点格式呢？因为性能。 NVIDIA Tesla P100（基于GP100 GPU）支持2路矢量半精度融合（FMA）指令(操作码)HFMA2)指令的发布速度和32位FMA指令相同。这意味着半精度算法的吞吐量是P单精度算法的两倍是双精度算法的四倍。具体来说，支持NVLink的P100（SXM2模块)半精度为21.2 Teraflop / s。如此巨大的性能优势值得研究如何使用它。 使用降低精度时要记住的一件事是，因为FP16规格化范围小，增加了生成次正规数(又称非正规数)的可能性。因此，重要的是，NVIDIA GPU性能必须低于正常水平FMA操作。有些处理器不会这样做，性能可能会受到影响。(注:使用刷新到零仍可能带来好处。请参阅“ CUDA Pro提示：放心刷新异常”。） 高性能、低精度整数 浮点数将高动态范围和高精度结合在一起，但在某些情况下，不需要动态范围，因此整数胜任。即使在某些应用程序中，处理数据的精度也很低，因此可以使用非常低精度的存储（例如C short或char / byte类型）。
图2：Tesla P4和P40 GPU中的新DP4A和DP2A32位整数累加的快速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。DP2AA是16位值的2元素向量，B是8位值的4元素向量DP2A不同形式的2选择高字节或低字节对。双向点积。这些灵活的指令可用于线性代数计算，如矩阵乘法和卷积。对于实现深度学习推理的8位整数卷积特别强大，这在深度神经网络中很常见，用于图像分类和目标检测。图3显示了在AlexNet上使用INT8卷积在Tesla P4 GPU提高电源效率。
图3：与上一代Tesla M4 GPU上的FP32相比，在Tesla P4上使用INT深度学习推理计算可以大大提高使用率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计算比较)在GPU6设置在时钟上.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（）”的声明。 __half __hadd（const __half a，const __half b）; __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 > = 530 int n2 = n / 2; half2 * x2 =（half2 ）x， y2 =（half2 *）y;
for（int i =开始； i <n2； i + =步幅） y2 [i] = 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] = 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版本，有符号和无符号两种形式： int __dp4a（int srcA，int srcB，int c）;int __dp4a （int srcA ，int srcB ，int c ）; int __dp4a（char4 srcA，char4 srcB，int c）;int __dp4a （char4 srcA ，char4 srcB ，int c ）; unsigned int __dp4a（unsigned int srcA，unsigned int srcB，unsigned int c）;unsigned int __dp4a （unsigned int srcA ，unsigned int srcB ，unsigned int c ）; 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] int __dp2a_lo（int srcA，int srcB，int c）;int __dp2a_lo （int srcA ，int srcB ，int c ）; 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] int __dp2a_lo（short2 srcA，char4 srcB，int c）;int __dp2a_lo （short2 srcA ，char4 srcB ，int c ）; unsigned int __dp2a_lo（ushort2 srcA，uchar4 srcB，unsigned int c）;unsigned int __dp2a_lo （ushort2 srcA ，uchar4 srcB ，unsigned int c ）;
//通用[_hi]//通用[_hi] int __dp2a_hi（int srcA，int srcB，int c）;int __dp2a_hi （int srcA ，int srcB ，int c ）; 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] int __dp2a_hi（short2 srcA，char4 srcB，int c）;int __dp2a_hi （short2 srcA ，char4 srcB ，int c ）; 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）上不可用。