CUDA內核中全局內存加載事務的計算

Question

我有一個關於CUDA內核中全局內存加載事務計算的問題，因爲配置文件的值與我的期望值不符。看看下面簡單的矩陣乘法的代碼，例如：

__global__ void matmul_kernel(float *A, float* B, float *C, int n) 
{ 
    int i, j, k; 
    float c; 

    i = blockIdx.x; 
    for(i=i; i < n; i += gridDim.x){ 
     j = threadIdx.x; 
     c = 0.0; 
     for(k = 0; k < n; k++) 
      c += A[i*n + k]*B[k*n + j]; 
     C[i*n + j] = c; 
    } 
} 

dim3 grid(1,1,1); 
dim3 block(128,1,1); 
n = 128; 
matmul_kernel<<<grid, block>>>(A, B, C, n); 

我用最簡單的矩陣乘法爲例。在上面的CUDA實現中，我將i循環迭代映射到塊索引，並將j循環映射到每個線程塊中的線程索引。線程塊和網格都是一維的。

請不要關注此實現的性能。我知道這並不高效，因爲我只是將它用於實驗目的。

在這個實現中，由於我在每個塊中分配了128個線程，所以j循環可以完全並行化。但我只爲i循環分配了1個塊，因此它將循環n次。下圖顯示了k=0時的執行狀態。在這種狀態下，128個線程訪問A的第一個元素和B的128個第一行元素。我在使用Kepler 40架構的Quadro K6000上執行此CUDA代碼，並打開L1緩存。由於對B的128次訪問被合併，因此加載的數量爲128*4/128 = 4（第一個128爲128個元素，第二個128爲L1高速緩存行大小（以字節爲單位），4爲浮點類型的字節數）。對於訪問A的128個訪問，由於它們訪問相同的元素，因此1個緩存行負載應該足夠。所以全球負荷的數量是4+1=5。但這只是k=0時的加載次數。 k將循環128次，i也將循環128次，所以負載總數爲5*128*128=81920。但是，概要的全局負載是131072。該值等於(4+4)*128*128。看起來在k=0的負載數量是4而不是1.任何人都可以解釋爲什麼配置值不符合我的預期值嗎？

Answer 1

只有你錯過了一個小點。全局內存訪問僅適用於經線內的線程（請參閱programming guide）。在你的情況下，有4個經紗。每個將需要一個內存交易的元素A.