This is an archive of the discontinued LLVM Phabricator instance.

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

Authored by skc7 on May 11 2022, 5:14 AM.

Details

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

Diff Detail

Event Timeline

skc7 created this revision.May 11 2022, 5:14 AM
Herald added a project: Restricted Project. · View Herald Transcript
skc7 requested review of this revision.May 11 2022, 5:14 AM
Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptMay 11 2022, 5:14 AM
skc7 updated this revision to Diff 428635.May 11 2022, 5:41 AM

clang-format

nikic requested changes to this revision.May 11 2022, 6:13 AM
nikic added a subscriber: 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.

This revision now requires changes to proceed.May 11 2022, 6:13 AM
cdevadas added inline comments.May 11 2022, 6:36 AM
clang/test/CodeGenHIP/shuffle-attr-verify.hip
34

Add a new line.

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

Ditto

jdoerfert requested changes to this revision.May 11 2022, 7:26 AM

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 added a comment.May 11 2022, 8:27 AM

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 added inline comments.May 11 2022, 8:38 AM
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?

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 requested changes to this revision.May 11 2022, 9:36 AM

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.

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.

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.