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中的发散

我正在研究GPGPU，我对支路发散很好奇。我知道，在处理GPU中的分支发散时，它使用SIMT堆栈，在相同的翘曲中选择的线程将被执行，而其他线程则不会执行。然后，如果没有选择相同经纱中的所有线程，那么指令是否会被执行？

例如，在上面的代码中，一半的线程将接受 if ，而另一半的则会接受else if。如果所有A都小于512个处于相同翘曲中的线程，那么如果要执行，是否需要执行有关的说明？或者干脆跳过它们。

Answer 1

例如，在上面的代码中，一半的线程将接受if，另一半的线程将接受else if。如果所有的A都小于512的线程是在相同的翘曲，是否有关的指令，如果要执行？或者干脆跳过它们。

直接跳过去。

如果同一翘曲中的所有线程都关闭了指令预测，则不会发出指令，也不会占用该翘曲的执行资源。

这就是关于幻灯片53 这里的声明/建议的原因。

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

实际上，如果条件边界与翘曲边界对齐，就不会出现“发散”。这被认为是一个“优化”。

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

·不同的翘曲可以执行不同的代码而不影响性能