CUDA编程优化(存储器访问优化,指令优化,参数优化,)

  1. 云栖社区>
  2. 博客>
  3. 正文

CUDA编程优化(存储器访问优化,指令优化,参数优化,)

cuda_study 2018-03-06 15:41:25 浏览2845
展开阅读全文

Chapter 1. 存储器访问优化

1、使用Pinned Memory

Pinned Memory又称页锁定存储器(Page-locked memory)。Pinned Memory由于“禁止”了系统的页交换功能,所以可以更快的在hostdevice之间传输。与一般GPU变量的空间分配不同,Pinned Memory通过特定的cudaHostAlloc函数分配的空间。如果想将已经分配的变量变成Pinned Memory,则需要通过cudaHostRegister函数。

通过Pinned Memory我们可以实现CUDA提供的多种高效的功能,其中对性能优化最有帮助的就是异步传输。Pinned Memory允许实现hostdevice之间数据的异步传输,这样程序将可以并行的处理计算与传输。在程序中以流水线(stream 流)作为实现方式,如下面这段程序:

 

size=N*sizeof(float)/nStreams;

for (i=0; i<nStreams; i++) {

offset = i*N/nStreams;

cudaMemcpyAsync(a_d+offset, a_h+offset, size, dir, stream[i]);

kernel<<<N/(nThreads*nStreams), nThreads, 0, stream[i]>>>(a_d+offset);

}

 

程序将需要计算的内容分为nStreams个部分,每个部分通过一个流来实现数据传输与计算。通过多个流的重叠,使得不同流之间的传输与计算可以重叠。

当然,Pinned Memory也是有缺点的,由于受到系统资源的限制,过多创建Pinned Memory会导致系统资源对其他程序不足,从而影响系统整体性能。

 

2global memory的访存优化

对属于计算密集型的CUDA程序,访存global memory的访存优化是十分必要的,但是在实际的实施过程中需要耗费较多的工作量,所以从实际角度出发,优化过程可以归结为:序号连续的线程应近可能访问地址连续的存储空间(即连续的数组元素)。

 

3shared memory的访存优化

Shared memory有多个等大小的内存模块组成,这些模块被称为bank。这些bank可以被同时访问。当若干个线程同时访问一个bank时,这些访存指令将会串行执行,这种情况被称为bank冲突。 计算能力1.X的设备Shared memory被分为16bank,计算能力2.X及以上的设备被分为32bank。下面以2.X设备为例介绍共享存储器的优化。假设有数组__shared__ float a[32][32],则数据在bank中的存储为:


e2a83f6d0d841961b6c9f4edf87c0a2538482113

warp中的线程以连续或者以交错而没有交集的方式读取一行数据,则这些访存指令不存在冲突。当这些线程按照如图所示的方式读取同属于bank 0的一列数据,则这些访存指令会串行执行。在发生冲突的情况下有一种特例,那就是当warp中所有线程同时访问同一个bank中的同一个元素时,会被自动优化成广播(broadcast)。在计算能力2.X以上的设备将这种情况进一步优化为,当warp中若干个线程访问同一个bank中的同一个元素时,这些访存指令只需进行一次,就可以使其他需要相同元素的线程也获得这个元素,这个过程称为多播(multicast)。

 

Chapter 2. 指令优化

1、如果程序对双精度没有要求,则应该使用单精度数(float)常量、变量和单精度计算函数。通过在编译时加入-use_fast_math选项,可以将核函数中使用的单精度计算函数替换为CUDA内部实现的高速版本(但是会影响计算精度)。

2、确保warp内的线程执行相同的指令,尽量减少在程序中使用分支语句。在线程执行时,同一个warp内的线程如果需要执行不同的指令,那么线程将顺序的执行所有的指令,并将线程中不需要执行的那一侧分支指令设置为无效。但是,不同warp之间执行不同的分支并不会使线程执行所有的指令。

3、在循环中不要使用__syncthreads()

 

Chapter 3. 参数设置

核函数参数的设置决定的运行时执行的方式和可以获得的资源。NVIDIA提供了如下参考规则以供。

1)每个block的线程数应该是warp的倍数。这样可以避免计算资源的浪费,同时有助于合并访存。

2)当运行包含多个block的核函数时,可以将线程数设置为64

3block中线程数的参考值是128256(实际依赖于具体物理硬件)。

4)对于调用__syncthreads()的核函数,可以将一个拥有较大线程数的block拆分成多个(34个)拥有较少线程的block

占有率是衡量参数设置正确与否的参考值,NVIDIATOOLKIT中提供了CUDA_Occupancy_Calculator.xls来计算占有率。但是高GPU占有率并不意味着核函数拥有最好的计算效率。


原文发布时间为:2016-6-24 10:21:58
原文由:NV开发者社区版主肖博士 发布,版权归属于原作者 
本文来自云栖社区合作伙伴NVIDIA,了解相关信息可以关注NVIDIA官方网站

网友评论

登录后评论
0/500
评论
cuda_study
+ 关注