CUDA軟體系統知識
阿新 • • 發佈：2018-11-15
本博文是根據中科大資訊學院譚立湘老師的課件加上自己的理解整理出來的
************************************************************************************
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多thread
cudaMemcpy(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;
}