這篇文章是對(duì) CUDA 的一個(gè)超級(jí)簡(jiǎn)單的介紹，這是一個(gè)流行的并行計(jì)算平臺(tái)和 NVIDIA 的編程模型。我在 2013 年給 CUDA 寫了一篇前一篇 “簡(jiǎn)單介紹” ，這幾年來非常流行。但是 CUDA 編程變得越來越簡(jiǎn)單， GPUs 也變得更快了，所以是時(shí)候更新（甚至更容易）介紹了。

CUDA C ++只是使用 CUDA 創(chuàng)建大規(guī)模并行應(yīng)用程序的一種方式。它讓您使用強(qiáng)大的 C ++編程語言來開發(fā)由數(shù)千個(gè)并行線程加速的高性能算法 GPUs 。許多開發(fā)人員已經(jīng)用這種方式加速了他們對(duì)計(jì)算和帶寬需求巨大的應(yīng)用程序，包括支持人工智能正在進(jìn)行的革命的庫和框架 深度學(xué)習(xí) 。

所以，您已經(jīng)聽說了 CUDA ，您有興趣學(xué)習(xí)如何在自己的應(yīng)用程序中使用它。如果你是 C 或 C ++程序員，這個(gè)博客應(yīng)該給你一個(gè)好的開始。接下來，您需要一臺(tái)具有 CUDA – 功能的 GPU 計(jì)算機(jī)（ Windows 、 Mac 或 Linux ，以及任何 NVIDIA GPU 都可以），或者需要一個(gè)具有 GPUs 的云實(shí)例（ AWS 、 Azure 、 IBM 軟層和其他云服務(wù)提供商都有）。您還需要安裝免費(fèi)的 CUDA 工具箱 。

我們開始吧！

從簡(jiǎn)單開始

我們將從一個(gè)簡(jiǎn)單的 C ++程序開始，它添加兩個(gè)數(shù)組的元素，每個(gè)元素有一百萬個(gè)元素。

#include 
#include  // function to add the elements of two arrays
void add(int n, float *x, float *y)
{ for (int i = 0; i < n; i++) y[i] = x[i] + y[i];
} int main(void)
{ int N = 1<<20; // 1M elements float *x = new float[N]; float *y = new float[N]; // initialize x and y arrays on the host for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } // Run kernel on 1M elements on the CPU add(N, x, y); // Check for errors (all values should be 3.0f) float maxError = 0.0f; for (int i = 0; i < N; i++) maxError = fmax(maxError, fabs(y[i]-3.0f)); std::cout << "Max error: " << maxError << std::endl; // Free memory delete [] x; delete [] y; return 0;
}

首先，編譯并運(yùn)行這個(gè) C ++程序。將代碼放在一個(gè)文件中，并將其保存為add.cpp，然后用 C ++編譯器編譯它。我在 Mac 電腦上，所以我用的是clang++，但你可以在 Linux 上使用g++，或者在 Windows 上使用 MSVC 。

> clang++ add.cpp -o add

然后運(yùn)行它：

> ./add Max error: 0.000000

（在 Windows 上，您可能需要命名可執(zhí)行文件添加. exe 并使用.dd運(yùn)行它。）

正如預(yù)期的那樣，它打印出求和中沒有錯(cuò)誤，然后退出?，F(xiàn)在我想讓這個(gè)計(jì)算在 GPU 的多個(gè)核心上運(yùn)行（并行）。其實(shí)邁出第一步很容易。

首先，我只需要將我們的add函數(shù)轉(zhuǎn)換成 GPU 可以運(yùn)行的函數(shù)，在 CUDA 中稱為內(nèi)核。要做到這一點(diǎn)，我所要做的就是把說明符__global__添加到函數(shù)中，它告訴 CUDA C ++編譯器，這是一個(gè)在 GPU 上運(yùn)行的函數(shù)，可以從 CPU 代碼調(diào)用。

// CUDA Kernel function to add the elements of two arrays on the GPU
__global__
void add(int n, float *x, float *y)
{ for (int i = 0; i < n; i++) y[i] = x[i] + y[i];
}

這些__global__函數(shù)被稱為果仁，在 GPU 上運(yùn)行的代碼通常稱為設(shè)備代碼，而在 CPU 上運(yùn)行的代碼是主機(jī)代碼。

CUDA 中的內(nèi)存分配

為了在 GPU 上計(jì)算，我需要分配 GPU 可訪問的內(nèi)存， CUDA 中的統(tǒng)一存儲(chǔ)器通過提供一個(gè)系統(tǒng)中所有 GPUs 和 CPU 都可以訪問的內(nèi)存空間，這使得這一點(diǎn)變得簡(jiǎn)單。要在統(tǒng)一內(nèi)存中分配數(shù)據(jù)，請(qǐng)調(diào)用cudaMallocManaged()，它返回一個(gè)指針，您可以從主機(jī)（ CPU ）代碼或設(shè)備（ GPU ）代碼訪問該指針。要釋放數(shù)據(jù)，只需將指針傳遞到cudaFree()。

我只需要將上面代碼中對(duì)new的調(diào)用替換為對(duì)cudaMallocManaged()的調(diào)用，并將對(duì)delete []的調(diào)用替換為對(duì)cudaFree.的調(diào)用

 // Allocate Unified Memory -- accessible from CPU or GPU float *x, *y; cudaMallocManaged(&x, N*sizeof(float)); cudaMallocManaged(&y, N*sizeof(float)); ... // Free memory cudaFree(x); cudaFree(y);

最后，我需要發(fā)射內(nèi)核，它在add()上調(diào)用它。 CUDA 內(nèi)核啟動(dòng)是使用三角括號(hào)語法指定的。我只需要在參數(shù)列表之前將它添加到對(duì) CUDA 的調(diào)用中。

add<<<1, 1>>>(N, x, y);

容易的！我很快將詳細(xì)介紹尖括號(hào)內(nèi)的內(nèi)容；現(xiàn)在您只需要知道這行代碼啟動(dòng)了一個(gè) GPU 線程來運(yùn)行add()。

還有一件事：我需要 CPU 等到內(nèi)核完成后再訪問結(jié)果（因?yàn)?CUDA 內(nèi)核啟動(dòng)不會(huì)阻塞調(diào)用的 CPU 線程）。為此，我只需在對(duì) CPU 進(jìn)行最后的錯(cuò)誤檢查之前調(diào)用cudaDeviceSynchronize()。

以下是完整的代碼：

#include 
#include 
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{ for (int i = 0; i < n; i++) y[i] = x[i] + y[i];
} int main(void)
{ int N = 1<<20; float *x, *y; // Allocate Unified Memory – accessible from CPU or GPU cudaMallocManaged(&x, N*sizeof(float)); cudaMallocManaged(&y, N*sizeof(float)); // initialize x and y arrays on the host for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } // Run kernel on 1M elements on the GPU add<<<1, 1>>>(N, x, y); // Wait for GPU to finish before accessing on host cudaDeviceSynchronize(); // Check for errors (all values should be 3.0f) float maxError = 0.0f; for (int i = 0; i < N; i++) maxError = fmax(maxError, fabs(y[i]-3.0f)); std::cout << "Max error: " << maxError << std::endl; // Free memory cudaFree(x); cudaFree(y); return 0;
}

CUDA 文件具有文件擴(kuò)展名；.cu。所以把代碼保存在一個(gè)名為

> nvcc add.cu -o add_cuda
> ./add_cuda
Max error: 0.000000

這只是第一步，因?yàn)檎缢鶎懙?，這個(gè)內(nèi)核只適用于一個(gè)線程，因?yàn)檫\(yùn)行它的每個(gè)線程都將在整個(gè)數(shù)組上執(zhí)行 add 。此外，還有一個(gè)競(jìng)爭(zhēng)條件，因?yàn)槎鄠€(gè)并行線程讀寫相同的位置。

注意：在 Windows 上，您需要確保在 Microsoft Visual Studio 中項(xiàng)目的配置屬性中將“平臺(tái)”設(shè)置為 x64 。

介紹一下！

我認(rèn)為找出運(yùn)行內(nèi)核需要多長(zhǎng)時(shí)間的最簡(jiǎn)單的方法是用nvprof運(yùn)行它，這是一個(gè)帶有 CUDA 工具箱的命令行 GPU 分析器。只需在命令行中鍵入nvprof ./add_cuda：

$ nvprof ./add_cuda
==3355== NVPROF is profiling process 3355, command: ./add_cuda
Max error: 0
==3355== Profiling application: ./add_cuda
==3355== Profiling result:
Time(%) Time Calls Avg Min Max Name
100.00% 463.25ms 1 463.25ms 463.25ms 463.25ms add(int, float*, float*)
...

上面是來自nvprof的截?cái)噍敵?，顯示了對(duì)add的單個(gè)調(diào)用。在 NVIDIA Tesla K80 加速器上需要大約半秒鐘的時(shí)間，而在我 3 歲的 Macbook Pro 上使用 NVIDIA GeForce GT 740M 大約需要半秒鐘的時(shí)間。

讓我們用并行來加快速度。

把線撿起來

既然你已經(jīng)用一個(gè)線程運(yùn)行了一個(gè)內(nèi)核，那么如何使它并行？鍵是在 CUDA 的<<<1, 1>>>語法中。這稱為執(zhí)行配置，它告訴 CUDA 運(yùn)行時(shí)要使用多少并行線程來啟動(dòng) GPU 。這里有兩個(gè)參數(shù)，但是讓我們從更改第二個(gè)參數(shù)開始：線程塊中的線程數(shù)。 CUDA GPUs 運(yùn)行內(nèi)核時(shí)使用的線程塊大小是 32 的倍數(shù)，因此 256 個(gè)線程是一個(gè)合理的選擇。

add<<<1, 256>>>(N, x, y);

如果我只在這個(gè)修改下運(yùn)行代碼，它將為每個(gè)線程執(zhí)行一次計(jì)算，而不是將計(jì)算分散到并行線程上。為了正確地執(zhí)行它，我需要修改內(nèi)核。 CUDA C ++提供了關(guān)鍵字，這些內(nèi)核可以讓內(nèi)核獲得運(yùn)行線程的索引。具體來說，threadIdx.x包含其塊中當(dāng)前線程的索引，blockDim.x包含塊中的線程數(shù)。我只需修改循環(huán)以使用并行線程跨過數(shù)組。

__global__
void add(int n, float *x, float *y)
{ int index = threadIdx.x; int stride = blockDim.x; for (int i = index; i < n; i += stride) y[i] = x[i] + y[i];
}

add函數(shù)沒有太大變化。事實(shí)上，將index設(shè)置為 0 ，stride設(shè)置為 1 會(huì)使其在語義上與第一個(gè)版本相同。

將文件另存為add_block.cu，然后再次在nvprof中編譯并運(yùn)行。在后面的文章中，我將只顯示輸出中的相關(guān)行。

Time(%) Time Calls Avg Min Max Name
100.00% 2.7107ms 1 2.7107ms 2.7107ms 2.7107ms add(int, float*, float*)

這是一個(gè)很大的加速（ 463 毫秒下降到 2 . 7 毫秒），但并不奇怪，因?yàn)槲覐?1 線程到 256 線程。 K80 比我的小 MacBookProGPU 快（ 3 . 2 毫秒）。讓我們繼續(xù)取得更高的表現(xiàn)。

走出街區(qū)

CUDA GPUs 有許多并行處理器組合成流式多處理器或 SMs 。每個(gè) SM 可以運(yùn)行多個(gè)并發(fā)線程塊。例如，基于 Tesla 的 Tesla P100帕斯卡 GPU 體系結(jié)構(gòu)有 56 個(gè)短消息，每個(gè)短消息能夠支持多達(dá) 2048 個(gè)活動(dòng)線程。為了充分利用所有這些線程，我應(yīng)該用多個(gè)線程塊啟動(dòng)內(nèi)核。

現(xiàn)在您可能已經(jīng)猜到執(zhí)行配置的第一個(gè)參數(shù)指定了線程塊的數(shù)量。這些平行線程塊一起構(gòu)成了所謂的網(wǎng)格。因?yàn)槲矣?code style="font-size:inherit;color:inherit;margin:0px;padding:0px;border:0px;font-style:inherit;font-variant:inherit;font-weight:inherit;line-height:inherit;vertical-align:baseline;background-color:rgb(244,244,244);">N元素要處理，每個(gè)塊有 256 個(gè)線程，所以我只需要計(jì)算塊的數(shù)量就可以得到至少 N 個(gè)線程。我只需將N除以塊大?。ㄗ⒁庠?code style="font-size:inherit;color:inherit;margin:0px;padding:0px;border:0px;font-style:inherit;font-variant:inherit;font-weight:inherit;line-height:inherit;vertical-align:baseline;background-color:rgb(244,244,244);">N不是blockSize的倍數(shù)的情況下向上取整）。

int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);

我還需要更新內(nèi)核代碼來考慮線程塊的整個(gè)網(wǎng)格。threadIdx.x提供了包含網(wǎng)格中塊數(shù)的gridDim.x和包含網(wǎng)格中當(dāng)前線程塊索引的blockIdx.x。圖 1 說明了使用 CUDA 、gridDim.x和threadIdx.x在 CUDA 中索引數(shù)組（一維）的方法。其思想是，每個(gè)線程通過計(jì)算到其塊開頭的偏移量（塊索引乘以塊大?。?code style="font-size:inherit;color:inherit;margin:0px;padding:0px;border:0px;font-style:inherit;font-variant:inherit;font-weight:inherit;line-height:inherit;vertical-align:baseline;background-color:rgb(244,244,244);">blockIdx.x * blockDim.x），并將線程的索引添加到塊內(nèi)（threadIdx.x）。代碼blockIdx.x * blockDim.x + threadIdx.x是慣用的 CUDA 。

__global__
void add(int n, float *x, float *y)
{ int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride) y[i] = x[i] + y[i];
}

更新的內(nèi)核還將stride設(shè)置為網(wǎng)格中的線程總數(shù)（blockDim.x * gridDim.x）。 CUDA 內(nèi)核中的這種類型的循環(huán)通常稱為柵格步幅循環(huán)。

將文件另存為&[EZX63 ；&[編譯并在&[EZX37 ；&]中運(yùn)行它]

Time(%) Time Calls Avg Min Max Name
100.00% 94.015us 1 94.015us 94.015us 94.015us add(int, float*, float*)

這是另一個(gè) 28 倍的加速，從運(yùn)行多個(gè)街區(qū)的所有短信 K80 ！我們?cè)?K80 上只使用了 2 個(gè) GPUs 中的一個(gè)，但是每個(gè) GPU 都有 13 條短信。注意，我筆記本電腦中的 GeForce 有 2 條（較弱的）短信，運(yùn)行內(nèi)核需要 680us 。

總結(jié)

下面是三個(gè)版本的add()內(nèi)核在 Tesla K80 和 GeForce GT 750M 上的性能分析。

如您所見，我們可以在 GPUs 上實(shí)現(xiàn)非常高的帶寬。這篇文章中的計(jì)算是非常有帶寬限制的，但是 GPUs 也擅長(zhǎng)于密集矩陣線性代數(shù)深度學(xué)習(xí)、圖像和信號(hào)處理、物理模擬等大量計(jì)算限制的計(jì)算。

關(guān)于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過 20 年的 GPUs 軟件開發(fā)經(jīng)驗(yàn)，從圖形和游戲到基于物理的模擬，到并行算法和高性能計(jì)算。當(dāng)他還是北卡羅來納大學(xué)的博士生時(shí)，他意識(shí)到了一種新生的趨勢(shì)，并為此創(chuàng)造了一個(gè)名字： GPGPU （圖形處理單元上的通用計(jì)算）。

審核編輯：郭婷