Chapter 1. 存储器访问优化
1、使用Pinned Memory
Pinned Memory又称页锁定存储器（Page-locked memory）。Pinned Memory由于“禁止”了系统的页交换功能，所以可以更快的在host和device之间传输。与一般GPU变量的空间分配不同，Pinned Memory通过特定的cudaHostAlloc函数分配的空间。如果想将已经分配的变量变成Pinned Memory，则需要通过cudaHostRegister函数。
通过Pinned Memory我们可以实现CUDA提供的多种高效的功能，其中对性能优化最有帮助的就是异步传输。Pinned Memory允许实现host和device之间数据的异步传输，这样程序将可以并行的处理计算与传输。在程序中以流水线（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会导致系统资源对其他程序不足，从而影响系统整体性能。
2、global memory的访存优化
对属于计算密集型的CUDA程序，访存global memory的访存优化是十分必要的，但是在实际的实施过程中需要耗费较多的工作量，所以从实际角度出发，优化过程可以归结为：序号连续的线程应近可能访问地址连续的存储空间（即连续的数组元素）。
3、shared memory的访存优化
Shared memory有多个等大小的内存模块组成，这些模块被称为bank。这些bank可以被同时访问。当若干个线程同时访问一个bank时，这些访存指令将会串行执行，这种情况被称为bank冲突。 计算能力1.X的设备Shared memory被分为16个bank，计算能力2.X及以上的设备被分为32个bank。下面以2.X设备为例介绍共享存储器的优化。假设有数组__shared__ float a[32][32]，则数据在bank中的存储为：
若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。
（3）block中线程数的参考值是128—256（实际依赖于具体物理硬件）。
（4）对于调用__syncthreads()的核函数，可以将一个拥有较大线程数的block拆分成多个（3、4个）拥有较少线程的block。
占有率是衡量参数设置正确与否的参考值，NVIDIA在TOOLKIT中提供了CUDA_Occupancy_Calculator.xls来计算占有率。但是高GPU占有率并不意味着核函数拥有最好的计算效率。
