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

[Johannes Doerfert via Phabricator via cfe-commits]

jdoerfert added a comment.

In D125378#3506383 <https://reviews.llvm.org/D125378#3506383>, @skc7 wrote:

> In D125378#3506215 <https://reviews.llvm.org/D125378#3506215>, @jdoerfert wrote:
>
>> In D125378#3506001 <https://reviews.llvm.org/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 <https://reviews.llvm.org/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 <https://reviews.llvm.org/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).


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D125378/new/

https://reviews.llvm.org/D125378