因为warp内32个线程必须锁步执行同一条指令，if分支导致部分线程空转、分时串行执行两路径，吞吐腰斩；仅当分支条件使同一warp内路径不一致时才发生发散。
为什么 if 分支在 warp 里会拖慢整个 warp
一个 warp 固定是 32 个线程，硬件要求它们在同一周期执行同一条指令。一旦出现 if (threadIdx.x % 2 == 0) 这类条件判断，warp 内部线程就可能分走不同分支：偶数 ID 的进 if 块，奇数的进 else 块。此时 GPU 不是“并行执行两个分支”，而是让一半线程空转，等另一半执行完再切回来——实际变成串行处理，吞吐直接腰斩。
注意：warp divergence 只发生在单个 warp 内部；warp 之间互不影响。所以问题不在“有没有分支”，而在“分支条件是否导致同一 warp 内线程路径不一致”。
threadIdx.x % 2 看似简单，但真正在 warp 0 里就分化了
假设 launch 配置是 >>，那么 warp 0 包含线程 ID 0–31。代入 threadIdx.x % 2，结果是 [0,1,0,1,...] 交替——warp 内严格一半走 if、一半走 else，必然 divergence。
- 别信“编译器会自动优化掉”：现代 NVCC 确实对极简分支（如上例）可能用
FSEL指令替换，但这依赖具体模式和 compute capability，不可控 - 调试时看 sass 指令不如直接看 profiler 数据：Nsight Compute 中的
sms__inst_executed_op_fadd.sum和sms__inst_executed_op_fadd.avg差距大，就是 divergence 信号 - 更稳妥的做法是主动消除分支依据，而不是赌编译器
用 selp 或三元运算符替代 if/else 是有效手段，但有前提
CUDA 的 selp（via PTX）或 C++ 三元运算符 a ? b : c 能被编译为谓词选择指令（如 FSEL），避免控制流跳转。但它只适用于分支体是纯表达式、无副作用的场景：
立即学习“C++免费学习笔记（深入）”；
- ✅ 安全：
C[idx] = (idx % 2 == 0) ? A[idx] + B[idx] : A[idx] - B[idx]; - ❌ 危险：
if (idx % 2 == 0) { atomicAdd(&sum, 1); }—— 三元无法包裹atomicAdd，强行写会丢原子性 - ⚠️ 注意：即使用了三元，若条件本身仍导致 warp 内不一致（如
idx % 32），只是把 divergence 换成了更隐蔽的 predication 开销，没根治
真正防 divergence 的硬核做法：让分支边界对齐 warp 尺寸
核心思路是让每个 warp 内所有线程做完全相同的事。常见模式：
- 把条件从线程粒度升到 warp 粒度：
int warp_id = threadIdx.x / 32;，然后用if (warp_id % 2 == 0)—— 这样 warp 0–1 全进 or 全不进，无分化 - 数据预处理阶段按 warp 对齐分块：例如规约时，让每个 warp 处理连续 32 个元素，而非跨步取样
- 用
__syncthreads()配合 shared memory 做 warp 内协作，把“分支逻辑”转成“统一计算+条件写入”，比如先算出两个候选值，再用谓词决定写哪个
最易被忽略的一点：launch 配置本身就得是 32 的整数倍。如果 >>，最后一个 warp（线程 96–99）只有 4 个有效线程，其余 28 个空转——这虽不算 divergence，但浪费资源，且可能干扰共享内存 bank 访问模式。
