In rGd08f34b592ff06ccb1f36da88ec09aa926427a4d#1100553, @kazu wrote:

In rGd08f34b592ff06ccb1f36da88ec09aa926427a4d#1100523, @jdoerfert wrote:

Where is the discussion and rational for this?

Please see https://reviews.llvm.org/D128131. Thanks!

In D128131#3594425, @dblaikie wrote:

Sounds good. Though it might be worth migrating/normalising uses to op* (rather than value()) and bool testing (where it do any have to be explicit) instead of has_value() when migrating callers.

Where is the discussion and rational for this?

LG. one nit.

I'm not so sure but phab things this is small, so no complaints:

LG

LG

In D128128#3604723, @nikic wrote:

Yeah, the non-constant case needs both nosync and norecurse.

Do we want JIT -> YES, but specalizing LLVM-IR JIT.
Do we want/need PTX, I do not, but I don't mind having it. Someone will ask for it eventually.

In rG083010312aa4a0ba0cd5299bd3b039af8fb6d58f#1098446, @eugenis wrote:

The leaks are still there, reverting.

^^^ so much for stored once

I think I misread it earlier. Constant & !thread_dependent *or* nosync. This makes more sense.

LG, I think :)

This breaks in the presence of synchronization.

In D127399#3581355, @kiranchandramohan wrote:

LG. I have one question. Is there a contract in place where the InsertionPointGuard is used to reset the IP back to the one at entry and would this change break it?

FWIW, I think we should have these attributes as spelled here, just w/ proper pass manager integration which then requires an RFC.
That said, I'm not opposed to this as an incremental step, albeit confusing until the PM support is integrated if we allow O1/2/3/fast

In D126984#3571573, @aeubanks wrote:

IIRC in the past there was a strong preference to not have the pass manager support this sort of thing
if you want to support this, there should be an RFC for how the optimization part of this will work as it may require invasive changes to the LLVM pass manager

(if this is purely a clang frontend thing then ignore me)

This seems to be uncontroversial. LGTM

Cool! Can we have that for begin declare variant too :D ?

LGTM

LG

Typos in the commit message.
s/new Clang/matching Clang/ ?

LG

LG, nice!

LG

In D125378#3506954, @aaron.ballman wrote:

In D125378#3506446, @jdoerfert wrote:

What I would suggest, if you want to use C/C++ attributes (which makes sense to me), is an attribute that avoids undef or introduces frozen:

I would prefer to avoid this approach for the moment. Such an attribute is highly focused on how the backend works and I'm not convinced that users in general would know when or how to properly use it. I believe all of the functions which need to be marked are known, and so we can keep an internal list of "interesting" functions that need special attention during clang codegen. If we later find that there's a wider need that cannot be handled transparently like this, we can reevaluate what to expose to users via an attribute at that point.

In D125378#3506383, @skc7 wrote:

In D125378#3506215, @jdoerfert wrote:

In D125378#3506001, @nikic wrote:

Please specify the semantics of the new LLVM attribute in LangRef -- though I don't really understand why you need an LLVM-side attribute at all.

+1. I doubt this patch is helpful.

Shuffle attribute has been added as per suggestions/comments from review: D124158

I failed to see where this was suggested.

---

My suggestion was and still is summarized in: https://reviews.llvm.org/D124158#3486110

@jdoerfert I meant to convey that, As a solution to review comments from D124158, introduced this shuffle attribute to identify __shfl_sync like APIs.

In D125378#3506001, @nikic wrote:

Please specify the semantics of the new LLVM attribute in LangRef -- though I don't really understand why you need an LLVM-side attribute at all.

LG

LG

I think we are lacking a good way to do what I want, let's go ahead with forward declares for now.

I'm not super happy about the forward decls. Should we augment begin declare tatrget with begin declare variant match(device={kind(nohost)}) instead?

LG

Not actually tested, is it?

Conceptually fine with me, @tra?

I though this is illegal to use align 4 loads for align 32 arguments.

Are there unresolved concerns we should address in this review?

Typo(s) in the commit message, otherwise this seems reasonable. LG

I think this is fine, we can argue about the name but it won't matter much. A test would be nice though.

In D124724#3486921, @tianshilei1992 wrote:

always use UIToFP because that is the only case that UpdateVal is a ConstantInt, where we created the constant integer for ++ or -- operation.

In D124158#3480566, @efriedma wrote:

In D124158#3480384, @arsenm wrote:

I'm thinking noundef is a bit of red herring here. The real problem seems to be arising from the assume call which is inserted, which now introduces the assumption that the lane ID must be 0

The optimizer is creating the llvm.assume call based on the violation of the noundef attribute.

LG, FWIW, you can commit these right away.

LG

In D124158#3477649, @efriedma wrote:

The issue you're describing sounds like it's specific to @__shfl_sync. In general, in C++, you aren't allowed to read from an uninitialized variable; see [basic.indet] in the standard. But if your testcase doesn't have undefined behavior, CUDA language rules must somehow allow this particular builtin function to take undef variables as input. (Is this documented somewhere?)

That isn't related to the "convergent" attribute; the transform you're describing doesn't break convergence rules.

For Ex: SimplifyCFG pass removes the branch leading to a BB which has an incoming value that will always trigger undefined behavior. This basically modifies the CFG and combines the basic blocks. This works for CPU execution. But on a GPU, there are intrinsics like "__shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize)", Where the exchange of variable occurs simultaneously for all active threads within the warp. So, here in the cuda/hip kernel, variable var in shuffl_sync may not be initialised, and LLVM IR treats it as undef. Currently all the arguments are tagged with noundef attribute and the above mentioned optimization by SimplifyCFG gets applied and the kernel execution becomes ambiguous. So, the proposed change is to skip adding noundef attribute to arguments when a function has been tagged with convergent attribute.

LG

Finally, I'd suggest to avoid the undef and UB.

Both argument promotion and lining are legal behavior. I think there's no undef/UB for single pass here.

Use set instead of map

Fix rebase issues, address comments

Fix rebase issues

Fix rebase errors, address comments

Address comments

LG