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

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

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

我們開始吧！

從簡單開始

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

#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;
}

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

> clang++ add.cpp -o add

然后運行它：

> ./add Max error: 0.000000

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

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

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

// 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 上運行的代碼通常稱為設備代碼，而在 CPU 上運行的代碼是主機代碼。

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

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

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

 // 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()上調用它。 CUDA 內(nèi)核啟動是使用三角括號語法指定的。我只需要在參數(shù)列表之前將它添加到對 CUDA 的調用中。

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

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

還有一件事：我需要 CPU 等到內(nèi)核完成后再訪問結果（因為 CUDA 內(nèi)核啟動不會阻塞調用的 CPU 線程）。為此，我只需在對 CPU 進行最后的錯誤檢查之前調用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 文件具有文件擴展名；.cu。所以把代碼保存在一個名為

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

這只是第一步，因為正如所寫的，這個內(nèi)核只適用于一個線程，因為運行它的每個線程都將在整個數(shù)組上執(zhí)行 add 。此外，還有一個競爭條件，因為多個并行線程讀寫相同的位置。

注意：在 Windows 上，您需要確保在 Microsoft Visual Studio 中項目的配置屬性中將“平臺”設置為 x64 。

介紹一下！

我認為找出運行內(nèi)核需要多長時間的最簡單的方法是用nvprof運行它，這是一個帶有 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的截斷輸出，顯示了對add的單個調用。在 NVIDIA Tesla K80 加速器上需要大約半秒鐘的時間，而在我 3 歲的 Macbook Pro 上使用 NVIDIA GeForce GT 740M 大約需要半秒鐘的時間。

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

把線撿起來

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

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

如果我只在這個修改下運行代碼，它將為每個線程執(zhí)行一次計算，而不是將計算分散到并行線程上。為了正確地執(zhí)行它，我需要修改內(nèi)核。 CUDA C ++提供了關鍵字，這些內(nèi)核可以讓內(nèi)核獲得運行線程的索引。具體來說，threadIdx.x包含其塊中當前線程的索引，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ù)沒有太大變化。事實上，將index設置為 0 ，stride設置為 1 會使其在語義上與第一個版本相同。

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

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

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

走出街區(qū)

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

現(xiàn)在您可能已經(jīng)猜到執(zhí)行配置的第一個參數(shù)指定了線程塊的數(shù)量。這些平行線程塊一起構成了所謂的網(wǎng)格。因為我有N元素要處理，每個塊有 256 個線程，所以我只需要計算塊的數(shù)量就可以得到至少 N 個線程。我只需將N除以塊大小（注意在N不是blockSize的倍數(shù)的情況下向上取整）。

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

我還需要更新內(nèi)核代碼來考慮線程塊的整個網(wǎng)格。threadIdx.x提供了包含網(wǎng)格中塊數(shù)的gridDim.x和包含網(wǎng)格中當前線程塊索引的blockIdx.x。圖 1 說明了使用 CUDA 、gridDim.x和threadIdx.x在 CUDA 中索引數(shù)組（一維）的方法。其思想是，每個線程通過計算到其塊開頭的偏移量（塊索引乘以塊大小：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設置為網(wǎng)格中的線程總數(shù)（blockDim.x * gridDim.x）。 CUDA 內(nèi)核中的這種類型的循環(huán)通常稱為柵格步幅循環(huán)。

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

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

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

總結

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

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

關于作者

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

審核編輯：郭婷