Cuda 7.5支持16位浮点变量。
任何人都可以提供示例代码来演示其用法吗?
最佳答案
首先要注意几件事:
float
)执行任何算术运算。计算能力为5.3的设备(当前为Tegra TX1)以及可能的将来的设备将支持“本机”半精度算术运算,但是这些当前通过诸如__hmul
之类的内在函数公开。在不支持本机操作的设备中,像__hmul
这样的内部函数将是未定义的。 cuda_fp16.h
。 考虑到以上几点,这是一个简单的代码,它采用一组
float
数量,将其转换为half
数量,并按比例因子进行缩放:$ cat t924.cu
#include <stdio.h>
#include <cuda_fp16.h>
#define DSIZE 4
#define SCF 0.5f
#define nTPB 256
__global__ void half_scale_kernel(float *din, float *dout, int dsize){
int idx = threadIdx.x+blockDim.x*blockIdx.x;
if (idx < dsize){
half scf = __float2half(SCF);
half kin = __float2half(din[idx]);
half kout;
#if __CUDA_ARCH__ >= 530
kout = __hmul(kin, scf);
#else
kout = __float2half(__half2float(kin)*__half2float(scf));
#endif
dout[idx] = __half2float(kout);
}
}
int main(){
float *hin, *hout, *din, *dout;
hin = (float *)malloc(DSIZE*sizeof(float));
hout = (float *)malloc(DSIZE*sizeof(float));
for (int i = 0; i < DSIZE; i++) hin[i] = i;
cudaMalloc(&din, DSIZE*sizeof(float));
cudaMalloc(&dout, DSIZE*sizeof(float));
cudaMemcpy(din, hin, DSIZE*sizeof(float), cudaMemcpyHostToDevice);
half_scale_kernel<<<(DSIZE+nTPB-1)/nTPB,nTPB>>>(din, dout, DSIZE);
cudaMemcpy(hout, dout, DSIZE*sizeof(float), cudaMemcpyDeviceToHost);
for (int i = 0; i < DSIZE; i++) printf("%f\n", hout[i]);
return 0;
}
$ nvcc -o t924 t924.cu
$ cuda-memcheck ./t924
========= CUDA-MEMCHECK
0.000000
0.500000
1.000000
1.500000
========= ERROR SUMMARY: 0 errors
$
如果您研究上述代码,则会注意到,除了cc5.3和更高版本的设备外,该算法是作为常规的
float
操作完成的。这与上面的注释3一致。外卖如下:
half
数据类型可能仍然有用,但主要是作为存储优化(以及相关地,也许是内存带宽优化,因为例如,给定的128位 vector 加载可能会加载 8 half
数量)。例如,如果您有一个大型神经网络,并且已确定权重可以容忍以半精度量存储(从而使存储密度增加一倍,或者大约可以将神经网络的大小增加一倍)。 GPU的存储空间),则可以将神经网络权重存储为半精度。然后,当您需要执行前向传递(推理)或后向传递(训练)时,可以从内存中加载权重,将其即时(使用内在函数)转换为float
数量,执行必要的操作(可能包括由于训练而调整权重),然后(如有必要)将权重再次存储为half
数量。 float
(可能还可以转换回half
),而是将所有数据保留在half
表示中,并直接进行必要的算术运算(例如使用__hmul
或__hadd
内部函数)。 尽管我在这里没有演示,但是
half
数据类型在主机代码中“可用”。就是说,我的意思是您可以为该类型的项目分配存储空间,并执行例如cudaMemcpy
操作就可以了。但是宿主代码对half
数据类型一无所知(例如,如何对其进行算术,打印输出或进行类型转换),并且内在函数在宿主代码中不可用。因此,如果您愿意(当然可以存储一组神经网络权重),则可以为大量的half
数据类型分配存储空间,但是只能从设备代码(而不是主机代码)轻松地直接操作该数据。其他一些评论:
half
数据。上面的描述应该对不同设备类型(即计算能力)在“幕后”可能发生的情况有所了解。 half
的一个相关问题是here。 关于cuda - 谁能提供示例代码来演示cuda中16位浮点的使用?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/32559619/