[Attribute] Introduce shuffle attribute to be used for __shfl_sync like cross-lane APIs

[skc7]

Summary

This change introduces shuffle as function attribute in clang and llvm IR. It is used to identify __shfl_sync like cross-lane APIs [allows exchange of variable across all active threads]. At clang codegen, noundef attribute is skipped to arguments and return types for functions with shuffle attribute.

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

[skc7]

clang-format

[skc7]

rebase

[nikic]

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.

[cdevadas]

clang/test/CodeGenHIP/shuffle-attr-verify.hip

34

Add a new line.

clang/test/CodeGenHIP/shuffle-noundef-attr.hip

49

Ditto

[jdoerfert]

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

[skc7]

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.

[aaron.ballman]

clang/include/clang/Basic/AttrDocs.td

1323–1344

This sounds highly specific to those language standards and like an expert-only feature that's going to be a footgun for everyone else. Is there a way these special functions can be identified at codegen time (or within the backend) so that we don't have to expose an attribute for this?

[jdoerfert]

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.

I'm confused. What is the Attribute in IR going to achieve? The problem is during clang codegen. I can see how the C++/Clang attribute part is useful but I would not call it "shuffle".
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:

__attribute__((maybe_undef(0))) int __shfl_sync(int var, int src_lane, int width = warpSize);
__attribute__((maybe_undef(var))) int __shfl_sync(int var, int src_lane, int width = warpSize);

*OR*

int __shfl_sync(int __attribute__((maybe_undef)) var, int src_lane, int width = warpSize);

And the attribute will cause clang either not to place noundef or place freeze for the respective argument(s).

[nhaehnle]

I'm sorry to say that this patch seems confused about semantics. It lacks clear definitions, and in particular for the shuffle attribute in LLVM IR, you almost certainly just want the already existing convergent instead.

If, separately from that, you still find a need to prevent Clang from labeling certain function arguments or return values as noundef, then presumably that should be done with an explicit "maybe undef" attribute instead, as @jdoerfert suggests.

[aaron.ballman]

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.

[jdoerfert]

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.

Works for me.