Nvidia CUDA初級教程6 CUDA編程一

視頻：https://www.bilibili.com/video/BV1kx411m7Fk?p=7
講師：周斌

GPU架構概覽

• GPU特別使用于：
  • 密集計算，高度可并行計算
  • 圖形學
• 晶體管主要被用于：
  • 執行計算
  • 而不是
    • 緩存數據
    • 控制指令流

圖中分別是CPU、GPU各個部件所占的芯片面積。可以看到，CPU芯片中大量部分是緩存和控制邏輯，而GPU中則絕大部分都是計算單元。

CUDA編程相關簡介

CUDA的一些信息

• 層次化線程集合
• 共享存儲
• 同步

CUDA術語

主機端和設備端

• HOST - 主機端，通常指CPU
  • 采用ANSI標準C語言編程
• Device - 設備端，通常指GPU（數據可并行）
  • 采用ANSI標準C的擴展語言編程 （CUDA C）
• HOST 和 Device 擁有各自的存儲器
• CUDA編程
  • 包括主機端和設備端兩部分代碼

核

• Kernel 數據并行處理函數
  • 通過調用 Kernel 函數在設備端創建輕量級的線程，線程由硬件負責創建并調度

類似于 OpenCL 的 shader？

• 核函數會在 N 個不同的 CUDA 線程上并行執行

  // 定義核函數 __global__ void VecAdd(float* a, float* B, float* C) {int i = threadIdx.x;C[i] = A[i] + B[i]; }int main() {// ...// 在N個線程上調用核函數VecAdd<<<1, N>>>(A, B, C); }

CUDA程序的執行

CUDA程序執行的流程大體上是這樣的：當我們在CPU端的代碼是串行執行的（這里簡單地認為指令在CPU上串行執行），當遇到需要并行大量處理數據時，會調用核函數在GPU上進行計算，計算完成后將結果返回給CPU。

線程層次

• Grid - 一維或多維線程塊（block）
  • 一維或二維
• Block - 一維線程
  • 一維，二維或三維
  • 一個 Grid 中的每個 Block 的線程數是一樣的
  • Block 內部的每個線程可以：
    • 同步 Synchronize
    • 訪問共享存儲器 shared memory

一個線程可以類比為一個員工，一個 block 是一個科室，grid 是整個公司。

線程ID

每一個線程都有一個索引：threadIdx

• 一維 Block Thread ID == Thread Index
• 二維 Block (Dx, Dy)
  • 索引為 (x, y) 的 Thread ID == x + yDy
• 三維 Block (Dx, Dy, Dz)
  • 索引為 (x, y) 的 Thread ID == x + yDy + zDxDy

代碼實例

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i = threadIdx.x;int j = threadIdx.y;C[i][j] = A[i][j] + B[i][j]; }int main() {int numBlocks = 1;dim3 threadsPerBlock(N, N);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); }

每個 Block 中的線程的索引是二維的，這在我們處理二維數據（矩陣）時可以很方便地進行對應。

線程數

Thread Block 線程塊

• 線程的的集合
  • G80 和 GT200：多達512個線程
  • Fermi：多達1024個線程
• 位于相同的處理器核（相同的SM）
• 共享所在核的存儲器

塊索引

• 塊索引：blockIdx
• 維度：blockDim
  • 一維，二維或三維

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {int i = blockIdx.x * blockDim.x + threadIdx.x;int j = blockIdx.y * blockDim.y + threadIdx.y;if (i < N && j < N)C[i][j] = A[i][j] + B[i][j]; }int main() {// ...dim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); }

例如 N = 32

• 每個塊有16x16個線程（跟N無關）
  • threadIdx([0, 15], [0, 15])
• Grid 里面有 2x2 個線程塊 block
  • blockIdx([0, 1], [0, 1])
  • blockDim = 16

int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y;

i = [0, 1] * 16 + [0, 15]

線程塊之間

線程塊彼此之間獨立執行

• 任意順序：并行或串行
• 被任意數量的處理器以任意順序調度
  • 處理器的數量具有可擴展性

一個塊內部的線程

一個塊內部的線程有一些重要的特性：

• 共享容量有限的低延遲存儲器 (shared memory)
• 同步執行
  • 合并訪存
  • __syncThreads()
    • barrier - 塊內線程一起等待所有的線程都
    • 輕量級線程

CUDA內存傳輸

主機端與設備端

CUDA內存傳輸

• device 端代碼可以：

  • 讀寫該線程的 registers
  • 讀寫該線程的local memory
  • 讀寫該線程所屬的塊的 shared memory
  • 讀寫grid的 global memory
  • 只讀grid的 constant memory

• host 端代碼可以：

  • 讀寫grid的 global memory 和 constant memory

• host 可以從 device 往返傳輸數據

  • global memory 全局存儲器
  • constant memory 常量存儲器

CUDA內存傳輸函數

• 在設備端分配 global memory：cudaMalloc()
• 釋放存儲空間 cudaFree()

float* Md; int size = Width * Width * sizof(float); cudaMalloc((void**)&Md, size); //... cudaFree(Md);

注意這里的指針 Md 是指向 device（GPU）上的存儲空間。

• 內存傳輸：cudaMemcpy()
  • host to host
  • host to device
  • device to host
  • device to device

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

示例：矩陣相乘 Matrix Multiply

矩陣相乘簡介

• 向量
• 點乘
• 行優先或列優先
• 每次點乘結果輸出一個元素

矩陣相乘CPU實現

void MatrixMulOnHost(float* M, float* N, float* P, int width) {for (int i=0; i<width; ++i) {for (int j=0; j<width; ++j) {float sum = 0;for (int k=0; k<width; ++k) {float a = M[i * width + k];float b = N[k * width + j]:sum += a * b;}P[i * width + j] = sum;}} }

CUDA算法框架

三步走：

int main(void) {// 1 分配device空間// 2 在GPU上，并行計算MatrixMulOnDevice(M, N, P, width);// 3 將結果拷貝回CPU，并釋放device空間return 0; }

偽代碼如下：

void MatrixMulOnDevice(float* M, float* N, float* P, int Width) {int size = Width * Width * sizeof(float);// 1cudaMalloc(Md, size);cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMalloc(Nd, size);cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);cudaMalloc(Pd, size);// 2 調用cuda核函數，并行計算// 3cudaMemcpy(P. Pd, size, cudaMemcpyDeivceToHost)cudaFree(Md); cudaFree(Nd); cudaFree(Pd); }

CUDA C 實現

矩陣相乘樣例

目前版本矩陣相乘的問題

• 在上述算法實現中最主要的性能問題是什么？
  • 訪存
• 主要的限制是什么？
  • 訪存帶寬

總結

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

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