不完整的全局內存訪問是否合併？

Question

它是否合併，如果n < warpSize？

// In kernel 
int x; 
if (threadId < n) 
    x = globalMem[threadId]; 

這種情況appers在循環的最後一次迭代，如果一些N是warpSize不可分割的。我應該運行關於這些sitatuations和分配設備內存只能由warpSize整除或它的合併，因爲它是？

Answer 1

如果threadId按照cuda programming guide - thread hierachy中的說明正確計算，則此訪問將被合併 - 這將是threadId = threadIdx.x的情況。

對於不同的計算體系結構，內存聯合會略有不同。更多詳情請見appendix G of cuda programming guide。

一般而言，如果您的線程從您的第一個線程訪問的元素地址開始抓取內存中的連續元素，則可以說全局內存訪問已合併。我們假設你有一個float數組。
float array[]
和你的內存存取權限看起來那樣

array[threadIdx.x == 0, threadId.x == 1, threadIdx.x == 2, ..., threadIdx.x == 31] 

比你的訪問將被coalesed。

但是，如果你以這種方式訪問​​內存（交錯例如）

array[threadIdx.x == 0, NONE, threadId.x == 1, NONE, threadIdx.x == 2, ..., NONE, threadIdx.x == 31] 

比你的訪問不會合並（NONE意味着這個數組元素是不會被任何線程訪問）

在第一種情況下，您可以獲取128個連續的字節內存在第二種情況下，您抓取256個字節。對於第二種情況，需要兩次彎曲從全局存儲器加載存儲器，而不是第一種情況的一次彎曲。但是在這兩種情況下，以下計算只需要32個浮點元素（即128個字節）。因此，在這種簡單的情況下，您的全球負載率將從1.0降至0.5。