任何人都可以提供示例代码,演示在cuda中使用16位浮点数吗？

Question

Cuda 7.5支持16位浮点变量.任何人都可以提供示例代码来证明它的使用吗？

Answer 1

预先有几点需要注意:

1. 请参阅半精度内在函数.
2. 请注意,大多数或所有这些内在函数仅在设备代码中受支持.(然而,@njuffa已经创建了一组的主机可用转换函数这里)
3. 请注意,5.2及以下计算能力的设备本身不支持半精度算术.这意味着要执行的任何算术运算都必须在某些受支持的类型上完成,例如float.计算能力设备5.3(目前的Tegra TX1)和可能未来的设备将支持"本机"半精度算术运算,但这些设备目前通过这样的内在函数暴露__hmul.在不支持本机操作的设备中,__hmul将不会定义内在类似内容.
4. 您应该包含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一致.

外卖如下:

1. 在cc5.2及更低版本的设备上,half数据类型可能仍然有用,但主要是作为存储优化(并且相关地,可能是存储器带宽优化,因为例如给定的128位向量负载可能加载8 half数量一次).例如,如果你有一个庞大的神经网络,并且你已经确定权重可以容忍被存储为半精度数量(从而使存储密度加倍,或者大约加倍可以在GPU的存储空间,然后您可以将神经网络权重存储为半精度.然后,当您需要执行正向传递(推理)或向后传递(训练)时,您可以从内存中加载权重,将它们即时(使用内在函数)转换为float数量,执行必要的操作(也许包括根据训练调整重量,然后(如果需要)再次将重量存储为half数量.
2. 对于cc5.3和未来的设备,如果算法能够容忍它,则可以执行与上面类似的操作,但不能转换为float(并且可能返回half),而是将所有数据保留在half表示中,并且执行必要的操作直接算术(使用eg __hmul或__haddintrinsics).

虽然我没有在这里演示,但half数据类型在主机代码中是"可用的".通过这种方式,我的意思是您可以为该类型的项目分配存储空间,并对其执行cudaMemcpy操作.但是主机代码对half数据类型一无所知(例如,如何对其进行算术运算,或将其打印出来,或进行类型转换),并且内部函数在主机代码中不可用.因此,half如果您希望(可能存储一组神经网络权重),您当然可以为大量数据类型分配存储空间,但您只能从设备代码中直接操作该数据,而不是主机代码.

还有一些评论:

1. CUBLAS库实现了矩阵 - 矩阵乘法,可直接处理half数据.上面的描述应该可以为不同的设备类型(即计算能力)提供一些有关"引擎盖下"的信息.

2. 关于half推力使用的相关问题在这里.