[PATCH] D16941: [NVPTX] Mark nvvm synchronizing intrinsics as convergent.

[Justin Lebar via llvm-commits]

jlebar added a comment.

Thank you for the quick review, jingyue.  I don't have an example of code that is miscompiled without this; I just went looking for optimizations that check convergent and not noduplicate.  I found LoopUnswitch, Sink, and InstructionCombining.  I think InstructionCombining is not relevant, because it appears that the only calls it's interested in are free()s.  LoopUnswitch seems OK because it checks Metrics.notDuplicatable, which checks for NoDuplicate calls.  But I don't immediately see how Sink is safe.  (This was the example majnemer	came up with off the top of his head when we talked.)  I also see something similar in MachineSink, although I'm not as sure that's relevant.

In http://reviews.llvm.org/D16941#345516, @jingyue wrote:

> According to PTX ISA, `bar.sync` should only be executed uniformly, so inlining won't introduce new divergence.


I don't know what this means; can you elaborate?


http://reviews.llvm.org/D16941