从一次翻车说起
去年给一个刚入职的小伙review代码，CUDA kernel里面一长串的if-else，我当时就笑了——这特么是SIMT的大忌啊。他一脸茫然：“不是能分支吗？”
能。当然能。但能得不优雅——这就是SIMT最骗人的地方。它让你以为每个线程都是独立的，实际上，硬件里一群线程被绑在同一辆战车上，战车却只有一个方向盘。
SIMT，全称Single Instruction, Multiple Threads，NVIDIA从G80时代就挂在嘴边的东西。可你要是真去翻PTX手册，会发现它跟SIMD根本就是堂兄弟。差别在哪？在程序员眼里，SIMT允许你写标量代码，每条线程走自己的逻辑；但在SM内部，那些线程被编成32个一组的warp，像小学生排队，一步都不能走散。一旦有人想拐弯——好嘛，全班都得停下来等它。
Warp调度：当指令撞上分支
说穿了，SIMT的核心算法就藏在那个叫“warp”的调度单元里。一个warp 32个线程，共享一个程序计数器。
没有分支时，大家齐步走，每个时钟周期发射一条指令，32个线程同时干一样的活，美得像阅兵式。可一旦遇到if-else，调度器就得掏出SIMT Stack——一个硬件结构，负责记录哪些线程走then，哪些走else。然后呢？先把走then的线程跑完，那些走else的线程在旁边干瞪眼（被mask掉），再反过来让else的跑，then的等。这就叫“serialization”，分支发散的代价。
我测过一次极端情况：在一个有1024线程的kernel里，故意让奇数线程走一条极长路径，偶数线程走短路。nsys profile显示，SM utilization直接从92%砸到27%。27%，简直是资源浪费的教科书。而如果那段代码改成warp内统一分支（比如用threadIdx.x % 32控制），利用率又回到85%以上。原因很简单——SIMT的调度粒度是warp，只要warp内32个线程路径一致，硬件就不用切换mask。所以写CUDA的人，永远要对“warp内分支”这件事保持病态般的敏感。
这背后的美学是什么？一种预约束的并行度。架构师在设计时就已经认了：为了简化指令发射和减少控制单元，牺牲单线程的完全灵活性，换来成百上千线程的吞吐。这个trade-off，我至今认为，是近代异构计算最性感的工程决策之一。
内存延迟大坑：从bank conflict到全局访存
分支只是开胃菜，真正让新手跪的是共享内存。
每个SM里的共享内存分成32个bank，同一周期可以服务不同地址的请求——前提是这些地址落在不同bank。一旦warp内两个线程访问同一个bank的不同地址，冲突就来了，请求被串行化。最惨的一次，我们团队写了一个规约kernel，天真地用stride=32来访问，结果bank conflict率100%，延迟直接上天。后来改成stride=33，加上一点点padding，冲突降到0，带宽利用飙升3倍。
坑点之一：bank conflict看似是硬件特性，实则是算法问题。你只需要记住：让warp内线程的访问地址对32取模后尽量分散。做法？加padding（__shared__ float smem[32][33]），或者用swizzling模式。这套路可复用，调一次，管半年。
还有全局内存，那个延迟高达几百个周期的地方。SIMT怎么扛？靠occupancy——让SM同时常驻多个warp，一个warp等内存的时候，立刻切到另一个warp执行。听起来完美，对吧？呵呵。
Occupancy：一个骗了多少人的指标
我刚玩CUDA那阵，迷信occupancy，觉得越高越好，寄存器能省就省，shared memory能分就分。结果有一次写计算密集型kernel，occupancy拉到100%，实测性能还不如50%的版本。
坑点之二：occupancy是延迟隐藏能力的指标，不是吞吐量的保证。如果你的kernel算术密集，指令级并行足够，减少warp数量反而能给每个warp留更多寄存器，避免溢出到local memory（那玩意慢得叫人想砸机器）。实测数据：一个矩阵乘kernel，限制每个SM 64个warp vs 32个warp，后者寄存器不溢出，IPC提高40%，整体时间缩短22%。所以别被NVIDIA的occupancy calculator洗脑了，它算的是理论上限，不是最优。
坑点之三：异步拷贝与流水线。从kepler开始，GPU支持async copy，用专门的Load/Store单元做全局到共享的搬运，同时SM继续算。可太多人写的kernel还是同步阻塞式的——把数据全搬完再算，SM就干等着，warp切换都救不了。正确的姿势是用pipeline，把数据切小块，搬一块算一块，让计算和访存重叠。一次图像卷积优化，就这么一改，吞吐翻倍。复用模板：cuda::pipeline配合shared memory double buffering，你的kernel至少快30%。
说到底，SIMT不是一个你打开手册就学得会的东西。它是那种——你得在nvprof里盯上三天三夜，亲手调过几十个kernel，才会对warp、bank、occupancy这些词建立肌肉记忆的硬核模型。
没别的捷径。去写代码吧，让profile成为你的另一双眼睛。
