CUDA中的GPU编程模型

一、相关名词

硬件相关：

SP(Streaming Processor,流处理器)：最基本的处理单元（等价于一个ALU），又叫做CUDA Core

SM(Streaming Multiprocessor,流多处理器)：又叫GPU Core，是多个SP和其他一些硬件资源组成的，下面以GTX 580为例讲述SM

一个GTX 580 的SM包含32个SP（CUDA core），2个取址译码单元（Fetch/Decode），和Shared Memory和Execution contexts组成。一个SM上驻扎block个数，warp个数，thread个数上限都是由硬件决定的，另外一个限制条件就是register和shared memory等资源的是否充足，不充足的情况下，驻扎个数达不到上限。

另外，这里SP的个数32与warp的尺寸相等只是巧合，在低版本架构中SP个数可能为8（一个warp要4次才能执行一遍），高版本Kepler架构的GTX 680中，一个SMX中有192个SP，还有48个SP的情况（这种情况下如何执行未深入了解，可能会让我们疑惑，根据规则：一个SM同一时刻只有一个warp在运行，且warp的尺寸为32，，SP小于等于32个我们都能理解，SP没有一个处于限制状态甚至很紧张，那么多于32个SP的SM中，其他SP难道都在闲置中？？？）。参考：http://15418.courses.cs.cmu.edu/spring2016/lecture/gpuarch/slide_056，上述规则可能并不适用，一个SM每个clock选择的warp数，跟Fetch/Decode单元的个数有关，上图中明明白白写着Two warps are selected each clock(decode, fetch and execute two warps in parallel)，SP总共32个，可能是两个warp轮换使用。

warp：warp是GPU执行程序时的调度单位，目前CUDA的warp大小为32，warp是由连续的32个thread组成的，warp可以理解成一种硬件机制，这种机制保证这32个连续thread能够天然同步执行（无时无刻不在执行相同的代码，无需软件同步，遇到divergent 则等待），warp之间切换时零开销的，因为每一个thread的相关数据，执行环境都有自己的专属空间（比如register file， execution contexts），切换warp只需要让sp关联到另一块thread的专属空间即可，没有重新加载的过程。另外，同一时刻一个SM上只能有一个warp在工作。

软件相关：

thread（线程）：每一个thread都有自己专属的register，执行上下文（还不确定，可能有误），local memory等硬件资源，每一个thread都对应一个SP，但是每一个sp则对应多个thread，thread是一个软件概念，它的硬件资源除了SP不是专属的，其他基本都是专属的，一个SM驻扎成千上万个thread，并不是说真的有成千上万个SP(也可能真的有这么多，但绝对不会有thread的数量多)，而是我们将寄存器register，local memory，执行上下文等硬件资源平均分配成 成千上万份，SP只需要在这些不同的thread所对应的专属资源间以warp为单位来回切换即可。

block（线程块）：多个thread组成一个block，block的维度和大小可以通过调用kernel函数时来指定，一个block中的所有thread必须驻扎在同一个SM上，且可以通过软件同步（通过调用__syncthreads()），但不同block之间的thread无法同步，并且核函数中开辟一块shared memory是一个block中所有thread共享的。

grid：多个block组成一个grid，尺寸在调用kernel函数时指定。

二、函数的声明

	Executed on the:	Only callable from the:
__global__ void KernelFunc()	device	host
__device__ float DeviceFunc()	device	device
__host__ float HostFunc()	host	host

__global__：返回值类型必须是void.

__global__和device函数：

    尽量不用递归（不鼓励）

    不要使用静态变量

    少用malloc（现在允许但不鼓励）

    小心通过指针实现的函数调用

三、内存模型

变量声明	存储器	作用域	生命期
必须是单独的自动变量而不能是数组	register	thread	kernel
自动变量数组	local	thread	kernel
__shared__ int sharedVar;	shared	block	kernel
__device__ int globalVar;	global	grid	application
__constant__ int constantVar;	constant	grid	application

四、  SIMD和SIMT的疑问

    

这里我们以这样一个只包含8个SP的SM为例，说明自己的看法，为仔细认证，可能有误：

上图中上图的SM可能驻扎了许多warp，但执行每一个warp时都需要轮换4次，因为warp有32个thread，而SM只有8个SP；那么在执行一个warp内的32个thread时，可以称之为SIMD，单指令多数据，但是一个GPU肯定有多个SM，有些芯片允许不同的SM允许执行的指令流不同（周斌老师说这是符合芯片的设计哲学的），但是在GPU中不同的SM上执行相同的指令流，可以称之为SIMT，单指令多线程。自己的理解，可能错的离谱。

参考：http://hustcat.github.io/gpu-architecture/