CUDA 的更簡單介紹

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

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

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

我們開始吧！

從簡單開始

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

#include <iostream>
#include <math.h> // 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 上，您可能需要命名可執行文件添加. exe 并使用 .dd 運行它。）

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

首先，我只需要將我們的 add 函數轉換成 GPU 可以運行的函數，在 CUDA 中稱為 內核 。要做到這一點，我所要做的就是把說明符 __global__ 添加到函數中，它告訴 CUDA C ++編譯器，這是一個在 GPU 上運行的函數，可以從 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__ 函數被稱為 果仁 ，在 GPU 上運行的代碼通常稱為 設備代碼 ，而在 CPU 上運行的代碼是 主機代碼 。

CUDA 中的內存分配

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

最后，我需要 發射 內核，它在 add() 上調用它。 CUDA 內核啟動是使用三角括號語法指定的。我只需要在參數列表之前將它添加到對 CUDA 的調用中。

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

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

還有一件事：我需要 CPU 等到內核完成后再訪問結果（因為 CUDA 內核啟動不會阻塞調用的 CPU 線程）。為此，我只需在對 CPU 進行最后的錯誤檢查之前調用 cudaDeviceSynchronize() 。

以下是完整的代碼：

#include <iostream>
#include <math.h>
// 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

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

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

介紹一下！

我認為找出運行內核需要多長時間的最簡單的方法是用 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 大約需要半秒鐘的時間。

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

把線撿起來

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

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

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

__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 函數沒有太大變化。事實上，將 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 毫秒）。讓我們繼續取得更高的表現。

走出街區

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

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

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

我還需要更新內核代碼來考慮線程塊的整個網格。 threadIdx.x 提供了包含網格中塊數的 gridDim.x 和包含網格中當前線程塊索引的 blockIdx.x 。圖 1 說明了使用 CUDA 、 gridDim.x 和 threadIdx.x 在 CUDA 中索引數組（一維）的方法。其思想是，每個線程通過計算到其塊開頭的偏移量（塊索引乘以塊大小： blockIdx.x * blockDim.x ），并將線程的索引添加到塊內（ 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];
}

更新的內核還將 stride 設置為網格中的線程總數（ blockDim.x * gridDim.x ）。 CUDA 內核中的這種類型的循環通常稱為 柵格步幅循環 。

將文件另存為&[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 倍的加速，從運行多個街區的所有短信 K80 ！我們在 K80 上只使用了 2 個 GPUs 中的一個，但是每個 GPU 都有 13 條短信。注意，我筆記本電腦中的 GeForce 有 2 條（較弱的）短信，運行內核需要 680us 。

總結

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

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

練習

為了讓你繼續前進，這里有幾件事你可以自己嘗試。請在下面的評論區發表你的經歷。

1. 瀏覽 工具包文件 。如果您還沒有安裝 CUDA ，請查看 快速入門指南 和安裝指南。然后瀏覽 編程指南 和 最佳實踐指南 。還有針對各種體系結構的調整指南。
2. 在內核中使用 printf() 進行實驗。嘗試打印出部分或所有線程的 threadIdx.x 和 blockIdx.x 的值。它們是按順序打印的嗎？為什么或者為什么不呢？
3. 在內核中打印 threadIdx.y 或 threadIdx.z （或 blockIdx.y ）的值。（同樣適用于 blockDim 和 gridDim ）。這些為什么存在？如何讓它們采用 0 以外的值（ 1 表示尺寸）？
4. 如果您可以訪問 基于 Pascal 的 GPU ，請嘗試在其上運行 add_grid.cu 。性能比 K80 的結果好還是差？為什么？（提示：閱讀關于 Pascal 的 PageMIG 定量引擎和 CUDA 8 統一內存 API 的信息）關于這個問題的詳細答案，請參閱 CUDA 初學者的統一內存 一文。

從這里到哪里去？

我希望這篇文章有助于提高 CUDA 的興趣，并且你有興趣在你自己的計算中學習更多的東西并應用 CUDA C ++。如果您有任何問題或意見，請使用下面的評論部分聯系您。

我計劃在這篇文章之后繼續提供更多的 CUDA 編程材料，但為了讓您暫時保持忙碌，您可以繼續閱讀一系列舊的介紹性文章（我計劃在將來根據需要進行更新/更換）：

• 如何在 CUDA C ++中實現性能度量
• 如何查詢 CUDA C ++中的設備屬性和處理錯誤
• 如何優化 CUDA C ++中的數據傳輸
• 如何在 CUDA C ++中重疊數據傳輸
• 如何在 CUDA C ++中高效訪問全局內存
• 在 CUDA C ++中使用共享內存
• CUDA C ++中的一種高效矩陣轉置
• CUDA C ++中的有限差分方法，第 1 部分
• CUDA C ++中的有限差分方法，第 2 部分
• 用 CUDA 在一個周末內加速光線追蹤

還有一系列的儀器。

您還有興趣從 Udacity 和 NVIDIA 注冊 CUDA 編程在線課程 。

關于 CUDA C ++和其他 GPU 計算主題，這里有很多關于 NVIDIA 并行 Forall 開發者博客 的內容，所以環顧四周！

?