Will threads in a warp merged again after divergence happened?

Accelerated Computing CUDA CUDA Programming and Performance
1

For example, I have this line of code:

// codeblock 1
if (threadIdx.x. < 16) { x = 3; }
// codeblock 2

Obviously divergence happened after if executed. My question is whether the diverged warp remerged again when running codeblock 2?

2

CUDA provides no guarantees for this, especially in the volta execution model.

3

Thank you for your reply. Could you give me more details? In what condition the diverged warp will be remerged again? Or under no condition would this happen? Or is there anything I can do in my code to force it happen?

4

If you know that it is legal to do so, you can issue a __syncwarp().

The decision about when to “remerge” a warp is a decision made by the compiler, and there are no published details that I am aware of. The compiler understands that generally, the machine runs most efficiently when warps are converged, so barring other considerations (which are unpublished/undocumented) I would expect the typical state of any warp that is not in a state of forced divergence would be that it would be converged. As far as we have covered here, this is not something I would suggest the CUDA programmer worry about.

If you need convergence for correctness, there are many ways to force reconvergence. For example, various warp collectives (e.g. shuffle) take a sync parameter which advises the machine to what level the warp needs to be reconverged.

Absent one of those situations, I would say that worrying generally about whether warps are diverged or reconverging quickly enough is not worth your time.