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

Posted 2023-03-23

【中文标题】任何人都可以提供示例代码来演示在 cuda 中使用 16 位浮点数吗？

【英文标题】：Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?

【发布时间】：2022-01-07 15:08:58

【问题描述】：

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

【问题讨论】：

在 cuda 中使用半浮点数（16 位浮点数）！！！

为什么会有人想要使用 64 位双精度以外的任何东西进行浮点计算？什么用例适合精度较低？

当您在使用相同内存的程序中需要更多变量时。此外，有些情况下 2 字节浮点数的精度就足够了。并非所有情况都相同

@duffymo - 16 位适用于许多应用程序（例如计算机视觉）。它也快了很多。目前也没有单个 GPU 可以访问超过 12GB，因此“购买更多”并不能解决所有问题。

@duffymo 在精度不如其他事物重要的地方：计算机图形学、进化算法……很多东西。与双打相比，缓存行上的一半是 4 倍，或者可以从 gmem 一次性获取。因此，性能增益和带宽增益。从 200k+ 会员那里看到这样的单边 cmets 真的很奇怪。

【参考方案1】：

有几点需要提前注意：

参考半精度intrinsics。 请注意，其中许多内在函数仅在设备代码中受支持。但是，在最近/当前的 CUDA 版本中，主机和设备代码都支持许多/大部分 conversion intrinsics。 （而且，@njuffa 创建了一组主机可用的转换函数here）因此，即使下面的代码示例显示了设备代码中的转换，相同类型的转换和内在函数（half->float、float->half ) 以相同的方式在宿主代码中使用和支持。 请注意，计算能力为 5.2 及以下的设备本机不支持半精度算术。这意味着要执行的任何算术运算都必须在某些受支持的类型上完成，例如float。计算能力为 5.3 的设备（目前为 Tegra TX1）和可能的未来设备将支持“本机”半精度算术运算，但这些目前通过 __hmul 等内在函数公开。在不支持本机操作的设备中，像 __hmul 这样的内部函数将未定义。 您应该在您打算在设备代码中使用这些类型和内在函数的任何文件中包含cuda_fp16.h。 half2 数据类型（向量类型）确实是压缩/批量半存储（例如向量或矩阵）的首选形式，因此您可能需要使用相关的half2 转换函数。

考虑到以上几点，下面是一个简单的代码，它采用一组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（也可能返回到 @ 987654342@)，而是将所有数据保留在half 表示中，并直接进行必要的算术运算（例如使用__hmul 或__hadd 内在函数）。

虽然我没有在这里演示，half 数据类型在主机代码中是“可用的”。我的意思是，您可以为该类型的项目分配存储空间，并执行例如cudaMemcpy 对其进行操作。但是主机代码对half 数据类型一无所知（例如，如何对其进行算术运算或将其打印出来），例如the arithmetic intrinsics 在主机代码中不可用。因此，您当然可以为大量 half（或者可能是 half2）数据类型分配存储空间（也许存储一组神经网络权重），但您只能使用任何从设备代码简化，而不是主机代码。

还有几个cmets：

CUBLAS 库implements a matrix-matrix multiply 旨在直接处理half 数据。上面的描述应该对不同设备类型（即计算能力）的“幕后”可能发生的事情提供一些见解。

关于在推力中使用half 的相关问题是here。

【讨论】：

所以在当前上下文中，我们不能将 cudaMalloc 用于“半”变量吗？或 cudaMemcpy 将主机中的半变量（使用库创建）复制到设备半变量？不是吗？

我在回答中特别说过，您可以将 cudaMalloc 和 cudaMemcpy 与 half 数据类型一起使用