[llvm] [NVPTX] Only run LowerUnreachable when necessary (PR #109868)

[Justin Fargnoli via llvm-commits]

justinfargnoli wrote:

> I'm confused. The comment you quoted in the description only says

Yes, this is my bad. I should have quoted places where @maleadt said things along the lines of:

> we only know about issues on Pascal and earlier because bar.sync is not allowed to be executed divergently there

---

> Citing from the relevant NVIDIA bug report (4078847):

Thank you for the bug number! I'm following up internally to confirm what release the fix was included in. 

---

> That would be a change in the instruction selection backend though, not disabling this pass altogether

I thought this pass added the `exit` instruction, not instruction selection. Assuming I'm missing something, where should we make this change within the instruction selection? 

> disabling this pass altogether ... mean(s) there's still a need for trap instructions to avoid bogus CFG edges.

This pass doesn't touch the trap instruction, so it should be okay to disable it, right? 

https://github.com/llvm/llvm-project/pull/109868