CUDA 編程入門

更好的閱讀體驗(yàn)

CUDA 概述

CUDA 是 NVIDIA 推出的用于其發(fā)布的 GPU 的并行計(jì)算架構(gòu)，使用 CUDA 可以利用 GPU 的并行計(jì)算引擎更加高效的完成復(fù)雜的計(jì)算難題。

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

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

CUDA 編程模型基礎(chǔ)

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

CUDA 模型時一個異構(gòu)模型，需要 CPU 和 GPU 協(xié)同工作，在 CUDA 中一般用 Host 指代 CPU 及其內(nèi)存，Device 指代 GPU 及其內(nèi)存。CUDA 程序中既包含在 Host 上運(yùn)行的程序，也包含在 Device 上運(yùn)行的程序，并且 Host 和 Device 之間可以進(jìn)行通信，如進(jìn)行數(shù)據(jù)拷貝等操作。一般的將需要串行執(zhí)行的程序放在 Host 上執(zhí)行，需要并行執(zhí)行的程序放在 Device 上進(jìn)行。

CUDA 程序一般的執(zhí)行流程：

分配 Host 內(nèi)存，并進(jìn)行數(shù)據(jù)初始化

分配 Device 內(nèi)存，并將 Host 上的數(shù)據(jù)拷貝到 Device 上

調(diào)用 CUDA Kernel 在 Device 上進(jìn)行并行運(yùn)算

將運(yùn)算結(jié)果從 Device 上拷貝到 Host 上，并釋放 Device 上對應(yīng)的內(nèi)存

并行運(yùn)算結(jié)束，Host 得到運(yùn)算結(jié)果，釋放 Host 上分配的內(nèi)存，程序結(jié)束

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

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

Kernel 的層次結(jié)構(gòu)

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

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

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

從線程的組織結(jié)構(gòu)可以得知，一個線程是由(blockIdx, threadIdx)來唯一標(biāo)識的，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)上運(yùn)行的，但是單個 SM 上的運(yùn)算核心(cuda core)有限，這導(dǎo)致線程塊中的線程數(shù)是有限制的，因此在設(shè)置 grid 和 block 的 shape 時需要根據(jù)所使用的 Device 來設(shè)計(jì)。

如果要知道一個線程在 block 中的全局 ID，就必須要根據(jù) block 的組織結(jié)構(gòu)來計(jì)算，對于一個 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 實(shí)現(xiàn)向量加法

查看 Device 基本信息

在進(jìn)行 CUDA 編程之前，需要先看一下自己的 Device 的配置，便于之后自己設(shè)定 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("設(shè)備名稱與型號: %s\n", deviceProp.name);printf("顯存大小: %d MB\n", (int)(deviceProp.totalGlobalMem / 1024 / 1024));printf("含有的SM數(shù)量: %d\n", deviceProp.multiProcessorCount);printf("CUDA CORE數(shù)量: %d\n", deviceProp.multiProcessorCount * 192);printf("計(jì)算能力: %d.%d\n", deviceProp.major, deviceProp.minor); }

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

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

實(shí)現(xiàn) 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); // 分配內(nèi)存cudaMalloc((void**)&device_B, sizeof(float) * size); // 分配內(nèi)存cudaMalloc((void**)&device_C, sizeof(float) * size); // 分配內(nèi)存const float perBlockThreads = 192.0;cudaMemcpy(device_A, A, sizeof(float) * size, cudaMemcpyHostToDevice); // 將數(shù)據(jù)從 Host 拷貝到 DevicecudaMemcpy(device_B, B, sizeof(float) * size, cudaMemcpyHostToDevice); // 將數(shù)據(jù)從 Host 拷貝到 DeviceadditionKernelVersion<<<ceil(size / perBlockThreads), perBlockThreads>>>(device_A, device_B, device_C, size); // 調(diào)用 Kernel 進(jìn)行并行計(jì)算cudaDeviceSynchronize();cudaMemcpy(device_C, C, sizeof(float) * size, cudaMemcpyDeviceToHost); // 將數(shù)據(jù)從 Device 拷貝到 HostcudaFree(device_A); // 釋放內(nèi)存cudaFree(device_B); // 釋放內(nèi)存cudaFree(device_C); // 釋放內(nèi)存 } __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

總結(jié)

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

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