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__可以同时使用（同时修饰同一个函数）
