cuda - 谁能提供示例代码来演示cuda中16位浮点的使用？

Question

Cuda 7.5支持16位浮点变量。
任何人都可以提供示例代码来演示其用法吗？

Answer 1

首先要注意几件事:

指的是半精度intrinsics。

注意，大多数或所有这些内在函数仅在设备代码中受支持。 (但是，@ njuffa创建了一组主机可用的转换函数here)

请注意，计算能力5.2及以下的设备本身不支持半精度算术。这意味着必须对某种受支持的类型(例如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一致。

外卖如下:

在cc5.2及更低版本的设备上，half数据类型可能仍然有用，但主要是作为存储优化(以及相关地，也许是内存带宽优化，因为例如，给定的128位 vector 加载可能会加载 8 half数量)。例如，如果您有一个大型神经网络，并且已确定权重可以容忍以半精度量存储(从而使存储密度增加一倍，或者大约可以将神经网络的大小增加一倍)。 GPU的存储空间)，则可以将神经网络权重存储为半精度。然后，当您需要执行前向传递(推理)或后向传递(训练)时，可以从内存中加载权重，将其即时(使用内在函数)转换为float数量，执行必要的操作(可能包括由于训练而调整权重)，然后(如有必要)将权重再次存储为half数量。

对于cc5.3及更高版本的设备，如果算法允许，则可以执行与上述类似的操作，但无需转换为float(可能还可以转换回half)，而是将所有数据保留在half表示中，并直接进行必要的算术运算(例如使用__hmul或__hadd内部函数)。

尽管我在这里没有演示，但是 half数据类型在主机代码中“可用”。就是说，我的意思是您可以为该类型的项目分配存储空间，并执行例如 cudaMemcpy操作就可以了。但是宿主代码对 half数据类型一无所知(例如，如何对其进行算术，打印输出或进行类型转换)，并且内在函数在宿主代码中不可用。因此，如果您愿意(当然可以存储一组神经网络权重)，则可以为大量的 half数据类型分配存储空间，但是只能从设备代码(而不是主机代码)轻松地直接操作该数据。

其他一些评论:

CUBLAS库implements a matrix-matrix multiply设计为直接处理half数据。上面的描述应该对不同设备类型(即计算能力)在“幕后”可能发生的情况有所了解。

有关在推力中使用half的一个相关问题是here。