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

Question

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

Answer 1

有几点需要提前注意:

引用半精度intrinsics .

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

请注意，计算能力为 5.2 及以下的设备本身不支持半精度算术。这意味着要执行的任何算术运算都必须在某些受支持的类型上完成，例如 float .计算能力为 5.3 的设备(目前为 Tegra TX1)和可能的 future 设备将支持“ native ”半精度算术运算，但这些目前通过诸如 __hmul 之类的内在函数公开。 .内在像 __hmul在不支持 native 操作的设备中将是未定义的。

你应该包括 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 位向量加载可以一次加载 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 .