In D69498#2728451, @JonChesterfield wrote:

I wonder if it is necessary for the exec mask to be implicit state, managed by a convergent/divergent abstraction.

We could reify it as an argument passed to kernels (where all bits are set), and adjusted by phi nodes at entry to basic blocks. Various intrinsics take and return that reified value. __ballot, a comparison op, possibly load/store.

At that point all the awkward control flow constraints are raised to data flow, and I think (but haven't really chased this into the dark corners) that solves the problems I know about. Notably, a uniform branch would be one where the exec mask value was unchanged, so the associated phi is constant folded.

Makes a mess of the human readable IR, but possibly at a lower overall complexity than continuing to handle divergence as control flow.

In D69498#2712706, @nhaehnle wrote:

In D69498#2711396, @foad wrote:

I don't have much to add to the conversation except to point out that this definition defines noconvergent in terms of divergent control flow, but the langref itself doesn't define what divergent control flow is, which makes it an incomplete spec. (Perhaps I'm just restating @arsenm's objections.) This seems unsatisfactory to me but I have no idea what to do about it. I agree with @sameerds that the current definition of convergent is too restrictive because in practice we really do want to be able to move convergent calls past uniform control flow.

That is one of the things that D85603 addresses. I suspect Sameer was assuming the language from there in the background.

I agree with Matt that it would be best to avoid too much TTI dependence. The existing situation with the intrinsics is already a bit of a strange one. I wonder if it is possible to move to a world where isSourceOfDivergence need not be target-dependent. (This requires e.g. new attributes on function arguments as well as new information about address spaces in the data layout, plus some new attributes to define intrinsic data flow. Likely beyond the scope of this patch.)

In D69498#2710675, @arsenm wrote:

In D69498#2709218, @sameerds wrote:

1. It is actually okay add control dependences to a convergent function as long as the new branches are uniform.
2. Some times we hack a transform to avoid doing things that cannot be described as "add new control dependence", and still talk about the function having the above named attribute.

There is another important bug that obfuscates the whole discussion: most places that use "isConvergent()" should in fact first check if it really matters to the surrounding control flow. There is too much loaded onto that one attribute, without sufficient context. The definition of "noconvergent" that I am proposing starts out by first fixing the definition of convergent itself. This definition is independent of target, and only talks about the properties of the control flow that reach the callsite. To repeat, this does not reinterpret the IR in a target-defined way. Like I said, in the new definition, all function calls are convergent even on CPUs, so I don't see where the target comes in. If you still insist on talking about interpretation of core attributes, please start by deconstructing the definition that I propose so I can see what I am missing.

Yes, the current concept of convergent is broken. I think this whole recent burst of conversation has been too focused on the current convergent situation and this patch in particular. I think we need to conceptually rebase this into the world that D85603 creates.

In D69498#2708990, @arsenm wrote:

You're still saying the same thing. This needs to be defined generically. Frontends don't *have* to to anything, they'll just get the assumed convergent behavior by default. Either frontends could always add noconvergent to avoid any possible optimization hit, or we could have a pass add noconvergent depending on the target. I don't want to change the interpretation of core attributes based on the target

In D69498#2706524, @arsenm wrote:

In D69498#2705441, @sameerds wrote:

I realize now that what @foad says above puts the idea in a clearer context. Essentially, any check for isConvergent() isn't happening in a vacuum. It is relevant only in the presence of divergent control flow, which in turn is relevant only when the target has divergence. Any standalone check for isConvergent() is merely making a pessimistic assumption that all the control flow incident on it is divergent. This has two consequences:

1. The meaning of the convergent attribute has always been target-dependent.
2. Dependence on TTI is not a real cost at all. We may eventually update every use of isConvergent() to depend on a check for divergence. The check for TTI is only the first step towards that.

The core IR semantics must *not* depend on target interpretation. convergent has never been target dependent, and was defined in an abstract way. Only certain targets will really care, but we cannot directly define it to mean what we want it to mean for particular targets

I realize now that what @foad says above puts the idea in a clearer context. Essentially, any check for isConvergent() isn't happening in a vacuum. It is relevant only in the presence of divergent control flow, which in turn is relevant only when the target has divergence. Any standalone check for isConvergent() is merely making a pessimistic assumption that all the control flow incident on it is divergent. This has two consequences:

In D69498#2704364, @foad wrote:

This recasts the whole question to be one about combining optimizations with target-specific information. The only changes required are in transforms that check CallInst::isConvergent(). These should now also check TTI, possibly adding a dependency on the TTI analysis where it didn't exist earlier.

@sameerds I agree with your conclusions but I would describe the situation a little differently. As I understand it, the optimizations that check isConvergent really only care about moving convergent calls past control flow that might be divergent. !hasBranchDivergence is a promise that there are no possible sources of divergence for a target, so you can run a divergence analysis if you like but it will just tell you that everything is uniform, so all control flow is uniform, so it's OK to move isConvergent calls around.

In practice the optimizations that check isConvergent don't seem to use divergence analysis, they just pessimistically assume that any control flow might be divergent (if hasBranchDivergence). But they could and perhaps should use divergence analysis, and then it would all just fall out in the wash with no need for an explicit hasBranchDivergence test.

The way I see it, the notion of convergence is relevant only to a certain class of targets (usually represented by GPUs) and it only affects certain optimizations. Then why not have only these optimizations check TTI to see if convergence matters? TTI.hasBranchDivergence() seems like a sufficient proxy for this information.

In D85603#2691257, @Anastasia wrote:

This needs to be inverted in the spirit of D69498. I would propose the following tweak:

1. By default, every call has an implicit convergencectrl bundle with a token returned by the @llvm.experimental.convergence.entry intrinsic from the entry block of the caller. This default is the most conservative setting within the semantics defined here.
2. A more informed frontend or a suitable transformation can replace this conservative token with one of the following:
   1. A token returned by any of the other intrinsics, which provides more specific information about convergence at this callsite.
   2. A predefined constant token (say none), which indicates complete freedom. This would be equivalent to the noconvergent attribute proposed in D69498.

Such a rewording would invert how we approach the spec. Instead of a representation that explicitly talks about special intrinsics that "need" convergence, the new semantics applies to all function calls. The redefined default is conservative instead of free, and the presence of the bundles relaxes the default instead of adding constraints.

Sounds good. If that would be acceptable to the wider community it might help to simplify the frontend design and improve the user interface and the coherence of the interfaces within the compiler stack too.

In D85603#2679362, @Anastasia wrote:

Sorry for not being clear - I was talking about two separate threads here (1) generalizing convergent attribute to non-uniform CF that is addressed by this patch and (2) inverting convergent attribute that is addressed in https://reviews.llvm.org/D69498. Just to provide more details regarding (2) - right now in clang we have a logic that adds convergent to every single function because when we parse the function we don't know whether it will call any function in a call tree that would use convergent operations. Therefore we need to be conservative to prevent incorrect optimizations but this is not ideal for multiple reasons. The optimiser can undo all or some of those convergent decorations if it can prove they are not needed. And for the uniform CF convergent operations this was the only "broken" functionality to my memory.

In D85603#2676631, @Anastasia wrote:

Regarding the HLL and frontend side, I believe this could be represented fairly similarly in different C/C++-based languages - considering that we already follow the same implementation for existing convergent semantics at least between CUDA and OpenCL. However, it isn't yet in its optimal state and perhaps we can attempt to refine this topic holistically for example also addressing the following rework that removes the need to make everything convergent: https://reviews.llvm.org/D69498. Otherwise, we will likely have to generate the convergent intrinsics absolutely everywhere, which is not ideal!

In D85603#2422804, @jholewinski wrote:

Few questions below. Please bear with me as I try to grok the proposal and the long stream of comments...

I believe that with this proposal, we can model this with the attributes we have by saying that subgroupShuffle is convergent readnone, while __shfl_sync is inaccessiblememonly.

subgroupShuffle would require convergentctrl and __shfl_sync would not, correct?

In D98958#2640544, @sameerds wrote:

Improved the readability of a potentially triple-negative boolean condition

Improved the readability of a potentially triple-negative boolean condition

This needs a deeper design. For now, the quick way forward is to disable loop unswitching on divergent non-trivial branches:

In D96754#2567456, @sameerds wrote:

That is not true. The flow in the legacy pass manager is more involved. The function LoopPass::preparePassManager() actually makes sure that if a loop transform T running inside a loop pass manager M invalidates analyses used by other passes in M, then T is split out into a separate loop pass manager M'. Thus every instance of loop pass manager is responsible for making sure that the used analyses are recomputed before starting any passes. The divergence analysis is not computed in the middle of loop unswitching like you seem to be suggesting here.

But now that I looked at it more closely, this does invalidate my original assumption that the legacy PM is providing a stale analysis. In that sense, my patch should not be seen as reproducing existing behaviour ... it's turning out to be more of a hack because the new PM lacks some functionality. There doesn't seem to be a way for a transform in the new PM to isolate itself from the effects of other transforms that invalidate analyses in the outer analysis manager.

In D96754#2566461, @asbirlea wrote:

This is absolutely not the right way resolve this.
First, the restriction to not allow getting a stale analysis is very much intentional and part of the design of the new pass manager. The API being added here must not exist.
Second, it is not safe to use the stale DA.

In D96615#2560885, @aeubanks wrote:

lgtm, thanks for doing this!

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".