CUDA Pro：通过向量化内存访问提高性能
CUDA Pro：通过向量化内存访问提高性能
许多CUDA内核受带宽限制，而新硬件中触发器与带宽的比率不断提高，导致带宽受限制的内核更多。这使得采取措施减轻代码中的带宽瓶颈非常重要。本文将展示如何在CUDA C / C ++中使用向量加载和存储，以帮助提高带宽利用率，同时减少已执行指令的数量。
本文参考文章链接
https://developer.nvidia.com/blog/cuda-pro-tip-increase-performance-with-vectorized-memory-access/
从以下简单的内存复制内核开始。
__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 = 128;
int blocks = min((N + threads-1) / threads, MAX_BLOCKS);
device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
代码使用的是网格跨度循环。图1显示了内核吞吐量（GB / s）与副本大小的关系。
图1：复制带宽与复制大小的关系。
可以使用CUDA Toolkit 附带的cuobjdump工具检查该内核的程序集。
％> cuobjdump -sass可执行文件
标量复制内核主体的SASS如下：
/ * 0058 * / IMAD R6.CC，R0，R9，c [0x0] [0x140]
/ * 0060 * / IMAD.HI.X R7，R0，R9，c [0x0] [0x144]
/ * 0068 * / IMAD R4.CC，R0，R9，c [0x0] [0x148]
/ * 0070 * / LD.E R2，[R6]
/ * 0078 * / IMAD.HI.X R5，R0，R9，c [0x0] [0x14c]
/ * 0090 * / ST.E [R4]，R2
可以看到总共六个与复制操作相关的指令。四个IMAD指令计算加载和存储地址和LD.E与ST.E负载位和32位来自这些地址存储。
可以使用向量化的加载和存储指令LD.E.{64,128}和来提高此操作的性能ST.E.{64,128}。这些操作也可以加载和存储数据，但可以64位或128位宽度进行加载和存储。使用矢量化负载减少了指令总数，减少了等待时间，并提高了带宽利用率。
使用矢量载荷的最简单的方法是使用在CUDA C / C ++标准头中定义的向量的数据类型，如int2，int4，或 float2。可以通过C / C ++中的类型转换轻松地使用这些类型。例如，在C ++可以重铸int指针d_in到一个int2使用指针reinterpret_cast<int2*>(d_in)。在C99中，可以使用强制转换运算符做相同的事情：(int2*(d_in))。
取消引用那些指针将导致编译器生成矢量化指令。但是，有一个重要警告：这些指令需要对齐的数据。设备分配的内存会自动对齐到数据类型大小的倍数，但是如果偏移指针，则偏移也必须对齐。例如reinterpret_cast<int2*>(d_in+1)，无效是因为d_in+1未与对齐sizeof(int2)。
如果使用“对齐”偏移量，则可以安全地偏移数组，如 reinterpret_cast<int2*>(d_in+2)中所示。也可以使用结构生成矢量化载荷，只要该结构的大小为2个字节即可。
struct Foo {int a，int b，double c}; // 16个字节
Foo * x，* y;
…
x [i] = y [i];
既然已经看到了如何生成向量化指令，那么让修改内存复制内核以使用向量加载。
__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 = 128;
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，看到以下内容。
/ * 0088 * / IMAD R10.CC，R3，R5，c [0x0] [0x140]
/ * 0090 * / IMAD.HI.X R11，R3，R5，c [0x0] [0x144]
/ * 0098 * / IMAD R8.CC，R3，R5，c [0x0] [0x148]
/ * 00a0 * / LD.E.64 R6，[R10]
/ * 00a8 * / IMAD.HI.X R9，R3，R5，c [0x0] [0x14c]
/ * 00c8 * / ST.E.64 [R8]，R6
编译器生成LD.E.64和ST.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 = 128;
int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);
device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}
相应的SASS是：
/*0090*/ IMAD R10.CC, R3, R13, c[0x0][0x140]
/*0098*/ IMAD.HI.X R11, R3, R13, c[0x0][0x144]
/*00a0*/ IMAD R8.CC, R3, R13, c[0x0][0x148]
/*00a8*/ LD.E.128 R4, [R10]
/*00b0*/ IMAD.HI.X R9, R3, R13, c[0x0][0x14c]
/*00d0*/ ST.E.128 [R8], R4
在这里可以看到生成的LD.E.128和ST.E.128。此版本的代码将指令数减少了4倍。可以在图2中看到所有3个内核的整体性能。
图2：矢量化内核的复制带宽与复制大小的关系。
在几乎所有情况下，矢量化载荷都优于标量载荷。但是请注意，使用矢量化负载会增加寄存器压力并降低总体并行度。因此，如果的内核已经受到寄存器限制或并行度很低，则可能需要坚持标量加载。同样，如前所述，如果指针未对齐或以字节为单位的数据类型大小不是2的幂，则不能使用矢量化加载。
向量化加载是应该尽可能使用的基本CUDA优化，因为它们会增加带宽，减少指令数量并减少延迟。本文展示了如何通过较少的更改就可以轻松地将向量化负载合并到现有内核中。
本文参考文章链接
https://developer.nvidia.com/blog/cuda-pro-tip-increase-performance-with-vectorized-memory-access/
