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显示的功能。