CUDA快速入门——学习总结

2024-03-25 08:10:04

**写在前面：**本文是作者总结的cuda基础部分重点内容，一些琐碎的重要细节部分没有提及，建议有需要的同学自行学习。

1. 基本概念

1.1 核函数、网格、线程块与线程

启动一个核函数，即调用一个网格（grid）

kernal<<<block, thread>>> (args);

grid中包含很多线程块（blocks），block中包含很多线程（threads），block和thread的数量在初始化核函数时配置，两者数量影响核函数的性能
每连续的32个threads组成线程束（warp）
当核函数启动时，所有blocks被分配到不同的流多处理器（SM）上
受限于硬件资源，SM的并发block数和线程束数存在上限

1.2 线程束

GPU执行核函数时，线程束warp是SM的最小执行单位
单指令多线程（SIMT）的执行方式：warp内所有的线程，同一时刻执行相同的指令，但处理的数据可以不同

1.3 线程束分化

产生分化的三个前提：

在同一个warp内，所有thread必须时刻同步
不同的thread可以处理不同的内存数据
存在判断语句（循环也是判断的一种）

表现：受限于同步要求，同一warp内的部分线程可能闲置
理想状态：同一个warp中，线程全员工作，不存在闲置线程

示例代码：

__global__ void test(arg1, arg2)
{
	int temp[blockDimx.x];
	tid=threadIdx.x;
	g_tid = threadIdx.x+blockDimx.x*blockIdx.x;
	if (g_tid % 2 == 0) temp[g_tid] = 0;
	else temp[tid]=1;
	
}

上述代码中，偶数线程号的线程需将对应数组位置元素置0，奇数则置为1

由于同一线程束中，所有线程是同步的，当偶数号线程执行写入0时，奇数号线程不能同时写入1。此时，所有线程都在执行判断语句 if (g_tid % 2 == 0)。偶数号线程判断为true，奇数号线程为false，但此时不能直接执行else，而是需要等待偶数号线程执行完写入0的操作后，所有线程共同进入else判断语句。

以2个线程束，共72个线程为例。奇数置1偶数置0，同一时刻中，两个warp中只有一半的线程执行写入内存操作，另一半处于闲置状态。即使处于闲置状态，也是耗费资源的。

我们希望warp中所有的线程要么都执行，要么都不执行。warp是基本执行单位，当一个warp中的线程均不运行时，整个warp休息，不占用资源。

处理方式：程序仅做示例，没有实际使用价值

__global__ void new_test(arg1, arg2)
{
	int temp[blockDimx.x];
	tid=threadIdx.x;
	g_tid = threadIdx.x+blockDimx.x*blockIdx.x;
	if (g_tid<32) temp[g_tid] = 0;
	else temp[tid]=1;	
}

改进后的程序，让前32个线程将数组置0，后32个置1，得到的最终结果和上一个代码一致——“32个0，32个1”，仅是数组中位置不同，后续使用时通过改变数组索引方式即可。
此时，同样考虑2个warp共72个线程，同一时刻，第一个warp内32个thread全员运行，后一个warp全员休息，避免了线程束分化。

2. cuda程序常见方法

2.1 建立线程与内存地址的对应关系

由于SIMT的执行方式，允许不同的thread处理不同的数据块，实现的关键在于建立thread和内存地址的对应关系
关键：

全局线程号和局部线程号的对应转化关系
二维向一维的转化关系
常用手段：

int a=blockIdx.x*blockDim.x+threadIdx.x;
//threadIdx.x和threadIdx.y是局部线程坐标，不同block块内可能重复
//threadIdx.x+blockDimx.x*blockIdx.x，将局部线程号转化为全局线程号，唯一且不存在重复
int ix=threadIdx.x+blockDim.x*blockIdx.x;
int iy=threadIdx.y+blockDim.y*blockIdx.y;
//将二维坐标转化为一维线程号
int transform_in_idx=iy*nx+ix

注：二维转化为一维的思维方式很重要。计算机存储信息均是以一维方式存储的，二维只是帮助人类理解的逻辑形式。

2.2 隐藏延迟

关键：请求访问内存和计算之间相互独立。在访问内存时，可同时进行计算，反之亦然。
通过增加线程数可实现：线程数足够多时，当一个线程执行某项工作时，SM可调度其他线程工作。（类似人可以在电饭煲煮饭的时候，同时进行炒菜；而不是非要等到饭煮好了再炒菜）
例如：线程A请求并执行访问内存的工作，此时A从请求到访问结束需要一段时间，SM在这段时间里可以调动线程B，给B下达其他的命令。显然，线程数越多，隐藏延迟的效果越好。
但是，线程数不是越多越好。单个block中线程数过多，会导致单个线程分得的硬件资源过少，反而会降低性能。

在资源允许的情况下，提升线程的数量

2.3 并行归约

这一节很重要，过段时间另写一篇文章。
关键点：

避免线程束分化
尽量连续地访问内存
在线程数小于32时，可通过展开线程进一步优化性能

2.4 常用优化方法

关键在于隐藏延迟、提高吞吐量和线程利用率

归约
展开线程
一个线程块处理多个数据块
流技术
使用共享内存代替全局内存

码农公寓