CUDA探查器報告全局內存訪問效率低下

Question

我有一個簡單的CUDA內核，我認爲它有效地訪問全局內存。然而，Nvidia profiler報告說我正在執行低效的全局內存訪問。我的內核代碼是：

__global__ void update_particles_kernel 
(
    float4 *pos, 
    float4 *vel, 
    float4 *acc, 
    float dt, 
    int numParticles 
) 
{ 
int index = threadIdx.x + blockIdx.x * blockDim.x; 
int offset = 0; 

while(index + offset < numParticles) 
{ 
    vel[index + offset].x += dt*acc[index + offset].x; // line 247 
    vel[index + offset].y += dt*acc[index + offset].y; 
    vel[index + offset].z += dt*acc[index + offset].z; 

    pos[index + offset].x += dt*vel[index + offset].x; // line 251 
    pos[index + offset].y += dt*vel[index + offset].y; 
    pos[index + offset].z += dt*vel[index + offset].z; 

    offset += blockDim.x * gridDim.x; 
} 

特別剖析報告如下：

從CUDA best practices guide它說：

「對於計算能力2.x的設備的，這些需求可以很容易地總結出來：一個warp線程的併發訪問將會合併成許多事務處理，這個事務處理的數量等同於服務所有經線。默認情況下，所有訪問都通過L1緩存，即128字節的行。對於分散訪問模式，爲了減少過度訪問，有時僅在L2中進行高速緩存是有用的，L2緩存了較短的32字節段（請參閱CUDA C編程指南）。

對於計算能力3.x的設備，對全局存儲器的訪問僅緩存在L2中; L1保留用於本地內存訪問。一些計算能力爲3.5,3.7或5.2的設備也允許在L1中選擇性地緩存全局變量。「

現在在我的內核的基礎上，我期望16個訪問需要服務一個32線程扭曲，因爲float4是16字節，並且在我的卡上（770米計算能力3.0）從L2緩存中讀取是以32字節塊（16字節* 32線程/ 32字節緩存線= 16訪問）執行的。 profiler報告我正在訪問16個數據，但我不明白的是，爲什麼profiler報告理想的訪問將涉及247行的每次訪問8個L2事務，其餘行的每次訪問只有4個L2事務。我在這裏失蹤了嗎？

Answer 1

我有一個簡單的CUDA內核，我認爲它可以高效地訪問全局內存。然而，Nvidia profiler報告說我正在執行低效的全局內存訪問。

舉一個例子，你float4vel陣列存儲在內存中是這樣的：

0.x 0.y 0.z 0.w 1.x 1.y 1.z 1.w 2.x 2.y 2.z 2.w 3.x 3.y 3.z 3.w ... 
^   ^   ^   ^   ... 
    thread0   thread1   thread2   thread3 

所以，當你這樣做：

vel[index + offset].x += ...; // line 247 

您正在訪問（存儲）在我在上面標記的位置（.x）。每個^標記之間的差距表明探測器指出的訪問模式效率低下。 （在下一行代碼中，無關緊要，您將存儲到.y位置。）

至少有2種解決方案，其中一種解決方案是經典的AoS - > SoA重組數據，通過適當的代碼調整。這是有據可查的（例如，here在cuda標籤和其他地方）就其含義而言，以及如何去做，所以我會讓你看看。

另一種典型的解決方案是在需要時爲每個線程加載float4數量，並根據需要存儲每個線程的float4數量。你的代碼可以平凡修改，以做到這一點，這應該給改進的分析結果：

//preceding code need not change 
while(index + offset < numParticles) 
{ 
    float4 my_vel = vel[index + offset]; 
    float4 my_acc = acc[index + offset]; 
    my_vel.x += dt*my_acc.x; 
    my_vel.y += dt*my_acc.y; 
    my_vel.z += dt*my_acc.z; 
    vel[index + offset] = my_vel; 

    float4 my_pos = pos[index + offset]; 
    my_pos.x += dt*my_vel.x; 
    my_pos.y += dt*my_vel.y; 
    my_pos.z += dt*my_vel.z; 
    pos[index + offset] = my_pos; 

    offset += blockDim.x * gridDim.x; 
} 

即使你可能會認爲這代碼是不是你的代碼「低效率」，因爲「出現」你的代碼只加載和存儲.x,.y,.z，而我的「出現」也加載和存儲.w，實際上基本上沒有區別，因爲GPU加載和存儲到全局存儲器/從全局存儲器存儲的方式。雖然您的代碼似乎未觸及.w，但在訪問相鄰元素的過程中，GPU將加載全局內存中的.w元素，並且還將（最終）將.w元素存儲回全局內存。

我不明白的是爲什麼剖析報告，理想的訪問將涉及每個接入8個L2交易線路247

對於你原來的代碼行247，您正在訪問一個float對於vel.x的加載操作，每個線程的每個線程的數量爲acc.x，每個線程的數量爲float。每個線程本身的數量應該需要128個字節用於一個warp，即4個32字節的L2高速緩存線。兩個負載一起需要8個L2高速緩存線負載。這是理想的情況，它假設數量很好地包裝在一起（SoA）。但這不是你所擁有的（你有AoS）。