CUDA 学习优化思路

参考http://m.blog.csdn.net/abcjennifer/article/details/42528569

CPU中memory<—>L3 Cache传输带宽为20GB/s, 除以64bytes/line得到传输记录速度约300M line/s，约为300M*8= 2.4G double/s. 一般地，浮点数操作需要两个输入+1个输出，那么loading 3个数（3 lines）的代价为 100Mflops。如果一个line中的全部8个variables都被用到，那么每秒浮点操作可以达到800Mflops。而CPU工作站典型为10 Gflops。这就要靠时间局部性来重用数据了。

Kepler GPU的cache line通常为128bytes（32个float or 16个double）。
数据传输带宽最高250GB/s
SMX的L2 cache统一1.5MB，L1 cache / shared memory有64KB
没有CPU中的全局缓存一致性，所以几乎没有两块block更新相同的全局数组元素。

GPU对浮点数的操作速度可达1Tflops。和上面CPU的计算类似，GPU中memory<—>L2Cache传输带宽为250GB/s, 除以128bytes/line得到传输记录速度约2G line/s，约为2G*16= 32G double/s. 一般地，浮点数操作需要两个输入+1个输出，那么loading 3个数（3 lines）的代价为 670Mflops。如果一个line中的全部16个variables都被用到，那么每秒浮点操作可以达到11Gflops。

这样的话每进行一次数据到device的传输需要45flops（45次浮点操作）才能达到500Gflops. 所以很多算法基本上不是卡在计算瓶颈，而是传输带宽。

5.1 global arrays

global arrays:

保存在/占用device memory由host code（非kernel部分code）声明一直存在，直到被host code释放因为所有block执行顺序不定，所以如果一个block修改了一个数组元素，其他block就不能再对该元素进行读写

5.2 global variables

声明前加标识符device，表示变量要放在device上了 e.g. device int reduction_lock=0;

shared（见4.6）和constant（见4.3）中至多有一个跟在device后面同时使用，标明用哪块memory空间，如果这两个都没写，则：

变量可以被grid内的所有threads读写与application同生死也可以定义为array，但是必须指定size可以在host code中通过以下函数读写：        1. cudaMemcpyToSymbol;             2. cudaMemcpyFromSymbol;        3. cudaMemcpy + cudaGetSymbolAddress

5.4 Register

默认一个kernel中的所有内部变量都存在register中64K 32-bit registers per SMXup to 63 registers per thread (up to 255 for K20 / K40)    这时有64K/63 = 1024个threads (256个threads for K20 / K40)up to 2048 threads (at most 1024 per thread block)    这时每个thread有32个registernot much difference between “fat” and “thin” threads如果程序需要更多的register呢？就“spill over”到L1 cache，这样访问速度就慢了，我们要尽量避免spill

5.5 Local Array

指kernel code中声明的数组。

简单情况下，编译器会将小数组float a[3]转换成3个标量registers：a0,a1,a2作处理复杂的情况，会将array放到L1（16KB），只能放4096个32-bit的变量，如果有1024个线程，每个线程只能分配放4个变量。

5.6 Shared Memory

前面加标识符shared e.g. shared int x_dim;

要占用thread block的shared memory space.要比global memory快很多,所以只要有机会就把global memory整成shared memory与block同生死thread block内所有threads共用（可读可写）啥时侯用呢？当所有threads访问都是同一个值的时候，这样就避免用register了

但是有问题就是，如果一个thread block有多个warp(上一篇blog中提到的概念，block中的thread每32个被分到一个warp，最后一个不足32个thread也没关系，同样形成一个warp)，各warp执行指令顺序是不定的，那么久需要线程同步机制，用指令__syncthreads(); 插入一个“barrier”，所有wrap执行到这个barrier之前没有thread/warp能够越过去。

Kepler GPU给L1 Cache + shared memory总共64KB，可以分为16+48，32+32，48+16；这个split可以通过cudaFuncSetCacheConfig（）或cudaDeviceSetCacheConfig（）设置，默认给shared memroy 48KB。这个具体情况看程序了。

参考：

1. CUDA C Programming Guide

2. different memory and variable types

3. CUDA 安装与配置

4. CUDA调试工具——CUDA GDB

5. GPU工作方式

6. Fermi 架构白皮书（GPU继承了Fermi的很多架构特点）

7. GTX460架构