In D85603#2379258, @nhaehnle wrote:

Roughly speaking, there are subgroup ops with "implicit" thread set (what Vinod calls textually aligned, and what this proposal mostly focuses on, because they require the most additional explanation) and subgroup ops with "explicit" thread set (sm70+).

In D85603#2361168, @jlebar wrote:

• Will this paint us into a corner wrt CUDA, and specifically sm70+?

/me summons @wash, who is probably a better person to speak to this than me.

My understanding is that the semantics of <sm70 convergent are pretty similar to what is described in these examples. But starting in sm70+, each sync operation takes an arg specifying which threads in the warp participate in the instruction.

I admit I do not fully understand what the purpose of this is. At one point in time I thought it was to let humans write (or compilers generate) code like this, where the identity of the convergent instruction does not matter.

// Warning, does not seem to work on sm75
if (cond)
  __syncwarp(FULL_MASK);
else
  __syncwarp(FULL_MASK);

but my testcase, https://gist.github.com/50d1b5fedc926c879a64436229c1cc05, dies with an illegal-instruction error (715) when I make cond have different values within the warp. So, guess not?

Anyway, clearly I don't fully understand the sm70+ convergence semantics. I'd ideally like someone from nvidia (hi, @wash) to speak to whether we can represent their convergent instruction semantics using this proposal. Then we should also double-check that clang can in fact generate the relevant LLVM IR.

In D85603#2370240, @vgrover99 wrote:

I believe what is described here about convergent, as best I can understand it, is the semantics of syncthreads in CUDA. This semantics is the same for <sm70 and sm70+. Not clear whether what is described here is a "textually aligned" semantics or unaligned. syncthreads is aligned, meaning that all threads in the threadblock must wait on the same lexical syncthreads().

LGTM. Thanks!

LGTM!

LGTM, to the extent of my high-level understanding of the analysis. I would recommend waiting a few days and give others a chance to finish the review. I do have some nitpicks, but feel free to decide which ones warrant a change.

https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md#warp-vote-and-ballot-functions

I think if the language interface insists on fixing the wave size, then I think the correct solution is to implement this in the header based on a wave size macro (which we're desperately missing). The library implementation should be responsible for inserting the extension to 64-bit for wave32

LGTM!

From the description: "This patch inserts the Lowerswitch pass in an appropriate place to ensure any dead blocks resulting from the transformation will be removed"

LGTM!

@cdevadas reported that other parts of the AMDGPU backend are also affected by the unreachable blocks being produced in the switch lowering. Instead of fixing each such pass separately, it seems the best way forward it to put the block elimination earlier in the flow. @cdevadas is already working on such a change.

In D83394#2142277, @foad wrote:

Rebase.
Fix silly mistake in checking for negative offsets.

The documentation for HIP __ballot seems to indicate that the user does not have to explicitly specify the warp size. How is that achieved with these new builtins? Can this be captured in a lit test here?

filed a bug; improved the description; added a failing testcase

@simoll, I've been staring at the code for some time now, and I am no longer sure that I have a proper grasp on the problem. It seems my fix works, but not for the reasons that I thought it should. Running the unmodified analysis on "join-at-loop-exit.ll" shows the following debug output:

In D81806#2095605, @madhur13490 wrote:

bb1: 
   %m.uni = <uni>
   %n.uni = <uni>
   br <uni>, %bb2, %bb3

bb2: [..]
   br <div>, %bb2, %bb3

bb3:
   %x.uni = phi [%m.uni, %bb2], [%n.uni %bb1]  <-- divergent loop exit that is not join divergent

I am new to this party. Why is %x marked as uni? Its value will differ amongst threads as branch in bb2 is divergent. Is my understanding incorrect?

eliminate DeferredTerminators; also a small refactor for readability

In D81806#2094869, @simoll wrote:

I am a bit surprised this is necessary. Actually, join divergence in bb3 should come about as a result of propagateLoopDivergence for the [bb1, bb2] loop and not immediately because it is a divergent exit.

Abandoned in favour of enabling the workarounds in the AMDGPU backend: https://reviews.llvm.org/D81211

This looks really good. Thanks a lot for bringing it this far! Please do check the #include for LoopInfo.h before submitting.

LGTM, but please wait for approval from @arsenm

In D80804#2063522, @saiislam wrote:

In D80804#2063451, @sameerds wrote:

Actually, the question really is about why inc/dec are needed as separate operations either as IR intrinsics or Clang builtins. Why not just expose a __builtin_amdgcn_atomicrmw that takes a scope, and map it to the LLVM atomicrmw? That would be way cleaner. The language can provide convenience functions for inc/dec that internally call the rmw builtin.

At the moment, atomic inc/dec exist along with atomicrmw. This patch only aims to devise a way to make it (inc/dec) accessible from the language level.

Actually, the question really is about why inc/dec are needed as separate operations either as IR intrinsics or Clang builtins. Why not just expose a __builtin_amdgcn_atomicrmw that takes a scope, and map it to the LLVM atomicrmw? That would be way cleaner. The language can provide convenience functions for inc/dec that internally call the rmw builtin.

The commit description needs fixing. These are not "llvm instructions" they are "AMDGCN intrinsics". They don't exist in the LangRef. Also, I would recommend inverting the first line of the description: "Introduce Clang builtins that are mapped to AMDGCN intrinsics".

LGTM!

LGTM!

In D79037#2054462, @ekatz wrote:

This is a second take of the solution, using SCCs instead of LoopInfo.
This algorithm can also handle irreducible loops, and, generally, any kind of cycle in the graph.

LGTM!

I think the patch is okay, but have a couple of doubts about the tests.

In D79037#2049573, @ekatz wrote:

In D79037#2040926, @sameerds wrote:

TBH, I spent a small amount of time visualizing the new testcase, but couldn't figure out the problem, nor the solution. I see that the check for currentLoopSize is moved up the control flow a bit, but I didn't see why.

We need to move up the check for *CurrentLoopSize, as in the case it is zero, then we no longer need to skip nodes, and just continue to the next node/loop.
The simplified, new test case, demonstrate a case where it is needed.

I support this enhancement. Thanks for doing it!

I am NOT familiar with this transformation. But I believe that the patch is well-formed, and I was able to follow the logic. The tests also demonstrate that it works as intended. LGTM.

TBH, I spent a small amount of time visualizing the new testcase, but couldn't figure out the problem, nor the solution. I see that the check for currentLoopSize is moved up the control flow a bit, but I didn't see why. A few LLVM_DEBUG lines would be very helpful to trace what the structurizer is doing. Also, can the new test case be reduced further? There are a couple of simple chains, and a few branches. Are all of them necessary? The main problem seems to be an inner loop with just two blocks where both blocks exit. I don't know whether both these properties matter. For example, will the problem happen with a single block that loops on itself?

Test update: Moved the convergent call to where it is actually relevant.

LGTM. It would be nice if you can move the following tests to the new file and demonstrate that that they are fixed. But that need not hold up this review any longer.

• test/Transforms/StructurizeCFG/AMDGPU/backedge-id-bug.ll
• test/Transforms/StructurizeCFG/AMDGPU/backedge-id-bug-xfail.ll

In D79037#2024745, @ekatz wrote:

In D79037#2024574, @sameerds wrote:

Actually, I take that back. Consider the following CFG:

A -> H1, H2
H1->B, C
B->H1
C->B, H2
H2->E
E->H2, F

A is the entry and F is the exit. The loops (H1, B, C) and (H2, E) are siblings with backedges B->H1 and E->H2 respectively. C is an exiting block in the first loop that branches to H2. One POT would be "B, F, E, H2, C, H1, A". When iterating in reverse, the loop with H1 is started first, but H2 from the second loop is encountered before the first loop is completed.

In this case, those loops are in different regions, and as this is a region pass, they won't appear in the same POT.

In D79037#2024482, @sameerds wrote:

In D79037#2022995, @ekatz wrote:

In D79037#2022768, @sameerds wrote:

I see. It sure does look like PR25378. Then the only thing left is one new test with two sibling loops. The intention is to cover the case where CurrentLoop is not the parent of L.

I don't think it is possible, although I might be wrong.
Because the way I see it, for loop A be interfered with another loop B, while POT, there must be at least one block that belongs to loop A that is dependent on at least one block in loop B. But then loop B must be finished before loop A, and actually inside it. Otherwise, it won't be recognized as a separate loop.

I hope I explained it clear enough. Otherwise, can you please give an example?

I agree. That means that if L exists, then CurrentLoop must be its parent, right? So instead of checking for this condition, the code will have to assert it.

In D79037#2022995, @ekatz wrote:

In D79037#2022768, @sameerds wrote:

I see. It sure does look like PR25378. Then the only thing left is one new test with two sibling loops. The intention is to cover the case where CurrentLoop is not the parent of L.

I don't think it is possible, although I might be wrong.
Because the way I see it, for loop A be interfered with another loop B, while POT, there must be at least one block that belongs to loop A that is dependent on at least one block in loop B. But then loop B must be finished before loop A, and actually inside it. Otherwise, it won't be recognized as a separate loop.

I hope I explained it clear enough. Otherwise, can you please give an example?

In D79037#2022612, @ekatz wrote:

The problem with those tests is something else, and not produced by the wrong ordering. You may refer to PR25378 for an example.
Still, this fix the ordering issue and fixes PR41509 (and its duplicates).
I think we should insert this patch, and attend to the other bug as a separate issue.

I have bad news. This fix does not help with the test needs-fr-ule.ll, i.e., the resulting CFG is incorret if I remove the unify-loop-exits pass. The fix-irreducible pass produces an inner loop with the header "irr.guard" and a latch "while.cond47". This is inside an outer loop with header "while.cond" and many backedges. When I run the structurizer with the current fix on this graph, the backedge of the inner loop is lost. Instead we see a spurious loop between two new blocks Flow5 and Flow4, which is governed by incorrect boolean constants when in fact they should have been %Pred variables. That information is lost along the edge Flow8 to Flow5.

The implementation looks okay to me. The new comments help a lot. Waiting for tests mentioned in the comments so far ... one new test (sibling loops where the DFS jumps from one to the other before finishing the first loop) and updates to existing tests.