GPU经典博客案例总结(Mark Harris的博客)

A. kernel里写成loop，比每个thread干一下就走要好：https://devblogs.nvidia.com/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/

扩展性：能适应更多的数据，尤其是当数据比Grid里threads个数多的时候；能tune grid/block大小（block数目尽量是SM个数的整数倍）；

性能：loop，能让线程创建和销毁的开销均摊开来，线程私有数据和shared-memory初始化等操作被均摊；

debug: 调成1个block1个thread, 就可以串行执行把结果作为baseline，并行版本对不对有了参照物，也便于看printf的输出；

移植性：便于移植到CPU版本；

B. 在app退出之前（或者不再调GPU计算和访存以后），应该调cudaDeviceReset()扫尾，否则profiling有可能得到不准确的数据（profiling数据被系统缓存在buffer里满了flush到磁盘，如果不调这个扫尾，可能app退出的时候buffer就扔了）

C. nvcc先将源代码编译成虚拟汇编（中间结果），再编译成特定卡型上的二进制码；后者可以编译阶段用户来指定，也可以运行时由CUDA来根据当时的卡型来编译(拖慢程序运行的启动速度）；

CUDA提供的两种降低启动时间开销的方式：1. nvcc允许用户指定第二阶段编译到多种卡型，运行的时候自动挑合适的（没有合适的只好当场编译）；2. 运行的时候编译的二进制码，会被cache到磁盘上，下次只要driver版本不变就直接加载cache里的来运行；

可能遇到的问题：1. 程序用了大部头的库的时候，可能造成代码超出cache默认大小导致不cache，每次启动都重新编译非常慢；此时可以nvcc指定第二阶段生成二进制码的卡型（就可以不在运行阶段生成了），或者把cache大小配置大些； 2. 二进制代码cache在默认目录可能慢，可以设成cache在更快的目录（比如用了SSD的磁盘）

D. CUDA for Python: Numba   https://devblogs.nvidia.com/numba-python-cuda-acceleration/

E. 命令行profiling工具nvprof: 可以profile自己的C-CUDA, Numba(python for CUDA), openACC; 可以remote profile, ssh到远程机器上运行nvprof, 输出日志文件，copy到本地，使用NVIDIA Visual Profiler来在UI里可视化的看。
F. Unified Memory: https://devblogs.nvidia.com/unified-memory-in-cuda-6/

cudaMallocManaged(); 分配完了CPU和GPU都可以拿来直接用，系统自动负责copy;  kernel启动之后要调cudaDeviceSynchronize()之后才能使用该内存，否者可能kernel还没干完活儿；

系统使用用的时候才copy策略，最大化性能；用户使用stream和cudaMemcpyAsync手工tune到极致的代码肯定更高性能，因为系统目前没那么聪明的知道数据什么时候被谁用；

和UVA不是一回事儿：UVA提供了一致的地址空间，指针的值不用改来改去，GPU上访问一个在host上开辟的内存的时候，是通过PCI总线访问的（但不整页copy到GPU主存里来）；UVA在cudaMemcpy的时候不用指定双方是在GPU/host上(host内存需要提前pin住)；

struct里带char*的，老方法需要cudaMemcpy三次（复制struct, 复制char*的内容, 填struct的char*指针）；现在直接分配struct和char*为managed内存即可；

链表等里面有指针的复杂数据结构，解放了；

C++的class，重载new和delete用managed内存来放自己，构造和析构函数里也用managed内存来放char*；则该class可在CPU和GPU之间无缝使用；

G. 用环境变量来控制app能够看到那些GPU卡：

CUDA_VISIBLE_DEVICES；用在这些情况下：1. 限制app只在某些卡上跑（或者程序里面一上来先查卡，使用跟自己兼容的卡）； 2.只有binary没有代码的时候，限制该app在兼容的卡上跑；3.程序用了Unified Memory，但是该机器不支持GPU卡的P2P访问，可以用该方式限制该程序只在某卡上跑；

H. warp内可以用shuffle来访问其他threads的register, 只需1个cycle，比shared-memory的3步要快（write, sync, read) https://devblogs.nvidia.com/cuda-pro-tip-kepler-shuffle/

I. 计算occupancy：https://devblogs.nvidia.com/cuda-pro-tip-occupancy-api-simplifies-launch-configuration/ 系统在runtime自动估算最优（近似最优）block大小和同时启动的block数, 可用在thrust等库编写者（无法知道用户传进来什么样的kernel函数）确定block大小，或者对性能不是追求极致的时候，使用。

J. 介绍CUDA相关的技术，语言，库，工具等：https://devblogs.nvidia.com/12-things-tesla-accelerated-computing-platform/

GPUDirect: 同节点上各个GPU以及host之间的数据传输，高速高效。

CUDA-aware MPI: （很多不同的MPI实现都支持CUDA)https://devblogs.nvidia.com/introduction-cuda-aware-mpi/

NVLink, https://devblogs.nvidia.com/nvlink-pascal-stacked-memory-feeding-appetite-big-data/

OpenACC: 类似OpenMP都是并行循环用的，OpenACC为底层GPU/CPU生成相应的代码（OpenMP只能在CPU上多线程拆分任务）http://www.gpuworld.cn/article/show/560.html

常用的Profiling和Debug工具：NVIDIA-Nsight集成到了VisualStudio和Eclipse里了；NVIDIA Visual Profiler是带UI的Profiler; nvprof是不带UI的命令行profiler; GUDA-GDB用来debug; CUDA-MEMCHECK用来debug显存；

K. NVLink/NVSwitch: https://devblogs.nvidia.com/how-nvlink-will-enable-faster-easier-multi-gpu-computing/ 

http://server.it168.com/a2018/0604/3206/000003206894.shtml

新型总线及其协议；带宽比PCIe高一个数量级；和Unified Memory是绝配；

文章前部关于异构超级计算机讲的不错：GPU是追求throughput, 隐藏而不是减少latency, 追求throughput的同时减少耗电量;  CPU是追求单线程尽量减少latency(分支预取，推测执行等），耗电量大不适合GPU; 

one size does not fit all: Parallel and serial segments of the workload execute on the best-suited processor—latency-optimized CPU or throughput-optimized GPU—delivering faster overall performance, greater efficiency, and lower energy and cost per unit of computation.

L. CUDA7新特性：https://devblogs.nvidia.com/cuda-7-release-candidate-feature-overview/

C++11的特性（lambda函数，auto类型等）可以在device代码里调用；

thrust支持在device代码里调用；可以指定thrust函数在哪个stream上运行；

cuSolver: cuSlover是基于cuBLAS和cuSPARES库的一个对矩阵操作的库，包括矩阵分解和求解方程等内容，它里面包含了三个独立的库文件

M. C++11/Thrust: https://devblogs.nvidia.com/power-cpp11-cuda-7/

thrust::seq 可以指定该函数只在本thread上执行（适合划分grid-block时候已经考虑了并行的情况）；thrust::device会使用Dynamic Parallel 【1,2,3】机制在kernel里再启动子kernel来并行执行；

N. 模板：https://devblogs.nvidia.com/cplusplus-11-in-cuda-variadic-templates/

O. 新特性：https://devblogs.nvidia.com/new-features-cuda-7-5/

Half-float(16bit-float)；Visual Profiler支持显示kernel里每行的执行时间（访存，同步，计算，...各占多少比例）

P. CUDA8新特性：https://devblogs.nvidia.com/cuda-8-features-revealed/

Unified Memory可以直接用malloc和free来开辟释放了；page交换机制使得可以开辟远超GP主存大小的Unified Memory; 讲的很细待研究；

支持16bit-float (存储空间比32bit-float少一半，传输数据耗时比32bit-float少一倍，计算速度比32bit-float快一倍，适用于深度学习）；神经网络正向预测的时候用8bit整数就够了；

NVIDIA Visual Profiler增加对critical path的高亮突出（CPU等GPU执行完才执行的情况）；可显示NVLink的架构和性能；可显示UnifiedMemory的page fault;

Q. 半精度计算/低精度整数计算（有例子）：https://devblogs.nvidia.com/mixed-precision-programming-cuda-8/

R. 专门为深度学习加速设计的服务器架构（2017年）：https://devblogs.nvidia.com/dgx-1-fastest-deep-learning-system/

S. Volta架构：https://devblogs.nvidia.com/inside-volta/

SM里有了专门计算矩阵(A点乘B+C)的专有计算单元(Tensor Cores)！比之前的有了1个数量级的速度提升; 可以用CUDA C++接口访问，cuBLAS和cuDNN也有相应的新接口;   (Tensor Cores可以加速元鹏的向量乘积！）

L1 cache和shared memory共用；kernel不需要shared memory的时候，全部用来做L1 cache使用；L1-cache较大可以让开发者少去折腾shared-memory（shared-memory其实相当于开发者自己去管理cache) (编程难度和性能之间的一个trade-off)

以前的架构是一个warp共享一个PC，有一组Mask来表示每个thread在当前指令上是否是“活”的，因此if-else的时候一半线程执行完才轮到另一半线程执行；现在的架构，每个thread都有自己的PC和callstack，if-else的时候可以交替执行（为什么还是不能并行执行？），可以齐步走的时候系统还是优化到齐步走(SIMT)，可以使用__syncwarp()来同步warp内的所有线程；

T. CUDA9新特性：https://devblogs.nvidia.com/cuda-9-features-revealed/

Cooperative Groups支持以下并行模式：producer-consumer parallelism, opportunistic parallelism, and global synchronization；一个group可以包括a few threads (smaller than a warp) to a whole thread block, to all thread blocks in a grid launch, to grids spanning multiple GPUs. Volta架构允许warp内的thread成group; 允许block/grid/multi-GPU-grid, 作为一个group, 里面用一维的thread_rank来表示自己（有点儿像MPI)

nvGraph, 支持BFS等多种图计算算法；cuSOLVER支持矩阵SVD分解，乔姆斯基分解，LU分解；cuBLAS的GEMM(矩阵乘矩阵）常用在RNN和FCN里

Visual Profiler新功能：1. 显示Unified Memory的page fault发生在哪行代码上；2. 显示3种事件：A. 一方在等另一方访存完毕; B.双方都在访问同一个页面导致页面搬来搬去；C.内存映射事件，即数据只存在一方那边（可能是为了避免A频繁发生，或者另一方没有空闲存储了）