在 CUDA C/C++ kernel中使用內(nèi)存
如何在主機(jī)和設(shè)備之間高效地移動(dòng)數(shù)據(jù)。本文將討論如何有效地從內(nèi)核中訪問(wèn)設(shè)備存儲(chǔ)器，特別是 全局內(nèi)存 。
在 CUDA 設(shè)備上有幾種內(nèi)存，每種內(nèi)存的作用域、生存期和緩存行為都不同。到目前為止，已經(jīng)使用了駐留在設(shè)備 DRAM 中的 全局內(nèi)存 ，用于主機(jī)和設(shè)備之間的傳輸，以及內(nèi)核的數(shù)據(jù)輸入和輸出。這里的名稱 global 是指作用域，因?yàn)樗梢詮闹鳈C(jī)和設(shè)備訪問(wèn)和修改。全局內(nèi)存可以像下面代碼片段的第一行那樣使用 device de Clara 說(shuō)明符在全局（變量）范圍內(nèi)聲明，或者使用 cudaMalloc() 動(dòng)態(tài)分配并分配給一個(gè)常規(guī)的 C 指針變量，如第 7 行所示。全局內(nèi)存分配可以在應(yīng)用程序的生命周期內(nèi)保持。根據(jù)設(shè)備的 計(jì)算能力 ，全局內(nèi)存可能被緩存在芯片上，也可能不在芯片上緩存。
device int globalArray[256];

void foo()
{
…
int *myDeviceMemory = 0;
cudaError_t result = cudaMalloc(&myDeviceMemory, 256 * sizeof(int));
…
}
在討論全局內(nèi)存訪問(wèn)性能之前，需要改進(jìn)對(duì) CUDA 執(zhí)行模型的理解。已經(jīng)討論了如何將 線程被分組為線程塊 分配給設(shè)備上的多處理器。在執(zhí)行過(guò)程中，有一個(gè)更精細(xì)的線程分組到 warps 。 GPU 上的多處理器以 SIMD （ 單指令多數(shù)據(jù) ）方式為每個(gè)扭曲執(zhí)行指令。所有當(dāng)前支持 CUDA – 的 GPUs 的翹曲尺寸（實(shí)際上是 SIMD 寬度）是 32 個(gè)線程。
全局內(nèi)存合并
將線程分組為扭曲不僅與計(jì)算有關(guān)，而且與全局內(nèi)存訪問(wèn)有關(guān)。設(shè)備 coalesces 全局內(nèi)存加載并存儲(chǔ)，由一個(gè) warp 線程發(fā)出的盡可能少的事務(wù)，以最小化 DRAM 帶寬（在計(jì)算能力小于 2 . 0 的老硬件上，事務(wù)合并在 16 個(gè)線程的一半扭曲內(nèi)，而不是整個(gè)扭曲中）。為了弄清楚 CUDA 設(shè)備架構(gòu)中發(fā)生聚結(jié)的條件，在三個(gè) Tesla 卡上進(jìn)行了一些簡(jiǎn)單的實(shí)驗(yàn)： a Tesla C870 （計(jì)算能力 1 . 0 ）、 Tesla C1060 （計(jì)算能力 1 . 3 ）和 Tesla C2050 （計(jì)算能力 2 . 0 ）。
運(yùn)行兩個(gè)實(shí)驗(yàn)，使用如下代碼（ GitHub 上也有 ）中所示的增量?jī)?nèi)核的變體，一個(gè)具有數(shù)組偏移量，這可能導(dǎo)致對(duì)輸入數(shù)組的未對(duì)齊訪問(wèn)，另一個(gè)是對(duì)輸入數(shù)組的跨步訪問(wèn)。
#include
#include

// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
if (result != cudaSuccess) {
fprintf(stderr, “CUDA Runtime Error: %sn”, cudaGetErrorString(result));
assert(result == cudaSuccess);
}
#endif
return result;
}

template
global void offset(T* a, int s)
{
int i = blockDim.x * blockIdx.x + threadIdx.x + s;
a[i] = a[i] + 1;
}

template
global void stride(T* a, int s)
{
int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
a[i] = a[i] + 1;
}

template
void runTest(int deviceId, int nMB)
{
int blockSize = 256;
float ms;

T *d_a;
cudaEvent_t startEvent, stopEvent;

int n = nMB10241024/sizeof(T);

// NB: d_a(33*nMB) for stride case
checkCuda( cudaMalloc(&d_a, n * 33 * sizeof(T)) );

checkCuda( cudaEventCreate(&startEvent) );
checkCuda( cudaEventCreate(&stopEvent) );

printf(“Offset, Bandwidth (GB/s):n”);

offset<<>>(d_a, 0); // warm up

for (int i = 0; i <= 32; i++) {
checkCuda( cudaMemset(d_a, 0.0, n * sizeof(T)) );

checkCuda( cudaEventRecord(startEvent,0) );
offset<<>>(d_a, i);
checkCuda( cudaEventRecord(stopEvent,0) );
checkCuda( cudaEventSynchronize(stopEvent) );checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
printf("%d, %fn", i, 2*nMB/ms);

}

printf(“n”);
printf(“Stride, Bandwidth (GB/s):n”);

stride<<>>(d_a, 1); // warm up
for (int i = 1; i <= 32; i++) {
checkCuda( cudaMemset(d_a, 0.0, n * sizeof(T)) );

checkCuda( cudaEventRecord(startEvent,0) );
stride<<>>(d_a, i);
checkCuda( cudaEventRecord(stopEvent,0) );
checkCuda( cudaEventSynchronize(stopEvent) );checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
printf("%d, %fn", i, 2*nMB/ms);

}

checkCuda( cudaEventDestroy(startEvent) );
checkCuda( cudaEventDestroy(stopEvent) );
cudaFree(d_a);
}

int main(int argc, char **argv)
{
int nMB = 4;
int deviceId = 0;
bool bFp64 = false;

for (int i = 1; i < argc; i++) {
if (!strncmp(argv[i], “dev=”, 4))
deviceId = atoi((char*)(&argv[i][4]));
else if (!strcmp(argv[i], “fp64”))
bFp64 = true;
}

cudaDeviceProp prop;

checkCuda( cudaSetDevice(deviceId) )
;
checkCuda( cudaGetDeviceProperties(&prop, deviceId) );
printf(“Device: %sn”, prop.name);
printf(“Transfer size (MB): %dn”, nMB);

printf("%s Precisionn", bFp64 ? “Double” : “Single”);

if (bFp64) runTest(deviceId, nMB);
else runTest(deviceId, nMB);
}
此代碼可以通過(guò)傳遞“ fp64 ”命令行選項(xiàng)以單精度（默認(rèn)值）或雙精度運(yùn)行偏移量?jī)?nèi)核和跨步內(nèi)核。每個(gè)內(nèi)核接受兩個(gè)參數(shù)，一個(gè)輸入數(shù)組和一個(gè)表示訪問(wèn)數(shù)組元素的偏移量或步長(zhǎng)的整數(shù)。內(nèi)核在一系列偏移和跨距的循環(huán)中被稱為。
未對(duì)齊的數(shù)據(jù)訪問(wèn)
下圖顯示了 Tesla C870 、 C1060 和 C2050 上的偏移內(nèi)核的結(jié)果。

設(shè)備內(nèi)存中分配的數(shù)組由 CUDA 驅(qū)動(dòng)程序與 256 字節(jié)內(nèi)存段對(duì)齊。該設(shè)備可以通過(guò) 32 字節(jié)、 64 字節(jié)或 128 字節(jié)的事務(wù)來(lái)訪問(wèn)全局內(nèi)存。對(duì)于 C870 或計(jì)算能力為 1 . 0 的任何其他設(shè)備，半線程的任何未對(duì)齊訪問(wèn)（或半扭曲線程不按順序訪問(wèn)內(nèi)存的對(duì)齊訪問(wèn)）將導(dǎo)致 16 個(gè)獨(dú)立的 32 字節(jié)事務(wù)。由于每個(gè) 32 字節(jié)事務(wù)只請(qǐng)求 4 個(gè)字節(jié)，因此可以預(yù)期有效帶寬將減少 8 倍，這與上圖（棕色線）中看到的偏移量（不是 16 個(gè)元素的倍數(shù)）大致相同，對(duì)應(yīng)于線程的一半扭曲。
對(duì)于計(jì)算能力為 1 . 2 或 1 . 3 的 Tesla C1060 或其他設(shè)備，未對(duì)準(zhǔn)訪問(wèn)的問(wèn)題較少。基本上，通過(guò)半個(gè)線程對(duì)連續(xù)數(shù)據(jù)的未對(duì)齊訪問(wèn)在幾個(gè)“覆蓋”請(qǐng)求的數(shù)據(jù)的事務(wù)中提供服務(wù)。由于未請(qǐng)求的數(shù)據(jù)正在傳輸，以及不同的半翹曲所請(qǐng)求的數(shù)據(jù)有些重疊，因此相對(duì)于對(duì)齊的情況仍然存在性能損失，但是這種損失遠(yuǎn)遠(yuǎn)小于 C870 。
計(jì)算能力為 2 . 0 的設(shè)備，如 Tesla C250 ，在每個(gè)多處理器中都有一個(gè) L1 緩存，其行大小為 128 字節(jié)。該設(shè)備將線程的訪問(wèn)合并到盡可能少的緩存線中，從而導(dǎo)致對(duì)齊，對(duì)跨線程順序內(nèi)存訪問(wèn)吞吐量的影響可以忽略不計(jì)。
快速內(nèi)存訪問(wèn)
步幅內(nèi)核的結(jié)果如下圖所示。

對(duì)于快速的全局內(nèi)存訪問(wèn)，有不同的看法。對(duì)于大步進(jìn)，無(wú)論架構(gòu)版本如何，有效帶寬都很差。這并不奇怪：當(dāng)并發(fā)線程同時(shí)訪問(wèn)物理內(nèi)存中相距很遠(yuǎn)的內(nèi)存地址時(shí)，硬件就沒(méi)有機(jī)會(huì)合并這些訪問(wèn)。從上圖中可以看出，在 Tesla C870 上，除 1 以外的任何步幅都會(huì)導(dǎo)致有效帶寬大幅降低。這是因?yàn)?compute capability 1 . 0 和 1 . 1 硬件需要跨線程進(jìn)行線性、對(duì)齊的訪問(wèn)以進(jìn)行合并，因此我們?cè)?offset 內(nèi)核中看到了熟悉的 1 / 8 帶寬。 Compute capability 1 . 2 及更高版本的硬件可以將訪問(wèn)合并為對(duì)齊的段（ CC 1 . 2 / 1 . 3 上為 32 、 64 或 128 字節(jié)段，在 CC 2 . 0 及更高版本上為 128 字節(jié)緩存線），因此該硬件可以產(chǎn)生平滑的帶寬曲線。
當(dāng)訪問(wèn)多維數(shù)組時(shí)，線程通常需要索引數(shù)組的更高維，因此快速訪問(wèn)是不可避免的。可以使用一種名為 共享內(nèi)存 的 CUDA 內(nèi)存來(lái)處理這些情況。共享內(nèi)存是一個(gè)線程塊中所有線程共享的片上內(nèi)存。共享內(nèi)存的一個(gè)用途是將多維數(shù)組的 2D 塊以合并的方式從全局內(nèi)存提取到共享內(nèi)存中，然后讓連續(xù)的線程繞過(guò)共享內(nèi)存塊。與全局內(nèi)存不同，對(duì)共享內(nèi)存的快速訪問(wèn)沒(méi)有懲罰。
概括
本文討論了如何從 CUDA 內(nèi)核代碼中有效地訪問(wèn)全局內(nèi)存的一些方面。設(shè)備上的全局內(nèi)存訪問(wèn)與主機(jī)上的數(shù)據(jù)訪問(wèn)具有相同的性能特征，即數(shù)據(jù)局部性非常重要。在早期的 CUDA 硬件中，內(nèi)存訪問(wèn)對(duì)齊和跨線程的局部性一樣重要，但在最近的硬件上，對(duì)齊并不是什么大問(wèn)題。另一方面，快速的內(nèi)存訪問(wèn)會(huì)損害性能，使用片上共享內(nèi)存可以減輕這種影響。

總結(jié)

以上是生活随笔為你收集整理的在 CUDA C/C++ kernel中使用内存的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問(wèn)題。

如果覺(jué)得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。