Note: This relands #140615 adding a ".count" suffix to the non-".all" variants. Our current intrinsic support for barrier intrinsics is confusing and incomplete, with multiple intrinsics mapping to the same instruction and intrinsic names not clearly conveying intrinsic semantics. Further, we lack support for some variants. This change unifies the IR representation to a single consistently named set of intrinsics. - llvm.nvvm.barrier.cta.sync.aligned.all(i32) - llvm.nvvm.barrier.cta.sync.aligned.count(i32, i32) - llvm.nvvm.barrier.cta.arrive.aligned.count(i32, i32) - llvm.nvvm.barrier.cta.sync.all(i32) - llvm.nvvm.barrier.cta.sync.count(i32, i32) - llvm.nvvm.barrier.cta.arrive.count(i32, i32) The following Auto-Upgrade rules are used to maintain compatibility with IR using the legacy intrinsics: * llvm.nvvm.barrier0 --> llvm.nvvm.barrier.cta.sync.aligned.all(0) * llvm.nvvm.barrier.n --> llvm.nvvm.barrier.cta.sync.aligned.all(x) * llvm.nvvm.bar.sync --> llvm.nvvm.barrier.cta.sync.aligned.all(x) * llvm.nvvm.barrier --> llvm.nvvm.barrier.cta.sync.aligned.count(x, y) * llvm.nvvm.barrier.sync --> llvm.nvvm.barrier.cta.sync.all(x) * llvm.nvvm.barrier.sync.cnt --> llvm.nvvm.barrier.cta.sync.count(x, y)
This reverts commit 735209c0688b10a66c24750422b35d8c2ad01bb5.
Our current intrinsic support for barrier intrinsics is confusing and incomplete, with multiple intrinsics mapping to the same instruction and intrinsic names not clearly conveying intrinsic semantics. Further, we lack support for some variants. This change unifies the IR representation to a single consistently named set of intrinsics. - llvm.nvvm.barrier.cta.sync.aligned.all(i32) - llvm.nvvm.barrier.cta.sync.aligned(i32, i32) - llvm.nvvm.barrier.cta.arrive.aligned(i32, i32) - llvm.nvvm.barrier.cta.sync.all(i32) - llvm.nvvm.barrier.cta.sync(i32, i32) - llvm.nvvm.barrier.cta.arrive(i32, i32) The following Auto-Upgrade rules are used to maintain compatibility with IR using the legacy intrinsics: * llvm.nvvm.barrier0 --> llvm.nvvm.barrier.cta.sync.aligned.all(0) * llvm.nvvm.barrier.n --> llvm.nvvm.barrier.cta.sync.aligned.all(x) * llvm.nvvm.bar.sync --> llvm.nvvm.barrier.cta.sync.aligned.all(x) * llvm.nvvm.barrier --> llvm.nvvm.barrier.cta.sync.aligned(x, y) * llvm.nvvm.barrier.sync --> llvm.nvvm.barrier.cta.sync.all(x) * llvm.nvvm.barrier.sync.cnt --> llvm.nvvm.barrier.cta.sync(x, y)
In case the same src BB targets to the same dest BB in different conditions/edges, such as switch-cases, we should use prob[SrcBB->SuccIndx] instead of prob[SrcBB->DstBB] to get probability.
In unreachable code, constant PHI nodes may appear and be replaced by their single value. As a result, instructions may become self-referencing. This commit adds checks to avoid going into infinite recursion when handling self-referencing compare instructions in `evaluateOnPredecessorEdge()`. This LLVM defect was identified via the AMD Fuzzing project.
These date back to when the non-intrinsic format of variable locations was still being tested and was behind a compile-time flag, so not all builds / bots would correctly run them. The solution at the time, to get at least some test coverage, was to have tests opt-in to non-intrinsic debug-info if it was built into LLVM. Nowadays, non-intrinsic format is the default and has been on for more than a year, there's no need for this flag to exist. (I've downgraded the flag from "try" to explicitly requesting non-intrinsic format in some places, so that we can deal with tests that are explicitly about non-intrinsic format in their own commit).
Although an unreachable BB is skipped by processBlock, its successor can still be handled by processBlock, and maybeMergeBasicBlockIntoOnlyPred may merge the two BBs and delete the unreachable BB. Then the garbage pointer is left in Unreachable set. This patch avoids merging a BB into unreachable predecessor.
We cannot infer more information from backedges in `solveBlockValueNonLocal`. However, since DT is unavailable in LVI, there is not a precise way to check whether a BB edge is a backedge. This patch only skips self loops to unblock the range analysis. The motivating case is extracted from https://github.com/llvm/llvm-project/pull/127663. Compile-time impact is high: https://llvm-compile-time-tracker.com/compare.php?from=84ddda58c870681dd12ed765e9d59d5e00567f94&to=af032f1351358f2f5b5d9f4e87c5601c23b9bd37&stat=instructions:u
This cleanup action only needs to be performed once when the entire optimization is converged. Doing it in every iteration has a very high time-complexity, as it queries every dbg value in a dense map Compare before and after for one internal source file with many basic blocks   >90% reduction in this extreme case.