[AMDGPU] Expose flat work group size, register and wave control attributes

[kzhuravl]

Summary

attribute((amdgpu_flat_work_group_size(<min>, <max>))) - request minimum and maximum flat work group size
attribute((amdgpu_waves_per_eu(<min>[, <max>]))) - request minimum and/or maximum waves per execution unit

[kzhuravl]

Update docs in AttrDocs.td

[kzhuravl]

Fix minor typos

[aaron.ballman]

include/clang/Basic/Attr.td

1054–1057

This FIXME is now detached from what "this" is referring to. Can you clarify by mentioning that the Subjects list is what should be modified?

1067

Looking at the documentation, are you sure this should be a VariadicUnsignedArgument? It seems like this should be an UnsignedArgument with the optional bit set. Or can you pass multiple Max values?

include/clang/Basic/AttrDocs.td

1145

There's an extra space between "can" and "hide".

1151

Comma after "executed".

1154

An error instead of The error.

include/clang/Basic/DiagnosticSemaKinds.td

2383

I think this should say the attribute argument is invalid, rather than the attribute parameter.

lib/CodeGen/TargetInfo.cpp

6958

Should be const auto *Attr.

6967–6969

Elide braces?

6972

const auto * here as well.

6983–6985

Elide braces?

6995

const auto * here as well.

lib/Sema/SemaDeclAttr.cpp

4932

I don't think this case is required (here, or elsewhere). getArgAsExpr() already returns an Expr *.

4944–4945

Please include this text in the diagnostic rather than hard-coding it here. You can use %select to differentiate between different text snippets in the diagnostic, and the diagnostic doesn't need to continue to try to be generic. Same comment applies elsewhere.

4973

This will change if you use the optional bit rather than a variadic argument, but this silently eats attributes that have more than two arguments, which is not a good user experience.

6033–6037

This list is getting to the point where we really need to start handling this in Attr.td soon. Are you planning to work on more AMDGPU attributes in the near future?

test/SemaOpenCL/amdgpu-attrs.cl

53

I'd like to see a test that shows amdgpu_waves_per_eu with > 2 arguments to make sure it's handled appropriately.

[kzhuravl]

Address review feedback

[kzhuravl]

include/clang/Basic/Attr.td

1067

You are right. Switched to UnsignedArgument since only one Max is allowed. Thanks.

lib/Sema/SemaDeclAttr.cpp

6033–6037

I agree, and yes, few more attributes will need to be added in the near future. Would it be ok if I change it to start handling in Attr.td after this change, but before other attributes are added?

[aaron.ballman]

lib/Sema/SemaDeclAttr.cpp

4952

Is it okay to supply 0, 0 as the min, max arguments?

4982

Is it okay to supply 0, 0 as the min, max arguments?

6033–6038

Yes, totally fine to be a follow-up patch. I was hoping it would look something like (we can bikeshed the name):

def SomeAttr {
  /* Blah */
}

def SomeOtherAttr {
  let RequiredCompanionAttributes = [SomeAttr];
}

[kzhuravl]

Mention 0, 0 case in the docs.

[kzhuravl]

lib/Sema/SemaDeclAttr.cpp

4952

Yes, I mentioned 0, 0 case in the docs.

4982

Yes, I mentioned 0, 0 case in the docs.

6033–6038

This seems like a good start. Thanks :)

[aaron.ballman]

LGTM, thank you!

[kzhuravl]

Thanks for the review Aaron!

Tom, would you be able to do a final glance over?

[• tstellarAMD]

LGTM.

include/clang/Basic/Attr.td

 // implementation of OpenCL be free to ignore all attributes and the
 // resulting executable binary will produce the same result."
 //
 // However, only AMD GPU targets will emit the corresponding IR
 // attribute.
 //
 // FIXME: This provides a sub-optimal error message if you attempt to
 // use this in CUDA, since CUDA does not use the same terminology.
-def AMDGPUNumVGPR : InheritableAttr {
-  let Spellings = [GNU<"amdgpu_num_vgpr">];
-  let Args = [UnsignedArgument<"NumVGPR">];
-  let Documentation = [AMDGPUNumVGPRDocs];
-
+//
 // FIXME: This should be for OpenCLKernelFunction, but is not to
 // workaround needing to see kernel attribute before others to know if
 // this should be rejected on non-kernels.
-  let Subjects = SubjectList<[Function], ErrorDiag,
-                             "ExpectedKernelFunction">;
+
+def AMDGPUFlatWorkGroupSize : InheritableAttr {
+  let Spellings = [GNU<"amdgpu_flat_work_group_size">];
+  let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max">];
+  let Documentation = [AMDGPUFlatWorkGroupSizeDocs];
+  let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">;
+}
+
+def AMDGPUWavesPerEU : InheritableAttr {
+  let Spellings = [GNU<"amdgpu_waves_per_eu">];
+  let Args = [UnsignedArgument<"Min">, VariadicUnsignedArgument<"Max">];
+  let Documentation = [AMDGPUWavesPerEUDocs];
+  let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">;
 }
 
 def AMDGPUNumSGPR : InheritableAttr {
   let Spellings = [GNU<"amdgpu_num_sgpr">];
   let Args = [UnsignedArgument<"NumSGPR">];
   let Documentation = [AMDGPUNumSGPRDocs];
-  let Subjects = SubjectList<[Function], ErrorDiag,
-                              "ExpectedKernelFunction">;
+  let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">;
+}
+
+def AMDGPUNumVGPR : InheritableAttr {
+  let Spellings = [GNU<"amdgpu_num_vgpr">];
+  let Args = [UnsignedArgument<"NumVGPR">];
+  let Documentation = [AMDGPUNumVGPRDocs];
+  let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">;
 }
 
 def NoSplitStack : InheritableAttr {
   let Spellings = [GCC<"no_split_stack">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [NoSplitStackDocs];
 }
 

include/clang/Basic/AttrDocs.td

 solely for use in an ``assert()`` statement, which makes the local variable
 unused when ``NDEBUG`` is defined.
 
 The attribute may be applied to the declaration of a class, a typedef, a
 variable, a function or method, a function parameter, an enumeration, an
 enumerator, a non-static data member, or a label.
 
 .. code-block: c++
-  #include <cassert>
-
-  [[maybe_unused]] void f([[maybe_unused]] bool thing1,
-                          [[maybe_unused]] bool thing2) {
-    [[maybe_unused]] bool b = thing1 && thing2;
-    assert(b);
+  #include <cassert>
+
+  [[maybe_unused]] void f([[maybe_unused]] bool thing1,
+                          [[maybe_unused]] bool thing2) {
+    [[maybe_unused]] bool b = thing1 && thing2;
+    assert(b);
   }
   }];
 }
 
 def WarnUnusedResultsDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "nodiscard, warn_unused_result, clang::warn_unused_result, gnu::warn_unused_result";
   let Content  = [{
 Clang supports the ability to diagnose when the results of a function call
 expression are discarded under suspicious circumstances. A diagnostic is
 generated when a function or its return type is marked with ``[[nodiscard]]``
 (or ``__attribute__((warn_unused_result))``) and the function call appears as a
 potentially-evaluated discarded-value expression that is not explicitly cast to
 `void`.
 
 .. code-block: c++
-  struct [[nodiscard]] error_info { /*...*/ };
-  error_info enable_missile_safety_mode();
-  
-  void launch_missiles();
-  void test_missiles() {
-    enable_missile_safety_mode(); // diagnoses
-    launch_missiles();
-  }
-  error_info &foo();
+  struct [[nodiscard]] error_info { /*...*/ };
+  error_info enable_missile_safety_mode();
+  
+  void launch_missiles();
+  void test_missiles() {
+    enable_missile_safety_mode(); // diagnoses
+    launch_missiles();
+  }
+  error_info &foo();
   void f() { foo(); } // Does not diagnose, error_info is a reference.
   }];
 }
 
 def FallthroughDocs : Documentation {
   let Category = DocCatStmt;
   let Heading = "fallthrough, clang::fallthrough";
   let Content = [{

 for the function.
 
 Example "subtarget features" from the x86 backend include: "mmx", "sse", "sse4.2",
 "avx", "xop" and largely correspond to the machine specific options handled by
 the front end.
 }];
 }
 
-def DocCatAMDGPURegisterAttributes :
-  DocumentationCategory<"AMD GPU Register Attributes"> {
-  let Content = [{
-Clang supports attributes for controlling register usage on AMD GPU
-targets. These attributes may be attached to a kernel function
-definition and is an optimization hint to the backend for the maximum
-number of registers to use. This is useful in cases where register
-limited occupancy is known to be an important factor for the
-performance for the kernel.
-
-The semantics are as follows:
-
-- The backend will attempt to limit the number of used registers to
-  the specified value, but the exact number used is not
-  guaranteed. The number used may be rounded up to satisfy the
-  allocation requirements or ABI constraints of the subtarget. For
-  example, on Southern Islands VGPRs may only be allocated in
-  increments of 4, so requesting a limit of 39 VGPRs will really
-  attempt to use up to 40. Requesting more registers than the
-  subtarget supports will truncate to the maximum allowed. The backend
-  may also use fewer registers than requested whenever possible.
-
-- 0 implies the default no limit on register usage.
-
-- Ignored on older VLIW subtargets which did not have separate scalar
-  and vector registers, R600 through Northern Islands.
+def DocCatAMDGPUAttributes : DocumentationCategory<"AMD GPU Attributes">;
 
-}];
+def AMDGPUFlatWorkGroupSizeDocs : Documentation {
+  let Category = DocCatAMDGPUAttributes;
+  let Content = [{}];
 }
 
-
-def AMDGPUNumVGPRDocs : Documentation {
-  let Category = DocCatAMDGPURegisterAttributes;
-  let Content = [{
-Clang supports the
-``__attribute__((amdgpu_num_vgpr(<num_registers>)))`` attribute on AMD
-Southern Islands GPUs and later for controlling the number of vector
-registers. A typical value would be between 4 and 256 in increments
-of 4.
-}];
+def AMDGPUWavesPerEUDocs : Documentation {
+  let Category = DocCatAMDGPUAttributes;
+  let Content = [{}];
 }
 
 def AMDGPUNumSGPRDocs : Documentation {
-  let Category = DocCatAMDGPURegisterAttributes;
-  let Content = [{
+  let Category = DocCatAMDGPUAttributes;
+  let Content = [{}];
+}
 
-Clang supports the
-``__attribute__((amdgpu_num_sgpr(<num_registers>)))`` attribute on AMD
-Southern Islands GPUs and later for controlling the number of scalar
-registers. A typical value would be between 8 and 104 in increments of
-8.
-
-Due to common instruction constraints, an additional 2-4 SGPRs are
-typically required for internal use depending on features used. This
-value is a hint for the total number of SGPRs to use, and not the
-number of user SGPRs, so no special consideration needs to be given
-for these.
-}];
+def AMDGPUNumVGPRDocs : Documentation {
+  let Category = DocCatAMDGPUAttributes;
+  let Content = [{}];
 }
 
 def DocCatCallingConvs : DocumentationCategory<"Calling Conventions"> {
   let Content = [{
 Clang supports several different calling conventions, depending on the target
 platform and architecture. The calling convention used for a function determines
 how parameters are passed, how results are returned to the caller, and other
 low-level details of calling a function.

 
 def MSABIDocs : Documentation {
   let Category = DocCatCallingConvs;
   let Content = [{
 On non-Windows x86_64 targets, this attribute changes the calling convention of
 a function to match the default convention used on Windows x86_64. This
 attribute has no effect on Windows targets or non-x86_64 targets.
   }];
 }
 
 def StdCallDocs : Documentation {
   let Category = DocCatCallingConvs;
   let Content = [{
 On 32-bit x86 targets, this attribute changes the calling convention of a
 function to clear parameters off of the stack on return. This convention does
 not support variadic calls or unprototyped functions in C, and has no effect on
 x86_64 targets. This calling convention is used widely by the Windows API and
 COM applications.  See the documentation for `__stdcall`_ on MSDN.
 
 .. _`__stdcall`: http://msdn.microsoft.com/en-us/library/zxk0tw93.aspx
   }];
 }
 
 def FastCallDocs : Documentation {
   let Category = DocCatCallingConvs;
   let Content = [{

include/clang/Basic/DiagnosticSemaKinds.td

 def err_swift_error_result_not_last : Error<
   "'swift_error_result' parameter must be last parameter of function">;
 def err_swift_error_result_not_after_swift_context : Error<
   "'swift_error_result' parameter must follow 'swift_context' parameter">;
 def err_swift_abi_parameter_wrong_type : Error<
   "'%0' parameter must have pointer%select{| to unqualified pointer}1 type; "
   "type here is %2">;
 
+def err_attribute_argument_invalid : Error<
+  "%0 attribute parameter is invalid: %1">;
 def err_attribute_argument_is_zero : Error<
   "%0 attribute must be greater than 0">;
 def warn_attribute_argument_n_negative : Warning<
   "%0 attribute parameter %1 is negative and will be ignored">,
   InGroup<CudaCompat>;
 def err_property_function_in_objc_container : Error<
   "use of Objective-C property in function nested in Objective-C "
   "container not supported, move function outside its container">;

lib/CodeGen/TargetInfo.cpp

public:
   unsigned getOpenCLKernelCallingConv() const override;
 };
 
 }
 
 static void appendOpenCLVersionMD (CodeGen::CodeGenModule &CGM);
 
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(
     const Decl *D,
     llvm::GlobalValue *GV,
     CodeGen::CodeGenModule &M) const {
   const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D);
   if (!FD)
     return;
 
-  if (const auto Attr = FD->getAttr<AMDGPUNumVGPRAttr>()) {
-    llvm::Function *F = cast<llvm::Function>(GV);
-    uint32_t NumVGPR = Attr->getNumVGPR();
-    if (NumVGPR != 0)
-      F->addFnAttr("amdgpu_num_vgpr", llvm::utostr(NumVGPR));
+
+  if (const auto Attr = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>()) {
+    unsigned Min = Attr->getMin();
+    unsigned Max = Attr->getMax();
+
+    if (Min != 0) {
+      assert(Min <= Max && "Min must be less than or equal Max");
+
+      std::string AttrVal = llvm::utostr(Min) + "," + llvm::utostr(Max);
+      F->addFnAttr("amdgpu-flat-work-group-size", AttrVal);
+    } else {
+      assert(Max == 0 && "Max must be zero");
+    }
+  }
+
+  if (const auto Attr = FD->getAttr<AMDGPUWavesPerEUAttr>()) {
+    unsigned Min = Attr->getMin();
+    unsigned Max = *Attr->max_begin();
+
+    if (Min != 0) {
+      assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max");
+
+      std::string AttrVal = llvm::utostr(Min);
+      if (Max != 0)
+        AttrVal = AttrVal + "," + llvm::utostr(Max);
+      F->addFnAttr("amdgpu-waves-per-eu", AttrVal);
+    } else {
+      assert(Max == 0 && "Max must be zero");
+    }
   }
 
   if (const auto Attr = FD->getAttr<AMDGPUNumSGPRAttr>()) {
-    llvm::Function *F = cast<llvm::Function>(GV);
     unsigned NumSGPR = Attr->getNumSGPR();
+
     if (NumSGPR != 0)
-      F->addFnAttr("amdgpu_num_sgpr", llvm::utostr(NumSGPR));
+      F->addFnAttr("amdgpu-num-sgpr", llvm::utostr(NumSGPR));
+  }
+
+  if (const auto Attr = FD->getAttr<AMDGPUNumVGPRAttr>()) {
+    uint32_t NumVGPR = Attr->getNumVGPR();
+
+    if (NumVGPR != 0)
+      F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR));
   }
 
   appendOpenCLVersionMD(M);
 }
 
 unsigned AMDGPUTargetCodeGenInfo::getOpenCLKernelCallingConv() const {
   return llvm::CallingConv::AMDGPU_KERNEL;
 }

lib/Sema/SemaDeclAttr.cpp

case llvm::Triple::x86_64:
     handleAnyX86InterruptAttr(S, D, Attr);
     break;
   default:
     handleARMInterruptAttr(S, D, Attr);
     break;
   }
 }
 
-static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D,
+static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D,
                                               const AttributeList &Attr) {
-  uint32_t NumRegs;
-  Expr *NumRegsExpr = static_cast<Expr *>(Attr.getArgAsExpr(0));
-  if (!checkUInt32Argument(S, Attr, NumRegsExpr, NumRegs))
+  uint32_t Min = 0;
+  Expr *MinExpr = static_cast<Expr*>(Attr.getArgAsExpr(0));
+  if (!checkUInt32Argument(S, Attr, MinExpr, Min))
+    return;
+
+  uint32_t Max = 0;
+  Expr *MaxExpr = static_cast<Expr*>(Attr.getArgAsExpr(1));
+  if (!checkUInt32Argument(S, Attr, MaxExpr, Max))
+    return;
+
+  if (Min == 0 && Max != 0) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_invalid)
+      << Attr.getName()
+      << "maximum flat work group size must be zero since minimum flat work "
+         "group size is zero";
+    return;
+  }
+  if (Min > Max) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_invalid)
+      << Attr.getName()
+      << "minimum flat work group size must not be greater than maximum flat "
+         "work group size";
     return;
+  }
+
+  D->addAttr(::new (S.Context)
+             AMDGPUFlatWorkGroupSizeAttr(Attr.getLoc(), S.Context, Min, Max,
+                                         Attr.getAttributeSpellingListIndex()));
+}
+
+static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D,
+                                       const AttributeList &Attr) {
+  uint32_t Min = 0;
+  Expr *MinExpr = static_cast<Expr*>(Attr.getArgAsExpr(0));
+  if (!checkUInt32Argument(S, Attr, MinExpr, Min))
+    return;
+
+  uint32_t Max = 0;
+  if (Attr.getNumArgs() == 2) {
+    Expr *MaxExpr = static_cast<Expr*>(Attr.getArgAsExpr(1));
+    if (!checkUInt32Argument(S, Attr, MaxExpr, Max))
+      return;
+  }
+
+  if (Min == 0 && Max != 0) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_invalid)
+      << Attr.getName()
+      << "maximum number of waves per execution unit must be zero since "
+         "minimum number of waves per execution unit is zero";
+    return;
+  }
+  if (Max != 0 && Min > Max) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_invalid)
+      << Attr.getName()
+      << "minimum number of waves per execution unit must not be greater than "
+         "maximum number of waves per execution unit";
+    return;
+  }
 
   D->addAttr(::new (S.Context)
-             AMDGPUNumVGPRAttr(Attr.getLoc(), S.Context,
-                               NumRegs,
+             AMDGPUWavesPerEUAttr(Attr.getLoc(), S.Context,
+                                  Min, &Max, sizeof(Max),
                                   Attr.getAttributeSpellingListIndex()));
 }
 
 static void handleAMDGPUNumSGPRAttr(Sema &S, Decl *D,
                                     const AttributeList &Attr) {
-  uint32_t NumRegs;
-  Expr *NumRegsExpr = static_cast<Expr *>(Attr.getArgAsExpr(0));
-  if (!checkUInt32Argument(S, Attr, NumRegsExpr, NumRegs))
+  uint32_t NumSGPR = 0;
+  Expr *NumSGPRExpr = static_cast<Expr*>(Attr.getArgAsExpr(0));
+  if (!checkUInt32Argument(S, Attr, NumSGPRExpr, NumSGPR))
     return;
 
   D->addAttr(::new (S.Context)
-             AMDGPUNumSGPRAttr(Attr.getLoc(), S.Context,
-                               NumRegs,
+             AMDGPUNumSGPRAttr(Attr.getLoc(), S.Context, NumSGPR,
+                               Attr.getAttributeSpellingListIndex()));
+}
+
+static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D,
+                                    const AttributeList &Attr) {
+  uint32_t NumVGPR = 0;
+  Expr *NumVGPRExpr = static_cast<Expr*>(Attr.getArgAsExpr(0));
+  if (!checkUInt32Argument(S, Attr, NumVGPRExpr, NumVGPR))
+    return;
+
+  D->addAttr(::new (S.Context)
+             AMDGPUNumVGPRAttr(Attr.getLoc(), S.Context, NumVGPR,
                                Attr.getAttributeSpellingListIndex()));
 }
 
 static void handleX86ForceAlignArgPointerAttr(Sema &S, Decl *D,
                                               const AttributeList& Attr) {
   // If we try to apply it to a function pointer, don't warn, but don't
   // do anything, either. It doesn't matter anyway, because there's nothing
   // special about calling a force_align_arg_pointer function.
case AttributeList::AT_DLLImport:
     break;
   case AttributeList::AT_Mips16:
     handleSimpleAttributeWithExclusions<Mips16Attr, MipsInterruptAttr>(S, D,
                                                                        Attr);
     break;
   case AttributeList::AT_NoMips16:
     handleSimpleAttribute<NoMips16Attr>(S, D, Attr);
     break;
-  case AttributeList::AT_AMDGPUNumVGPR:
-    handleAMDGPUNumVGPRAttr(S, D, Attr);
+  case AttributeList::AT_AMDGPUFlatWorkGroupSize:
+    handleAMDGPUFlatWorkGroupSizeAttr(S, D, Attr);
+    break;
+  case AttributeList::AT_AMDGPUWavesPerEU:
+    handleAMDGPUWavesPerEUAttr(S, D, Attr);
     break;
   case AttributeList::AT_AMDGPUNumSGPR:
     handleAMDGPUNumSGPRAttr(S, D, Attr);
     break;
+  case AttributeList::AT_AMDGPUNumVGPR:
+    handleAMDGPUNumVGPRAttr(S, D, Attr);
+    break;
   case AttributeList::AT_IBAction:
     handleSimpleAttribute<IBActionAttr>(S, D, Attr);
     break;
   case AttributeList::AT_IBOutlet:
     handleIBOutlet(S, D, Attr);
     break;
   case AttributeList::AT_IBOutletCollection:
     handleIBOutletCollection(S, D, Attr);
if (Attr *A = D->getAttr<ReqdWorkGroupSizeAttr>()) {
       Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
       D->setInvalidDecl();
     } else if (Attr *A = D->getAttr<WorkGroupSizeHintAttr>()) {
       Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
       D->setInvalidDecl();
     } else if (Attr *A = D->getAttr<VecTypeHintAttr>()) {
       Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
       D->setInvalidDecl();
-    } else if (Attr *A = D->getAttr<AMDGPUNumVGPRAttr>()) {
+    } else if (Attr *A = D->getAttr<AMDGPUFlatWorkGroupSizeAttr>()) {
+      Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
+        << A << ExpectedKernelFunction;
+      D->setInvalidDecl();
+    } else if (Attr *A = D->getAttr<AMDGPUWavesPerEUAttr>()) {
       Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
         << A << ExpectedKernelFunction;
       D->setInvalidDecl();
     } else if (Attr *A = D->getAttr<AMDGPUNumSGPRAttr>()) {
       Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
         << A << ExpectedKernelFunction;
       D->setInvalidDecl();
+    } else if (Attr *A = D->getAttr<AMDGPUNumVGPRAttr>()) {
+      Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
+        << A << ExpectedKernelFunction;
+      D->setInvalidDecl();
     }
   }
 }
 
 // Annotation attributes are the only attributes allowed after an access
 // specifier.
 bool Sema::ProcessAccessDeclAttributeList(AccessSpecDecl *ASDecl,
                                           const AttributeList *AttrList) {

test/CodeGenOpenCL/amdgpu-attrs.cl

+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s
+
+__attribute__((amdgpu_flat_work_group_size(0, 0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0() {}
+__attribute__((amdgpu_waves_per_eu(0))) // expected-no-diagnostics
+kernel void waves_per_eu_0() {}
+__attribute__((amdgpu_waves_per_eu(0, 0))) // expected-no-diagnostics
+kernel void waves_per_eu_0_0() {}
+__attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void num_sgpr0() {}
+__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void num_vgpr0() {}
+
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_num_sgpr_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_num_vgpr_0() {}
+__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void waves_per_eu_0_num_sgpr_0() {}
+__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void waves_per_eu_0_num_vgpr_0() {}
+__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void waves_per_eu_0_0_num_sgpr_0() {}
+__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void waves_per_eu_0_0_num_vgpr_0() {}
+__attribute__((amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void num_sgpr_0_num_vgpr_0() {}
+
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_num_sgpr_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_num_vgpr_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_sgpr_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_vgpr_0() {}
+
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_num_sgpr_0_num_vgpr_0() {}
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_sgpr_0_num_vgpr_0() {}
+
+__attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics
+kernel void waves_per_eu_2() {
+// CHECK: define amdgpu_kernel void @waves_per_eu_2() [[WAVES_PER_EU_2:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2, 4))) // expected-no-diagnostics
+kernel void waves_per_eu_2_4() {
+// CHECK: define amdgpu_kernel void @waves_per_eu_2_4() [[WAVES_PER_EU_2_4:#[0-9]+]]
+}
+__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
+kernel void num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @num_sgpr_32() [[NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @num_vgpr_64() [[NUM_VGPR_64:#[0-9]+]]
+}
+
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64_waves_per_eu_2() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64_waves_per_eu_2_4() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_sgpr(32))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64_num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32))) // expected-no-diagnostics
+kernel void waves_per_eu_2_num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @waves_per_eu_2_num_sgpr_32() [[WAVES_PER_EU_2_NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void waves_per_eu_2_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @waves_per_eu_2_num_vgpr_64() [[WAVES_PER_EU_2_NUM_VGPR_64:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32))) // expected-no-diagnostics
+kernel void waves_per_eu_2_4_num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @waves_per_eu_2_4_num_sgpr_32() [[WAVES_PER_EU_2_4_NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void waves_per_eu_2_4_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @waves_per_eu_2_4_num_vgpr_64() [[WAVES_PER_EU_2_4_NUM_VGPR_64:#[0-9]+]]
+}
+__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void num_sgpr_32_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @num_sgpr_32_num_vgpr_64() [[NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]]
+}
+
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32)))
+kernel void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64)))
+kernel void flat_work_group_size_32_64_waves_per_eu_2_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_VGPR_64:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32)))
+kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64)))
+kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64:#[0-9]+]]
+}
+
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]]
+}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() {
+// CHECK: define amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]]
+}
+
+// Make sure this is silently accepted on other targets.
+// X86-NOT: "amdgpu-flat-work-group-size"
+// X86-NOT: "amdgpu-waves-per-eu"
+// X86-NOT: "amdgpu-num-vgpr"
+// X86-NOT: "amdgpu-num-sgpr"
+
+// CHECK-NOT: "amdgpu-flat-work-group-size"="0,0"
+// CHECK-NOT: "amdgpu-waves-per-eu"="0"
+// CHECK-NOT: "amdgpu-waves-per-eu"="0,0"
+// CHECK-NOT: "amdgpu-num-sgpr"="0"
+// CHECK-NOT: "amdgpu-num-vgpr"="0"
+
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = { nounwind "amdgpu-flat-work-group-size"="32,64"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = { nounwind "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2_4]] = { nounwind "amdgpu-waves-per-eu"="2,4"
+// CHECK-DAG: attributes [[NUM_SGPR_32]] = { nounwind "amdgpu-num-sgpr"="32"
+// CHECK-DAG: attributes [[NUM_VGPR_64]] = { nounwind "amdgpu-num-vgpr"="64"
+
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-waves-per-eu"="2,4"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-vgpr"="64"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_SGPR_32]] = { nounwind "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2_NUM_VGPR_64]] = { nounwind "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_SGPR_32]] = { nounwind "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4"
+// CHECK-DAG: attributes [[WAVES_PER_EU_2_4_NUM_VGPR_64]] = { nounwind "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"
+// CHECK-DAG: attributes [[NUM_SGPR_32_NUM_VGPR_64]] = { nounwind "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64"
+
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_VGPR_64]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-waves-per-eu"="2,4"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"
+
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2"
+// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64]] = { nounwind "amdgpu-flat-work-group-size"="32,64" "amdgpu-num-sgpr"="32" "amdgpu-num-vgpr"="64" "amdgpu-waves-per-eu"="2,4"

test/CodeGenOpenCL/amdgpu-num-gpr-attr.cl

-// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s
-
-// Make sure this is silently accepted on other targets.
-
-__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
-kernel void test_num_vgpr64() {
-// CHECK: define amdgpu_kernel void @test_num_vgpr64() [[ATTR_VGPR64:#[0-9]+]]
-}
-
-__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
-kernel void test_num_sgpr32() {
-// CHECK: define amdgpu_kernel void @test_num_sgpr32() [[ATTR_SGPR32:#[0-9]+]]
-}
-
-__attribute__((amdgpu_num_vgpr(64), amdgpu_num_sgpr(32))) // expected-no-diagnostics
-kernel void test_num_vgpr64_sgpr32() {
-// CHECK: define amdgpu_kernel void @test_num_vgpr64_sgpr32() [[ATTR_VGPR64_SGPR32:#[0-9]+]]
-
-}
-
-__attribute__((amdgpu_num_sgpr(20), amdgpu_num_vgpr(40))) // expected-no-diagnostics
-kernel void test_num_sgpr20_vgpr40() {
-// CHECK: define amdgpu_kernel void @test_num_sgpr20_vgpr40() [[ATTR_SGPR20_VGPR40:#[0-9]+]]
-}
-
-__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
-kernel void test_num_vgpr0() {
-}
-
-__attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics
-kernel void test_num_sgpr0() {
-}
-
-__attribute__((amdgpu_num_vgpr(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
-kernel void test_num_vgpr0_sgpr0() {
-}
-
-
-// X86-NOT: "amdgpu_num_vgpr"
-// X86-NOT: "amdgpu_num_sgpr"
-
-// CHECK-NOT: "amdgpu_num_vgpr"="0"
-// CHECK-NOT: "amdgpu_num_sgpr"="0"
-// CHECK-DAG: attributes [[ATTR_VGPR64]] = { nounwind "amdgpu_num_vgpr"="64"
-// CHECK-DAG: attributes [[ATTR_SGPR32]] = { nounwind "amdgpu_num_sgpr"="32"
-// CHECK-DAG: attributes [[ATTR_VGPR64_SGPR32]] = { nounwind "amdgpu_num_sgpr"="32" "amdgpu_num_vgpr"="64"
-// CHECK-DAG: attributes [[ATTR_SGPR20_VGPR40]] = { nounwind "amdgpu_num_sgpr"="20" "amdgpu_num_vgpr"="40"

test/SemaCUDA/amdgpu-attrs.cu

+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+
+// expected-error@+2 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64)))
+__global__ void flat_work_group_size_32_64() {}
+
+// expected-error@+2 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2)))
+__global__ void waves_per_eu_2() {}
+
+// expected-error@+2 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2, 4)))
+__global__ void waves_per_eu_2_4() {}
+
+// expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_num_sgpr(32)))
+__global__ void num_sgpr_32() {}
+
+// expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_num_vgpr(64)))
+__global__ void num_vgpr_64() {}
+
+
+// expected-error@+3 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2() {}
+
+// expected-error@+3 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_4() {}
+
+// expected-error@+3 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_sgpr(32)))
+__global__ void flat_work_group_size_32_64_num_sgpr_32() {}
+
+// expected-error@+3 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_vgpr(64)))
+__global__ void flat_work_group_size_32_64_num_vgpr_64() {}
+
+// expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32)))
+__global__ void waves_per_eu_2_num_sgpr_32() {}
+
+// expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64)))
+__global__ void waves_per_eu_2_num_vgpr_64() {}
+
+// expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32)))
+__global__ void waves_per_eu_2_4_num_sgpr_32() {}
+
+// expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64)))
+__global__ void waves_per_eu_2_4_num_vgpr_64() {}
+
+// expected-error@+3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
+__global__ void num_sgpr_32_num_vgpr_64() {}
+
+
+// expected-error@+4 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32() {}
+
+// expected-error@+4 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_num_vgpr_64() {}
+
+// expected-error@+4 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32() {}
+
+// expected-error@+4 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+3 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() {}
+
+
+// expected-error@+5 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+4 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() {}
+
+// expected-error@+5 {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+// fixme-expected-error@+4 {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+// fixme-expected-error@+3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+// fixme-expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
+__global__ void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() {}

test/SemaCUDA/amdgpu-num-gpr-attr.cu

-// RUN: %clang_cc1 -fsyntax-only -verify %s
-
-#include "Inputs/cuda.h"
-
-__attribute__((amdgpu_num_vgpr(64)))
-__global__ void test_num_vgpr() { } // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
-
-__attribute__((amdgpu_num_sgpr(32)))
-__global__ void test_num_sgpr() { } // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
-
-// fixme-expected-error@+3 {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
-// expected-error@+2 {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
-__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
-__global__ void test_num_vgpr_num_sgpr() { }

test/SemaOpenCL/amdgpu-attrs.cl

+// RUN: %clang_cc1 -triple amdgcn-- -verify -fsyntax-only %s
+
+typedef __attribute__((amdgpu_flat_work_group_size(32, 64))) struct struct_flat_work_group_size_32_64 { // expected-error {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+  int x;
+  float y;
+} struct_flat_work_group_size_32_64;
+typedef __attribute__((amdgpu_waves_per_eu(2))) struct struct_waves_per_eu_2 { // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+  int x;
+  float y;
+} struct_waves_per_eu_2;
+typedef __attribute__((amdgpu_waves_per_eu(2, 4))) struct struct_waves_per_eu_2_4 { // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+  int x;
+  float y;
+} struct_waves_per_eu_2_4;
+typedef __attribute__((amdgpu_num_sgpr(32))) struct struct_num_sgpr_32 { // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+  int x;
+  float y;
+} struct_num_sgpr_32;
+typedef __attribute__((amdgpu_num_vgpr(64))) struct struct_num_vgpr_64 { // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+  int x;
+  float y;
+} struct_num_vgpr_64;
+
+__attribute__((amdgpu_flat_work_group_size(32, 64))) void func_flat_work_group_size_32_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2))) void func_waves_per_eu_2() {} // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+__attribute__((amdgpu_waves_per_eu(2, 4))) void func_waves_per_eu_2_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
+__attribute__((amdgpu_num_sgpr(32))) void func_num_sgpr_32() {} // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_num_vgpr(64))) void func_num_vgpr_64() {} // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
+
+__attribute__((amdgpu_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, "ABC"))) kernel void kernel_flat_work_group_size_32_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("ABC", 64))) kernel void kernel_flat_work_group_size_ABC_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
+__attribute__((amdgpu_waves_per_eu("ABC"))) kernel void kernel_waves_per_eu_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, "ABC"))) kernel void kernel_waves_per_eu_2_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
+__attribute__((amdgpu_waves_per_eu("ABC", 4))) kernel void kernel_waves_per_eu_ABC_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
+__attribute__((amdgpu_num_sgpr("ABC"))) kernel void kernel_num_sgpr_ABC() {} // expected-error {{'amdgpu_num_sgpr' attribute requires an integer constant}}
+__attribute__((amdgpu_num_vgpr("ABC"))) kernel void kernel_num_vgpr_ABC() {} // expected-error {{'amdgpu_num_vgpr' attribute requires an integer constant}}
+
+__attribute__((amdgpu_flat_work_group_size(4294967296, 4294967296))) kernel void kernel_flat_work_group_size_L_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_flat_work_group_size(32, 4294967296))) kernel void kernel_flat_work_group_size_32_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_flat_work_group_size(4294967296, 64))) kernel void kernel_flat_work_group_size_L_64() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_waves_per_eu(4294967296))) kernel void kernel_waves_per_eu_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_waves_per_eu(2, 4294967296))) kernel void kernel_waves_per_eu_2_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_waves_per_eu(4294967296, 4))) kernel void kernel_waves_per_eu_L_4() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_num_sgpr(4294967296))) kernel void kernel_num_sgpr_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_num_vgpr(4294967296))) kernel void kernel_num_vgpr_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+
+__attribute__((amdgpu_flat_work_group_size(0, 64))) kernel void kernel_flat_work_group_size_0_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute parameter is invalid: maximum flat work group size must be zero since minimum flat work group size is zero}}
+__attribute__((amdgpu_waves_per_eu(0, 4))) kernel void kernel_waves_per_eu_0_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute parameter is invalid: maximum number of waves per execution unit must be zero since minimum number of waves per execution unit is zero}}
+
+__attribute__((amdgpu_flat_work_group_size(64, 32))) kernel void kernel_flat_work_group_size_64_32() {} // expected-error {{'amdgpu_flat_work_group_size' attribute parameter is invalid: minimum flat work group size must not be greater than maximum flat work group size}}
+__attribute__((amdgpu_waves_per_eu(4, 2))) kernel void kernel_waves_per_eu_4_2() {} // expected-error {{'amdgpu_waves_per_eu' attribute parameter is invalid: minimum number of waves per execution unit must not be greater than maximum number of waves per execution unit}}
+
+__attribute__((amdgpu_flat_work_group_size(0, 0))) kernel void kernel_flat_work_group_size_0_0() {}
+__attribute__((amdgpu_waves_per_eu(0))) kernel void kernel_waves_per_eu_0() {}
+__attribute__((amdgpu_waves_per_eu(0, 0))) kernel void kernel_waves_per_eu_0_0() {}
+__attribute__((amdgpu_num_sgpr(0))) kernel void kernel_num_sgpr_0() {}
+__attribute__((amdgpu_num_vgpr(0))) kernel void kernel_num_vgpr_0() {}
+
+kernel __attribute__((amdgpu_flat_work_group_size(32, 64))) void kernel_flat_work_group_size_32_64() {}
+kernel __attribute__((amdgpu_waves_per_eu(2))) void kernel_waves_per_eu_2() {}
+kernel __attribute__((amdgpu_waves_per_eu(2, 4))) void kernel_waves_per_eu_2_4() {}
+kernel __attribute__((amdgpu_num_sgpr(32))) void kernel_num_sgpr_32() {}
+kernel __attribute__((amdgpu_num_vgpr(64))) void kernel_num_vgpr_64() {}

test/SemaOpenCL/amdgpu-num-register-attrs.cl

-// RUN: %clang_cc1 -triple r600-- -verify -fsyntax-only %s
-
-typedef __attribute__((amdgpu_num_vgpr(128))) struct FooStruct { // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
-  int x;
-  float y;
-} FooStruct;
-
-
-__attribute__((amdgpu_num_vgpr("ABC"))) kernel void foo2() {} // expected-error {{'amdgpu_num_vgpr' attribute requires an integer constant}}
-__attribute__((amdgpu_num_sgpr("ABC"))) kernel void foo3() {} // expected-error {{'amdgpu_num_sgpr' attribute requires an integer constant}}
-
-
-__attribute__((amdgpu_num_vgpr(40))) void foo4() {} // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
-__attribute__((amdgpu_num_sgpr(64))) void foo5() {} // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
-
-__attribute__((amdgpu_num_vgpr(40))) kernel void foo7() {}
-__attribute__((amdgpu_num_sgpr(64))) kernel void foo8() {}
-__attribute__((amdgpu_num_vgpr(40), amdgpu_num_sgpr(64))) kernel void foo9() {}
-
-// Check 0 VGPR is accepted.
-__attribute__((amdgpu_num_vgpr(0))) kernel void foo10() {}
-
-// Check 0 SGPR is accepted.
-__attribute__((amdgpu_num_sgpr(0))) kernel void foo11() {}
-
-// Check both 0 SGPR and VGPR is accepted.
-__attribute__((amdgpu_num_vgpr(0), amdgpu_num_sgpr(0))) kernel void foo12() {}
-
-// Too large VGPR value.
-__attribute__((amdgpu_num_vgpr(4294967296))) kernel void foo13() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
-
-__attribute__((amdgpu_num_sgpr(4294967296))) kernel void foo14() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
-
-__attribute__((amdgpu_num_sgpr(4294967296), amdgpu_num_vgpr(4294967296))) kernel void foo15() {} // expected-error 2 {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
-
-
-// Make sure it is accepted with kernel keyword before the attribute.
-kernel __attribute__((amdgpu_num_vgpr(40))) void foo16() {}
-
-kernel __attribute__((amdgpu_num_sgpr(40))) void foo17() {}