CUDA程序优化
1. 确定任务中的串行和并行部分,选择合适的算法,将问题分为几个步骤看那些可以用并行来实现,确定要使用的算法
2. 按照算法确定任务和数据的划分方式,将每个需要并行实现的步骤映射为一个满足CUDA两层并行模式的内核函数,要尽量让每个sm上拥有至少六个活动的warp 核至少两个活动的线程块。
3. 编写一个能正常运行的程序作为优化的起点
4. 优化显存的访问,避免显存的带宽成为瓶颈。在显存带宽优化前,其他的优化不会产生明显的结果。
a) 显存优化的主要方法
b) 将可以采用相同的block和grid的维度实现几个kernel的合并减少对显存的访问
c) 尽量不要让私有变量分配到local memory
d) 为满足合并访问,采用cudaMallocPitch()或者cudaMalloc3D()分配显存。
e) 为满足合并访问,将数据进行对齐(使用__align)
f) 为满足合并访问,保证访问的首地址从16的整数倍开始,如果可能,尽量让每个线程一次读的数据字长都为32bit;
g) 在数据只会被访问一次的,并且满足合并访问时,考虑使用zerocopy
h) 在某些情况下,考虑存储器控制器负责不均衡造成的风情冲突
i) 使用拥有换成的常数存储器和纹理存储器
5. 优化指令流
a) 如果只需要少量的线程进行操作,一定要使用 类似‘if threaded < N’的方式,避免多个线程同时运行占用更多的时间或者产生错误的结果
b) 在不会出现不可接受的误差的情况下,采用CUDA算术指令集中的快速指令
c) 使用#unroll让编译器有效的展开循环
d) 采用原子函数实现更复杂的算法,并保证结果的正确性
e) 避免多余的同步
f) 如果不产生bank conflict的算法不会造成算法的效率的的下降或者非合并访问,究应该避免bank conflict
6. 资源均衡
a) 调整shared memory的使用量和register的使用量,保证更高的SM占用率,调整每个线程的数据处理量,shared memory,register的使用量
i. 使用括号明确变量的生存周期,使用shared memory存储变量
b)
7. 与主机之间的通信的优化
a) 尽量减少通信
b) 使用cudaMallocHost()分配主机的内存,获得更大的带宽
c) 一次缓存较多的数据,在一并传输,可以获得较大的带宽
d) 用流和异步处理来隐藏通信时间
e) 使用zerocopy 和write-combined memory提高可用带宽
任务划分原则
1. 按照输入划分
a) 输入很多而输出很少时,如规约,直方图
b) 输入满足合并条件,但是输出的位置是随机的,或者输出时要进行显存原子操作
2. 按照输出划分
a) 输入参数很少,但是输出很多时,如函数随机数发生器,block内每个线程的输入与其他的线程公用,比如卷积,滤波 每个线程的输入和周围的输入有公共的部分,此时应先用合并访问的形式,将数据放入shared memory,在由每个线程计算一定数量的输出
b) 输入数据在存储器中的位置是随机的,而输出数据时可以满足合并访问
3. 一个block可以进行一维带状划分也可以进行二维棋盘划分 但是要保证每个block里的线程是32的整数倍
block和grid的维度设计
1. 每个SM中要只是保持6个active warp 才能有效的隐藏延时。每个SM中的shared memory 的容量是16KB
2. ??????每个block中到底最大能运行512 条线程还是768个线程????丫的
3. 最后让每个block保持在64~~256之间,block的设计主要是为了避免整数的除法和求模运算应该让blockDim.x为16或16的整数倍,提高global memory和shared memory的访问效率。
4. 先确定block在确定grid grid中的数量应该是SM数量的几倍,让每个SM满负荷运转
存储器访问优化
1. 主机----设备访问优化
a) PCI双向8GB/S
b) Pinned memory
c) 异步执行(内核函数和存储器拷贝函数有同步和异步两个版本)
d) 使用不同的流之间的异步执行,使流之间的传输和运算能够同时执行。更好的利用GPU资源
全局存储器优化
1. half-warp对内存进行装载或者存储操作的时候,如果满足访问条件只要一次读写就可以满足访问。Global是一个每行有64byte对其的段
2. half-warp 访问段的首地址必须是每个线程所访问字长的64倍,如果在此half-warp种有几个线程不需要读取数据,也会进行并行访问
3. 计算能力为1.0和1.1的设备上访问字长满足32bit,64bit,128bit在1.2和1.3的设备上字长可以为8,16,32,64,128bit
4. 连续但是没有对其的访问,连续对其的范围
5. 矩阵是按行存储的,当一个half-warp没有按照行去访问元素时,性能就会严重下降,在1.0和1.1的设备上当stride不为1时性能严重下降,在1.2和1.3的设备上当stride=2时性能折半
共享存储器的访问优化
1. 共享存储器和bank conflict
a) 为了满足并行访问的需要,共享存储器被划分为大小相等的能够被同时访问的模块,称为bank 可以互不干扰的工作,对于n个bank上的n个地址的访问能够同时进行,也就增加了带宽
b) 每个bank大小32bit 访问时间一个周期
c) 每个warp被划分为两个,SM的共享存储器被划分为16个bank,只有在同一个half-warp中的线程间才有可能出现bank conflict
d) 和全局存储器不同在shared memory中只要地址不共用,线性、乱序随机访问都可以并行,当stride=3时也可以,只要stride和16没有公约数就可以,实际写的过程中可以自己画一下
e) 当16个线程去访问一个地址时不会出现bank conflict 数据会被广播到16个线程中
并行程序思想:
(1)并行程序主 要是将循环进行拆解,CUDA里面的每一条线程可以对应为每一次循环,在这里,要求每次循环都不依赖于上一次循环的结果。举个例子:将有100个元素的数 组中的每个元素加上1,再赋值到另一个数组中去的时候,串行就一个for循环,并行的话就一条语句 B[tid] = A[tid]+1, tid为线程号,每个线程都取一条数据,加上1之后赋值到B中对应的位置。其实并行就这么简单!
(2)关于编CUDA程序建议大家都先了解一下GPU的存储结构,《GPU高性能计算之CUDA》这本书用来入门还是不错的。我写程序的时候那个存储结构图一直就贴在电脑旁边。
(3)当对程序有些理解之后,下一步一定要搞明白的是并行里面常用的规约算法和前缀和算法,这两个算法几乎是随处可用。一定要理解这两个算法。
(4)关于调试,没有什么好的调试方法,最好的就是不断的输出,将显存里面觉得有问题的空间拷贝出来看看
(5)一定要理解GPU中线程是怎样的顺序执行的。要理解Wrap是个什么东西。
最后最后,一定不要用串行的头脑去考虑并行的问题,一旦程序并行,什么执行顺序都是可能的,该同步的地方同步,该原子操作的地方原子操作。先保证程序正确了,然后再考虑程序的改进~
cuda核心思想之规约:
//求32*128矩阵的每行之和
//kernel
__global__ void RowSum(float* A, float* B){
int bid = blockIdx.x;
int tid = threadIdx.x;
__shared__ s_data[128];
//read data to shared memory
s_data[tid] = A[bid*128 + tid];
__synctheads(); //sync
for(int i=64; i>0; i/=2){
if(tid<i)
s_data[tid] = s_data[tid] + s_data[tid+i] ;
__synctheads();
}
if(tid==0)
B[bid] = s_data[0];
}
int main(){
RowSum<<<32,128>>>(A,B);
}
规约算法(reduction)用来求连加、连乘、最值等,应用广泛。每次循环参加运算的线程减少一半,上面的代码还要优化。
main中的调用表示线程的组织是32个block,每个一维block有128个thread。GPU架构中每个block的线程数和shared memory的大小是受限的。
转自http://hi.baidu.com/superkiki1989/blog/item/1429ce2a290825ee98250a2c.html
联系客服