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...
}

Divergence in GPGPU

我正在研究 GPGPU，我开始对分支分歧感到好奇。 我知道在处理 GPU 中的分支分歧时，它使用 SIMT 堆栈，并且将执行同一 warp 中的选定线程，而其他线程则不会。那么，如果同一个warp中的所有线程都没有被选中，指令会执行还是不执行？

例如， 在上面的代码中，一半线程将采用 if，而另一半线程将采用 else if。如果在同一个warp中的线程中所有A都小于512，是否会执行关于else if的指令？或者直接跳过它们。

Answer 1

例如，在上面的代码中，一半的线程会使用 if，而另一半会使用 else if。如果在同一个warp中的所有线程中的所有A都小于512，那么执行有关else if的指令吗？或者直接跳过它们。

跳过。

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

这就是幻灯片 53 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

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

也就是说，从上面的同一张幻灯片中：

• 不同的 warp 可以执行不同的代码，而不会影响 性能