CUDA on Platform 学习笔记3--线程层次

CV温故知新

已于 2022-07-10 17:28:34 修改

阅读量1.2k

点赞数 1

分类专栏： CUDA教程文章标签：学习人工智能

于 2022-07-06 15:23:10 首次发布

本文链接：https://blog.csdn.net/lansebingxuan/article/details/125636045

版权

CUDA教程专栏收录该内容

6 篇文章 1 订阅

订阅专栏

本文是参加2022 CUDA on Platform 线上训练营学习笔记，感谢NVIDIA各位老师的精彩讲解!

1. CUDA线程层次

GPU在管理线程的时候是以block为单元调度到SM上执行，每个block中以warp作为一次执行的单位，每个warp包括32个线程。

3个线程层次

Thread: sequential execution unit
1）所有线程执行相同的核函数
2）并行执行
3）thread是最基本单元，32个thread组成一个warp，一个 warp 对应一条指令流。
Thread Block: a group of threads
1）执行在一个Streaming Multiprocessor (SM)
2）同一个Block中的线程可以协作
3）block内部的线程可以共享存储单元，SM是硬件层次，一个硬件SM可以执行多个blook，一个block只能在一个SM中执行
Thread Grid: a collection of thread blocks
1）一个Grid当中的Block可以在多个SM中执行
2）线程网格是由多个线程块组成，每个线程块又包含若干个线程
三者的关系图如下图所示

布局设置

在CUDA中我们通常采用dim3这个数据类型来指定grid或者block的大小，它实际上是一个结构体，有x, y, z三个变量, 分别代表三个维度上的大小，这里至少指定一个变量x，其他变量若不指定，默认是1。例如下图中，线程网格在水平和竖直方向上分别有3个和2个block，总共有6个block；线程块在水平和竖直方向上分别有5个和3个线程，因此线程网格和block的维度可以表示为：dim3 grid(3,2,1), block(5,3,1)
在这里插入图片描述
实际执行时，变量的含义和设置如下，
Built-in variables:

threadIdx.[x y z]，如上图中Thread(0,0)
是执行当前kernel函数的线程在block中的索引值
blockIdx.[x y z]
是指执行当前kernel函数的线程所在block，在grid中的索引值，如上图中block(1,1)
blockDim.[x y z]
表示一个block中包含多少个线程
gridDim.[x y z]
表示一个grid中包含多少个block
在程序中有时会看到类似下面的设置，尖括号结构中第一个？表示grid的索引，第二个？表示block的索引。

HelloFromGPU <<<?, ?>>>();
<<<grid, block>>>

程序示例

__global__ void add( int *a, int *b, int *c ) {
c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}
add<<<1,4>>>( a, b, c);

上图是一段程序，程序实际在设备上运行如下
在这里插入图片描述
需注意：GPU和CPU多线程意义不同，CPU上的多线程通常执行不同的操作，而GPU上的多个线程执行相同的指令，并且同时运行。从GPU显存读取数据时用不同的线程索引区分。

软件和硬件对应图

在这里插入图片描述
GPU硬件中存在三个层次——core、SM、Device，CUDA分3个软件层次thread，block，grid正好和硬件一一对应。grid是GPU的调度单位，block是SM的调度单位，thread/warp是CUDA core的调度单位。

设置block、grid层次优点

增加block和grid层次是由GPU的硬件架构决定的，若没有block，当同时进行同步，通信时，所有的核芯都要等待，即使增加硬件核芯数量计算效率也不会提升。而以 block 的形式组织 warp能实现可扩展性。全局同步开销太大，但是如果让一个 block 中的 warp 通过 barrier 同步来通信开销就小得多，这样设计出来的程序性能才有可能随着 ALU 数量的增加线性增长。

CUDA实际执行流程

加载核函数
将Grid分配到一个Device
根据<<<…>>>内的执行设置的第一个参数，Giga threads engine将block分配到SM中。一个Block内的线程一定会在同一个SM内，一个SM可以有很多个Block。
根据<<<…>>>内的执行设置的第二个参数，Warp调度器会调用线程。
Warp调度器为了提高运行效率，会将每32个线程分为一组，称作一个warp。
在某个时刻，每个Warp指令会被SM内部的某些单元执行(Warp具体分配给多少个SP是不确定的，可以根据deviceQuery查询，例如计算能力7.5可能给1组（16个）SP，连续2个周期执行)

2. CUDA线程索引

索引计算

在这里插入图片描述
如上图所示，第一行是一个warp，32个thread，block中将其分为4组，每组8个，threadIdx.x代表组内的索引，blockIdx.x代表组索引，blockDim.x代表每组线程个数，本例中是8个，因此，计算的索引如下：

int index = threadIdx.x + blockIdx.x * blockDim.x;
	          =      5      +     2      * 8;
	          = 21;

CPU执行改为GPU执行

下图是CPU执行一段代码的示例，在for循环中串行执行加法。
在这里插入图片描述
在GPU执行同样功能代码时，可以多个线程并行执行加法，提高效率，下图是在不同线程上分配不同数据的示例：

如下图所示，改写成GPU执行的代码分5个步骤：

Allocate GPU Memories（图中紫色框标注1）
Copy data from CPU to GPU（图中蓝色框标注2）
Invoke the CUDA Kernel（图中绿色框标注3）
Copy result from GPU to CPU（图中橙色框标注4）
Release GPU Memories（图中绿色框标注5）
这样，在CPU中循环n次执行的代码在GPU的多个线程中同时分别执行一次，效率提升。

3. 线程分配

是否有最优值？
这个值的设置没有一个最优值，需要根据实际运行的程序进行调整，例如设置
block_size = 128;
grid_size = (N + block_size - 1) / block_size;
grid_size需要确保分配的大小是充足的。
可以申请的最大线程数？
每个block可以申请的最大线程数是有限制的，比如下图在devicequery中可以查到

比如上图中，倒数第三行每个block可以设置的最大线程数是1024，倒数第二行一个block的每个维度可以设置的最大线程数分别是1024,1024,64，但是3个维度线程数的乘积不能大于1024，比如x维度设置为1024，则后两个维度只能设置成1,1。
实际中应该申请多少个线程？
前面提过，32个thread组成一个warp，实际在GPU中，硬件也是以32为单位组成一个warp，warp是一个SM执行的基本单元。因此，申请1个线程和申请32个线程硬件都会分配一个warp。因此，申请的线程数最好是32的倍数，这样可以最有效的利用线程，减少不必要的浪费。
若数据过大，线程不够用怎么办？
例如下图，一共分配2*4=8个线程，实际有32个数据

可以参考下面的代码执行，用for循环。蓝色框是一次循环时8个线程执行的数据，红色框代表索引值为0的线程处理的数据，循环4次，可以完成32个数据运算。

__global__ add(const double *x, const double *y, double *z, int n)
{
	int index = blockDim.x * blockIdx.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
	for(; index <n; index +=stride)
		z[index] = x[index] + y[index];
}