如何正確合併從全局內存寫入全局內存？

Question

請理解我，但我不懂英文。

我的計算環境是

• CPU：英特爾至強X5690 3.46GHz的* 2EA
• OS：CentOS的5.8
• VGA：NVIDIA公司的GeForce GTX580（CC是2.0）

我請閱讀CUDA C編程指南中關於「合併內存訪問」的文檔。 但我不能將它們應用於我的情況。

我有32x32塊/網格和16x16線程/塊。 這意味着如下代碼。

dim3 grid(32, 32); 
dim3 block(16,16); 

kernel<<<grid, block>>>(...); 

然後，我如何使用聚結內存訪問？

我在下面的內核中使用了代碼。

int i = blockIdx.x*16 + threadIdx.x; 
int j = blockIdx.y*16 + threadIdx.y; 

... 

global_memory[i*512+j] = ...; 

我使用了常量512，因爲線程總量是512x512個線程：它是grid_size x block_size。

但是，我從Visual Profiler看到「低全局內存存儲效率[9.7％平均值，對於計算100％的內核而言]」。

幫手說使用合併內存訪問。 但是，我不知道我應該使用內存的索引上下文。

用於詳細信息代碼的更多信息，The result of an experiment different from CUDA Occupancy Calculator

Answer 1

聚結存儲器加載和存儲在CUDA是一個非常簡單的概念 - 在相同的經紗需要從/加載或存儲線程成合適地對準，在連續的字記憶。

CUDA中的變形大小爲32，並且由相同塊內的線程形成變形，從而使得threadIdx.{xyz}的x尺寸變化最快，y次最快，並且z最慢（從功能上來說，這是與數組中的列主要排序相同）。

您發佈的代碼沒有實現合併內存存儲，因爲同一個warp內的線程以512字的間距存儲，而不是所需的32個連續字。

一個簡單的黑客攻擊，提高凝聚將是解決在列優先的順序記憶，所以：

int i = blockIdx.x*16 + threadIdx.x; 
int j = blockIdx.y*16 + threadIdx.y; 

... 

global_memory[i+512*j] = ...; 

二維塊和網格上較普遍的辦法，以實現在你表現出什麼樣的精神凝聚在的問題是這樣的：

tid_in_block = threadIdx.x + threadIdx.y * blockDim.x; 
    bid_in_grid = blockIdx.x + blockIdx.y * gridDim.x; 
    threads_per_block = blockDim.x * blockDim.y; 

    tid_in_grid = tid_in_block + thread_per_block * bid_in_grid; 

    global_memory[tid_in_grid] = ...; 

最合適的解決方案將取決於你有沒有描述的代碼和數據的其他細節。