本博文是根據中科大信息學院譚立湘老師的課件加上自己的理解整理出來的

************************************************************************************

NVIDIA在2007年推出CUDA這個統一計算架構

CUDA的基本思想是支持大量的線程級并行，并在硬件中動態地調度和執行這些線程

?

CUDA軟件體系可以分為三層結構

• CUDA函數庫（CUDA Library）

• CUDA運行時API（Runtime API）https://blog.csdn.net/qq_41598072/article/details/81030272

• CUDA驅動API（Driver API）

Difference between the driver and runtime APIs

https://docs.nvidia.com/cuda/cuda-driver-api/driver-vs-runtime-api.html#driver-vs-runtime-api

CUDA軟件環境：

CUDA支持Windows、Linux、MacOS三種主流操作系統，支持CUDA C及CUDA Fortran等多種語言。無論使用何種語言或接口，指令最終都會被驅動程序轉換成PTX（ParallelThread Execution，并行線程執行，CUDA架構中的指令集，類似于匯編語言）代碼，交由GPU核心計算。CUDA最主要的包含兩個方面：ISA指令集架構與硬件計算引擎；實際上是硬件和指令集。見下圖中的綠色部分，CUDA 架構的組件組成是：
(1)NVIDIA GPU中的并行計算引擎；
(2)對硬件初始化、配置的OS內核級支持；
(3)用戶模式的驅動，為開發者的PTX 指令集架構(ISA，Instructionset architecture)

Kernel

Kernel函數：

Kernel函數是指為GPU設備編譯的一個函數。也就是一個編譯好的、在GPU上并行運行的計算函數。Kernel在GPU上以多個線程的方式被執行。
運行在GPU上的CUDA并行計算函數稱為kernel函數（內核函數）。一個完整的CUDA程序是由一系列的設備端kernel函數并行部分和主機端的串行處理部分共同組成的。這些處理步驟會按照程序中相應語句的順序依次執行，滿足順序一致性。

CUDA編程中的術語:

• Host：宿主，指CPU，系統的CPU。負責啟動應用程序，運行程序的串行部分，將程序的并行、計算密集的部分offload到GPU上運行，并最終返回程序的運行結果。

• Device：設備，指GPU，CPU的協處理器。負責程序的并行、計算密集部分的處理，并將處理結果返回給Host。

Block：線程塊
——執行Kernel的一組線程組成一個線程塊。(一個Kernel只做同一件事)
一個線程塊最多可包含1024個并行執行的線程，線程之間通過共享內存有效地共享數據，并實現線程的通信和柵欄同步。
線程ID：線程在線程塊中的線程號（唯一標識）
基于線程ID的復雜尋址，應用程序可以將線程塊指定為任意大小的二維或三維數組，并使用2個或3個索引來標識每個線程。

• 對于大小是（Dx，Dy）的二維線程塊，索引為（x，y）的線程的線程ID為（x+y*Dx）

• 對于大小為（Dx，Dy，Dz）的三維線程塊，索引為（x，y，z）的線程的線程ID為：

（x+y*Dx+z*Dx*Dy）

Grid：線程塊組成的線程網格（最多2^32 個blocks）
執行相同Kernel、具有相同維數和大小的線程塊可以組合到一個網格中。這樣單個Kernel調用中啟動的線程數就可以很大。同一網格中的不同線程塊中的線程不能互相通信和同步。
Grid 是一個線程塊陣列，執行相同的內核，從全局內存讀取輸入數據，將計算結果寫入全局內存。

Block ID：線程塊ID
線程塊ID是線程塊在Grid中的塊號。實現基于塊ID的復雜尋址，應用程序可以將Grid指定為任意大小的二維數組，并用2個索引來標識每個線程塊。對于大小為（Dx，Dy）的二維線程塊，索引為（x，y）的線程塊的ID為（x+y*Dx）。現已支持三維
Wrap：線程束
一個線程塊中連續的固定數量（32）的線程組。
將線程塊中的線程劃分成wrap的方式是：每個wrap包含線程ID連續遞增的32個線程，從線程0開始遞增到線程31。

Stream：
CUDA的一個Stream表示一個按特定順序執行的GPU操作序列。諸如kernel啟動、內存拷貝、事件啟動和停止等操作可以排序放置到一個Stream中。
一個Stream包含了一系列Grids，并且可以多個Stream并行執行。

在CUDA 架構下，GPU芯片執行時的最小單位是thread。
若干個thread可以組成一個線程塊（block）。一個block中的thread能存取同一塊共享內存，可以快速進行同步和通信操作。
每一個block 所能包含的thread 數目是有限的。執行相同程序的block，可以組成grid。不同block 中的thread 無法存取同一共享內存，因此無法直接通信或進行同步。
不同的grid可以執行不同的程序（kernel）。

Grid是由線程塊組成的網格。每個線程都執行該kernel，應用程序指定了Grid和線程塊的維數，Grid的布局可以是一維、二維或三維的。
每個線程塊有一個唯一的線程塊ID，線程塊中的每個線程具有唯一的線程ID。同一個線程塊中的線程可以協同訪問共享內存，實現線程之間的通信和同步。
每個線程塊最多可以包含的線程的個數為1024個，線程塊中的線程以32個線程為一組的Wrap的方式進行分時調度。每個線程在數據的不同部分并行地執行相同的操作。

CUDA處理流程：

在CUDA 的架構下，一個程序分為兩個部份：Host 端和Device 端。Host 端是指在CPU 上執行的部份，而Device 端則是在GPU上執行的部份。Device端的程序又稱為kernel函數。
通常Host 端程序會將數據準備好后，復制到GPU的內存中，再由GPU執行Device 端程序，完成后再由Host 端程序將結果從GPU的內存中取回。
CPU 存取GPU 內存時只能通過PCI-E 接口，速度有限。

• 1)從系統內存中復制數據到GPU內存

• 2)CPU指令驅動GPU運行；

• 3)GPU 的每個CUDA核心并行處理

• 4)GPU 將CUDA處理的最終結果返回到系統的內存

?

CUDA編程模型：

• CPU作為主機端只能有一個

• GPU作為設備端可以有多個

• CPU主要負責邏輯處理

• GPU負責密集型的并行計算

? ? ? ??完整的CUDA程序包括主機端和設備端兩部分代碼，主機端代碼在CPU上執行。
設備端代碼（kernel函數）運行在GPU上。其中一個kernel函數對應一個grid，每個grid根據需要配置不同的block數量和thread數量。

CUDA包含兩個并行邏輯層：block層和thread層。
在執行時block映射到SM
thread映射到SP（Core）
如何在實際應用程序中高效地開發這兩個層次的并行是CUDA編程與優化的關鍵之一。
Stream > Grid > Block > Warp > Thread
學校? ? ? ? 年級? ? ? 班級? ? ? ?小組? ? ? ? 學生

?

Kernel的啟動參數

• cuda程序執行流程：

• 單顯卡只需要考慮紅色的，多顯卡要七步曲

• 1)cudaSetDevice(0); //獲取設備；只有一個GPU時或默認使用0號GPU時可以省略
  2)cudaMalloc((void**) &d_a,sizeof(float)*n); //分配顯存
  3)cudaMemcpy(d_a,a,sizeof(float)*n,cudaMemcpyHostToDevice); //數據傳輸
  4)gpu_kernel<<<blocks,threads>>>(***); //kernel函數
  5)cudaMemcpy(a,d_a,sizeof(float)*n,cudaMemcpyDeviceToHost); //數據傳輸
  6)cudaFree(d_a); //釋放顯存空間
  7)cudaDeviceReset( ); //重置設備；可以省略

完整的向量點積CUDA程序

/*
a = [a1, a2, …an], b = [b1, b2, …bn]
a*b = a1*b1 + a2*b2 + … + an*bn
*/#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#include <malloc.h>
#define N 10
__global__ void Dot(int *a, int *b, int *c) //聲明kernel函數
{__shared__ int temp[N]; // 聲明在共享存儲中的變量temp[threadIdx.x] = a[threadIdx.x] * b[threadIdx.x];//__syncthreads();if (0 == threadIdx.x){//Kernel函數中利用threadIdx.x 獲得線程索引號//threadIdx是內建變量，它指定block內thread索引號int sum = 0;for (int i = 0; i < N; i++)sum += temp[i];*c = sum;printf("sum Calculated on Device:%d\n", *c);}
}void random_ints(int *a, int n)
{for (int i = 0; i< n; i++)*(a + i) = rand() % 10;
}int main()
{int *a, *b, *c;int *d_a, *d_b, *d_c;int size = N * sizeof(int);cudaMalloc((void **)&d_a, size);cudaMalloc((void **)&d_b, size);cudaMalloc((void **)&d_c, sizeof(int));a = (int *)malloc(size); random_ints(a, N);b = (int *)malloc(size); random_ints(b, N);c = (int *)malloc(sizeof(int));printf("Array a[N]:\n");for (int i = 0; i < N; i++) printf("%d ", a[i]);printf("\n");printf("Array b[N]:\n");for (int i = 0; i < N; i++) printf("%d ", b[i]);printf("\n\n");cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);Dot << <1, N >> >(d_a, d_b, d_c); //單block多threadcudaMemcpy(c, d_c, sizeof(int), cudaMemcpyDeviceToHost);int sumHost = 0;for (int i = 0; i < N; i++)sumHost += a[i] * b[i];printf("sum Calculated on Host=%d\n", sumHost);printf("Device to Host: a*b=%d\n", *c);free(a); free(b); free(c);cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);return 0;
}

?

?

?

?

?

?

?

?

總結

以上是生活随笔為你收集整理的CUDA软件系统知识的全部內容，希望文章能夠幫你解決所遇到的問題。

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