GPGPU 中的分支分歧

Question

__golbal__ void function(){

... some codes... 

if(0<A<512){

... a few codes...
}

else if (512<=A){

... a lot of codes...
}

... some codes...
}




int  main(){
some code for GPU...
...
for(int i=0; i<1024 ; i++){
A[i] = i;
}
...
some code for GPU...
}

我在研究GPGPU，对分支发散产生了好奇。我知道在 GPU 中处理分支发散时，它使用 SIMT 堆栈，并且将执行同一 warp 中的选定线程，而其他线程则不会。然后，如果同一个warp中的所有线程都没有被选中，指令是否执行？

例如，在上面的代码中，一半的线程将使用 if，另一半的 half 将使用 else if。如果在同一个warp的threads中所有A都小于512，执行else if的指令？或者跳过它们。

Answer 1

for example, In the above code, half of threads will take if and the other half will take else if. If all A are less than 512 in threads that are in the same warp, dose the instructions about else if will be executed? or just skip them.

直接跳过。

如果同一个 warp 中的所有线程都有一条指令 predicated 关闭，则不会发出该指令，也不会消耗执行资源 该 warp。

这就是幻灯片 53 statement/suggestion here 的原因：

Avoid diverging within a warp
– Example with divergence:
• if (threadIdx.x > 2) {...} else {...}
• Branch granularity < warp size
– Example without divergence:
• if (threadIdx.x / WARP_SIZE > 2) {...} else {...}
• Branch granularity is a whole multiple of warp size

如果条件边界与扭曲边界对齐，实际上“没有分歧”。这被认为是“优化”。

也就是说，来自与上面相同的幻灯片：

• Different warps can execute different code with no impact on performance

GPGPU 中的分支分歧

Branch divergence in GPGPU

cuda

gpgpu