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）上不可用。