有几点需要提前注意：
- 参考半精度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。