CUDA学习笔记(四)——CUDA性能

转自:http://blog.sina.com.cn/s/blog_48b9e1f90100fm5h.html

四、CUDA性能

       CUDA中的block被划分成一个个的warp,在GeForce8800GTX上,一个warp有32个线程。若不够32个线程,则padding相应数目的线程。Warp中的线程ID是连续且递增的。对于二维组织的线程来说,先把threadIdx.y为0的线程按照threadIdx.x从小到大排,然后把threadIdx.y为1的线程按照threadIdx.x从小到大的顺序排列成warp。对于三维组织的线程来说,先排列threadIdx.z为0的二维线程,再排列threadIdx.z为1的二维线程,以此类推。

       任何时刻,硬件都只能一次选择执行一个warp。

       下面2个图是执行元素总和的操作,不同的算法实现,其效率不同。第二种方法使得线程没有分支。      

 CUDA学习笔记(四)鈥斺擟UDA性能

 

CUDA学习笔记(四)鈥斺擟UDA性能

      

 

 

 

全局内存性能

       Memory coalescing技术能够大大提高全局内存访问的速度。硬件会检测线程所访问的位置是否是在全局内存中的相邻位置,如果是相邻位置,硬件会把这些访问合并成一个相邻访问。

CUDA学习笔记(四)鈥斺擟UDA性能

CUDA学习笔记(四)鈥斺擟UDA性能

 

 

 

SM资源的动态划分

1.      thread限制:一个SM最多768个

2.      block限制:一个block最多512个thread,如果block有256个thread,根据线程数目的限制,一个SM最多只能有3个block了。

3.      register限制:一个SM最多8192个register

资源之间的分配和权衡对性能影响非常大。

例如,假设在矩阵乘中,block事16*16线程的,每个thread使用10个register,那么一个block需要使用2560个register,那么3个block就要用8120个register,小于8192的寄存器数目限制。但是如果再加一个block的话,寄存器数目就到达10240了,超过了寄存器数目的限制。假设每个thread使用11个register,那么一个block需要使用16*16*11=2816个寄存器,3个block需要8448个寄存器,超过了寄存器限制。所以一个SM只能容纳2个block了。那么SM上的线程数目就从768减少为512了。线程并行度减少了1/3!

另一个例子:假设在全局内存load和使用之间有四条独立的指令,在G80中,每条指令要4cyle,所以4条指令需要16个cycle。而200内存延迟需要调度200/16=14个warps来保证执行单元的利用率。若把指令数从4增加到8,那么就只需要200/(4*8)=7个warps。也就是说,即使我们把block数目从3减少到2,因此warps数目从24(768/32)减少到16(512/32),我们有足够的warps来完全地利用执行单元,性能最后还提高了!参考[RyooCGO08]

 

数据预取

       Prefetch techniques are often combined with tiling to simultaneously address the problems of limited bandwidth and long latency

循环展开

       减少分支指令和循环计数器的更新。

线程粒度

       线程粒度通常是程序调优中一个重要的标准。

 

2009-11-9

做了一些实验,问题:

如果不设置足够多的thread,处理器是怎样处理的,实验结果是否会出错?

若thread数目不足以计算矩阵乘,那么剩下的就不会计算

 

编程:

gridDim.x, gridDim.y, gridDim.z表示Grid的三个维度

blockDim.x, blockDim.y, blockDim.z表示block的三个维度

对于GeForce8800来说,由于一个block的shared memory大小为 16K字节,所以对于float型(8个字节)的矩阵乘来说,tile 块最大不能超过2048

 

posted @ 2013-10-22 17:13  qingsun_ny  阅读(971)  评论(0编辑  收藏  举报