Index: include/clang/Basic/Attr.td =================================================================== --- include/clang/Basic/Attr.td +++ include/clang/Basic/Attr.td @@ -1323,45 +1323,50 @@ def AMDGPUFlatWorkGroupSize : InheritableAttr { let Spellings = [GNU<"amdgpu_flat_work_group_size">, CXX11<"","hc_flat_workgroup_size", 201511>]; - let Args = [UnsignedArgument<"Min">, - UnsignedArgument<"Max", 1>, + let Args = [ExprArgument<"Min">, + ExprArgument<"Max", 1>, StringArgument<"ISA", 1>]; let Documentation = [AMDGPUFlatWorkGroupSizeDocs]; let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">; + let TemplateDependent = 1; } def AMDGPUWavesPerEU : InheritableAttr { let Spellings = [GNU<"amdgpu_waves_per_eu">, CXX11<"", "hc_waves_per_eu", 201511>]; - let Args = [UnsignedArgument<"Min">, - UnsignedArgument<"Max", 1>, + let Args = [ExprArgument<"Min">, + ExprArgument<"Max", 1>, StringArgument<"ISA", 1>]; let Documentation = [AMDGPUWavesPerEUDocs]; let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">; + let TemplateDependent = 1; } def AMDGPUNumSGPR : InheritableAttr { let Spellings = [GNU<"amdgpu_num_sgpr">]; - let Args = [UnsignedArgument<"NumSGPR">]; + let Args = [ExprArgument<"NumSGPR">]; let Documentation = [AMDGPUNumSGPRNumVGPRDocs]; let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">; + let TemplateDependent = 1; } def AMDGPUNumVGPR : InheritableAttr { let Spellings = [GNU<"amdgpu_num_vgpr">]; - let Args = [UnsignedArgument<"NumVGPR">]; + let Args = [ExprArgument<"NumVGPR">]; let Documentation = [AMDGPUNumSGPRNumVGPRDocs]; let Subjects = SubjectList<[Function], ErrorDiag, "ExpectedKernelFunction">; + let TemplateDependent = 1; } -def AMDGPUMaxWorkGroupDim : InheritableParamAttr { +def AMDGPUMaxWorkGroupDim : InheritableAttr { let Spellings = [CXX11<"","hc_max_workgroup_dim", 201511>]; - let Args = [IntArgument<"X">, - IntArgument<"Y", 1>, - IntArgument<"Z", 1>, + let Args = [ExprArgument<"X">, + ExprArgument<"Y">, + ExprArgument<"Z">, StringArgument<"ISA", 1>]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [Undocumented]; + let TemplateDependent = 1; } def NoSplitStack : InheritableAttr { Index: lib/CodeGen/TargetInfo.cpp =================================================================== --- lib/CodeGen/TargetInfo.cpp +++ lib/CodeGen/TargetInfo.cpp @@ -7662,6 +7662,18 @@ }; } +namespace +{ + inline + llvm::APSInt getConstexprInt(const Expr *E, const ASTContext& Ctx) + { + llvm::APSInt r{32, 0}; + if (E) E->EvaluateAsInt(r, Ctx); + + return r; + } +} + void AMDGPUTargetCodeGenInfo::setTargetAttributes( const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M, ForDefinition_t IsForDefinition) const { @@ -7678,8 +7690,11 @@ FD->getAttr() : nullptr; const auto *FlatWGS = FD->getAttr(); if (ReqdWGS || FlatWGS) { - unsigned Min = FlatWGS ? FlatWGS->getMin() : 0; - unsigned Max = FlatWGS ? FlatWGS->getMax() : 0; + llvm::APSInt min = getConstexprInt(FlatWGS->getMin(), FD->getASTContext()); + llvm::APSInt max = getConstexprInt(FlatWGS->getMax(), FD->getASTContext()); + + unsigned Min = min.getZExtValue(); + unsigned Max = std::max(min, max).getZExtValue(); if (ReqdWGS && Min == 0 && Max == 0) Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim(); @@ -7695,8 +7710,13 @@ } if (const auto *Attr = FD->getAttr()) { - unsigned Min = Attr->getMin(); - unsigned Max = Attr->getMax(); + llvm::APSInt min = getConstexprInt(Attr->getMin(), FD->getASTContext()); + llvm::APSInt max = getConstexprInt(Attr->getMax(), FD->getASTContext()); + + Attr->getMin()->dump(); + Attr->getMax()->dump(); + unsigned Min = min.getZExtValue(); + unsigned Max = std::max(min, max).getZExtValue(); if (Min != 0) { assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max"); @@ -7711,23 +7731,31 @@ } if (const auto *Attr = FD->getAttr()) { - unsigned NumSGPR = Attr->getNumSGPR(); + llvm::APSInt sgprs = + getConstexprInt(Attr->getNumSGPR(), FD->getASTContext()); + unsigned NumSGPR = sgprs.getZExtValue(); if (NumSGPR != 0) F->addFnAttr("amdgpu-num-sgpr", llvm::utostr(NumSGPR)); } if (const auto *Attr = FD->getAttr()) { - uint32_t NumVGPR = Attr->getNumVGPR(); + llvm::APSInt vgprs = + getConstexprInt(Attr->getNumVGPR(), FD->getASTContext()); + unsigned NumVGPR = vgprs.getZExtValue(); if (NumVGPR != 0) F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR)); } if (const auto *Attr = FD->getAttr()) { - unsigned X = Attr->getX(); - unsigned Y = Attr->getY(); - unsigned Z = Attr->getZ(); + llvm::APSInt x = getConstexprInt(Attr->getX(), FD->getASTContext()); + llvm::APSInt y = getConstexprInt(Attr->getY(), FD->getASTContext()); + llvm::APSInt z = getConstexprInt(Attr->getZ(), FD->getASTContext()); + + unsigned X = x.getZExtValue(); + unsigned Y = y.getZExtValue(); + unsigned Z = z.getZExtValue(); std::string AttrVal = llvm::utostr(X) + "," + llvm::utostr(Y) + "," + llvm::utostr(Z); F->addFnAttr("amdgpu-max-work-group-dim", AttrVal); Index: lib/Sema/SemaDeclAttr.cpp =================================================================== --- lib/Sema/SemaDeclAttr.cpp +++ lib/Sema/SemaDeclAttr.cpp @@ -5580,17 +5580,39 @@ } }; +namespace +{ + inline + bool checkAllAreIntegral(const AttributeList &Attr, Sema &S) { + for (auto i = 0u; i != Attr.getNumArgs(); ++i) { + auto e = Attr.getArgAsExpr(i); + if (e && !e->getType()->isIntegralOrEnumerationType()) { + S.Diag(getAttrLoc(Attr), diag::err_attribute_argument_n_type) + << getAttrName(Attr) << i << AANT_ArgumentIntegerConstant + << e->getSourceRange(); + + return false; + } + } + + return true; + } +} + static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D, const AttributeList &Attr) { uint32_t Min = 0; Expr *MinExpr = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, MinExpr, Min)) + if (MinExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, MinExpr, Min)) return; uint32_t Max = 0; + Expr *MaxExpr = MinExpr; if (Attr.getNumArgs() > 1 ) { - Expr *MaxExpr = Attr.getArgAsExpr(1); - if (!checkUInt32Argument(S, Attr, MaxExpr, Max)) + MaxExpr = Attr.getArgAsExpr(1); + if (MaxExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, MaxExpr, Max)) return; } @@ -5609,21 +5631,27 @@ StringRef ISA; if (VC.checkAMDGPUISAVersion(Attr, 2, ISA)) D->addAttr(::new (S.Context) - AMDGPUFlatWorkGroupSizeAttr(Attr.getLoc(), S.Context, Min, Max, - ISA, Attr.getAttributeSpellingListIndex())); + AMDGPUFlatWorkGroupSizeAttr(Attr.getLoc(), S.Context, MinExpr, + MaxExpr, ISA, Attr.getAttributeSpellingListIndex())); } static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (!checkAllAreIntegral(Attr, S)) + return; + uint32_t Min = 0; Expr *MinExpr = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, MinExpr, Min)) + if (MinExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, MinExpr, Min)) return; uint32_t Max = 0; + Expr *MaxExpr = MinExpr; if (Attr.getNumArgs() > 1) { - Expr *MaxExpr = Attr.getArgAsExpr(1); - if (!checkUInt32Argument(S, Attr, MaxExpr, Max)) + MaxExpr = Attr.getArgAsExpr(1); + if (MaxExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, MaxExpr, Max)) return; } @@ -5642,59 +5670,73 @@ StringRef ISA; if (VC.checkAMDGPUISAVersion(Attr, 2, ISA)) D->addAttr(::new (S.Context) - AMDGPUWavesPerEUAttr(Attr.getLoc(), S.Context, Min, Max, ISA, - Attr.getAttributeSpellingListIndex())); + AMDGPUWavesPerEUAttr(Attr.getLoc(), S.Context, MinExpr, MaxExpr, + ISA, Attr.getAttributeSpellingListIndex())); } static void handleAMDGPUNumSGPRAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (!checkAllAreIntegral(Attr, S)) + return; + uint32_t NumSGPR = 0; Expr *NumSGPRExpr = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, NumSGPRExpr, NumSGPR)) + if (NumSGPRExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, NumSGPRExpr, NumSGPR)) return; D->addAttr(::new (S.Context) - AMDGPUNumSGPRAttr(Attr.getLoc(), S.Context, NumSGPR, + AMDGPUNumSGPRAttr(Attr.getLoc(), S.Context, NumSGPRExpr, Attr.getAttributeSpellingListIndex())); } static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (!checkAllAreIntegral(Attr, S)) + return; + uint32_t NumVGPR = 0; Expr *NumVGPRExpr = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, NumVGPRExpr, NumVGPR)) + if (NumVGPRExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, NumVGPRExpr, NumVGPR)) return; D->addAttr(::new (S.Context) - AMDGPUNumVGPRAttr(Attr.getLoc(), S.Context, NumVGPR, + AMDGPUNumVGPRAttr(Attr.getLoc(), S.Context, NumVGPRExpr, Attr.getAttributeSpellingListIndex())); } static void handleAMDGPUMaxWorkGroupDimAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (!checkAllAreIntegral(Attr, S)) + return; if (!checkAttributeAtLeastNumArgs(S, Attr, 3)) return; uint32_t X = 0; Expr *XExpr = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, XExpr, X)) + if (XExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, XExpr, X)) return; uint32_t Y = 0; Expr *YExpr = Attr.getArgAsExpr(1); - if (!checkUInt32Argument(S, Attr, YExpr, Y)) + if (YExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, YExpr, Y)) return; uint32_t Z = 0; Expr *ZExpr = Attr.getArgAsExpr(2); - if (!checkUInt32Argument(S, Attr, ZExpr, Z)) + if (ZExpr->isEvaluatable(S.Context) && + !checkUInt32Argument(S, Attr, ZExpr, Z)) return; AMDGPUISAVersionChecker VC(S); StringRef ISA; if (VC.checkAMDGPUISAVersion(Attr, 3, ISA)) D->addAttr(::new (S.Context) - AMDGPUMaxWorkGroupDimAttr(Attr.getLoc(), S.Context, X, Y, Z, ISA, + AMDGPUMaxWorkGroupDimAttr(Attr.getLoc(), S.Context, XExpr, YExpr, + ZExpr, ISA, Attr.getAttributeSpellingListIndex())); }