Index: include/clang/Basic/Attr.td =================================================================== --- include/clang/Basic/Attr.td +++ include/clang/Basic/Attr.td @@ -1484,14 +1484,14 @@ def AMDGPUFlatWorkGroupSize : InheritableAttr { let Spellings = [Clang<"amdgpu_flat_work_group_size", 0>]; - let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max">]; + let Args = [ExprArgument<"Min">, ExprArgument<"Max">]; let Documentation = [AMDGPUFlatWorkGroupSizeDocs]; let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; } def AMDGPUWavesPerEU : InheritableAttr { let Spellings = [Clang<"amdgpu_waves_per_eu", 0>]; - let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max", 1>]; + let Args = [ExprArgument<"Min">, ExprArgument<"Max", 1>]; let Documentation = [AMDGPUWavesPerEUDocs]; let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; } Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -8674,6 +8674,16 @@ void AddXConsumedAttr(Decl *D, SourceRange SR, unsigned SpellingIndex, RetainOwnershipKind K, bool IsTemplateInstantiation); + /// addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size + /// attribute to a particular declaration. + void addAMDGPUFlatWorkGroupSizeAttr(SourceRange AttrRange, Decl *D, Expr *Min, + Expr *Max, unsigned SpellingListIndex); + + /// addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a + /// particular declaration. + void addAMDGPUWavesPerEUAttr(SourceRange AttrRange, Decl *D, Expr *Min, + Expr *Max, unsigned SpellingListIndex); + bool checkNSReturnsRetainedReturnType(SourceLocation loc, QualType type); //===--------------------------------------------------------------------===// Index: lib/CodeGen/TargetInfo.cpp =================================================================== --- lib/CodeGen/TargetInfo.cpp +++ lib/CodeGen/TargetInfo.cpp @@ -7797,8 +7797,16 @@ const auto *FlatWGS = FD->getAttr(); if (ReqdWGS || FlatWGS) { - unsigned Min = FlatWGS ? FlatWGS->getMin() : 0; - unsigned Max = FlatWGS ? FlatWGS->getMax() : 0; + unsigned Min = 0; + unsigned Max = 0; + if (FlatWGS) { + Min = FlatWGS->getMin() + ->EvaluateKnownConstInt(M.getContext()) + .getExtValue(); + Max = FlatWGS->getMax() + ->EvaluateKnownConstInt(M.getContext()) + .getExtValue(); + } if (ReqdWGS && Min == 0 && Max == 0) Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim(); @@ -7812,8 +7820,12 @@ } if (const auto *Attr = FD->getAttr()) { - unsigned Min = Attr->getMin(); - unsigned Max = Attr->getMax(); + unsigned Min = + Attr->getMin()->EvaluateKnownConstInt(M.getContext()).getExtValue(); + unsigned Max = Attr->getMax() ? Attr->getMax() + ->EvaluateKnownConstInt(M.getContext()) + .getExtValue() + : 0; if (Min != 0) { assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max"); Index: lib/Sema/SemaDeclAttr.cpp =================================================================== --- lib/Sema/SemaDeclAttr.cpp +++ lib/Sema/SemaDeclAttr.cpp @@ -245,11 +245,11 @@ !Expr->isIntegerConstantExpr(I, S.Context)) { if (Idx != UINT_MAX) S.Diag(getAttrLoc(AI), diag::err_attribute_argument_n_type) - << AI << Idx << AANT_ArgumentIntegerConstant + << &AI << Idx << AANT_ArgumentIntegerConstant << Expr->getSourceRange(); else S.Diag(getAttrLoc(AI), diag::err_attribute_argument_type) - << AI << AANT_ArgumentIntegerConstant << Expr->getSourceRange(); + << &AI << AANT_ArgumentIntegerConstant << Expr->getSourceRange(); return false; } @@ -261,7 +261,7 @@ if (StrictlyUnsigned && I.isSigned() && I.isNegative()) { S.Diag(getAttrLoc(AI), diag::err_attribute_requires_positive_integer) - << AI << /*non-negative*/ 1; + << &AI << /*non-negative*/ 1; return false; } @@ -5853,57 +5853,115 @@ } } -static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D, - const ParsedAttr &AL) { +static bool +checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr, + const AMDGPUFlatWorkGroupSizeAttr &Attr) { + // Accept template arguments for now as they depend on something else. + // We'll get to check them when they eventually get instantiated. + if (MinExpr->isValueDependent() || MaxExpr->isValueDependent()) + return false; + uint32_t Min = 0; - Expr *MinExpr = AL.getArgAsExpr(0); - if (!checkUInt32Argument(S, AL, MinExpr, Min)) - return; + if (!checkUInt32Argument(S, Attr, MinExpr, Min, 0)) + return true; uint32_t Max = 0; - Expr *MaxExpr = AL.getArgAsExpr(1); - if (!checkUInt32Argument(S, AL, MaxExpr, Max)) - return; + if (!checkUInt32Argument(S, Attr, MaxExpr, Max, 1)) + return true; if (Min == 0 && Max != 0) { - S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 0; - return; + S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid) + << &Attr << 0; + return true; } if (Min > Max) { - S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1; - return; + S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid) + << &Attr << 1; + return true; } - D->addAttr(::new (S.Context) - AMDGPUFlatWorkGroupSizeAttr(AL.getLoc(), S.Context, Min, Max, - AL.getAttributeSpellingListIndex())); + return false; } -static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) { - uint32_t Min = 0; - Expr *MinExpr = AL.getArgAsExpr(0); - if (!checkUInt32Argument(S, AL, MinExpr, Min)) +void Sema::addAMDGPUFlatWorkGroupSizeAttr(SourceRange AttrRange, Decl *D, + Expr *MinExpr, Expr *MaxExpr, + unsigned SpellingListIndex) { + AMDGPUFlatWorkGroupSizeAttr TmpAttr(AttrRange, Context, MinExpr, MaxExpr, + SpellingListIndex); + + if (checkAMDGPUFlatWorkGroupSizeArguments(*this, MinExpr, MaxExpr, TmpAttr)) return; + D->addAttr(::new (Context) AMDGPUFlatWorkGroupSizeAttr( + AttrRange, Context, MinExpr, MaxExpr, SpellingListIndex)); +} + +static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D, + const ParsedAttr &AL) { + Expr *MinExpr = AL.getArgAsExpr(0); + Expr *MaxExpr = AL.getArgAsExpr(1); + + S.addAMDGPUFlatWorkGroupSizeAttr(AL.getRange(), D, MinExpr, MaxExpr, + AL.getAttributeSpellingListIndex()); +} + +static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr, + Expr *MaxExpr, + const AMDGPUWavesPerEUAttr &Attr) { + if (S.DiagnoseUnexpandedParameterPack(MinExpr) || + (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr))) + return true; + + // Accept template arguments for now as they depend on something else. + // We'll get to check them when they eventually get instantiated. + if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent())) + return false; + + uint32_t Min = 0; + if (!checkUInt32Argument(S, Attr, MinExpr, Min, 0)) + return true; + uint32_t Max = 0; - if (AL.getNumArgs() == 2) { - Expr *MaxExpr = AL.getArgAsExpr(1); - if (!checkUInt32Argument(S, AL, MaxExpr, Max)) - return; - } + if (MaxExpr && !checkUInt32Argument(S, Attr, MaxExpr, Max, 1)) + return true; if (Min == 0 && Max != 0) { - S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 0; - return; + S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid) + << &Attr << 0; + return true; } if (Max != 0 && Min > Max) { - S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1; - return; + S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid) + << &Attr << 1; + return true; } - D->addAttr(::new (S.Context) - AMDGPUWavesPerEUAttr(AL.getLoc(), S.Context, Min, Max, - AL.getAttributeSpellingListIndex())); + return false; +} + +void Sema::addAMDGPUWavesPerEUAttr(SourceRange AttrRange, Decl *D, + Expr *MinExpr, Expr *MaxExpr, + unsigned SpellingListIndex) { + AMDGPUWavesPerEUAttr TmpAttr(AttrRange, Context, MinExpr, MaxExpr, + SpellingListIndex); + + if (checkAMDGPUWavesPerEUArguments(*this, MinExpr, MaxExpr, TmpAttr)) + return; + + D->addAttr(::new (Context) AMDGPUWavesPerEUAttr(AttrRange, Context, MinExpr, + MaxExpr, SpellingListIndex)); +} + +static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + if (!checkAttributeAtLeastNumArgs(S, AL, 1) || + !checkAttributeAtMostNumArgs(S, AL, 2)) + return; + + Expr *MinExpr = AL.getArgAsExpr(0); + Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr; + + S.addAMDGPUWavesPerEUAttr(AL.getRange(), D, MinExpr, MaxExpr, + AL.getAttributeSpellingListIndex()); } static void handleAMDGPUNumSGPRAttr(Sema &S, Decl *D, const ParsedAttr &AL) { Index: lib/Sema/SemaTemplateInstantiateDecl.cpp =================================================================== --- lib/Sema/SemaTemplateInstantiateDecl.cpp +++ lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -344,6 +344,51 @@ Attr.getRange()); } +static void instantiateDependentAMDGPUFlatWorkGroupSizeAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const AMDGPUFlatWorkGroupSizeAttr &Attr, Decl *New) { + // Both min and max expression are constant expressions. + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + + ExprResult Result = S.SubstExpr(Attr.getMin(), TemplateArgs); + if (Result.isInvalid()) + return; + Expr *MinExpr = Result.getAs(); + + Result = S.SubstExpr(Attr.getMax(), TemplateArgs); + if (Result.isInvalid()) + return; + Expr *MaxExpr = Result.getAs(); + + S.addAMDGPUFlatWorkGroupSizeAttr(Attr.getLocation(), New, MinExpr, MaxExpr, + Attr.getSpellingListIndex()); +} + +static void instantiateDependentAMDGPUWavesPerEUAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const AMDGPUWavesPerEUAttr &Attr, Decl *New) { + // Both min and max expression are constant expressions. + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + + ExprResult Result = S.SubstExpr(Attr.getMin(), TemplateArgs); + if (Result.isInvalid()) + return; + Expr *MinExpr = Result.getAs(); + + Expr *MaxExpr = nullptr; + if (auto Max = Attr.getMax()) { + Result = S.SubstExpr(Max, TemplateArgs); + if (Result.isInvalid()) + return; + MaxExpr = Result.getAs(); + } + + S.addAMDGPUWavesPerEUAttr(Attr.getLocation(), New, MinExpr, MaxExpr, + Attr.getSpellingListIndex()); +} + void Sema::InstantiateAttrsForDecl( const MultiLevelTemplateArgumentList &TemplateArgs, const Decl *Tmpl, Decl *New, LateInstantiatedAttrVec *LateAttrs, @@ -437,6 +482,18 @@ continue; } + if (const AMDGPUFlatWorkGroupSizeAttr *AMDGPUFlatWorkGroupSize = + dyn_cast(TmplAttr)) { + instantiateDependentAMDGPUFlatWorkGroupSizeAttr( + *this, TemplateArgs, *AMDGPUFlatWorkGroupSize, New); + } + + if (const AMDGPUWavesPerEUAttr *AMDGPUFlatWorkGroupSize = + dyn_cast(TmplAttr)) { + instantiateDependentAMDGPUWavesPerEUAttr(*this, TemplateArgs, + *AMDGPUFlatWorkGroupSize, New); + } + // Existing DLL attribute on the instantiation takes precedence. if (TmplAttr->getKind() == attr::DLLExport || TmplAttr->getKind() == attr::DLLImport) { Index: test/SemaCUDA/amdgpu-attrs.cu =================================================================== --- test/SemaCUDA/amdgpu-attrs.cu +++ test/SemaCUDA/amdgpu-attrs.cu @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsyntax-only -verify %s +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s #include "Inputs/cuda.h" @@ -78,3 +78,119 @@ // expected-error@+2{{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel function}} __attribute__((intel_reqd_sub_group_size(64))) __global__ void intel_reqd_sub_group_size_64() {} + +// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}} +__attribute__((amdgpu_flat_work_group_size("32", 64))) +__global__ void non_int_min_flat_work_group_size_32_64() {} +// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}} +__attribute__((amdgpu_flat_work_group_size(32, "64"))) +__global__ void non_int_max_flat_work_group_size_32_64() {} + +int nc_min = 32, nc_max = 64; +// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}} +__attribute__((amdgpu_flat_work_group_size(nc_min, 64))) +__global__ void non_cint_min_flat_work_group_size_32_64() {} +// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}} +__attribute__((amdgpu_flat_work_group_size(32, nc_max))) +__global__ void non_cint_max_flat_work_group_size_32_64() {} + +const int c_min = 16, c_max = 32; +__attribute__((amdgpu_flat_work_group_size(c_min * 2, 64))) +__global__ void cint_min_flat_work_group_size_32_64() {} +__attribute__((amdgpu_flat_work_group_size(32, c_max * 2))) +__global__ void cint_max_flat_work_group_size_32_64() {} + +// expected-error@+3{{'T' does not refer to a value}} +// expected-note@+1{{declared here}} +template +__attribute__((amdgpu_flat_work_group_size(T, 64))) +__global__ void template_class_min_flat_work_group_size_32_64() {} +// expected-error@+3{{'T' does not refer to a value}} +// expected-note@+1{{declared here}} +template +__attribute__((amdgpu_flat_work_group_size(32, T))) +__global__ void template_class_max_flat_work_group_size_32_64() {} + +template +__attribute__((amdgpu_flat_work_group_size(a, b))) +__global__ void template_flat_work_group_size_32_64() {} +template __global__ void template_flat_work_group_size_32_64<32, 64>(); + +template +__attribute__((amdgpu_flat_work_group_size(a + b, b + c))) +__global__ void template_complex_flat_work_group_size_32_64() {} +template __global__ void template_complex_flat_work_group_size_32_64<16, 16, 48>(); + +unsigned ipow2(unsigned n) { return n == 0 ? 1 : 2 * ipow2(n - 1); } +constexpr unsigned ce_ipow2(unsigned n) { return n == 0 ? 1 : 2 * ce_ipow2(n - 1); } + +__attribute__((amdgpu_flat_work_group_size(ce_ipow2(5), ce_ipow2(6)))) +__global__ void cexpr_flat_work_group_size_32_64() {} +// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}} +__attribute__((amdgpu_flat_work_group_size(ipow2(5), 64))) +__global__ void non_cexpr_min_flat_work_group_size_32_64() {} +// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}} +__attribute__((amdgpu_flat_work_group_size(32, ipow2(6)))) +__global__ void non_cexpr_max_flat_work_group_size_32_64() {} + +// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu("2"))) +__global__ void non_int_min_waves_per_eu_2() {} +// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu(2, "4"))) +__global__ void non_int_max_waves_per_eu_2_4() {} + +// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu(nc_min))) +__global__ void non_cint_min_waves_per_eu_2() {} +// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu(2, nc_max))) +__global__ void non_cint_min_waves_per_eu_2_4() {} + +__attribute__((amdgpu_waves_per_eu(c_min / 8))) +__global__ void cint_min_waves_per_eu_2() {} +__attribute__((amdgpu_waves_per_eu(c_min / 8, c_max / 8))) +__global__ void cint_min_waves_per_eu_2_4() {} + +// expected-error@+3{{'T' does not refer to a value}} +// expected-note@+1{{declared here}} +template +__attribute__((amdgpu_waves_per_eu(T))) +__global__ void cint_min_waves_per_eu_2() {} +// expected-error@+3{{'T' does not refer to a value}} +// expected-note@+1{{declared here}} +template +__attribute__((amdgpu_waves_per_eu(2, T))) +__global__ void cint_min_waves_per_eu_2_4() {} + +template +__attribute__((amdgpu_waves_per_eu(a))) +__global__ void template_waves_per_eu_2() {} +template __global__ void template_waves_per_eu_2<2>(); + +template +__attribute__((amdgpu_waves_per_eu(a, b))) +__global__ void template_waves_per_eu_2_4() {} +template __global__ void template_waves_per_eu_2_4<2, 4>(); + +template +__attribute__((amdgpu_waves_per_eu(a + b, c - b))) +__global__ void template_complex_waves_per_eu_2_4() {} +template __global__ void template_complex_waves_per_eu_2_4<1, 1, 5>(); + +// expected-error@+2{{expression contains unexpanded parameter pack 'Args'}} +template +__attribute__((amdgpu_waves_per_eu(Args))) +__global__ void template_waves_per_eu_2() {} +template __global__ void template_waves_per_eu_2<2, 4>(); + +__attribute__((amdgpu_waves_per_eu(ce_ipow2(1)))) +__global__ void cexpr_waves_per_eu_2() {} +__attribute__((amdgpu_waves_per_eu(ce_ipow2(1), ce_ipow2(2)))) +__global__ void cexpr_waves_per_eu_2_4() {} +// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu(ipow2(1)))) +__global__ void non_cexpr_waves_per_eu_2() {} +// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu(2, ipow2(2)))) +__global__ void non_cexpr_waves_per_eu_2_4() {} Index: test/SemaOpenCL/amdgpu-attrs.cl =================================================================== --- test/SemaOpenCL/amdgpu-attrs.cl +++ test/SemaOpenCL/amdgpu-attrs.cl @@ -27,12 +27,12 @@ __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_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be 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 parameter 1 to be 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 parameter 0 to be an integer constant}} +__attribute__((amdgpu_waves_per_eu("ABC"))) kernel void kernel_waves_per_eu_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 0 to be 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 parameter 1 to be 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 parameter 0 to be 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}}