CUDA翻譯：How to Access Global Memory Efficiently in CUDA

在先前的兩篇博客里，我們展示了如何高效地在Host端和Device端交換數(shù)據(jù)，在本篇博客我們著重討論如何高效地訪問顯存，特別是Global Memory

CUDA設備上有好幾種顯存，具有各自的作用域，生命周期，緩存機制。這個系列里我們以駐留在設備DRAM的GlobalMemory為例，它將被用于host和device端數(shù)據(jù)交換，以及核函數(shù)的輸入輸出數(shù)據(jù)交換。

名字中的global表示他的作用域，他可以在host，device端來訪問，修改其內容。GlobalMemory可以使用device關鍵字在global作用域聲明，或者使用cudaMalloc配合指針來動態(tài)分配：

????__device__?int?globalArray[256];
?????
????void?foo()
????{
????????...
????????int?*myDeviceMemory?=?0;
????????cudaError_t?result?=?cudaMalloc(&myDeviceMemory,?256?*?sizeof(int));
????????...
????}

根據(jù)設備的compute capability，Global Memory 可能會在片上進行緩存。

在我們深入探討GlobalMemory訪存性能之前，我們需要完善對CUDA執(zhí)行模型的理解。我們已經討論了線程是如何被組織進線程塊的，這些線程塊又被分配給設備上的Multi Processor。在執(zhí)行過程中，我們又把線程進一步分組為線程束。GPU上的Multi Processor以SIMD的方式給每個線程束執(zhí)行指令，線程束則是由32個線程組成。

譯者注：我覺得這里原文寫錯了，應該是以SIMT的方式執(zhí)行指令，歡迎感興趣的讀者討論

全局內存合并

將線程組成一個線程束不僅跟計算有關，還跟訪存有關。設備將Global Memory的加載和存儲盡可能合并成較少的內存事務，以最小化DRAM帶寬(在Compute Capability<2.0 更老的設備上，事務會在半個線程束內合并)。為了弄清楚CUDA設備架構上發(fā)生訪存合并的條件，我們分別在

• Tesla C870 (compute capability 1.0)
• Tesla C1060 (compute capability 1.3)
• Tesla C2050 (compute capability 2.0). 這上面進行測試

我們使用兩個不同的Kernel，其中一個是帶有offset，導致對數(shù)組訪問是不對齊的，而另外一個則是跨步長來訪問數(shù)組：

具體代碼參考：https://github.com/NVIDIA-developer-blog/code-samples/blob/master/series/cuda-cpp/coalescing-global/coalescing.cu

????template?<typename?T>
????__global__?void?offset(T*?a,?int?s)
????{
??????int?i?=?blockDim.x?*?blockIdx.x?+?threadIdx.x?+?s;
??????a[i]?=?a[i]?+?1;
????}
????
????template?<typename?T>
????__global__?void?stride(T*?a,?int?s)
????{
??????int?i?=?(blockDim.x?*?blockIdx.x?+?threadIdx.x)?*?s;
??????a[i]?=?a[i]?+?1;
????}

通過設置fp64命令行選項，我們可以讓代碼在單精度/雙精度下運行。每個kernel一共有兩個輸入，第一個是輸入數(shù)組，第二個則表示訪問數(shù)組的stride/offset，我們在一個stride/offset的范圍內進行循環(huán)調用核函數(shù)

不對齊的數(shù)據(jù)訪問

offset kernel的測試結果如下：

分配到設備內存上的數(shù)組由CUDA驅動程序對齊到256-byte內存段，設備可以通過32/64/128-byte大小的內存事務來訪問Global Memory，這些內存事務和他們的大小對齊。

對于C870這樣的Compute Cabability<1.0的設備，任意 半個線程束 內的不對齊訪問(或半個線程束不按順序訪問內存的對齊訪問)，將會把256-byte訪問拆分成16個32-byte內存訪問事務

譯者注：這里為什么是16個32-byte，因為是半個線程束，所以是16。因為沒有對齊訪問，所以每一次訪問一個元素，是要用單獨一個最小內存事務來完成，而內存事務最小為32-byte。因此總的就是16x32-byte

在每一次32-byte大小的內存事務中，我們只取了其中的4byte(譯者注這里應該是float32情況下)。那么相當于有效帶寬為原先的1/8。對應我們上圖棕褐色曲線中，offset不為16倍數(shù)時對應的帶寬大小。

對于C1060這些Compute Cabability為1.2/1.3的設備，不對齊訪問帶來的問題則比較小。基本上由半個線程束帶來的不對齊訪問是在其中幾個覆蓋了所需要的數(shù)據(jù)的事務中進行。但因為傳輸了一些不需要的數(shù)據(jù)，以及不同的半個線程束內數(shù)據(jù)有所重疊，仍會帶來性能損失，但損失相較C870少的多

而在Compute Cabability為2.0的設備，如C205，它在每個MultiProecssor上都有一個linesize為128byte的L1 Cache。設備將盡可能把線程束訪問合并到盡可能少的cache line中，從而使跨線程的內存訪問對吞吐量影響可忽略不計

跨Stride內存訪問

跨Stride訪問的Kernel結果如下圖所示：

我們有一張完全不一樣的數(shù)據(jù)圖。對于比較大的Strides，無論架構版本如何，其有效帶寬都很低。這也無需感到驚訝，當線程并發(fā)地訪問相距較遠的內存地址時，硬件就沒有機會訪問?？梢栽趫D中看到C870，當stride〉1后，其有效帶寬大幅度降低，這是因為Compute Cabability為1.0/1.1的硬件要求線性，對齊的內存訪問以合并 ，所以我們也能在offset kernel那里看到性能降低至1/8。Compute capability 1.2及更高版本的設備可以合并成段對齊的訪問（CC 1.2/1.3上的32、64或128字節(jié)段，以及CC 2.0及更高版本上的128bytes的cache line），因此這些設備的帶寬曲線更加平滑。

當我們訪問高維數(shù)組的時候，線程通常要索引數(shù)組中更高的維度，所以stirded access是不可避免的。我們可以借助shared memory來處理這種情形，shared memory是在一個線程塊內被所有線程共享的片上顯存。Shared Memory的一個用途就是以合并的方式，從Global Memory取一個二維的分塊到Shared Memory中，然后連續(xù)的線程經過Shared Memory分塊跨Stride訪問。與Global Memory不同的是， Shared Memory跨Stride訪問并不會有任何懲罰，我們將在下篇博客詳細講解Shared Memory。

總結

在本篇博客中，我們討論了CUDA Kernel中如何高效訪問Global Memory，設備上的Shared Memory訪問與Host端數(shù)據(jù)訪問具有相同的特性，數(shù)據(jù)局部性非常重要。在早期的CUDA設備中中，內存訪問對齊與線程間的局部性一樣重要，但在最近的設備上，內存訪問對齊并不是什么大問題。另一方面，跨Stride訪問顯存可能會影響性能，使用Shared Memory可以緩解這一問題。

在下一篇文章中，我們將詳細探討Shared Memory，并展示如何使用Shared Memory，以避免在矩陣轉置期間進行跨Stride的Global Memory訪問。