许多 CUDA 内核受带宽限制，新硬件中 FLOPS 与带宽的比例不断增加，导致更多内核受带宽限制。因此，采取措施缓解代码中的带宽瓶颈非常重要。在这篇文章中，我将向您展示如何在 CUDA C++ 中使用向量加载和存储来帮助提高带宽利用率，同时减少执行指令的数量。
首先，让我们来看一下以下简单的内存复制内核。
__global__ void device_copy_scalar_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
d_out[i] = d_in[i];
}
}
void device_copy_scalar(int* d_in, int* d_out, int N)
{
int threads = 256;
int blocks = min((N + threads-1) / threads, MAX_BLOCKS);
device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
在此代码中，我使用了网格步长循环，如之前的 CUDA Pro Tip 帖子中所述。图 1 显示了内核的吞吐量（单位：GB/s）与复制大小的关系。
我们可以使用 CUDA 工具包中包含的 cuobjdump 工具来检查此内核的汇编。
%> cuobjdump -sass executable
标量复制内核的 SASS 包括以下指令：
...
LDG.E R3, desc[UR6][R2.64] ;
...
STG.E desc[UR6][R4.64], R3 ;
...
LDG.E 和 STG.E 指令分别从全局内存中加载和存储 32 位。
我们可以通过使用矢量化加载和存储指令 LDG.E 来提高此操作的性能。{64,128} 和 STG.E.{64,128}。这些操作也会加载和存储数据，但会以 64 或 128 位的宽度进行。使用矢量化加载可以减少指令总数、降低延迟并提高带宽利用率。使用矢量化加载的最简单方法是使用 CUDA C++ 标准头文件中定义的矢量数据类型，例如 int2、int4 或 float2、float4。这些类型表示多个值打包到单个数据单元中。您可以通过 C++ 中的类型转换轻松使用这些类型。例如，在 C++ 中，您可以使用 reinterpret_cast<int2*>(d_in) 将 int 指针 d_in 重铸为 int2 指针，该指针将一对“int”值视为一个单元。在 C99 中，您可以使用 casting 运算符来执行相同的操作： (int2* (d_in)) 。
取消引用这些指针将导致编译器生成矢量化指令。
int2* int2Ptr = reinterpret_cast<int2*>(d_in);
int2 data = int2Ptr[0]; // Loads the first two int values as one int2
但是，有一个重要的注意事项：这些指令需要对齐的数据。设备分配的内存会自动对齐到数据类型大小的倍数，但如果偏移指针，则偏移量也必须对齐。例如，reinterpret_cast<int2*>(d_in+1) 是无效的，因为 d_in+1 未对齐到 int2 的倍数。
如果使用“对齐”偏移量，则可以安全地偏移数组，如 reinterpret_cast<int2*>(d_in+2) 中所示。
您还可以使用结构生成矢量化负载，只要结构的大小是 2 字节的幂即可。
struct Foo {int a, int b, double c}; // 16 bytes in size
Foo *x, *y;
…
x[i]=y[i];
非 2 的幂次大小可能会导致内存对齐效率低下，从而可能导致编译器自动添加填充以在典型架构上正确对齐数据。
现在，我们已经了解了如何生成矢量化指令，接下来让我们修改内存复制内核以使用矢量加载。
__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
}
// in only one thread, process final element (if there is one)
if (idx==N/2 && N%2==1)
d_out[N-1] = d_in[N-1];
}
void device_copy_vector2(int* d_in, int* d_out, int n) {
threads = 256;
blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS);
device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
此内核只有少数更改。首先，由于每次迭代处理两个元素，因此循环现在只执行 N/2 次。第二，我们在复制中使用上述的投射技术。第三，我们处理 N 不能被 2 整除时可能出现的任何剩余元素。最后，我们启动了与标量内核相同数量的一半线程。
检查 SASS 后，我们看到以下更改：
...
LDG.E.64 R2, desc[UR4][R2.64] ;
...
STG.E.64 desc[UR4][R4.64], R2 ;
...
请注意，现在编译器生成了 LDG.E.64 和 STG.E.64。所有其他指令均相同。但是，需要注意的是，由于循环只执行 N/2 次，因此执行的指令数量将减半。在指令受限或延迟受限的内核中，指令数量的 2 倍提升非常重要。
我们还可以编写 vector4 版本的复制内核。
__global__ void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
}
// in only one thread, process final elements (if there are any)
int remainder = N%4;
if (idx==N/4 && remainder!=0) {
while(remainder) {
int idx = N - remainder--;
d_out[idx] = d_in[idx];
}
}
}
void device_copy_vector4(int* d_in, int* d_out, int N) {
int threads = 256;
int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);
device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
相应的 SASS 更改包括：
...
LDG.E.128 R4, desc[UR4][R4.64] ;
...
STG.E.128 desc[UR4][R8.64], R4 ;
...
在这里，我们可以看到生成的 LDG.E.128 和 STG.E.128。此版本的代码将指令数量减少了四倍。您可以在图 2 中看到所有三个内核的整体性能。
在几乎所有情况下，矢量化加载优于标量加载。不过，请注意，使用矢量化加载会增加寄存器压力并降低整体并行性。因此，如果内核已经受到寄存器限制或并行度非常低，您可能需要坚持使用标量加载。此外，如前所述，如果指针未对齐或数据类型大小（以字节为单位）不是 2 的幂，则无法使用矢量化加载。
矢量化负载是基本的 CUDA 优化，在可能的情况下，您应该使用矢量化负载，因为它们可以增加带宽、减少指令数量并降低延迟。在这篇文章中，我展示了如何通过相对较少的更改轻松地将矢量化负载集成到现有内核中。
此博客的版本于 2013 年 12 月 4 日发布。该示例已更新，以反映当前 GPU 上的行为。
