CUDA 編程入門

更好的閱讀體驗

CUDA 概述

CUDA 是 NVIDIA 推出的用于其發布的 GPU 的并行計算架構，使用 CUDA 可以利用 GPU 的并行計算引擎更加高效的完成復雜的計算難題。

在目前主流使用的馮·諾依曼體系結構的計算機中，GPU 屬于一個外置設備，因此即便在利用 GPU 進行并行計算的時候也無法脫離 CPU，需要與 CPU 協同工作。因此當我們在說 GPU 并行計算時，其實指的是基于 CPU+GPU 的異構計算架構。在異構計算架構中，CPU 和 GPU 通過 PCI-E 總線連接在一起進行協同工作，所以 CPU 所在位置稱為 Host，GPU 所在位置稱為 Device，如下圖所示。

從上圖可以看到，GPU 中有著更多的運算核心，非常適合數據并行的計算密集型任務，比如大型的矩陣計算。

CUDA 編程模型基礎

在了解了 CUDA 的基本概念之后，還需要了解 CUDA 編程模型的基本概念以便于之后利用 CUDA 編寫并行計算程序。

CUDA 模型時一個異構模型，需要 CPU 和 GPU 協同工作，在 CUDA 中一般用 Host 指代 CPU 及其內存，Device 指代 GPU 及其內存。CUDA 程序中既包含在 Host 上運行的程序，也包含在 Device 上運行的程序，并且 Host 和 Device 之間可以進行通信，如進行數據拷貝等操作。一般的將需要串行執行的程序放在 Host 上執行，需要并行執行的程序放在 Device 上進行。

CUDA 程序一般的執行流程：

分配 Host 內存，并進行數據初始化

分配 Device 內存，并將 Host 上的數據拷貝到 Device 上

調用 CUDA Kernel 在 Device 上進行并行運算

將運算結果從 Device 上拷貝到 Host 上，并釋放 Device 上對應的內存

并行運算結束，Host 得到運算結果，釋放 Host 上分配的內存，程序結束

在第 3 步中，CUDA Kernel 指的是在 Device 線程上并行執行的函數，在程序中利用 __global__ 符號聲明，在調用時需要用 <<<grid, block>>> 來指定 Kernel 執行的線程數量，在 CUDA 中每一個線程都要執行 Kernel 函數，并且每個線程會被分配到一個唯一的 Thread ID，這個 ID 值可以通過 Kernel 的內置變量 threadIdx 來獲得。

__gloabl__ vectorAddition(float* device_a, float* device_b, float* device_c); // 定義 Kernel int main() {/*some codes*/vectorAddition<<<10, 32>>>(parameters); // 調用 Kernel 并指定 grid 為 10, block 為 32/*some codes*/ }

Kernel 的層次結構

Kernel 在 Device 執行的時候實際上是啟動很多線程，這些線程都執行 Kernel 這個函數。其中，由這個 Kernel 啟動的所有線程稱為一個 grid，同一個 grid 中的線程共享相同的 Global memory，grid 是線程結構的第一個層次。一個 grid 又可以劃分為多個 block，每一個 block 包含多個線程，其中的所有線程又共享 Per-block shared memory，block 是線程結構的第二個層次。最后，每一個線程(thread)有著自己的 Per-thread local memory。

下圖是一個線程兩層組織結構的示意圖，其中 grid 和 block 均為 2-dim 的線程組織。grid 和 block 都是定義為 dim3 類型的變量，dim3 可以看成是包含三個無符號整數(x, y, z)成員的結構體變量，在定義時，缺省值初始化為1。

dim3 grid(3, 2); dim3 block(5, 3); kernel<<<grid, block>>>(parameters);

從線程的組織結構可以得知，一個線程是由(blockIdx, threadIdx)來唯一標識的，blockIdx 和 threadIdx 都是 dim3 類型的變量，其中 blockIdx 指定線程所在 block 在 grid 中的位置，threadIdx 指定線程在 block 中的位置，如圖中的 Thread(2,1) 滿足：

threadIdx.x = 2; threadIdx.y = 1; blockIdx.x = 1; blockIdx.y = 1;

一個 block 是放在同一個流式多處理器(SM)上運行的，但是單個 SM 上的運算核心(cuda core)有限，這導致線程塊中的線程數是有限制的，因此在設置 grid 和 block 的 shape 時需要根據所使用的 Device 來設計。

如果要知道一個線程在 block 中的全局 ID，就必須要根據 block 的組織結構來計算，對于一個 2-dim 的 block($D_x$, $D_y$)，線程($x$, $y$)的 ID 值為 $x+y?D_x$，如果是 3-dim 的 block($D_x$, $D_y$, $D_z$)，線程($x$, $y$, $z$)的 ID 值為 $x+y?D_x+z?D_x?D_y$。

CUDA 實現向量加法

查看 Device 基本信息

在進行 CUDA 編程之前，需要先看一下自己的 Device 的配置，便于之后自己設定 grid 和 block 更好的利用 GPU。

#include <stdio.h> #include "cuda_runtime.h" #include "device_launch_parameters.h" int main() {cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, 0);printf("Device 0 information:\n");printf("設備名稱與型號: %s\n", deviceProp.name);printf("顯存大小: %d MB\n", (int)(deviceProp.totalGlobalMem / 1024 / 1024));printf("含有的SM數量: %d\n", deviceProp.multiProcessorCount);printf("CUDA CORE數量: %d\n", deviceProp.multiProcessorCount * 192);printf("計算能力: %d.%d\n", deviceProp.major, deviceProp.minor); }

Device 0 information:
設備名稱與型號: Tesla K20c
顯存大小: 4743 MB
含有的SM數量: 13
CUDA CORE數量: 2496
計算能力: 3.5
Device 1 information:
設備名稱與型號: Tesla K20c
顯存大小: 4743 MB
含有的SM數量: 13
CUDA CORE數量: 2496
計算能力: 3.5

其中第 12 行乘 192 的原因是我所使用的設備為 Tesla K20，而 Tesla K 系列均采用 Kepler 架構，該架構下每個 SM 中的 cuda core 的數量為 192。

實現 Vector Addition

#include <stdio.h> #include <time.h> #include <math.h> #include "cuda_runtime.h" #include "device_launch_parameters.h"const int LENGTH = 5e4; clock_t start, end; void vectorAdditionOnDevice(float*, float*, float*, const int); __global__ void additionKernelVersion(float*, float*, float*, const int); int main() {start = clock();float A[LENGTH], B[LENGTH], C[LENGTH] = {0};for (int i = 0; i < LENGTH; i ++) A[i] = 6, B[i] = 5;vectorAdditionOnDevice(A, B, C, LENGTH); //calculation on GPUend = clock();printf("Calculation on GPU version1 use %.8f seconds.\n", (float)(end - start) / CLOCKS_PER_SEC); } void vectorAdditionOnDevice(float* A, float* B, float* C, const int size) {float* device_A = NULL;float* device_B = NULL;float* device_C = NULL;cudaMalloc((void**)&device_A, sizeof(float) * size); // 分配內存cudaMalloc((void**)&device_B, sizeof(float) * size); // 分配內存cudaMalloc((void**)&device_C, sizeof(float) * size); // 分配內存const float perBlockThreads = 192.0;cudaMemcpy(device_A, A, sizeof(float) * size, cudaMemcpyHostToDevice); // 將數據從 Host 拷貝到 DevicecudaMemcpy(device_B, B, sizeof(float) * size, cudaMemcpyHostToDevice); // 將數據從 Host 拷貝到 DeviceadditionKernelVersion<<<ceil(size / perBlockThreads), perBlockThreads>>>(device_A, device_B, device_C, size); // 調用 Kernel 進行并行計算cudaDeviceSynchronize();cudaMemcpy(device_C, C, sizeof(float) * size, cudaMemcpyDeviceToHost); // 將數據從 Device 拷貝到 HostcudaFree(device_A); // 釋放內存cudaFree(device_B); // 釋放內存cudaFree(device_C); // 釋放內存 } __global__ void additionKernelVersion(float* A, float* B, float* C, const int size) {// 此處定義用于向量加法的 Kernelint i = blockIdx.x * blockDim.x + threadIdx.x;C[i] = A[i] + B[i]; }

Calculation on GPU version1 use 0.14711700 seconds.

參考資料

CUDA編程入門極簡教程

CUDA C Programming Guide

總結

以上是生活随笔為你收集整理的CUDA 编程入门的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。