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

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

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

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

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

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

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

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

全局內(nèi)存合并

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

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

我們使用兩個(gè)不同的Kernel，其中一個(gè)是帶有offset，導(dǎo)致對(duì)數(shù)組訪問(wèn)是不對(duì)齊的，而另外一個(gè)則是跨步長(zhǎng)來(lái)訪問(wèn)數(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;
????}

通過(guò)設(shè)置fp64命令行選項(xiàng)，我們可以讓代碼在單精度/雙精度下運(yùn)行。每個(gè)kernel一共有兩個(gè)輸入，第一個(gè)是輸入數(shù)組，第二個(gè)則表示訪問(wèn)數(shù)組的stride/offset，我們?cè)谝粋€(gè)stride/offset的范圍內(nèi)進(jìn)行循環(huán)調(diào)用核函數(shù)

不對(duì)齊的數(shù)據(jù)訪問(wèn)

offset kernel的測(cè)試結(jié)果如下：

分配到設(shè)備內(nèi)存上的數(shù)組由CUDA驅(qū)動(dòng)程序?qū)R到256-byte內(nèi)存段，設(shè)備可以通過(guò)32/64/128-byte大小的內(nèi)存事務(wù)來(lái)訪問(wèn)Global Memory，這些內(nèi)存事務(wù)和他們的大小對(duì)齊。

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

譯者注：這里為什么是16個(gè)32-byte，因?yàn)槭前雮€(gè)線程束，所以是16。因?yàn)闆](méi)有對(duì)齊訪問(wèn)，所以每一次訪問(wèn)一個(gè)元素，是要用單獨(dú)一個(gè)最小內(nèi)存事務(wù)來(lái)完成，而內(nèi)存事務(wù)最小為32-byte。因此總的就是16x32-byte

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

對(duì)于C1060這些Compute Cabability為1.2/1.3的設(shè)備，不對(duì)齊訪問(wèn)帶來(lái)的問(wèn)題則比較小?；旧嫌砂雮€(gè)線程束帶來(lái)的不對(duì)齊訪問(wèn)是在其中幾個(gè)覆蓋了所需要的數(shù)據(jù)的事務(wù)中進(jìn)行。但因?yàn)閭鬏斄艘恍┎恍枰臄?shù)據(jù)，以及不同的半個(gè)線程束內(nèi)數(shù)據(jù)有所重疊，仍會(huì)帶來(lái)性能損失，但損失相較C870少的多

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

跨Stride內(nèi)存訪問(wèn)

跨Stride訪問(wèn)的Kernel結(jié)果如下圖所示：

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

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

總結(jié)

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

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