[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

1319–1340

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.

clang/include/clang/Basic/Attr.td

 
 def Convergent : InheritableAttr {
   let Spellings = [Clang<"convergent">];
   let Subjects = SubjectList<[Function]>;
   let Documentation = [ConvergentDocs];
   let SimpleHandler = 1;
 }
 
+def Shuffle : InheritableAttr {
+  let Spellings = [Clang<"shuffle">];
+  let Subjects = SubjectList<[Function]>;
+  let Documentation = [ShuffleDocs];
+  let SimpleHandler = 1;
+}
+
 def NoInline : DeclOrStmtAttr {
   let Spellings = [GCC<"noinline">, CXX11<"clang", "noinline">,
                    C2x<"clang", "noinline">, Declspec<"noinline">];
   let Accessors = [Accessor<"isClangNoInline", [CXX11<"clang", "noinline">,
                                                 C2x<"clang", "noinline">]>];
   let Documentation = [NoInlineDocs];
   let Subjects = SubjectList<[Function, Stmt], WarnDiag,
                              "functions and statements">;

clang/include/clang/Basic/AttrDocs.td

.. code-block:: c
 
   void convfunc(void) __attribute__((convergent));
   // Setting it as a C++11 attribute is also valid in a C++ program.
   // void convfunc(void) [[clang::convergent]];
 
   }];
 }
 
+def ShuffleDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+The ``shuffle`` attribute can be placed on a function declaration. It indicates
+that the call instructions of a function with this attribute can take undef
+arguments and is still valid.
+
+In languages HIP or CUDA, there are APIs like 
+T __shfl_sync(unsigned mask,T var, int srcLane, int width=warpSize);
+etc which permit exchanging of a variable between threads within a warp without 
+use of shared memory. These APIs allow variable var to be uninitialised in the program.
+Noundef analysis on such APIs can lead to ambiguous kernel execution. 
+So shuffle attribute on a function is used to skip adding noundef attribute to such APIs.
+
+Sample usage:
+.. code-block:: c
+
+  void shufflefunc(void) __attribute__((shuffle));
+  // Setting it as a C++11 attribute is also valid in a C++ program.
+  // void shufflefunc(void) [[clang::shuffle]];
+
+  }];
+}
+
 def NoSplitStackDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 The ``no_split_stack`` attribute disables the emission of the split stack
 preamble for a particular function. It has no effect if ``-fsplit-stack``
 is not specified.
   }];
 }

clang/lib/CodeGen/CGCall.cpp

if (const MatrixType *Matrix = dyn_cast<MatrixType>(QTy))
     return DetermineNoUndef(Matrix->getElementType(), Types, DL, AI, false);
   if (const ArrayType *Array = dyn_cast<ArrayType>(QTy))
     return DetermineNoUndef(Array->getElementType(), Types, DL, AI, false);
 
   // TODO: Some structs may be `noundef`, in specific situations.
   return false;
 }
 
+static bool DetermineNoUndefForShuffle(const Decl *TargetDecl) {
+  if (!TargetDecl)
+    return true;
+
+  // Function has shuffle attribute. 
+  // Skip adding noundef in this case.
+  if (TargetDecl->hasAttr<ShuffleAttr>()) {
+    return false;
+  }
+
+  return true;
+}
+
 /// Construct the IR attribute list of a function or call.
 ///
 /// When adding an attribute, please consider where it should be handled:
 ///
 ///   - getDefaultFunctionAttributes is for attributes that are essentially
 ///     part of the global target configuration (but perhaps can be
 ///     overridden on a per-function basis).  Adding attributes there
 ///     will cause them to also be set in frontends that build on Clang's
if (TargetDecl) {
     if (TargetDecl->hasAttr<ColdAttr>())
       FuncAttrs.addAttribute(llvm::Attribute::Cold);
     if (TargetDecl->hasAttr<HotAttr>())
       FuncAttrs.addAttribute(llvm::Attribute::Hot);
     if (TargetDecl->hasAttr<NoDuplicateAttr>())
       FuncAttrs.addAttribute(llvm::Attribute::NoDuplicate);
     if (TargetDecl->hasAttr<ConvergentAttr>())
       FuncAttrs.addAttribute(llvm::Attribute::Convergent);
+    if (TargetDecl->hasAttr<ShuffleAttr>())
+      FuncAttrs.addAttribute(llvm::Attribute::Shuffle);
 
     if (const FunctionDecl *Fn = dyn_cast<FunctionDecl>(TargetDecl)) {
       AddAttributesFromFunctionProtoType(
           getContext(), FuncAttrs, Fn->getType()->getAs<FunctionProtoType>());
       if (AttrOnCallSite && Fn->isReplaceableGlobalAllocationFunction()) {
         // A sane operator new returns a non-aliasing pointer.
         auto Kind = Fn->getDeclName().getCXXOverloadedOperator();
         if (getCodeGenOpts().AssumeSaneOperatorNew &&
void CodeGenModule::ConstructAttributeList(StringRef Name,
   HasStrictReturn &= getCodeGenOpts().StrictReturn ||
                      !MayDropFunctionReturn(getContext(), RetTy) ||
                      getLangOpts().Sanitize.has(SanitizerKind::Memory) ||
                      getLangOpts().Sanitize.has(SanitizerKind::Return);
 
   // Determine if the return type could be partially undef
   if (CodeGenOpts.EnableNoundefAttrs && HasStrictReturn) {
     if (!RetTy->isVoidType() && RetAI.getKind() != ABIArgInfo::Indirect &&
-        DetermineNoUndef(RetTy, getTypes(), DL, RetAI))
+        DetermineNoUndef(RetTy, getTypes(), DL, RetAI) &&
+        DetermineNoUndefForShuffle(TargetDecl)) {
         RetAttrs.addAttribute(llvm::Attribute::NoUndef);
     }
+  }
 
   switch (RetAI.getKind()) {
   case ABIArgInfo::Extend:
     if (RetAI.isSignExt())
       RetAttrs.addAttribute(llvm::Attribute::SExt);
     else
       RetAttrs.addAttribute(llvm::Attribute::ZExt);
     LLVM_FALLTHROUGH;
if (IRFunctionArgs.hasPaddingArg(ArgNo)) {
             llvm::AttributeSet::get(
                 getLLVMContext(),
                 llvm::AttrBuilder(getLLVMContext()).addAttribute(llvm::Attribute::InReg));
       }
     }
 
     // Decide whether the argument we're handling could be partially undef
     if (CodeGenOpts.EnableNoundefAttrs &&
-        DetermineNoUndef(ParamType, getTypes(), DL, AI)) {
+        DetermineNoUndef(ParamType, getTypes(), DL, AI) &&
+        DetermineNoUndefForShuffle(TargetDecl)) {
         Attrs.addAttribute(llvm::Attribute::NoUndef);
     }
 
     // 'restrict' -> 'noalias' is done in EmitFunctionProlog when we
     // have the corresponding parameter variable.  It doesn't make
     // sense to do it here because parameters are so messed up.
     switch (AI.getKind()) {
     case ABIArgInfo::Extend:
       if (AI.isSignExt())

clang/lib/Headers/__clang_cuda_intrinsics.h

#define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask,    \
   inline __device__ long long __FnName(long long __val, __Type __offset,       \
                                        int __width = warpSize) {               \
     struct __Bits {                                                            \
       int __a, __b;                                                            \
     };                                                                         \
     _Static_assert(sizeof(__val) == sizeof(__Bits));                           \
     _Static_assert(sizeof(__Bits) == 2 * sizeof(int));                         \
     __Bits __tmp;                                                              \
-    memcpy(&__tmp, &__val, sizeof(__val));                                \
+    memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp.__a = ::__FnName(__tmp.__a, __offset, __width);                      \
     __tmp.__b = ::__FnName(__tmp.__b, __offset, __width);                      \
     long long __ret;                                                           \
     memcpy(&__ret, &__tmp, sizeof(__tmp));                                     \
     return __ret;                                                              \
   }                                                                            \
   inline __device__ long __FnName(long __val, __Type __offset,                 \
                                   int __width = warpSize) {                    \

 __MAKE_SHUFFLES(__shfl_xor, __nvvm_shfl_bfly_i32, __nvvm_shfl_bfly_f32, 0x1f,
                 int);
 #pragma pop_macro("__MAKE_SHUFFLES")
 
 #endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300
 
 #if CUDA_VERSION >= 9000
 #if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 300)
+#define __shuffle __attribute__((shuffle))
 // __shfl_sync_* variants available in CUDA-9
 #pragma push_macro("__MAKE_SYNC_SHUFFLES")
 #define __MAKE_SYNC_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic,       \
                              __Mask, __Type)                                   \
-  inline __device__ int __FnName(unsigned int __mask, int __val,               \
+  inline __device__ __shuffle int __FnName(unsigned int __mask, int __val,     \
                                  __Type __offset, int __width = warpSize) {    \
     return __IntIntrinsic(__mask, __val, __offset,                             \
                           ((warpSize - __width) << 8) | (__Mask));             \
   }                                                                            \
-  inline __device__ float __FnName(unsigned int __mask, float __val,           \
+  inline __device__ __shuffle float __FnName(unsigned int __mask, float __val, \
                                    __Type __offset, int __width = warpSize) {  \
     return __FloatIntrinsic(__mask, __val, __offset,                           \
                             ((warpSize - __width) << 8) | (__Mask));           \
   }                                                                            \
-  inline __device__ unsigned int __FnName(unsigned int __mask,                 \
+  inline __device__ __shuffle unsigned int __FnName(unsigned int __mask,       \
                                           unsigned int __val, __Type __offset, \
                                           int __width = warpSize) {            \
     return static_cast<unsigned int>(                                          \
         ::__FnName(__mask, static_cast<int>(__val), __offset, __width));       \
   }                                                                            \
-  inline __device__ long long __FnName(unsigned int __mask, long long __val,   \
+  inline __device__ __shuffle long long __FnName(unsigned int __mask,          \
+                                      long long __val,                         \
                                        __Type __offset,                        \
                                        int __width = warpSize) {               \
     struct __Bits {                                                            \
       int __a, __b;                                                            \
     };                                                                         \
     _Static_assert(sizeof(__val) == sizeof(__Bits));                           \
     _Static_assert(sizeof(__Bits) == 2 * sizeof(int));                         \
     __Bits __tmp;                                                              \
     memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp.__a = ::__FnName(__mask, __tmp.__a, __offset, __width);              \
     __tmp.__b = ::__FnName(__mask, __tmp.__b, __offset, __width);              \
     long long __ret;                                                           \
     memcpy(&__ret, &__tmp, sizeof(__tmp));                                     \
     return __ret;                                                              \
   }                                                                            \
-  inline __device__ unsigned long long __FnName(                               \
+  inline __device__ __shuffle unsigned long long __FnName(                     \
       unsigned int __mask, unsigned long long __val, __Type __offset,          \
       int __width = warpSize) {                                                \
     return static_cast<unsigned long long>(::__FnName(                         \
         __mask, static_cast<unsigned long long>(__val), __offset, __width));   \
   }                                                                            \
-  inline __device__ long __FnName(unsigned int __mask, long __val,             \
+  inline __device__ __shuffle long __FnName(unsigned int __mask, long __val,   \
                                   __Type __offset, int __width = warpSize) {   \
     _Static_assert(sizeof(long) == sizeof(long long) ||                        \
                    sizeof(long) == sizeof(int));                               \
     if (sizeof(long) == sizeof(long long)) {                                   \
       return static_cast<long>(::__FnName(                                     \
           __mask, static_cast<long long>(__val), __offset, __width));          \
     } else if (sizeof(long) == sizeof(int)) {                                  \
       return static_cast<long>(                                                \
           ::__FnName(__mask, static_cast<int>(__val), __offset, __width));     \
     }                                                                          \
   }                                                                            \
-  inline __device__ unsigned long __FnName(                                    \
+  inline __device__ __shuffle unsigned long __FnName(                          \
       unsigned int __mask, unsigned long __val, __Type __offset,               \
       int __width = warpSize) {                                                \
     return static_cast<unsigned long>(                                         \
         ::__FnName(__mask, static_cast<long>(__val), __offset, __width));      \
   }                                                                            \
-  inline __device__ double __FnName(unsigned int __mask, double __val,         \
+  inline __device__ __shuffle double __FnName(unsigned int __mask,             \
+                                    double __val,                              \
                                     __Type __offset, int __width = warpSize) { \
     long long __tmp;                                                           \
     _Static_assert(sizeof(__tmp) == sizeof(__val));                            \
     memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp = ::__FnName(__mask, __tmp, __offset, __width);                      \
     double __ret;                                                              \
     memcpy(&__ret, &__tmp, sizeof(__ret));                                     \
     return __ret;                                                              \

clang/lib/Sema/SemaDeclAttr.cpp

case ParsedAttr::AT_AMDGPUWavesPerEU:
     handleAMDGPUWavesPerEUAttr(S, D, AL);
     break;
   case ParsedAttr::AT_AMDGPUNumSGPR:
     handleAMDGPUNumSGPRAttr(S, D, AL);
     break;
   case ParsedAttr::AT_AMDGPUNumVGPR:
     handleAMDGPUNumVGPRAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_Shuffle:
+    handleSimpleAttribute<ShuffleAttr>(S, D, AL);
+    break;
   case ParsedAttr::AT_AVRSignal:
     handleAVRSignalAttr(S, D, AL);
     break;
   case ParsedAttr::AT_BPFPreserveAccessIndex:
     handleBPFPreserveAccessIndexAttr(S, D, AL);
     break;
   case ParsedAttr::AT_BTFDeclTag:
     handleBTFDeclTagAttr(S, D, AL);

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

+// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm  %s \
+// RUN:   -o - | FileCheck %s
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+#define __shuffle __attribute__((shuffle))
+#define HYPRE_WARP_SIZE 64
+
+static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE;
+
+__device__ static inline unsigned int __lane_id() {
+    return  __builtin_amdgcn_mbcnt_hi(
+        -1, __builtin_amdgcn_mbcnt_lo(-1, 0));
+}
+
+// CHECK: define linkonce_odr i32 @_Z11__shfl_synciii(i32 %var, i32 %src_lane, i32 %width) #[[attr1:[0-9]+]]
+__device__
+inline
+int __shuffle __shfl_sync(int var, int src_lane, int width = warpSize) {
+    int self = __lane_id();
+    int index = src_lane + (self & ~(width-1));
+    return __builtin_amdgcn_ds_bpermute(index<<2, var);
+}
+
+__global__ void
+shufflekernel()
+{
+    int t;
+    int res, res1;
+    res = __shfl_sync(HYPRE_WARP_SIZE, t, 0);
+}
+
+// CHECK-DAG: attributes #[[attr1]] = { {{[^}]*}}shuffle{{[^}]*}} }
 No newline at end of file

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

+// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip -fcuda-is-device -emit-llvm  %s \
+// RUN:   -o - | FileCheck %s
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+#define __shuffle __attribute__((shuffle))
+#define HYPRE_WARP_SIZE 64
+
+static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE;
+
+__device__ static inline unsigned int __lane_id() {
+    return  __builtin_amdgcn_mbcnt_hi(
+        -1, __builtin_amdgcn_mbcnt_lo(-1, 0));
+}
+
+__device__
+inline
+int __shfl(int var, int src_lane, int width = warpSize) {
+    int self = __lane_id();
+    int index = src_lane + (self & ~(width-1));
+    return __builtin_amdgcn_ds_bpermute(index<<2, var);
+}
+
+template <typename T>
+static __device__
+T __shuffle __shfl_sync(unsigned mask, T val, int src_line, int width=HYPRE_WARP_SIZE)
+{
+   return __shfl(val, src_line, width);
+}
+
+template <typename T>
+static __device__
+T __shfl_sync_1(unsigned mask, T val, int src_line, int width=HYPRE_WARP_SIZE)
+{
+   return __shfl(val, src_line, width);
+}
+
+// CHECK-LABEL: @_Z13shufflekernelv(
+// CHECK: call i32 @_ZL11__shfl_syncIiET_jS0_ii(i32 64, i32 %0, i32 0, i32 64)
+// CHECK: call noundef i32 @_ZL13__shfl_sync_1IiET_jS0_ii(i32 noundef 64, i32 noundef %1, i32 noundef 0, i32 noundef 64)
+__global__ void
+shufflekernel()
+{
+    int t;
+    int res, res1;
+    res = __shfl_sync(HYPRE_WARP_SIZE, t, 0);
+    res1 = __shfl_sync_1(HYPRE_WARP_SIZE, t, 0);
+}
 No newline at end of file

clang/test/Misc/pragma-attribute-supported-attributes-list.test

 // CHECK-NEXT: Restrict (SubjectMatchRule_function)
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
 // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
 // CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member)
+// CHECK-NEXT: Shuffle (SubjectMatchRule_function)
 // CHECK-NEXT: SpeculativeLoadHardening (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: StandaloneDebug (SubjectMatchRule_record)
 // CHECK-NEXT: SwiftAsync (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: SwiftAsyncContext (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: SwiftAsyncError (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: SwiftAsyncName (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: SwiftAttr ()
 // CHECK-NEXT: SwiftBridgedTypedef (SubjectMatchRule_type_alias)

llvm/include/llvm/Bitcode/LLVMBitCodes.h

enum AttributeKindCodes {
   ATTR_KIND_VSCALE_RANGE = 74,
   ATTR_KIND_SWIFT_ASYNC = 75,
   ATTR_KIND_NO_SANITIZE_COVERAGE = 76,
   ATTR_KIND_ELEMENTTYPE = 77,
   ATTR_KIND_DISABLE_SANITIZER_INSTRUMENTATION = 78,
   ATTR_KIND_NO_SANITIZE_BOUNDS = 79,
   ATTR_KIND_ALLOC_ALIGN = 80,
   ATTR_KIND_ALLOCATED_POINTER = 81,
+  ATTR_KIND_SHUFFLE = 82,
 };
 
 enum ComdatSelectionKindCodes {
   COMDAT_SELECTION_KIND_ANY = 1,
   COMDAT_SELECTION_KIND_EXACT_MATCH = 2,
   COMDAT_SELECTION_KIND_LARGEST = 3,
   COMDAT_SELECTION_KIND_NO_DUPLICATES = 4,
   COMDAT_SELECTION_KIND_SAME_SIZE = 5,

llvm/include/llvm/IR/Attributes.td

 def WriteOnly : EnumAttr<"writeonly", [FnAttr, ParamAttr]>;
 
 /// Zero extended before/after call.
 def ZExt : EnumAttr<"zeroext", [ParamAttr, RetAttr]>;
 
 /// Function is required to make Forward Progress.
 def MustProgress : EnumAttr<"mustprogress", [FnAttr]>;
 
+/// Function is a __shfl_sync like API.
+def Shuffle : EnumAttr<"shuffle", [FnAttr]>;
+
 /// Target-independent string attributes.
 def LessPreciseFPMAD : StrBoolAttr<"less-precise-fpmad">;
 def NoInfsFPMath : StrBoolAttr<"no-infs-fp-math">;
 def NoNansFPMath : StrBoolAttr<"no-nans-fp-math">;
 def ApproxFuncFPMath : StrBoolAttr<"approx-func-fp-math">;
 def NoSignedZerosFPMath : StrBoolAttr<"no-signed-zeros-fp-math">;
 def UnsafeFPMath : StrBoolAttr<"unsafe-fp-math">;
 def NoJumpTables : StrBoolAttr<"no-jump-tables">;

llvm/include/llvm/IR/Function.h

public:
   bool willReturn() const { return hasFnAttribute(Attribute::WillReturn); }
   void setWillReturn() { addFnAttr(Attribute::WillReturn); }
 
   /// Get what kind of unwind table entry to generate for this function.
   UWTableKind getUWTableKind() const {
     return AttributeSets.getUWTableKind();
   }
 
+  /// Determine if the function is __shfl_sync like.
+  bool isShuffle() const {
+    return hasFnAttribute(Attribute::Shuffle);
+  }
+  void setShuffle() { addFnAttr(Attribute::Shuffle); }
+
   /// True if the ABI mandates (or the user requested) that this
   /// function be in a unwind table.
   bool hasUWTable() const {
     return getUWTableKind() != UWTableKind::None;
   }
   void setUWTableKind(UWTableKind K) {
     addFnAttr(Attribute::getWithUWTableKind(getContext(), K));
   }

llvm/lib/Bitcode/Reader/BitcodeReader.cpp

static Attribute::AttrKind getAttrFromCode(uint64_t Code) {
   case bitc::ATTR_KIND_NOUNDEF:
     return Attribute::NoUndef;
   case bitc::ATTR_KIND_BYREF:
     return Attribute::ByRef;
   case bitc::ATTR_KIND_MUSTPROGRESS:
     return Attribute::MustProgress;
   case bitc::ATTR_KIND_HOT:
     return Attribute::Hot;
+  case bitc::ATTR_KIND_SHUFFLE:
+    return Attribute::Shuffle;
   }
 }
 
 Error BitcodeReader::parseAlignmentValue(uint64_t Exponent,
                                          MaybeAlign &Alignment) {
   // Note: Alignment in bitcode files is incremented by 1, so that zero
   // can be used for default alignment.
   if (Exponent > Value::MaxAlignmentExponent + 1)

llvm/lib/Bitcode/Writer/BitcodeWriter.cpp

case Attribute::MustProgress:
     return bitc::ATTR_KIND_MUSTPROGRESS;
   case Attribute::EndAttrKinds:
     llvm_unreachable("Can not encode end-attribute kinds marker.");
   case Attribute::None:
     llvm_unreachable("Can not encode none-attribute.");
   case Attribute::EmptyKey:
   case Attribute::TombstoneKey:
     llvm_unreachable("Trying to encode EmptyKey/TombstoneKey");
+  case Attribute::Shuffle:
+    return bitc::ATTR_KIND_SHUFFLE;
   }
 
   llvm_unreachable("Trying to encode unknown attribute");
 }
 
 void ModuleBitcodeWriter::writeAttributeGroupTable() {
   const std::vector<ValueEnumerator::IndexAndAttrSet> &AttrGrps =
       VE.getAttributeGroups();

llvm/lib/IR/Attributes.cpp

if (ASK & ASK_UNSAFE_TO_DROP)
       Incompatible.addAttribute(Attribute::Nest)
           .addAttribute(Attribute::SwiftError)
           .addAttribute(Attribute::Preallocated)
           .addAttribute(Attribute::InAlloca)
           .addAttribute(Attribute::ByVal)
           .addAttribute(Attribute::StructRet)
           .addAttribute(Attribute::ByRef)
           .addAttribute(Attribute::ElementType)
-          .addAttribute(Attribute::AllocatedPointer);
+          .addAttribute(Attribute::AllocatedPointer)
+          .addAttribute(Attribute::Shuffle);
   }
 
     // Attributes that only apply to pointers or vectors of pointers.
   if (!Ty->isPtrOrPtrVectorTy()) {
     if (ASK & ASK_SAFE_TO_DROP)
       Incompatible.addAttribute(Attribute::Alignment);
   }
 

llvm/lib/Transforms/Utils/CodeExtractor.cpp

if (Attr.isStringAttribute()) {
       case Attribute::StackProtectReq:
       case Attribute::StackProtectStrong:
       case Attribute::StrictFP:
       case Attribute::UWTable:
       case Attribute::VScaleRange:
       case Attribute::NoCfCheck:
       case Attribute::MustProgress:
       case Attribute::NoProfile:
+      case Attribute::Shuffle:
         break;
       // These attributes cannot be applied to functions.
       case Attribute::Alignment:
       case Attribute::AllocatedPointer:
       case Attribute::AllocAlign:
       case Attribute::ByVal:
       case Attribute::Dereferenceable:
       case Attribute::DereferenceableOrNull: