Nvidia CUDA初級教程5 CUDA/GPU編程模型
視頻：https://www.bilibili.com/video/BV1kx411m7Fk?p=6
講師：周斌
本節內容：
- CPU和GPU互動模式
- GPU線程組織模型（需要不停強化）
- GPU存儲模型
- 基本的編程問題
CPU與GPU交互
- 各自的物理內存空間
- 通過PCIE總線互連
- 交互開銷較大
GPU的存儲器層次架構（硬件）
下圖是OpenCL的圖
GPU存儲層次訪存速度
| Registers | Dedicated HW | Single cycle |
| Shared Memory | Dedicated HW | Single cycle |
| Local Memory | DRAM | *slow* |
| Global Memory | DRAM | *slow* |
| Constant Memory | DRAM, cached | 1-10s-100s cycles, depending on cache locality |
| Texture Memory | DRAM, cached | 1-10s-100s cycles, depending on cache locality |
| Instruction Memory (invisible) | DRAM, cached |
GPU線程組織模型
GPU線程組織模型圖示
線程組織架構說明
- 一個 Kernel 具有大量線程
- 線程被劃分為線程塊 blocks
- 一個 block 內部的線程共享 shared memory
- 可以同步 _syncthreads()
- Kernel 可以啟動一個 grid，包含若干線程塊
- 用戶設定
- 線程和線程塊具有唯一標識
GPU線程映射關系
GPU和內存線程等關系
線程和存儲器的整體關系
編程模型
常規意義的GPU用于處理圖形圖像
操作于每個像素，每個像素的操作都類似，可以應用 SIMD
SIMD
SIMD (single instruction multiple data)，可以理解為是數據的并行分割。
SIMT
single instruction multiple thread
- GPU版本的SIMD
- 大量線程模型獲得高度并行
- 線程切換獲得延遲掩藏
- 多個線程執行相同的指令流
- GPU上大量線程承載和調度
CUDA編程模式：Extended C
可以理解為就是一個擴展的C語言，增加了一些關鍵字、API、函數調用、存儲位置聲明等
Declspecs
global, device, shared, local, constant
實例：
__device__ float filter[N];
__global__ void convolve(float* image) {}
__shared__ float region(M);
關鍵詞
threadIdx, blockIdx
實例：
region[threadIdx] = image[i];
Intrinsics
__syncthreads
實例：
__syncthreads()
運行時API
Memory, symbol, execution, management
實例：
// 分配GPU顯存
void* myimage = cudaMalloc(bytes)
函數調用
// 100個block，每個block 10個thread
convolve<<<100, 10>>> (myimage);
CUDA函數聲明
| 執行位置 | 調用位置 | |
|---|---|---|
__device__ float DeviceFunc() | device | device |
__global__ void KernelFunc() | device | host |
__host__ float HostFunc() | host | host |
__global__定義一個 kernel 函數- 入口函數，CPU上調用，GPU上執行
- 必須返回 void
__device__和__host__可以同時使用（同時修飾同一個函數）
)
