Industry news

CUDA基本优化方法

一、基于编程模型和执行模型的优化方法

1.选取合适的gridDim和blockDim

blockDim最好为32的整数倍:因为执行指令的基本单位为线程束,线程束内的所有线程统一执行广播下来的命令,而线程束的线程数量基本为32。当block被分到SM中去,其会被划分为多个线程束,若blockDim!=线程束内线程数整数倍,则会造成线程的浪费。

2.减少存在分支的if

因为线程束中的所有线程执行同一条的指令,若出现存在出现分支的if,即线程束内的部分线程符合if条件,部分符合else条件,因cuda编译不具备分支预测能力,则会在一个指令执行if,下一个指令执行else,两个指令中都有部分不符合条件的线程空闲,从而造成浪费,如下述代码所示:

 

线程束内奇数线程(threadIdx.x为奇数)会执行else,偶数线程执行if,分化很严重。对此,我们可以更改if条件,使整个线程束内都符合if或else即可:

 

可以得到相同但是错乱的结果C,这个顺序是无所谓的,可以后期调整。

3.循环展开

减少循环判断的次数,本质上也是减少分支

 
 

最新的cuda应该不需要手动展开了,在for前加上#pragma unroll 循环次数命令即可。

二、基于内存的优化方法

内存结构

1.合并访问和对齐访问

核函数运行时需要从全局内存(DRAM)中读取数据,只有两种粒度:128字节、32字节。粒度可以理解为最小单位,也就是核函数运行时每次读内存,哪怕是读一个字节的变量,也要读128字节,或者32字节。
合并访问是指所有线程访问连续的对齐的内存块,对于L1 cache,内存块大小支持32字节、64字节以及128字节为单位读取数据。前提是,访问必须连续,并且访问的地址是以32字节对齐。
我们把一次内存请求——也就是从内核函数发起请求,到硬件响应返回数据这个过程称为一个内存事务。当一个内存事务的首个访问地址是缓存粒度(32或128字节)的偶数倍的时候:比如二级缓存32字节的偶数倍64,128字节的偶数倍256的时候,这个时候被称为对齐内存访问
在这里插入图片描述
规约中可以通过设计累加的位置来实现合并对齐访问,提高效率
在这里插入图片描述

 
2.利用缓存来进行交叉访问(非合并的访问)

在某些情况下,无法避免进行交叉的读或写,如求矩阵的转置。这种情况下可以利用从缓存中读取的高效性,进行交叉内存访问(虽然按照列读是不合并的,但是使用一级缓存加载过来的数据在后面会被使用,虽然一级缓存一次读取128字节的数据,其中只有一个单位是有用的,但是剩下的并不会被马上覆盖,粒度是128字节,但是一级缓存的大小有几k或是更大,这些数据很有可能不会被替换,所以,我们按列读取数据,虽然第一行只用了一个,但是下一列的时候,理想情况是所有需要读取的元素都在一级缓存中,这时候,数据直接从缓存里面读取),从而让写结果到全局内存的过程是合并访问(写的过程不走缓存)。

 
3.基于共享内存的优化

共享内存是在他所属的线程块被执行时建立,线程块执行完毕后共享内存释放,线程块和他的共享内存有相同的生命周期。
共享内存读写速度较快,且可编程,故可用它来存储中间变量,或者先将内存中的数据读到共享内存中来,再做后续的操作。

 

共享内存按列读取不会导致交叉访问那么严重的延迟,代码中申请了共享内存来进行交叉访问,避免了交叉的全局内存访问,还对申请的动态内存做了填充,防止读取冲突(有疑问)。

CONTACT US

Contact: 蓝狮-蓝狮娱乐-蓝狮在线站

Phone: 13800000000

Tel: 400-123-4567

E-mail: admin@youweb.com

Add: Here is your company address

平台注册入口