CUDA快速入门——学习总结
作者:互联网
**写在前面:**本文是作者总结的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 常用优化方法
关键在于隐藏延迟、提高吞吐量和线程利用率
- 归约
- 展开线程
- 一个线程块处理多个数据块
- 流技术
- 使用共享内存代替全局内存
标签:总结,入门,32,threadIdx,warp,线程,CUDA,tid,执行 来源: https://blog.csdn.net/weixin_43688140/article/details/118327701