CUDA聯合訪問二維塊

Question

對於一維情況，我非常瞭解CUDA中全局內存的整合訪問需求。

但是我對二維情況（即我們有一個2D網格，由2D塊組成）有點卡住了。

假設我有一個向量in_vector，並且在我的內核中，我想以合併的方式訪問它。像這樣：

__global__ void my_kernel(float* out_matrix, float* in_vector, int size) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    int j = blockIdx.y * blockDim.y + threadIdx.y; 
    // ... 
    float vx = in_vector[i]; // This is good. Here we have coalesced access 
    float vy = in_vector[j]; // Not sure about this. All threads in my warp access the same global address. (See explanation) 
    // ... 
    // Do some calculations... Obtain result 
} 

在我對這種2D情況的理解中，塊內的線程以列主要方式「排列」。例如：假設一（threadIdx.x，threadIdx.y）表示法：

• 第一經線將是：（0,0），（1,0），（2,0），...，（ （0,1），（1,1），（2,1），...，（31,1），
• 等等。

在這種情況下，調用in_vector[i]給了我們一個合併訪問，因爲同一warp中的每個連續線程都將訪問連續的地址。然而，調用in_vector[j]似乎是一個壞主意，因爲每個連續的線程將訪問全局內存中的相同地址（例如warp 0中的所有線程將訪問in_vector [0]，這會給我們32個不同的全局內存請求）

我對此有正確的理解嗎？如果是這樣，我怎樣才能使用in_vector[j]聯合訪問全局內存？

Answer 1

您在問題中顯示的內容僅適用於特定的塊大小。你的「聚結的」訪問：

int i = blockIdx.x * blockDim.x + threadIdx.x; 
float vx = in_vector[i]; 

將導致從全局存儲器的in_vector聚結的訪問僅當blockDim.x大於或等於32即使在所述聚結的情況下，共享相同threadIdx.x塊內的每個線程值從全局內存中讀取相同的單詞，這似乎是違反直覺和浪費的。

正確的方式，以確保讀取每線程獨特的合併是計算塊內的線程數，涉及的電力網，也許像中的偏移量：

int tid = threadIdx.x + blockDim.x * threadIdx.y; // must use column major order 
int bid = blockIdx.x + gridDim.x * blockDim.y; // can either use column or row major 
int offset = (blockDim.x * blockDim.y) * bid; // block id * threads per block 
float vx = in_vector[tid + offset]; 

如果你的意圖真的是不閱讀每個線程的唯一值，那麼你就可以節省大量的內存帶寬和實現使用共享內存，這樣的合併：

__shared__ float vx[32], vy[32]; 

int tid = threadIdx.x + blockDim.x * threadIdx.y; 

if (tid < 32) { 
    vx[tid] = in_vector[blockIdx.x * blockDim.x + tid]; 
    vy[tid] = in_vector[blockIdx.y * blockDim.y + tid]; 
} 
__syncthread(); 

，你會得到一個經獨特的閱讀值共享一次。其他線程可以從共享內存中讀取值，而不需要進一步的全局內存訪問。請注意，在上面的示例中，我遵循了代碼的約定，即使這樣讀取in_vector兩次也沒有多大意義。