gpu编程注意事项（二） 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中对应的位置。其实并行就这么简单！ cuda核心思想之规约： //求32*128矩阵的每行之和 __shared__ s_data[128]; for(int i=64; i>0; i/=2){ int main(){ 规约算法(reduction)用来求连加、连乘、最值等，应用广泛。每次循环参加运算的线程减少一半，上面的代码还要优化。 main中的调用表示线程的组织是32个block，每个一维block有128个thread。GPU架构中每个block的线程数和shared memory的大小是受限的。 转自http://hi.baidu.com/superkiki1989/blog/item/1429ce2a290825ee98250a2c.html |
|
|
来自： yangshiquan > 《CUDA》
