物理概念:

streaming processor(sp): 最基本的处理单元。GPU进行并行计算,也就是很多个sp同时做处理。现在SP的术语已经有点弱化了,而是直接使用thread来代替。一个SP对应一个thread
Warp:warp是SM调度和执行的基础概念,通常一个SM中的SP(thread)会分成几个warp(也就是SP在SM中是进行分组的,物理上进行的分组),一般每一个WARP中有32个thread.这个WARP中的32个thread(sp)是一起工作的,执行相同的指令,如果没有这么多thread需要工作,那么这个WARP中的一些thread(sp)是不工作的
每一个线程都有自己的寄存器内存和local memory,一个warp中的线程是同时执行的,也就是当进行并行计算时,线程数尽量为32的倍数,如果线程数不上32的倍数的话;假如是1,则warp会生成一个掩码,当一个指令控制器对一个warp单位的线程发送指令时,32个线程中只有一个线程在真正执行,其他31个 进程会进入静默状态。)
streaming multiprocessor(sm):多个sp加上其他的一些资源组成一个sm, 其他资源也就是存储资源,共享内存,寄储器等。可见,一个SM中的所有SP是先分成warp的,是共享同一个memory和instruction unit(指令单元)。从硬件角度讲,一个GPU由多个SM组成(当然还有其他部分),一个SM包含有多个SP(以及还有寄存器资源,shared memory资源,L1cache,scheduler,SPU,LD/ST单元等等)

软件概念:

thread–>block–>grid:
CUDA在执行的时候是让host里面的一个一个的kernel按照线程网格(Grid)的概念在显卡硬件(GPU)上执行。当要执行这些任务的时候,每一个Grid又把任务分成一部分一部分的block,block再分线程来完成。每个Grid中的任务是一定的。二维线程块的索引关系为如下:
unsigned int xIndex = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int yIndex = blockDim.y * blockIdx.y + threadIdx.y;

GPU中的几个基本概念_数据

cuda内存模型

GPU中的几个基本概念_并行计算_02


每个 thread 都有自己的一份 register 和 local memory 的空间。

一组thread构成一个 block,这些thread 则共享有一份shared memory。

所有的 thread(包括不同 block 的 thread)都共享一份global memory、constant memory、和 texture memory。

不同的 grid 则有各自的 global memory、constant memory 和 texture memory。

每一个时钟周期内,warp(一个block里面一起运行的thread,其中各

个线程对应的数据资源不同(指令相同但是数据不同)包含的thread数量是有限的,现在的规定是32个。一个block中含有16个warp。所以一个block中最多含有512个线程.每次Device(就是GPU)只