概述
恭喜自己完成第三章内容，第三章内容主要讲了：
- 通过配置文件驱动的方法优化内核
- 理解线程束执行的本质
- 增大GPU的并行性
- 掌握网格和线程块的启发式配置
- 学习多种CUDA的性能指标和事件
- 了解动态并行与嵌套执行
相对于上一章无目的的修改grid和block的size，这章从原理进行了分析，如何去提高运行性能，个人认为这章比较重要的是：
理解GPU的wrap、以及好几种提高GPU计算性能的方法。
但是由于核函数多种多样，各种芯片资源也不同，需要能够通过性能分析工具如nsys、ncu进行内核函数的指标分析，进而思考核函数的优化方法。
1关键概念
首先是几个关键的概念，都跟性能的提升有关：
- 线程束和线程块：线程束是SM中基本的执行单元。当一个线程块的网格被启动后，网格中的线程块分布在SM中。一旦线程块被调度到一个SM上，线程块中的线程会被进一步划分为线程束。一个线程束由32个连续的线程组成，在一个线程束中，所有的线程按照单指令多线程（SIMT）方式执行；也就是说，所有线程都执行相同的指令，每个线程在私有数据上进行操作。
- 线程束分化：假设在一个线程束中有16个线程执行这段代码，cond为true，但对于其他16个来说cond为false。一半的线程束需要执行if语句块中的指令，而另一半需要执行else语句块中的指令，但是同一wrap只能执行相同的指令，因此else语句的线程只能在wrap执行if语句的线程时停下，等其执行完毕再执行else，这降低了wrap的并行能力。在同一线程束中的线程执行不同的指令，被称为线程束分化。
- 资源分配：线程束的本地执行上下文主要由以下资源组成：·程序计数器·寄存器·共享内存，因此block的大小要合理划分，避免blocksize过大引起SM的资源没有充分利用，导致占用率较低；或者过小导致SM资源不足，无法充分发挥wrap的并行能力。
- 延迟隐藏：在指令发出和完成之间的时钟周期被定义为指令延迟，分为两种：算术指令和内存指令，其中算术指令的延迟为10～20个周期，内存指令延迟为400~800个周期。在线程进行这两个指令时会浪费很多时间，wrap会在这些指令执行后，跳转到其他就绪的wrap上运行，待这个wrap的所有指令完成后再来处理，因此需要足够多的活跃wrap来进行跳转以延迟隐藏。常见的一个量化指标是：
所需线程束数量＝延迟×吞吐量
假设在内核里一条指令的平均延迟是5个周期。为了保持在每个周期内执行6个线程束的吞吐量，则至少需要30个未完成的线程束。 - 占用率occupancy：在每个CUDA核心里指令是顺序执行的。当一个线程束阻塞时，SM切换执行其他符合条件的线程束。理想情况下，我们想要有足够的线程束占用设备的核心。占用率是每个SM中活跃的线程束占最大线程束数量的比值：
换句话讲就是希望让SM中的所有wrap全力运行。 - 同步 synchronize：可分为两种·系统级：等待主机和设备完成所有的工作,常用
cudaDeviceSynchronize()，·块级：在设备执行过程中等待一个线程块中所有线程到达同一点,常用__syncthread() - 可扩展性：可扩展性比性能更重要，使得代码能够在不同的GPU上运行，运算效率可以随着计算资源的增加而提高。
2使用性能检测工具
这部分放我常用的nsys和ncu指令：
nsys：因为我使用的是WSL2环境，因此不能直接运行ncu的ui，但是能使用nsys的ui，这个工具看程序运行的timeline很方便。
//使用nsys工具运行cuda_file并保存为报告nsys profile-o nsys_report_file cuda_path_flie//使用nsys-ui打开报告nsys-ui report_flie_name一般打开的界面是这样的:
ncu：这个工具可以分析很多性能相关的信息，如SM占用率，内存效率，寄存器使用情况等等，进行性能调优的时候推荐使用。但是wsl上不能直接打开报告文件,使用稍微麻烦一点
//使用ncu运行cuda_file并保存计算负载分析报告,注意顺序ncu--section ComputeWorkloadAnsysis-o ncu_report_file cuda_path_file//使用ncu运行cuda_file并保存内存负载分析报告,注意顺序ncu--section MemoryWorkloadAnsysis-o ncu_report_file cuda_path_file//运行的话要使用windows上的ncu打开保存的报告打开的报告界面是这样的:
还可以看详细的参数:
还能给你提供改进建议,非常好用啊就是说
3关于性能提高的优化方法
这章主要讲的是两方面:
- 避免线程束分化
- 展开循环
这里拿reduce sum简要说明一下:
1尽量优化减少if语句，如相对于相邻规约，使用交错配对的规约方法，可以大大提高访存效率
2.进行循环展开，在规约前处理多个连续相邻的数据块，可以很好的让内存带宽饱和和隐藏加载/存储延迟；若已知规约迭代次数，可以将规约循环完全展开
3.对最后的32个线程束进行循环展开，使用串行规约，在一个数据块里进行处理，避免了执行循环控制和线程同步逻辑。注意，短时间内多次执行计算语句需要加上volatile修饰符，告诉编译器每次赋值时必须讲值存回全局内存里。
4.使用模板化内核如template <int iblocksize>指定块大小作为模板参数，使得在编译器在编译阶段对核函数的分支进行优化
关于动态并行,现代的GPU上逐渐取消了这个功能,所以暂时跳过这部分的内容
