其他分享
首页 > 其他分享> > CUDA快速入门——学习总结

CUDA快速入门——学习总结

作者:互联网

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

1. 基本概念

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

  1. 启动一个核函数,即调用一个网格(grid)
kernal<<<block, thread>>> (args);
  1. grid中包含很多线程块(blocks),block中包含很多线程(threads),block和thread的数量在初始化核函数时配置,两者数量影响核函数的性能
  2. 每连续的32个threads组成线程束(warp)
  3. 当核函数启动时,所有blocks被分配到不同的流多处理器(SM)上
  4. 受限于硬件资源,SM的并发block数和线程束数存在上限

1.2 线程束

  1. GPU执行核函数时,线程束warp是SM的最小执行单位
  2. 单指令多线程(SIMT)的执行方式:warp内所有的线程,同一时刻执行相同的指令,但处理的数据可以不同

1.3 线程束分化

产生分化的三个前提:

  1. 在同一个warp内,所有thread必须时刻同步
  2. 不同的thread可以处理不同的内存数据
  3. 存在判断语句(循环也是判断的一种)

表现:受限于同步要求,同一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和内存地址的对应关系
关键:

  1. 全局线程号和局部线程号的对应转化关系
  2. 二维向一维的转化关系
    常用手段:
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 并行归约

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

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

2.4 常用优化方法

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

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

标签:总结,入门,32,threadIdx,warp,线程,CUDA,tid,执行
来源: https://blog.csdn.net/weixin_43688140/article/details/118327701