Nvidia CUDA初级教程5 CUDA/GPU编程模型

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