CUDA Pro：通過向量化內存訪問提高性能
許多CUDA內核受帶寬限制，而新硬件中觸發器與帶寬的比率不斷提高，導致帶寬受限制的內核更多。這使得采取措施減輕代碼中的帶寬瓶頸非常重要。本文將展示如何在CUDA C / 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 = 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優化，因為它們會增加帶寬，減少指令數量並減少延遲。本文展示了如何通過較少的更改就可以輕松地將向量化負載合並到現有內核中。