Add attributes for AMD GPU Tools

[• kzhuravl-AMD]

Summary

Add following kernel function attributes for AMD GPU Tools (debugger, profiler):

• amdgpu_tools_insert_nops - insert two nop instructions for each high level source statement
• amdgpu_tools_num_reserved_vgpr(<num>) - reserve <num> vector registers and do not use throughout kernel execution
• amdgpu_tools_num_reserved_sgpr(<num>) - reserve <num> scalar registers and do not use throughout kernel execution

Also add similar options that cause to insert attributes for each kernel, options take precedence

+ Updated docs, added tests

[arsenm]

include/clang/Basic/Attr.td

993–998

I didn't envision these as being user code facing attributes, and only the function emission would add them. Is there a use for these being used manually?

include/clang/Driver/Options.td

355–365

These are user facing options, not cc1 flags? I wouldn't expect these to be exposed to users, and they would just be implied by -g. Is there a need for this? Even if there is some need for these, I don't see anything testing for these with only -g

lib/CodeGen/CGCall.cpp

1601

I've been trying to switch to consistently using the '-' separates words conventions most function attributes use.

test/CodeGenOpenCL/amdgpu-tools-attrs.cl

2

I all of the run lines should use amdhsa triple

[aaron.ballman]

Some quick thoughts below.

include/clang/Basic/Attr.td

997

I don't see any code that checks whether the function is actually a kernel function, nor sema tests that diagnose if this is attached to a non-kernel function.

It may make sense to make a SubsetSubject that does the checking for you, and use that in all of these attributes.

include/clang/Basic/AttrDocs.td

943

This reads a bit strangely, can it be "debuggers" and "profilers" (plural)?

951

"supports the __attribute__" (insert the).

952

"it causes" (insert it).

954

"before the first ISA" (insert the, capitalize acronym).
"of the high level" (insert the).

955

"after the last ISA" (insert the, capitalize acronym).
"of the high level" (insert the).

959

"function in module" -> "function in the translation unit"? At least insert "the", but uncertain what you mean by "module", which I assume to be a translation unit.

Also add "the" before the command line option.

964

Same sentence comments apply to the following attributes as above.

test/SemaOpenCL/amdgpu-tools-attrs.cl

6

You can make the struct empty in these tests, but more importantly, it would be good to see a test that ensures the following diagnoses:

__attribute__((amdgpu_tools_insert_nops)) void not_a_kernel_func() {}

[• kzhuravl-AMD]

Review Feedback

include/clang/Basic/Attr.td

993–998

Tools team expressed a desire to have the ability to specify them manually as well.

include/clang/Driver/Options.td

355–365

These should be independent of -g, user facing options. One of the use cases is profiler, where debug information is needed (-g), but no registers need to be reserved, no nops inserted.

lib/CodeGen/CGCall.cpp

1601

Do you mean use '-' instead of '_' in attribute spellings?

[• kzhuravl-AMD]

Review Feedback - Updated diff

[arsenm]

lib/CodeGen/CGCall.cpp

1601

Yes

[• tstellarAMD]

include/clang/Basic/AttrDocs.td

969–973

What does the value of this attribute mean? If I pass 16, does it reserve VGPRS 0 - 15 or
VGPRS 240 - 255 ? Can you clarify this in the documentation?

[• kzhuravl-AMD]

Review Feedback - Updated diff

[aaron.ballman]

lib/Sema/SemaDeclAttr.cpp

5653

All of these should call D->setInvalidDecl(), similar to the above kernel attributes.

[• kzhuravl-AMD]

Aaron's Review Comments

[aaron.ballman]

Attribute side of things LGTM, but I don't have the expertise to comment on the OpenCL functionality itself.

[• kzhuravl-AMD]

After recent discussions we decided to use target specific options instead

include/clang/Basic/Attr.td

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">;
 }
 
+def AMDGPUToolsInsertNops : InheritableAttr {
+  let Spellings = [GNU<"amdgpu_tools_insert_nops">];
+  let Documentation = [AMDGPUToolsInsertNopsDocs];
+  let Subjects = SubjectList<[Function], ErrorDiag,
+                             "ExpectedKernelFunction">;
+}
+
+def AMDGPUToolsNumReservedVGPR : InheritableAttr {
+  let Spellings = [GNU<"amdgpu_tools_num_reserved_vgpr">];
+  let Args = [UnsignedArgument<"NumReservedVGPR">];
+  let Documentation = [AMDGPUToolsNumReservedVGPRDocs];
+  let Subjects = SubjectList<[Function], ErrorDiag,
+                             "ExpectedKernelFunction">;
+}
+
+def AMDGPUToolsNumReservedSGPR : InheritableAttr {
+  let Spellings = [GNU<"amdgpu_tools_num_reserved_sgpr">];
+  let Args = [UnsignedArgument<"NumReservedSGPR">];
+  let Documentation = [AMDGPUToolsNumReservedSGPRDocs];
+  let Subjects = SubjectList<[Function], ErrorDiag,
+                             "ExpectedKernelFunction">;
+}
+
 def NoSplitStack : InheritableAttr {
   let Spellings = [GCC<"no_split_stack">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [NoSplitStackDocs];
 }
 
 def NonNull : InheritableAttr {
   let Spellings = [GCC<"nonnull">];

include/clang/Basic/AttrDocs.td

 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 DocCatAMDGPUToolsAttributes :
+  DocumentationCategory<"AMD GPU Tools Attributes"> {
+  let Content = [{
+
+Clang supports following AMD GPU attributes for tools, such as debuggers and
+profilers:
+  }];
+}
+
+def AMDGPUToolsInsertNopsDocs : Documentation {
+  let Category = DocCatAMDGPUToolsAttributes;
+  let Content = [{
+
+Clang supports the ``__attribute__((amdgpu_tools_insert_nops))`` attribute on
+AMD Southern Islands GPUs and later. If specified, it causes AMD GPU Backend to
+insert two nop instructions for each high level source statement: one nop
+instruction is inserted before the first ISA instruction of the high level
+source statement, and one nop instruction is inserted after the last ISA
+instruction of the high level source statement.
+
+In addition to specifying this attribute manually, clang can add this attribute
+for each kernel function in the translation unit if the
+``--amdgpu-tools-insert-nops`` clang command line option is specified.
+  }];
+}
+
+def AMDGPUToolsNumReservedVGPRDocs : Documentation {
+  let Category = DocCatAMDGPUToolsAttributes;
+  let Content = [{
+
+Clang supports the ``__attribute__((amdgpu_tools_num_reserved_vgpr(<num>)))``
+attribute on AMD Southern Islands GPUs and later. If specified, it causes AMD
+GPU Backend to reserve ``<num>`` number of vector registers and do not use those
+registers throughout kernel function execution. Index of the first reserved
+vector register is recorded in the ``amd_kernel_code_t``.
+
+In addition to specifying this attribute manually, clang can add this attribute
+for each kernel function in the translation unit if the
+``--amdgpu-tools-num-reserved-vgpr=<num>`` clang command line option is
+specified. In this case, the option takes precedence over the attribute.
+  }];
+}
+
+def AMDGPUToolsNumReservedSGPRDocs : Documentation {
+  let Category = DocCatAMDGPUToolsAttributes;
+  let Content = [{
+
+Clang supports the ``__attribute__((amdgpu_tools_num_reserved_sgpr(<num>)))``
+attribute on AMD Southern Islands GPUs and later. If specified, it causes AMD
+GPU Backend to reserve ``<num>`` number of scalar registers and do not use those
+registers throughout kernel function execution. Index of the first reserved
+scalar register is recorded in the ``amd_kernel_code_t``.
+
+In addition to specifying this attribute manually, clang can add this attribute
+for each kernel function in the translation unit if the
+``--amdgpu-tools-num-reserved-sgpr=<num>`` clang command line option is
+specified. In this case, the option takes precedence over the attribute.
+  }];
+}
+
 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.
   }];
 }

include/clang/Driver/Options.td

 def Xpreprocessor : Separate<["-"], "Xpreprocessor">,
   HelpText<"Pass <arg> to the preprocessor">, MetaVarName<"<arg>">;
 def X_Flag : Flag<["-"], "X">;
 def X_Joined : Joined<["-"], "X">;
 def Z_Flag : Flag<["-"], "Z">;
 def Z_Joined : Joined<["-"], "Z">;
 def all__load : Flag<["-"], "all_load">;
 def allowable__client : Separate<["-"], "allowable_client">;
+def amdgpu_tools_insert_nops :
+  Flag<["--"], "amdgpu-tools-insert-nops">, Flags<[CC1Option, HelpHidden]>,
+  HelpText<"Insert two nop instructions for each high level source statement">;
+def amdgpu_tools_num_reserved_vgpr :
+  Joined<["--"], "amdgpu-tools-num-reserved-vgpr=">,
+  Flags<[CC1Option, HelpHidden]>,
+  HelpText<"Reserve <num> vector registers">, MetaVarName<"<num>">;
+def amdgpu_tools_num_reserved_sgpr :
+  Joined<["--"], "amdgpu-tools-num-reserved-sgpr=">,
+  Flags<[CC1Option, HelpHidden]>,
+  HelpText<"Reserve <num> scalar registers">, MetaVarName<"<num>">;
 def ansi : Flag<["-", "--"], "ansi">;
 def arch__errors__fatal : Flag<["-"], "arch_errors_fatal">;
 def arch : Separate<["-"], "arch">, Flags<[DriverOption]>;
 def arch__only : Separate<["-"], "arch_only">;
 def a : Joined<["-"], "a">;
 def bind__at__load : Flag<["-"], "bind_at_load">;
 def bundle__loader : Separate<["-"], "bundle_loader">;
 def bundle : Flag<["-"], "bundle">;

include/clang/Frontend/CodeGenOptions.def

 CODEGENOPT(StackRealignment  , 1, 0) ///< Control whether to force stack
                                      ///< realignment.
 CODEGENOPT(UseInitArray      , 1, 0) ///< Control whether to use .init_array or
                                      ///< .ctors.
 VALUE_CODEGENOPT(StackAlignment    , 32, 0) ///< Overrides default stack 
                                             ///< alignment, if not 0.
 VALUE_CODEGENOPT(StackProbeSize    , 32, 4096) ///< Overrides default stack
                                                ///< probe size, even if 0.
+
+/// \brief Control setting of ``amdgpu_tools_insert_nops`` attribute
+CODEGENOPT(AMDGPUToolsInsertNopsOpt, 1, 0)
+/// \brief Control setting of ``amdgpu_tools_num_reserved_vgpr(<num>)``
+/// attribute
+VALUE_CODEGENOPT(AMDGPUToolsNumReservedVGPROpt, 32, 0)
+/// \brief Control setting of ``amdgpu_tools_num_reserved_sgpr(<num>)``
+/// attribute
+VALUE_CODEGENOPT(AMDGPUToolsNumReservedSGPROpt, 32, 0)
+
 CODEGENOPT(DebugColumnInfo, 1, 0) ///< Whether or not to use column information
                                   ///< in debug info.
 
 CODEGENOPT(DebugTypeExtRefs, 1, 0) ///< Whether or not debug info should contain
                                    ///< external references to a PCH or module.
 
 CODEGENOPT(DebugExplicitImport, 1, 0)  ///< Whether or not debug info should 
                                        ///< contain explicit imports for 

lib/CodeGen/CGCall.cpp

if (FD && FD->hasAttr<TargetAttr>()) {
         std::sort(Features.begin(), Features.end());
         FuncAttrs.addAttribute(
             "target-features",
             llvm::join(Features.begin(), Features.end(), ","));
       }
     }
   }
 
+  if (getTarget().getTriple().getArch() == llvm::Triple::amdgcn) {
+    // Add AMDGPU Tools attributes if needed.
+    if (CodeGenOpts.AMDGPUToolsInsertNopsOpt)
+      FuncAttrs.addAttribute("amdgpu_tools_insert_nops");
+    if (CodeGenOpts.AMDGPUToolsNumReservedVGPROpt)
+      FuncAttrs.addAttribute(
+        "amdgpu_tools_num_reserved_vgpr",
+        llvm::utostr(CodeGenOpts.AMDGPUToolsNumReservedVGPROpt));
+    if (CodeGenOpts.AMDGPUToolsNumReservedSGPROpt)
+      FuncAttrs.addAttribute(
+        "amdgpu_tools_num_reserved_sgpr",
+        llvm::utostr(CodeGenOpts.AMDGPUToolsNumReservedSGPROpt));
+  }
+
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     // Conservatively, mark all functions and calls in CUDA as convergent
     // (meaning, they may call an intrinsically convergent op, such as
     // __syncthreads(), and so can't have certain optimizations applied around
     // them).  LLVM will remove this attribute where it safely can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
   }
 

lib/CodeGen/TargetInfo.cpp

void AMDGPUTargetCodeGenInfo::setTargetAttributes(
   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));
+      F->addFnAttr(Attr->getSpelling(), llvm::utostr(NumVGPR));
   }
 
   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(Attr->getSpelling(), llvm::utostr(NumSGPR));
+  }
+
+  if (const auto Attr = FD->getAttr<AMDGPUToolsInsertNopsAttr>()) {
+    llvm::Function *F = cast<llvm::Function>(GV);
+    if (!F->hasFnAttribute(Attr->getSpelling()))
+      F->addFnAttr(Attr->getSpelling());
+  }
+
+  if (const auto Attr = FD->getAttr<AMDGPUToolsNumReservedVGPRAttr>()) {
+    llvm::Function *F = cast<llvm::Function>(GV);
+    unsigned NumReservedVGPR = Attr->getNumReservedVGPR();
+    if (!F->hasFnAttribute(Attr->getSpelling()) && NumReservedVGPR != 0)
+      F->addFnAttr(Attr->getSpelling(), llvm::utostr(NumReservedVGPR));
+  }
+
+  if (const auto Attr = FD->getAttr<AMDGPUToolsNumReservedSGPRAttr>()) {
+    llvm::Function *F = cast<llvm::Function>(GV);
+    unsigned NumReservedSGPR = Attr->getNumReservedSGPR();
+    if (!F->hasFnAttribute(Attr->getSpelling()) && NumReservedSGPR != 0)
+      F->addFnAttr(Attr->getSpelling(), llvm::utostr(NumReservedSGPR));
   }
 }
 
 
 //===----------------------------------------------------------------------===//
 // SPARC v9 ABI Implementation.
 // Based on the SPARC Compliance Definition version 2.4.1.
 //

lib/Driver/Tools.cpp

if (Args.hasArg(options::OPT_mstack_probe_size)) {
     StringRef Size = Args.getLastArgValue(options::OPT_mstack_probe_size);
 
     if (!Size.empty())
       CmdArgs.push_back(Args.MakeArgString("-mstack-probe-size=" + Size));
     else
       CmdArgs.push_back("-mstack-probe-size=0");
   }
 
+  // Translate AMDGPU Tools arguments.
+  if (Args.hasArg(options::OPT_amdgpu_tools_insert_nops))
+    CmdArgs.push_back("--amdgpu-tools-insert-nops");
+  if (Args.hasArg(options::OPT_amdgpu_tools_num_reserved_vgpr)) {
+    StringRef Size =
+      Args.getLastArgValue(options::OPT_amdgpu_tools_num_reserved_vgpr);
+
+    if (!Size.empty())
+      CmdArgs.push_back(
+        Args.MakeArgString("--amdgpu-tools-num-reserved-vgpr=" + Size));
+  }
+  if (Args.hasArg(options::OPT_amdgpu_tools_num_reserved_sgpr)) {
+    StringRef Size =
+      Args.getLastArgValue(options::OPT_amdgpu_tools_num_reserved_sgpr);
+
+    if (!Size.empty())
+      CmdArgs.push_back(
+        Args.MakeArgString("--amdgpu-tools-num-reserved-sgpr=" + Size));
+  }
+
   switch (getToolChain().getArch()) {
   case llvm::Triple::aarch64:
   case llvm::Triple::aarch64_be:
   case llvm::Triple::arm:
   case llvm::Triple::armeb:
   case llvm::Triple::thumb:
   case llvm::Triple::thumbeb:
     CmdArgs.push_back("-fallow-half-arguments-and-returns");

lib/Frontend/CompilerInvocation.cpp

static bool ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args, InputKind IK,
 
   if (Arg *A = Args.getLastArg(OPT_mstack_probe_size)) {
     StringRef Val = A->getValue();
     unsigned StackProbeSize = Opts.StackProbeSize;
     Val.getAsInteger(0, StackProbeSize);
     Opts.StackProbeSize = StackProbeSize;
   }
 
+  // Set up AMDGPU Tools arguments.
+  Opts.AMDGPUToolsInsertNopsOpt = Args.hasArg(OPT_amdgpu_tools_insert_nops);
+  if (Arg *A = Args.getLastArg(OPT_amdgpu_tools_num_reserved_vgpr)) {
+    StringRef Val = A->getValue();
+    unsigned AMDGPUToolsNumReservedVGPROpt = Opts.AMDGPUToolsNumReservedVGPROpt;
+    Val.getAsInteger(0, AMDGPUToolsNumReservedVGPROpt);
+    Opts.AMDGPUToolsNumReservedVGPROpt = AMDGPUToolsNumReservedVGPROpt;
+  }
+  if (Arg *A = Args.getLastArg(OPT_amdgpu_tools_num_reserved_sgpr)) {
+    StringRef Val = A->getValue();
+    unsigned AMDGPUToolsNumReservedSGPROpt = Opts.AMDGPUToolsNumReservedSGPROpt;
+    Val.getAsInteger(0, AMDGPUToolsNumReservedSGPROpt);
+    Opts.AMDGPUToolsNumReservedSGPROpt = AMDGPUToolsNumReservedSGPROpt;
+  }
+
   if (Arg *A = Args.getLastArg(OPT_fobjc_dispatch_method_EQ)) {
     StringRef Name = A->getValue();
     unsigned Method = llvm::StringSwitch<unsigned>(Name)
       .Case("legacy", CodeGenOptions::Legacy)
       .Case("non-legacy", CodeGenOptions::NonLegacy)
       .Case("mixed", CodeGenOptions::Mixed)
       .Default(~0U);
     if (Method == ~0U) {

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,
-                                    const AttributeList &Attr) {
-  uint32_t NumRegs;
-  Expr *NumRegsExpr = static_cast<Expr *>(Attr.getArgAsExpr(0));
-  if (!checkUInt32Argument(S, Attr, NumRegsExpr, NumRegs))
+template <typename AMDGPUAttrType>
+static void handleAMDGPUUInt32Attr(Sema &S, Decl *D,
+  uint32_t UInt32Arg;
+  Expr *UInt32ArgExpr = static_cast<Expr *>(Attr.getArgAsExpr(0));
+  if (!checkUInt32Argument(S, Attr, UInt32ArgExpr, UInt32Arg))
     return;
 
   D->addAttr(::new (S.Context)
-             AMDGPUNumVGPRAttr(Attr.getLoc(), S.Context,
-                               NumRegs,
-                               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))
-    return;
-
-  D->addAttr(::new (S.Context)
-             AMDGPUNumSGPRAttr(Attr.getLoc(), S.Context,
-                               NumRegs,
-                               Attr.getAttributeSpellingListIndex()));
+             AMDGPUAttrType(Attr.getLoc(), S.Context,
+                            UInt32Arg,
 }
 
 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.
   ValueDecl *VD = dyn_cast<ValueDecl>(D);
static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   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);
+    handleAMDGPUUInt32Attr<AMDGPUNumVGPRAttr>(S, D, Attr);
     break;
   case AttributeList::AT_AMDGPUNumSGPR:
-    handleAMDGPUNumSGPRAttr(S, D, Attr);
+    handleAMDGPUUInt32Attr<AMDGPUNumSGPRAttr>(S, D, Attr);
+    break;
+  case AttributeList::AT_AMDGPUToolsInsertNops:
+    handleSimpleAttribute<AMDGPUToolsInsertNopsAttr>(S, D, Attr);
+    break;
+  case AttributeList::AT_AMDGPUToolsNumReservedVGPR:
+    handleAMDGPUUInt32Attr<AMDGPUToolsNumReservedVGPRAttr>(S, D, Attr);
+    break;
+  case AttributeList::AT_AMDGPUToolsNumReservedSGPR:
+    handleAMDGPUUInt32Attr<AMDGPUToolsNumReservedSGPRAttr>(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:
if (!D->hasAttr<OpenCLKernelAttr>()) {
     } else if (Attr *A = D->getAttr<AMDGPUNumVGPRAttr>()) {
       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<AMDGPUToolsInsertNopsAttr>()) {
+      Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
+        << A << ExpectedKernelFunction;
+    } else if (Attr *A = D->getAttr<AMDGPUToolsNumReservedVGPRAttr>()) {
+      Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
+        << A << ExpectedKernelFunction;
+    } else if (Attr *A = D->getAttr<AMDGPUToolsNumReservedSGPRAttr>()) {
+      Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
+        << A << ExpectedKernelFunction;
     }
   }
 }
 
 // Annotation attributes are the only attributes allowed after an access
 // specifier.
 bool Sema::ProcessAccessDeclAttributeList(AccessSpecDecl *ASDecl,
                                           const AttributeList *AttrList) {

test/CodeGenOpenCL/amdgpu-tools-attrs-opts-precedence.cl

+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=4 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=PREC1 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-sgpr=8 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=PREC2 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=4 --amdgpu-tools-num-reserved-sgpr=8 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=PREC3 %s
+
+__attribute__((amdgpu_tools_num_reserved_vgpr(8))) // expected-no-diagnostics
+kernel void test_prec1() {
+// PREC1: define void @test_prec1() [[ATTR_PREC1:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_num_reserved_sgpr(4))) // expected-no-diagnostics
+kernel void test_prec2() {
+// PREC2: define void @test_prec2() [[ATTR_PREC2:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) // expected-no-diagnostics
+kernel void test_prec3() {
+// PREC3: define void @test_prec3() [[ATTR_PREC3:#[0-9]+]]
+}
+
+// PREC1-DAG: attributes [[ATTR_PREC1]] = { nounwind "amdgpu_tools_num_reserved_vgpr"="4"
+// PREC2-DAG: attributes [[ATTR_PREC2]] = { nounwind "amdgpu_tools_num_reserved_sgpr"="8"
+// PREC3-DAG: attributes [[ATTR_PREC3]] = { nounwind "amdgpu_tools_num_reserved_sgpr"="8" "amdgpu_tools_num_reserved_vgpr"="4"

test/CodeGenOpenCL/amdgpu-tools-attrs.cl

+// RUN: %clang_cc1 -triple amdgcn--amdhsa -target-cpu kaveri -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_tools_insert_nops)) // expected-no-diagnostics
+kernel void test_tools_insert_nops() {
+// CHECK: define void @test_tools_insert_nops() [[ATTR_NOP:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_num_reserved_vgpr(8))) // expected-no-diagnostics
+kernel void test_tools_num_reserved_vgpr8() {
+// CHECK: define void @test_tools_num_reserved_vgpr8() [[ATTR_VGPR8:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_vgpr(8))) // expected-no-diagnostics
+kernel void test_tools_insert_nops_num_reserved_vgpr8() {
+// CHECK: define void @test_tools_insert_nops_num_reserved_vgpr8() [[ATTR_NOP_VGPR8:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_num_reserved_sgpr(4))) // expected-no-diagnostics
+kernel void test_tools_num_reserved_sgpr4() {
+// CHECK: define void @test_tools_num_reserved_sgpr4() [[ATTR_SGPR4:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_sgpr(4))) // expected-no-diagnostics
+kernel void test_tools_insert_nops_num_reserved_sgpr4() {
+// CHECK: define void @test_tools_insert_nops_num_reserved_sgpr4() [[ATTR_NOP_SGPR4:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) // expected-no-diagnostics
+kernel void test_tools_num_reserved_vgpr8_sgpr4() {
+// CHECK: define void @test_tools_num_reserved_vgpr8_sgpr4() [[ATTR_VGPR8_SGPR4:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) // expected-no-diagnostics
+kernel void test_tools_insert_nops_num_reserved_vgpr8_sgpr4() {
+// CHECK: define void @test_tools_insert_nops_num_reserved_vgpr8_sgpr4() [[ATTR_NOP_VGPR8_SGPR4:#[0-9]+]]
+}
+
+__attribute__((amdgpu_tools_num_reserved_vgpr(0))) // expected-no-diagnostics
+kernel void test_tools_num_reserved_vgpr0() {}
+
+__attribute__((amdgpu_tools_num_reserved_sgpr(0))) // expected-no-diagnostics
+kernel void test_tools_num_reserved_sgpr0() {}
+
+__attribute__((amdgpu_tools_num_reserved_vgpr(0), amdgpu_tools_num_reserved_sgpr(0))) // expected-no-diagnostics
+kernel void test_tools_num_reserved_vgpr0_sgpr0() {}
+
+// CHECK-DAG: attributes [[ATTR_NOP]] = { nounwind "amdgpu_tools_insert_nops"
+// CHECK-DAG: attributes [[ATTR_VGPR8]] = { nounwind "amdgpu_tools_num_reserved_vgpr"="8"
+// CHECK-DAG: attributes [[ATTR_NOP_VGPR8]] = { nounwind "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_vgpr"="8"
+// CHECK-DAG: attributes [[ATTR_SGPR4]] = { nounwind "amdgpu_tools_num_reserved_sgpr"="4"
+// CHECK-DAG: attributes [[ATTR_NOP_SGPR4]] = { nounwind "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="4"
+// CHECK-DAG: attributes [[ATTR_VGPR8_SGPR4]] = { nounwind "amdgpu_tools_num_reserved_sgpr"="4" "amdgpu_tools_num_reserved_vgpr"="8"
+// CHECK-DAG: attributes [[ATTR_NOP_VGPR8_SGPR4]] = { nounwind "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="4" "amdgpu_tools_num_reserved_vgpr"="8"
+// CHECK-NOT: "amdgpu_tools_num_reserved_vgpr"="0"
+// CHECK-NOT: "amdgpu_tools_num_reserved_sgpr"="0"
+// X86-NOT: "amdgpu_tools_insert_nops"
+// X86-NOT: "amdgpu_tools_num_reserved_vgpr"
+// X86-NOT: "amdgpu_tools_num_reserved_sgpr"

test/CodeGenOpenCL/amdgpu-tools-opts.cl

+// RUN: %clang_cc1 --amdgpu-tools-insert-nops -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=8 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=VGPR8 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-vgpr=8 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP_VGPR8 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-sgpr=4 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-sgpr=4 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=8 --amdgpu-tools-num-reserved-sgpr=4 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=VGPR8_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-vgpr=8 --amdgpu-tools-num-reserved-sgpr=4 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP_VGPR8_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=0 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=VGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-vgpr=0 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP_VGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-sgpr=0 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=SGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-sgpr=0 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP_SGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=0 --amdgpu-tools-num-reserved-sgpr=0 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=VGPR0_SGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-vgpr=0 --amdgpu-tools-num-reserved-sgpr=0 -triple amdgcn--amdhsa -target-cpu kaveri -O0 -emit-llvm -o - %s | FileCheck -check-prefix=NOP_VGPR0_SGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_NOP %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=8 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_VGPR8 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-vgpr=8 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_NOP_VGPR8 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-sgpr=4 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-sgpr=4 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_NOP_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=8 --amdgpu-tools-num-reserved-sgpr=4 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_VGPR8_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-insert-nops --amdgpu-tools-num-reserved-vgpr=8 --amdgpu-tools-num-reserved-sgpr=4 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_NOP_VGPR8_SGPR4 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-vgpr=0 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_VGPR0 %s
+// RUN: %clang_cc1 --amdgpu-tools-num-reserved-sgpr=0 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86_SGPR0 %s
+
+kernel void foo() { // expected-no-diagnostics
+// NOP: define void @foo() [[ATTR_NOP:#[0-9]+]]
+// VGPR8: define void @foo() [[ATTR_VGPR8:#[0-9]+]]
+// NOP_VGPR8: define void @foo() [[ATTR_NOP_VGPR8:#[0-9]+]]
+// SGPR4: define void @foo() [[ATTR_SGPR4:#[0-9]+]]
+// NOP_SGPR4: define void @foo() [[ATTR_NOP_SGPR4:#[0-9]+]]
+// VGPR8_SGPR4: define void @foo() [[ATTR_VGPR8_SGPR4:#[0-9]+]]
+// NOP_VGPR8_SGPR4: define void @foo() [[ATTR_NOP_VGPR8_SGPR4:#[0-9]+]]
+}
+
+kernel void another_foo() { // expected-no-diagnostics
+// NOP: define void @another_foo() [[ATTR_SECOND_NOP:#[0-9]+]]
+}
+
+// NOP-DAG: attributes [[ATTR_NOP]] = { nounwind "amdgpu_tools_insert_nops"
+// NOP-DAG: attributes [[ATTR_SECOND_NOP]] = { nounwind "amdgpu_tools_insert_nops"
+// VGPR8-DAG: attributes [[ATTR_VGPR8]] = { nounwind "amdgpu_tools_num_reserved_vgpr"="8"
+// NOP_VGPR8-DAG: attributes [[ATTR_NOP_VGPR8]] = { nounwind "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_vgpr"="8"
+// SGPR4-DAG: attributes [[ATTR_SGPR4]] = { nounwind "amdgpu_tools_num_reserved_sgpr"="4"
+// NOP_SGPR4-DAG: attributes [[ATTR_NOP_SGPR4]] = { nounwind "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="4"
+// VGPR8_SGPR4-DAG: attributes [[ATTR_VGPR8_SGPR4]] = { nounwind "amdgpu_tools_num_reserved_sgpr"="4" "amdgpu_tools_num_reserved_vgpr"="8"
+// NOP_VGPR8_SGPR4-DAG: attributes [[ATTR_NOP_VGPR8_SGPR4]] = { nounwind "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="4" "amdgpu_tools_num_reserved_vgpr"="8"
+// VGPR0-NOT: "amdgpu_tools_num_reserved_vgpr"="0"
+// NOP_VGPR0-NOT: "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_vgpr"="0"
+// SGPR0-NOT: "amdgpu_tools_num_reserved_sgpr"="0"
+// NOP_SGPR0-NOT: "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="0"
+// VGPR0_SGPR0-NOT: "amdgpu_tools_num_reserved_vgpr"="0" "amdgpu_tools_num_reserved_sgpr"="0"
+// NOP_VGPR0_SGPR0-NOT: "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="0" "amdgpu_tools_num_reserved_vgpr"="0"
+// X86_NOP-NOT: "amdgpu_tools_insert_nops"
+// X86_VGPR8-NOT: "amdgpu_tools_num_reserved_vgpr"
+// X86_NOP_VGPR8-NOT: "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_vgpr"="8"
+// X86_SGPR4-NOT: "amdgpu_tools_num_reserved_sgpr"
+// X86_NOP_SGPR4-NOT: "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="4"
+// X86_VGPR8_SGPR4-NOT: "amdgpu_tools_num_reserved_sgpr"="4" "amdgpu_tools_num_reserved_vgpr"="8"
+// X86_NOP_VGPR8_SGPR4-NOT: "amdgpu_tools_insert_nops" "amdgpu_tools_num_reserved_sgpr"="4" "amdgpu_tools_num_reserved_vgpr"="8"
+// X86_VGPR0-NOT: "amdgpu_tools_num_reserved_vgpr"
+// X86_SGPR0-NOT: "amdgpu_tools_num_reserved_sgpr"

test/SemaOpenCL/amdgpu-tools-attrs.cl

+// RUN: %clang_cc1 -triple amdgcn--amdhsa -verify -fsyntax-only %s
+
+// Check attributes with non kernel function
+__attribute__((amdgpu_tools_insert_nops)) void not_kernel0() {} // expected-error {{'amdgpu_tools_insert_nops' attribute only applies to kernel functions}}
+__attribute__((amdgpu_tools_num_reserved_vgpr(8))) void not_kernel1() {} // expected-error {{'amdgpu_tools_num_reserved_vgpr' attribute only applies to kernel functions}}
+__attribute__((amdgpu_tools_num_reserved_sgpr(4))) void not_kernel2() {} // expected-error {{'amdgpu_tools_num_reserved_sgpr' attribute only applies to kernel functions}}
+typedef __attribute__((amdgpu_tools_insert_nops)) struct foo0_s { // expected-error {{'amdgpu_tools_insert_nops' attribute only applies to kernel functions}}
+  int x;
+  int y;
+} foo0_t;
+typedef __attribute__((amdgpu_tools_num_reserved_vgpr(8))) struct foo1_s { // expected-error {{'amdgpu_tools_num_reserved_vgpr' attribute only applies to kernel functions}}
+  int x;
+  int y;
+} foo1_t;
+typedef __attribute__((amdgpu_tools_num_reserved_sgpr(4))) struct foo2_s { // expected-error {{'amdgpu_tools_num_reserved_sgpr' attribute only applies to kernel functions}}
+  int x;
+  int y;
+} foo2_t;
+
+// Check non integer attribute values.
+__attribute__((amdgpu_tools_num_reserved_vgpr("ABC"))) kernel void foo3() {} // expected-error {{'amdgpu_tools_num_reserved_vgpr' attribute requires an integer constant}}
+__attribute__((amdgpu_tools_num_reserved_sgpr("DEF"))) kernel void foo4() {} // expected-error {{'amdgpu_tools_num_reserved_sgpr' attribute requires an integer constant}}
+
+// Check large attribute values.
+__attribute__((amdgpu_tools_num_reserved_vgpr(4294967296))) kernel void foo5() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_tools_num_reserved_sgpr(4294967296))) kernel void foo6() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__attribute__((amdgpu_tools_num_reserved_vgpr(4294967296), amdgpu_tools_num_reserved_sgpr(4294967296))) kernel void foo7() {} // expected-error 2 {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+
+// Check valid attributes.
+__attribute__((amdgpu_tools_insert_nops)) kernel void foo8() {}
+__attribute__((amdgpu_tools_num_reserved_vgpr(8))) kernel void foo9() {}
+__attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_vgpr(8))) kernel void foo10() {}
+__attribute__((amdgpu_tools_num_reserved_sgpr(4))) kernel void foo11() {}
+__attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_sgpr(4))) kernel void foo12() {}
+__attribute__((amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) kernel void foo13() {}
+__attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) kernel void foo14() {}
+// Make sure 0 VGPRs is accepted.
+__attribute__((amdgpu_tools_num_reserved_vgpr(0))) kernel void foo15() {}
+// Make sure 0 SGPRs is accepted.
+__attribute__((amdgpu_tools_num_reserved_sgpr(0))) kernel void foo16() {}
+// Make sure 0 VGPRs and 0 SGPRs is accepted.
+__attribute__((amdgpu_tools_num_reserved_vgpr(0), amdgpu_tools_num_reserved_sgpr(0))) kernel void foo17() {}
+// Make sure kernel keyword can go before __attribute__ keyword.
+kernel __attribute__((amdgpu_tools_insert_nops)) void foo18() {}
+kernel __attribute__((amdgpu_tools_num_reserved_vgpr(8))) void foo19() {}
+kernel __attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_vgpr(8))) void foo20() {}
+kernel __attribute__((amdgpu_tools_num_reserved_sgpr(4))) void foo21() {}
+kernel __attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_sgpr(4))) void foo22() {}
+kernel __attribute__((amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) void foo23() {}
+kernel __attribute__((amdgpu_tools_insert_nops, amdgpu_tools_num_reserved_vgpr(8), amdgpu_tools_num_reserved_sgpr(4))) void foo24() {}