- Parallel Reduction 并行规约
- Warp 分割
- Memory Coalescing 访存合并
- Bank 冲突
- SM 资源动态分割
- 数据预读
- 指令混合
- 循环展开
有效的数据并行算法 + 针对GPU架构特性的优化 = 最优性能
Parallel Reduction 并行规约
如果我们改进一下，也就是减少空闲线程，修改索引方式。
代码如下：
第二种好处在于可以释放多余资源。
Warp 分割
线程块内如何划分warp
通晓：warp分割有助于：减少分支发散，让warp尽早完工
（1）warp是连续的32个线程。
（2）warp是基本的调度单元。
（3）warp执行相同的指令。(SIMT)
（4）每个线程只能执行自己的代码路径。
（5）warp内的线程存在分歧时，warp的线程两种分支都要执行。
（6）warp间线程互不干扰。大量的warp在一起可以隐藏访存延时。
block中thread编号：
二维block始于theadIdx.y==0 分割为一维block，重复增长threadIdx.y
三维block始于theadIdx.z==0 分割为二维block，重复增长threadIdx.z
warp存在分支发散：
部分线程执行黄色部分，部分线程执行白色部分。(代码中的分支逻辑导致的)
CPU-GPU数据传输最小化
访存合并 Coalescing
访问global memory延时：400-800cycles 几乎是最重要的性能影响因子
按照步长访问global、memory可以通过shared memory避免。 一般步长>>18的时候，访存的开销会比较大，这是可以通过将数据加载到shared memory中进行处理。
shared memory
shared memory关于bank的问题
bank0：对应地址0 地址32 地址64等，每周期只能相应一个地址。
bank conflict
下图产生了bank冲突
所有线程读取同一个地址的时候，通过广播机制不会发生bank confict，只有同时访问一个bank的时候才会发生冲突。
举例：矩阵转置
问题：shared memory的矩阵转置中，bank conflict是如何发生的？
CUDA的Texture纹理
SM资源分割
block的数量首先于硬件配置，每个block的线程越多，由于一个SM的ALU有限，最多容纳有限数量的线程，每个block分的线程越多，那么blockd的数量就会变小。这种变化有点离散。
performance cliff：增加资源用量后导致并行性急速下降
Kernel 执行配置
Grid Size试探法 (1)软件的block数大于等于SM的个数 保证都在执行一个block (2) 每个SM上的block数量应该大于2 保证多个blocks可以在SM上并发执行 如果一个block在等待同步，启动另一个block (3) 每个SM上的block数量大于100 对支持未来的设备拥有很好的伸缩性
块大小的设置 (1)块大小必须是32的倍数 (2)需要尽量多的warp尽可能地隐藏延时 (3)最少：64，通常采用128 or 256.视程序而定。 (4)也依赖于问题本身，多实验。
Latency Hiding 延时隐藏
关键点： 指令按顺序发布 一个线程的任一操作数没有预备好时将阻塞等待 延时在切换线程时被隐藏 总的来说就是需要多一些线程来隐藏延时
Occupancy 占用率
定义：一个SM里面激活的warp与最大可容纳的warp数目的比值 最大warp数目：32inTesla 48 in Fermi ??(ampere?)
假设一个指令需要2个时钟周期。
26 / 48(SM最多驻留48个warp) ≈ 54%
数据预读 Data Prefecthing
指令优化
循环展开
不用循环 用手写
编译器自动实现循环展开
循环展开的缺点：
