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
联系客服
