CUDA7.5开始,支持16位浮点数的存储和计算,添加了half
和half2
两种数据类型,并内置了用来操作它们的函数。16位“半精度”浮点类型在应用程序中很有用,这些应用程序可以处理更大的数据集,也可以通过存储和操作更低精度的数据来获得性能。例如对一些规模比较大的神经网络模型来说,它们可能会受限于有限的GPU存储;一些信号处理内核(signal processing kernels 如FFTs)受限于存储的带宽。
许多应用都会得益于使用半精度来储存数据,然后用32位的单精度来处理这些数据。Pascal架构的GPU将会全面支持这种“混合精度”的计算,使用FP16计算将会获得比FP32和FP64更高的吞吐量。
CUDA7.5提供了3个FP16的特性:
1.一个新的头文件 cuda_fp16.h
:定义了half
和half2
两个数据类型,__half2float()
和__float2half()
两个函数,这两个函数用于FP16和FP32之间的相互转换;
2.一个新的程序 cublasSgemmEx()
:使用FP16数据作为输入,执行混合精度的矩阵乘法,也可在32位精度下执行所有的计算。这使得GPU上能够计算比原来矩阵大两倍的乘法;
3.函数支持 :cuda_fp16.h
定义了16位的计算和比较。cuCLAS中也包含了cublasHgemm()
(半精度矩阵乘)程序。
NVIDIA GPU实现了IEEE 754浮点数标准[2008],半精度浮点数的定义为:
符号:1 bit
指数位:5 bits
精度位:10 bits
半精度数的范围大约是5.96×10^-8~6.55×10^4。half2结构在一个32位字里存储两个half值: