[PATCH] D16941: [NVPTX] Mark nvvm synchronizing intrinsics as convergent.
Justin Lebar via llvm-commits
llvm-commits at lists.llvm.org
Fri Feb 5 18:22:46 PST 2016
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?
More information about the llvm-commits