GPU

CUDA-aware MPI

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#changes-from-previous-version

用处：透明的把Host主存orGPU内存里的数据发送到对方Host主存orGPU内存; （省去了手动在Host和GPU之间Copy）

UVA，统一地址空间，例如Host主存映射到0~32GB, GPU主存映射到32GB~48GB，所以MPI调用无需对地址来自哪里的参数，从地址上进行区分的；

 

CUDA的Context这个概念 资料

lazy创建的；所以要想不被首次创建所拖慢（影响计时统计），就预先调用cudaSetDevice() 提前创建context

 

卷积操作，如何用GPU加速？

 

很好的知识总结

 

1.

GPU比CPU的计算能力高接近1个数量级；存储带宽也高接近1个数量级；

快的原因：GPU把钱花在了计算单元ALU上；CPU把钱花在了cache和control上；

GPU设计目标就是计算密集，高度并行计算的任务；

尤其适合：

1. 数据并行计算(data-parallel computation)的任务--相同的程序，并行的，执行在很多数据上；因此不需要太多的控制逻辑；

2. 计算密集--计算比访存频繁；因此访存延迟可以隐藏在计算时间的背后，从而不需要复杂而大的cache;

如何利用好这么多计算单元，又不让编程难度加大？C语言，3个基本抽象：1. 线程组层次；2.shared memory; 3. barrier同步；

子问题--block; 同一个block里的threads们有shared memory; 无论GPU有多少SM, 都能被充分利用（每个SM串行处理多个blocks, SM们并行处理），SM多则每个SM少算几个block，SM少则每个SM多算几个block；user不需要关心GPU有多少SM, GPU简单扩展SM个数即可使符合该编程模型的程序加速。

2.

Grid-->Block-->Thread; Grid和Block可以是1~3维的；block里的threads数量上限，受制于一个SM上的shared-memory，因为同一个block只能被调度到同一个SM上；

block们的执行顺序是不确定的，程序员必须保证block顺序不影响最终计算结果；

一个block内，threads访问shared-memory(对标CPU的L1/L2 cache), 通过__syncthreads()来barrier (block-wise)

内存架构：1. thread local memory; 2. 每个block内被所有threads共享的shared memory; 3. global memory (还有texture memory, constant memory)（对标DRAM)

生命周期：1: thread; 2. block执行期间保证一直有效；3. 同一个App执行期间，保证一直有效（不同kernel的生死不影响他）

device相当于host的协处理器；两者各自维护各自的内存空间(GPU的是global/texture/constant meory)；"Unified Memory"提供一致内存空间的功能

GPU的compute capability version（硬件），每个主版本号对应一个名字(3:Kepler, 4:Maxwell, ...); 和CUDA（软件平台）的版本号不是一回事儿

3.

3.2 Device memory: 老三样:cudaMalloc(), cudaFree(), cudaMemcpy(); 分配2D/3D内存时建议使用cudaMallocPitch()和cudaMalloc3D()拷贝时建议使用cudaMemcpy2D()和cudaMemcpy3D()，考虑了padding以获取最优访问性能；cudaMemcpyToSymbol是往device的global内存里的变量赋值的操作（例如声明的变量__device__ float devData;）；

3.2.3 分块矩阵乘法；每个block负责一个目标矩阵的块，每个thread负责这个块的一个元素的计算；kernel里每次从global memory加载进来两个子矩阵进到shared memory(每个thread搬A的一个元素和B的一个元素），sync一下，计算乘法，sync一下；好几轮之后子矩阵计算完了，就写入结果（一个数）到global memory；注意：2个sync都有用，第二个如果没有的话，快的thread去下一轮搬数据了，慢的thread还在上一轮计算数据就乱了！

3.2.4 Page-locked host memory(又叫pinned memory): cudaHostAlloc(),cudaFreeHost(),cudaHostRegister(); 3大好处：1. 主机的这种内存和GPU之间的copy，可以和kernel计算并行执行；2. Mapped-memory, 主机的这种内存，可以映射到GPU内存地址空间里；3. 在有front-side总线的机子里，这种内存和GPU之间的copy更快，可设成Write-combining内存（不进行L1和L2 cache, 写的更快，读的更慢）

3.2.5 并行

- 同一个GPU上可以多个Kernel并行执行；（属于不同CUDA context的kernel之间不能并行）

- 主存和GPU内存之间的copy，和kernel执行，可以并行（需要主存开成pinned); GPU内存之间的copy，和kernel执行，可以并行；GPU内存之间的copy，和主存和GPU内存之间的copy，可以并行

- 主存->GPU内存， GPU内存->主存，可以并行（需要主存开成pinned)

- Stream:一组命令序列，GPU确保同一Stream的命令能顺序执行；默认是所有线程使用默认stream(0号stream)，如果加编译选项或者宏就每个线程默认使用自己单独的stream; 

把默认stream和"stream们各个操作提交顺序造成的性能差异"讲的很透彻的文章： https://devblogs.nvidia.com/how-overlap-data-transfers-cuda-cc/   默认stream(NULL)就是他要执行必须之前提交的所有stream的操作都完成，他执行的时候后面所有stream提交的操作都不能开始；对于host而言，cudaMemcpy是阻塞的，kernel提交是非阻塞的（可以和接下来的CPU计算并行），对device而言，kernel执行也是阻塞的，他执行完，后面的操作才能开始（同stream内的）；想copy和计算并行，必须用cudaMemcpyAsync版本，且pinned host内存；提交完一个stream的3个操作再提交下一个stream的3个操作那个例子，讲的很清楚，按提交顺序queue到不同的engine队列里（必须按先入先出来执行），且满足同一个stream内的顺序；Hyper-Q之后这些trick没用了。

 

Local Memory就是当Register不够的时候不得已而使用的，其实就是Global Memory的一部分；即Kernel里的变量，优先往Register里放，放不下就只好退化到Local Memory里；