《CUDA编程基础与实践》读书笔记
先吐槽一下，184页的数69元，着实比较贵。而且184页入门内容偏少
以下是读后总结，不全仅供参考。
4. CUDA程序的错误检查
4.1. 一个检测CUDA运行错误的宏函数 —— 写一个宏函数捕获函数返回值，打印错误信息。 4.2. 核函数可在其后使用CHECK(cudaGetLastError())函数捕获最近一次错误代码。 4.3. 检查内存错误可使用CUDA-MEMCHECK工具集。
5. 获得GPU加速的关键
5.1. 用CUDA事件计时——cudaEvent_t
5.2. CUDA工具箱nvprof进行性能剖析
5.3. 影响GPU加速的关键因素——数据传输比例、算法强度和并行规模
- 减少主机与设备之间的数据传输
- 提高核函数的算法强度
- 增大核函数的并行规模
6. CUDA的内存组织
6.1. CUDA中不同类型的内存
全局内存——所有线程均可访问，可读可写，包括静态全局内存；
常亮内存——有常亮缓存的全局内存，数量有限，仅64KB，可读不可写；
纹理内存和表面内存——是一种有缓存的全局内存；
寄存器——线程私有内存，数量有限；
局部内存——是全局内存的一部分，所以延迟也很高；
共享内存——线程块内共享，速度仅次于寄存器，数量也有限；
L1和L2缓存——用来缓存全局内存和局部内存的访问，减少延迟。
6.2. SM及其占有率
寄存器和共享内存使用量很少的情况下，SM占有率受线程束的数量影响，尽量为32的倍数；
有限的寄存器数量对于占有率的约束是若单线程占用过多寄存器，会导致某些线程没有寄存器，从而降低SM占用率；
有限共享内存对占有率的约束是若单线程需要共享内存过大，导致某些线程无法分配共享内存，从而降低SM占有率。
7. 全局内存的合理使用
7.1. 全局内存的合并与非合并访问
合并访问指的是一个线程束对全局内存的一次访问请求导致最少数量的数据传输（以32字节为单位和起始的顺序传输；
在不能满足读取和写入都是合并的情况下，一般来说应尽量做到合并地写入。
8. 共享内存的合理使用
8.1. 函数规约
共享内存使用__shared__限定符定义，使用共享内存在内存访问高的情况下比全局内存加速效果明显，否则相当；
动态共享内存——在调用核函数是通过第三个参数动态申请动态共享内存，且共享内存申明时使用extern限定词。
8.2. 矩阵转置
通过全局内存 -> 共享内存 ->全局内存可以实现矩阵转置中读写都合并访问，其结果比读合并、写非合并要快（没有读非合并、写合并快）；
避免共享内存的bank冲突——共享内存在物理上被分为32个同样宽度、能被同时访问的内存bank。在每个bank中又可以对其中的内存地址从0开始编号（相当于二位矩阵，每列都可以同时访问）。同一线程束内的多个线程同时访问同一个bank中不同层的数据时，则会发生bank冲突。所以尽量让同一线程束中多个线程同时访问不同bank的数据。
9. 原子函数的合理使用
9.1. 原子函数
原子函数对它的第一个参数指向的数据进行一次“读-改-写”的原子操作，避免了多线程同时读写同一内存导致的冲突问题；
原子函数的种类——加法、减法、交换、最小值、自增、自减、比较-交换、按位与、按位或、按位异或；
不同的原子函数支持的数据类型也不同，使用前需要查阅资料；
10. 线程束基本函数与协作组（本章没有理解，待用到时再研究补充）
10.1. 单指令-多线程执行模式 —— 注意分支分散问题（同一时刻，一个线程束的线程只能执行一个共同的指令或者闲置，当一部分线程执行A分支时，B分支闲置。再B分支执行，A分支闲置；
10.2. 线程束内的线程同步函数 —— __syncwarp()比syncthreads()函数更加高效；
10.3. 更多线程束内的基本函数（并未理解所有函数，待补充）
10.4. 协作组 —— 可在线程块中组更小的组，实现更灵活的并行；
| 机器和方法 | 计算结果 | 计算时间/ms | 单次加速比 | 累积加速比 |
|---|---|---|---|---|
| CPU（循环累加） | 33554432.0 | 100 | 1 | 1 |
| GPU（只用全局内存） | 123633392.0 | 5.8 | 17 | 17 |
| GPU（使用静态共享内存） | 123633392.0 | 5.8 | 1 | 17 |
| GPU（使用动态共享内存） | 123633392.0 | 5.8 | 1 | 17 |
| GPU（使用原子函数） | 123633392.0 | 3.8 | 1.5 | 26 |
| GPU（使用束内同步函数） | 123633392.0 | 3.4 | 1.1 | 29 |
| GPU（使用洗牌函数） | 123633392.0 | 2.8 | 1.2 | 36 |
| GPU（使用协作组） | 123633392.0 | 2.8 | 1 | 36 |
| GPU（增大线程利用率） | 123000064.0 | 2.0 | 1.4 | 50 |
| GPU（使用静态全局内存） | 123000064.0 | 1.5 | 1.3 | 67 |
注：计算的精确结果为123000000.0
11. CUDA流
11.1. CUDA流实现核函数外部的并行
- 核函数计算与数据传输之间的并行；
- 主机计算与数据传输之间的并行；
- 不同的数据传输之间的并行；
- 核函数计算与主机计算之间的并行；
- 不同核函数之间的并行。
11.2. CUDA流概念
- 一个CUDA流指的是由主机发出的在一个设备中执行的CUDA操作序列；
- CUDA流分为默认流（空流）和非默认流（非空流）；
- CUDA流的类型是：cudaStraeam_t,创建函数：cudaError_t cudaStreamCreate(cudaSteam_t*),析构函数：cdudaError_t cudaStreamDestroy(cudaStream_t)。
11.3. 在默认流中重叠主机和设备计算
- cudaMemcpy函数是个阻塞函数，主机调用后需等待函数返回；
- 核函数是非阻塞的，可在调用核函数后调用主机计算函数，实现主机和设备计算的重叠。
11.4. 用非默认CUDA流重叠多个核函数的执行
- 使用my_kernal«<N_grid, N_block, N_shared, stream_id»>的方式调用核函数实现不同流的调用；
- 不同GPU架构支持并发执行的核函数个数上限不同，达到上限后加速比会下降。
11.5. 用非默认CUDA流重叠核函数的执行与数据传递
- 并发的数据传输需要使用cudaMemcpyAsync()函数；
- 异步的数据传输需要将主机内存定义为不可分页内存或者固定内存，使用cudaError_t cudaMallocHost(void** ptr, size_t size)或cudaError_t cudaHostAlloc(void** ptr, size_t size, size_t flags)两种函数申请内存。
- 将执行和数据传输并发的方法是将数据分割并在不同的CUDA流中执行。极限加速比为3，即内存H2D、核函数计算和内存D2H三者速度相当时。
12. 使用统一内存编程
12.1. 统一内存简介
- 统一内存是一种逻辑上的概念，是一种系统中任何处理器（CPU或GPU）都可以访问，并能保证一致性的虚拟存储器。
- 使用统一内存的硬件要求——1. GPU架构不低于开普勒架构，主机程序必须是64位；2. 新功能需要帕斯卡以上架构，Linux系统；3. 部分系统才支持访问全部的主机内存。
- 统一内存的优势——1. 是CUDA编程更加简单；2. 可能提供比手工移动数据更好的性能；3. 允许GPU进行超量分配。
12.2. 统一内存的基本使用方法
- 必须在主机端定义或分配内存。
- 动态统一内存用cudaMallocManaged()函数申请；
- 静态统一内存用__device__ __managed__申请。
12.3. 使用统一内存申请超量内存
- 由于统一内存的试剂分配发生在主机或者设备第一次访问预留内存时。
12.2. 优化使用统一内存的程序
- 为了在使用统一内存时获得较高性能，需要避免缺页异常、保持数据的局部性（让相关数据尽量靠近对应的处理器），但避免内存抖动（即频繁地在不同的处理器之间传输数据）。
- 在使用统一内存时，要尽可能多地使用cudaMemPreferchAsync()函数，将缺页异常的次数最小化。
13. 编程实践
14. CUDA标准库的使用
14.1. 目前有20多个库，部分如下：
- Thrust——类似于C++的标准模版库
- cuBLAS——基本线性代数子程序
- cuFFT——快速傅里叶变换
- cuSPARSE——稀疏矩阵
- cuRAND——随机数生成器
- cuSolver——稠密矩阵和稀疏矩阵计算库
- cuDNN——深度神经网络