[PATCH] D12246: [NVPTX] change threading intrinsics from noduplicate to convergent

[Owen Anderson via llvm-commits]

resistor added a comment.

I don’t think the example code here is legal under any SPMD models I am aware of.  It’s generally not legal to have barrier operations under divergent control flow, such as divergent trip-count loops.

From the CUDA docs:

> __syncthreads() is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects.


—Owen


http://reviews.llvm.org/D12246