warp scheduler
sm 调度器是核心，并没有对外透露任何文档，有人在做逆向。
用几个 SM
在激活 kernel 的时候传入的 block 数决定了 kernel 用几个 SM。
什么时候被调度？
在当前的 warp 可以发射的指令发完应该会触发调度，发完指的是，和顺序向量一样，最前面的那个被依赖堵住了，然后触发调度。 这意味着，在 warp 内部指令排布的方式要跟顺序向量一样，最好依赖无关来先把 ILP 打上去。然后交给调度器来盖掉延时， 这里面最好是合理到多个 warp 的安排能够充分的吧延时盖掉。
分支
遇到悬而未决分支怎么办？还是靠调度，分支不会两边都执行。在引入了线程独立执行之后，分支会在出结果之后，按照结果 各个线程独自向下，但是汇聚的时候要 syncwarp() 同步。分支悬而未决的延迟由调度来盖。
依赖管理
软硬件协同的计分牌，整个过程和 cpu 的投机调度有点像。
首先在指令中就有软件计分板，等于对于简单的算数指令，能直接推算出延迟的，延迟直接放在软件里，在发射的时候把延迟读出来 放在特定的地方记录，每个周期数字减1，减到0就结束了。后续的指令想发射得等这个减到0。
长延时的就硬件做记录，比如 load，专门在计分牌写，等写回来了再擦掉。对于异步的要等这个擦掉再继续向下。
同步
别人写你读的时候就可能出竟态条件。寄存器之间的同步的场景用 __shfl_sync 等。和内存相关的应该要显式通过原语同步， 比如 __syncthreads() 等。不做同步就锁不住，默认向下，可能出错。
WS
即一个 kernel 内不同的 warp 做不同的事，本身cuda没有给出 warpid 这种概念，想要用自己对线程进行划分。 形成一种生产者消费者模式，但是 WS 不局限于这种。一旦分工，产生的思想就是各司其职了，不能在一些专用的 warp 里面干一些通用的事，这样就有点矛盾。
SIMT 和 SIMD
本身来说，SIMT 本身是用线程抽象的方式驱动向量，但是现在加了 DSA 一样的操作后，感觉这种所有线程做一件事 的通用性被弱化了。同时按照之间的分析，即使 SIMT 的前端能把延时盖掉，但是尽力做指令排序还是很重要的。
SIMT 的模式有点想代码被复制多份，通过切换的时候同样的指令被发出，这个我能想到的好处是软件预取、向内存的 请求被直接打出去了，不用做更多的考虑。
任务安排
激活时候传入的 block 被调度到各个 SM 上，并且就在一个 SM 上被一直驻留，不会离开。同时 SM 安排的策略是有 套路的，不可控的，总体来说：
- 考虑寄存器压力：SM 上所有 block 的所有 thread 合起来的寄存器没炸掉。
- smem：所有的合起来 SMEM 不超标。
- 总 block 数不超标：总 block 不超 32.
- 总线程不超：最多 2048 thread。
执行完 thread 后，完事一个释放相关的寄存器，然后补一个上来。就这样直到全做完为止。 一般来说 block 内的 thread 数为 256/512 等感觉比较固定，或许不是一个太值得深思的指标， 等研究着看吧。