diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index aa5a33a67875e..e4e7d7c667896 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2182,17 +2182,25 @@ class Sema final { SYCLIntelFPGAIVDepAttr * BuildSYCLIntelFPGAIVDepAttr(const AttributeCommonInfo &CI, Expr *Expr1, Expr *Expr2); - template - FPGALoopAttrT *BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, - Expr *E = nullptr); - LoopUnrollHintAttr *BuildLoopUnrollHintAttr(const AttributeCommonInfo &A, Expr *E); OpenCLUnrollHintAttr * BuildOpenCLLoopUnrollHintAttr(const AttributeCommonInfo &A, Expr *E); SYCLIntelFPGALoopCountAttr * - BuildSYCLIntelFPGALoopCount(const AttributeCommonInfo &CI, Expr *E); + BuildSYCLIntelFPGALoopCountAttr(const AttributeCommonInfo &CI, Expr *E); + SYCLIntelFPGAInitiationIntervalAttr * + BuildSYCLIntelFPGAInitiationIntervalAttr(const AttributeCommonInfo &CI, + Expr *E); + SYCLIntelFPGAMaxConcurrencyAttr * + BuildSYCLIntelFPGAMaxConcurrencyAttr(const AttributeCommonInfo &CI, Expr *E); + SYCLIntelFPGAMaxInterleavingAttr * + BuildSYCLIntelFPGAMaxInterleavingAttr(const AttributeCommonInfo &CI, Expr *E); + SYCLIntelFPGASpeculatedIterationsAttr * + BuildSYCLIntelFPGASpeculatedIterationsAttr(const AttributeCommonInfo &CI, + Expr *E); + SYCLIntelFPGALoopCoalesceAttr * + BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E); bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc); @@ -13527,51 +13535,6 @@ void Sema::AddOneConstantPowerTwoValueAttr(Decl *D, D->addAttr(::new (Context) AttrType(Context, CI, E)); } -template -FPGALoopAttrT *Sema::BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, - Expr *E) { - if (!E && !(A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGALoopCoalesce)) - return nullptr; - - if (E && !E->isInstantiationDependent()) { - Optional ArgVal = E->getIntegerConstantExpr(getASTContext()); - - if (!ArgVal) { - Diag(E->getExprLoc(), diag::err_attribute_argument_type) - << A.getAttrName() << AANT_ArgumentIntegerConstant - << E->getSourceRange(); - return nullptr; - } - - int Val = ArgVal->getSExtValue(); - - if (A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGAInitiationInterval || - A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGALoopCoalesce) { - if (Val <= 0) { - Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << A.getAttrName() << /* positive */ 0; - return nullptr; - } - } else if (A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGAMaxConcurrency || - A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving || - A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations || - A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGALoopCount) { - if (Val < 0) { - Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << A.getAttrName() << /* non-negative */ 1; - return nullptr; - } - } else { - llvm_unreachable("unknown sycl fpga loop attr"); - } - } - - return new (Context) FPGALoopAttrT(Context, A, E); -} - /// RAII object that enters a new expression evaluation context. class EnterExpressionEvaluationContext { Sema &Actions; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 0486afbda2fec..5083b2bd757a1 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -552,11 +552,11 @@ MDNode *LoopInfo::createMetadata( } // Setting max_concurrency attribute with number of threads - if (Attrs.SYCLMaxConcurrencyEnable) { - Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_concurrency.count"), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLMaxConcurrencyNThreads))}; + if (Attrs.SYCLMaxConcurrencyNThreads) { + Metadata *Vals[] = { + MDString::get(Ctx, "llvm.loop.max_concurrency.count"), + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxConcurrencyNThreads))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -582,11 +582,11 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLMaxInterleavingEnable) { + if (Attrs.SYCLMaxInterleavingNInvocations) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLMaxInterleavingNInvocations))}; + *Attrs.SYCLMaxInterleavingNInvocations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -596,16 +596,16 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLSpeculatedIterationsEnable) { + if (Attrs.SYCLSpeculatedIterationsNIterations) { Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), ConstantAsMetadata::get( ConstantInt::get(llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLSpeculatedIterationsNIterations))}; + *Attrs.SYCLSpeculatedIterationsNIterations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - for (auto &VC : Attrs.SYCLIntelFPGAVariantCount) { + for (const auto &VC : Attrs.SYCLIntelFPGAVariantCount) { Metadata *Vals[] = {MDString::get(Ctx, VC.first), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), VC.second))}; @@ -622,15 +622,12 @@ LoopAttributes::LoopAttributes(bool IsParallel) UnrollAndJamEnable(LoopAttributes::Unspecified), VectorizePredicateEnable(LoopAttributes::Unspecified), VectorizeWidth(0), VectorizeScalable(LoopAttributes::Unspecified), InterleaveCount(0), - SYCLIInterval(0), SYCLMaxConcurrencyEnable(false), - SYCLMaxConcurrencyNThreads(0), SYCLLoopCoalesceEnable(false), + SYCLIInterval(0), SYCLLoopCoalesceEnable(false), SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false), - SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), - SYCLSpeculatedIterationsEnable(false), - SYCLSpeculatedIterationsNIterations(0), UnrollCount(0), - UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), - PipelineDisabled(false), PipelineInitiationInterval(0), - SYCLNofusionEnable(false), MustProgress(false) {} + UnrollCount(0), UnrollAndJamCount(0), + DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), + PipelineInitiationInterval(0), SYCLNofusionEnable(false), + MustProgress(false) {} void LoopAttributes::clear() { IsParallel = false; @@ -640,15 +637,12 @@ void LoopAttributes::clear() { GlobalSYCLIVDepInfo.reset(); ArraySYCLIVDepInfo.clear(); SYCLIInterval = 0; - SYCLMaxConcurrencyEnable = false; - SYCLMaxConcurrencyNThreads = 0; + SYCLMaxConcurrencyNThreads.reset(); SYCLLoopCoalesceEnable = false; SYCLLoopCoalesceNLevels = 0; SYCLLoopPipeliningDisable = false; - SYCLMaxInterleavingEnable = false; - SYCLMaxInterleavingNInvocations = 0; - SYCLSpeculatedIterationsEnable = false; - SYCLSpeculatedIterationsNIterations = 0; + SYCLMaxInterleavingNInvocations.reset(); + SYCLSpeculatedIterationsNIterations.reset(); SYCLIntelFPGAVariantCount.clear(); UnrollCount = 0; UnrollAndJamCount = 0; @@ -679,14 +673,12 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.VectorizeScalable == LoopAttributes::Unspecified && Attrs.InterleaveCount == 0 && !Attrs.GlobalSYCLIVDepInfo.hasValue() && Attrs.ArraySYCLIVDepInfo.empty() && Attrs.SYCLIInterval == 0 && - Attrs.SYCLMaxConcurrencyEnable == false && + !Attrs.SYCLMaxConcurrencyNThreads && Attrs.SYCLLoopCoalesceEnable == false && Attrs.SYCLLoopCoalesceNLevels == 0 && Attrs.SYCLLoopPipeliningDisable == false && - Attrs.SYCLMaxInterleavingEnable == false && - Attrs.SYCLMaxInterleavingNInvocations == 0 && - Attrs.SYCLSpeculatedIterationsEnable == false && - Attrs.SYCLSpeculatedIterationsNIterations == 0 && + !Attrs.SYCLMaxInterleavingNInvocations && + !Attrs.SYCLSpeculatedIterationsNIterations && Attrs.SYCLIntelFPGAVariantCount.empty() && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && @@ -1025,39 +1017,42 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, IntelFPGAIVDep->getArrayDecl()); if (const auto *IntelFPGAII = - dyn_cast(A)) - setSYCLIInterval(IntelFPGAII->getIntervalExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + dyn_cast(A)) { + const auto *CE = cast(IntelFPGAII->getIntervalExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + setSYCLIInterval(ArgVal.getSExtValue()); + } if (const auto *IntelFPGAMaxConcurrency = dyn_cast(A)) { - setSYCLMaxConcurrencyEnable(); - setSYCLMaxConcurrencyNThreads(IntelFPGAMaxConcurrency->getNThreadsExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + const auto *CE = + cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + setSYCLMaxConcurrencyNThreads(ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCountAvg = dyn_cast(A)) { - unsigned int Count = IntelFPGALoopCountAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue(); + const auto *CE = + cast(IntelFPGALoopCountAvg->getNTripCount()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); const char *Var = IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" : "llvm.loop.intel.loopcount_avg"; - setSYCLIntelFPGAVariantCount(Var, Count); + setSYCLIntelFPGAVariantCount(Var, ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = dyn_cast(A)) { - if (auto *LCE = IntelFPGALoopCoalesce->getNExpr()) - setSYCLLoopCoalesceNLevels( - LCE->getIntegerConstantExpr(Ctx)->getSExtValue()); - else + if (const auto *LCE = IntelFPGALoopCoalesce->getNExpr()) { + const auto *CE = cast(LCE); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + setSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); + } else { setSYCLLoopCoalesceEnable(); + } } if (isa(A)) @@ -1065,19 +1060,17 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGAMaxInterleaving = dyn_cast(A)) { - setSYCLMaxInterleavingEnable(); - setSYCLMaxInterleavingNInvocations(IntelFPGAMaxInterleaving->getNExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + setSYCLMaxInterleavingNInvocations(ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = dyn_cast(A)) { - setSYCLSpeculatedIterationsEnable(); - setSYCLSpeculatedIterationsNIterations( - IntelFPGASpeculatedIterations->getNExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + const auto *CE = + cast(IntelFPGASpeculatedIterations->getNExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + setSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); } if (isa(A)) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 317972a34ebeb..f48cfb248c3cb 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -15,6 +15,7 @@ #define LLVM_CLANG_LIB_CODEGEN_CGLOOPINFO_H #include "llvm/ADT/ArrayRef.h" +#include "llvm/ADT/Optional.h" #include "llvm/ADT/SmallVector.h" #include "llvm/IR/DebugLoc.h" #include "llvm/IR/Value.h" @@ -111,11 +112,8 @@ struct LoopAttributes { /// Value for llvm.loop.ii.count metadata. unsigned SYCLIInterval; - /// Flag for llvm.loop.max_concurrency.count metadata. - bool SYCLMaxConcurrencyEnable; - /// Value for llvm.loop.max_concurrency.count metadata. - unsigned SYCLMaxConcurrencyNThreads; + llvm::Optional SYCLMaxConcurrencyNThreads; /// Value for count variant (min/max/avg) and count metadata. llvm::SmallVector, 2> @@ -130,17 +128,11 @@ struct LoopAttributes { /// Flag for llvm.loop.intel.pipelining.enable, i32 0 metadata. bool SYCLLoopPipeliningDisable; - /// Flag for llvm.loop.max_interleaving.count metadata. - bool SYCLMaxInterleavingEnable; - /// Value for llvm.loop.max_interleaving.count metadata. - unsigned SYCLMaxInterleavingNInvocations; - - /// Flag for llvm.loop.intel.speculated.iterations.count metadata. - bool SYCLSpeculatedIterationsEnable; + llvm::Optional SYCLMaxInterleavingNInvocations; /// Value for llvm.loop.intel.speculated.iterations.count metadata. - unsigned SYCLSpeculatedIterationsNIterations; + llvm::Optional SYCLSpeculatedIterationsNIterations; /// llvm.unroll. unsigned UnrollCount; @@ -363,12 +355,7 @@ class LoopInfoStack { /// Set value of an initiation interval for the next loop pushed. void setSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } - /// Set flag of max_concurrency for the next loop pushed. - void setSYCLMaxConcurrencyEnable() { - StagedAttrs.SYCLMaxConcurrencyEnable = true; - } - - /// Set value of threads for the next loop pushed. + /// Set value of max_concurrency for the next loop pushed. void setSYCLMaxConcurrencyNThreads(unsigned C) { StagedAttrs.SYCLMaxConcurrencyNThreads = C; } @@ -388,22 +375,12 @@ class LoopInfoStack { StagedAttrs.SYCLLoopPipeliningDisable = true; } - /// Set flag of max_interleaving for the next loop pushed. - void setSYCLMaxInterleavingEnable() { - StagedAttrs.SYCLMaxInterleavingEnable = true; - } - /// Set value of max interleaved invocations for the next loop pushed. void setSYCLMaxInterleavingNInvocations(unsigned C) { StagedAttrs.SYCLMaxInterleavingNInvocations = C; } - /// Set flag of speculated_iterations for the next loop pushed. - void setSYCLSpeculatedIterationsEnable() { - StagedAttrs.SYCLSpeculatedIterationsEnable = true; - } - - /// Set value of concurrent speculated iterations for the next loop pushed. + /// Set value of speculated iterations for the next loop pushed. void setSYCLSpeculatedIterationsNIterations(unsigned C) { StagedAttrs.SYCLSpeculatedIterationsNIterations = C; } diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 5f5e40df21ae9..623bd6c6b5461 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -68,12 +68,149 @@ static Attr *handleSuppressAttr(Sema &S, Stmt *St, const ParsedAttr &A, S.Context, A, DiagnosticIdentifiers.data(), DiagnosticIdentifiers.size()); } -template -static Attr *handleIntelFPGALoopAttr(Sema &S, Stmt *St, const ParsedAttr &A) { +SYCLIntelFPGAMaxConcurrencyAttr * +Sema::BuildSYCLIntelFPGAMaxConcurrencyAttr(const AttributeCommonInfo &CI, + Expr *E) { + if (!E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return nullptr; + E = Res.get(); + + // This attribute requires a non-negative value. + if (ArgVal < 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; + return nullptr; + } + } + + return new (Context) SYCLIntelFPGAMaxConcurrencyAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGAMaxConcurrencyAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { + S.CheckDeprecatedSYCLAttributeSpelling(A); + + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGAMaxConcurrencyAttr(A, E); +} + +SYCLIntelFPGAInitiationIntervalAttr * +Sema::BuildSYCLIntelFPGAInitiationIntervalAttr(const AttributeCommonInfo &CI, + Expr *E) { + if (!E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return nullptr; + E = Res.get(); + + // This attribute requires a strictly positive value. + if (ArgVal <= 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*positive*/ 0; + return nullptr; + } + } + + return new (Context) SYCLIntelFPGAInitiationIntervalAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGAInitiationIntervalAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { + S.CheckDeprecatedSYCLAttributeSpelling(A); + + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGAInitiationIntervalAttr(A, E); +} + +SYCLIntelFPGAMaxInterleavingAttr * +Sema::BuildSYCLIntelFPGAMaxInterleavingAttr(const AttributeCommonInfo &CI, + Expr *E) { + if (!E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return nullptr; + E = Res.get(); + + // This attribute requires a non-negative value. + if (ArgVal < 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; + return nullptr; + } + } + + return new (Context) SYCLIntelFPGAMaxInterleavingAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGAMaxInterleavingAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - return S.BuildSYCLIntelFPGALoopAttr( - A, A.getNumArgs() ? A.getArgAsExpr(0) : nullptr); + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGAMaxInterleavingAttr(A, E); +} + +SYCLIntelFPGALoopCoalesceAttr * +Sema::BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, + Expr *E) { + if (E && !E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return nullptr; + E = Res.get(); + + // This attribute requires a strictly positive value. + if (ArgVal <= 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*positive*/ 0; + return nullptr; + } + } + + return new (Context) SYCLIntelFPGALoopCoalesceAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGALoopCoalesceAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { + S.CheckDeprecatedSYCLAttributeSpelling(A); + + Expr *E = A.isArgExpr(0) ? A.getArgAsExpr(0) : nullptr; + return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); +} + +SYCLIntelFPGASpeculatedIterationsAttr * +Sema::BuildSYCLIntelFPGASpeculatedIterationsAttr(const AttributeCommonInfo &CI, + Expr *E) { + if (!E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) + return nullptr; + E = Res.get(); + + // This attribute requires a non-negative value. + if (ArgVal < 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; + return nullptr; + } + } + + return new (Context) SYCLIntelFPGASpeculatedIterationsAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGASpeculatedIterationsAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { + S.CheckDeprecatedSYCLAttributeSpelling(A); + + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGASpeculatedIterationsAttr(A, E); } static Attr *handleSYCLIntelFPGADisableLoopPipeliningAttr(Sema &S, Stmt *, @@ -270,27 +407,32 @@ CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, } } -static SYCLIntelFPGALoopCountAttr * -handleIntelFPGALoopCountAttr(Sema &S, Stmt *St, const ParsedAttr &A) { - Expr *E = A.getArgAsExpr(0); - if (E && !E->isInstantiationDependent()) { - Optional ArgVal = - E->getIntegerConstantExpr(S.getASTContext()); - - if (!ArgVal) { - S.Diag(E->getExprLoc(), diag::err_attribute_argument_type) - << A << AANT_ArgumentIntegerConstant << E->getSourceRange(); +SYCLIntelFPGALoopCountAttr * +Sema::BuildSYCLIntelFPGALoopCountAttr(const AttributeCommonInfo &CI, Expr *E) { + if (!E->isValueDependent()) { + llvm::APSInt ArgVal; + ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); + if (Res.isInvalid()) return nullptr; - } + E = Res.get(); - if (ArgVal->getSExtValue() < 0) { - S.Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << A << /* non-negative */ 1; + // This attribute requires a non-negative value. + if (ArgVal < 0) { + Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << CI << /*non-negative*/ 1; return nullptr; } } - return new (S.Context) - SYCLIntelFPGALoopCountAttr(S.Context, A, A.getArgAsExpr(0)); + + return new (Context) SYCLIntelFPGALoopCountAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGALoopCountAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { + S.CheckDeprecatedSYCLAttributeSpelling(A); + + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGALoopCountAttr(A, E); } static Attr *handleIntelFPGANofusionAttr(Sema &S, Stmt *St, @@ -731,21 +873,19 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A, case ParsedAttr::AT_SYCLIntelFPGAIVDep: return handleIntelFPGAIVDepAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGAInitiationInterval: - return handleIntelFPGALoopAttr(S, St, - A); + return handleSYCLIntelFPGAInitiationIntervalAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGAMaxConcurrency: - return handleIntelFPGALoopAttr(S, St, A); + return handleSYCLIntelFPGAMaxConcurrencyAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGALoopCoalesce: - return handleIntelFPGALoopAttr(S, St, A); + return handleSYCLIntelFPGALoopCoalesceAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGADisableLoopPipelining: return handleSYCLIntelFPGADisableLoopPipeliningAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving: - return handleIntelFPGALoopAttr(S, St, A); + return handleSYCLIntelFPGAMaxInterleavingAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations: - return handleIntelFPGALoopAttr(S, St, - A); + return handleSYCLIntelFPGASpeculatedIterationsAttr(S, St, A); case ParsedAttr::AT_SYCLIntelFPGALoopCount: - return handleIntelFPGALoopCountAttr(S, St, A); + return handleSYCLIntelFPGALoopCountAttr(S, St, A); case ParsedAttr::AT_OpenCLUnrollHint: case ParsedAttr::AT_LoopUnrollHint: return handleLoopUnrollHint(S, St, A, Range); diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index d8d6203fe5861..f0cee17d30ec3 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1541,9 +1541,8 @@ TemplateInstantiator::TransformSYCLIntelFPGAInitiationIntervalAttr( const SYCLIntelFPGAInitiationIntervalAttr *II) { Expr *TransformedExpr = getDerived().TransformExpr(II->getIntervalExpr()).get(); - return getSema() - .BuildSYCLIntelFPGALoopAttr( - *II, TransformedExpr); + return getSema().BuildSYCLIntelFPGAInitiationIntervalAttr(*II, + TransformedExpr); } const SYCLIntelFPGAMaxConcurrencyAttr * @@ -1551,33 +1550,29 @@ TemplateInstantiator::TransformSYCLIntelFPGAMaxConcurrencyAttr( const SYCLIntelFPGAMaxConcurrencyAttr *MC) { Expr *TransformedExpr = getDerived().TransformExpr(MC->getNThreadsExpr()).get(); - return getSema().BuildSYCLIntelFPGALoopAttr( - *MC, TransformedExpr); + return getSema().BuildSYCLIntelFPGAMaxConcurrencyAttr(*MC, TransformedExpr); } const SYCLIntelFPGALoopCoalesceAttr * TemplateInstantiator::TransformSYCLIntelFPGALoopCoalesceAttr( const SYCLIntelFPGALoopCoalesceAttr *LC) { Expr *TransformedExpr = getDerived().TransformExpr(LC->getNExpr()).get(); - return getSema().BuildSYCLIntelFPGALoopAttr( - *LC, TransformedExpr); + return getSema().BuildSYCLIntelFPGALoopCoalesceAttr(*LC, TransformedExpr); } const SYCLIntelFPGAMaxInterleavingAttr * TemplateInstantiator::TransformSYCLIntelFPGAMaxInterleavingAttr( const SYCLIntelFPGAMaxInterleavingAttr *MI) { Expr *TransformedExpr = getDerived().TransformExpr(MI->getNExpr()).get(); - return getSema().BuildSYCLIntelFPGALoopAttr( - *MI, TransformedExpr); + return getSema().BuildSYCLIntelFPGAMaxInterleavingAttr(*MI, TransformedExpr); } const SYCLIntelFPGASpeculatedIterationsAttr * TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI) { Expr *TransformedExpr = getDerived().TransformExpr(SI->getNExpr()).get(); - return getSema() - .BuildSYCLIntelFPGALoopAttr( - *SI, TransformedExpr); + return getSema().BuildSYCLIntelFPGASpeculatedIterationsAttr(*SI, + TransformedExpr); } const SYCLIntelFPGALoopCountAttr * @@ -1585,8 +1580,7 @@ TemplateInstantiator::TransformSYCLIntelFPGALoopCountAttr( const SYCLIntelFPGALoopCountAttr *LCA) { Expr *TransformedExpr = getDerived().TransformExpr(LCA->getNTripCount()).get(); - return getSema().BuildSYCLIntelFPGALoopAttr( - *LCA, TransformedExpr); + return getSema().BuildSYCLIntelFPGALoopCountAttr(*LCA, TransformedExpr); } const LoopUnrollHintAttr *TemplateInstantiator::TransformLoopUnrollHintAttr( diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index a0dfafdbcd92f..06c0d80d7988c 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -12,8 +12,10 @@ // CHECK: br label %for.cond13, !llvm.loop ![[MD_LC_3:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_MI:[0-9]+]] // CHECK: br label %for.cond2, !llvm.loop ![[MD_MI_2:[0-9]+]] +// CHECK: br label %for.cond13, !llvm.loop ![[MD_MI_3:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_SI:[0-9]+]] // CHECK: br label %for.cond2, !llvm.loop ![[MD_SI_2:[0-9]+]] +// CHECK: br label %for.cond13, !llvm.loop ![[MD_SI_3:[0-9]+]] // CHECK: br label %for.cond, !llvm.loop ![[MD_LCA:[0-9]+]] // CHECK: br label %for.cond2, !llvm.loop ![[MD_LCA_1:[0-9]+]] // CHECK: br label %for.cond13, !llvm.loop ![[MD_LCA_2:[0-9]+]] @@ -86,7 +88,7 @@ void loop_coalesce() { a[i] = 0; } -template +template void max_interleaving() { int a[10]; // CHECK: ![[MD_MI]] = distinct !{![[MD_MI]], ![[MP]], ![[MD_max_interleaving:[0-9]+]]} @@ -97,9 +99,15 @@ void max_interleaving() { // CHECK-NEXT: ![[MD_max_interleaving_2]] = !{!"llvm.loop.max_interleaving.count", i32 2} [[intel::max_interleaving(2)]] for (int i = 0; i != 10; ++i) a[i] = 0; + + // CHECK: ![[MD_MI_3]] = distinct !{![[MD_MI_3]], ![[MP]], ![[MD_max_interleaving_3:[0-9]+]]} + // CHECK-NEXT: ![[MD_max_interleaving_3]] = !{!"llvm.loop.max_interleaving.count", i32 0} + [[intel::max_interleaving(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + } -template +template void speculated_iterations() { int a[10]; // CHECK: ![[MD_SI]] = distinct !{![[MD_SI]], ![[MP]], ![[MD_speculated_iterations:[0-9]+]]} @@ -110,6 +118,11 @@ void speculated_iterations() { // CHECK-NEXT: ![[MD_speculated_iterations_2]] = !{!"llvm.loop.intel.speculated.iterations.count", i32 5} [[intel::speculated_iterations(5)]] for (int i = 0; i != 10; ++i) a[i] = 0; + + // CHECK: ![[MD_SI_3]] = distinct !{![[MD_SI_3]], ![[MP]], ![[MD_speculated_iterations_3:[0-9]+]]} + // CHECK-NEXT: ![[MD_speculated_iterations_3]] = !{!"llvm.loop.intel.speculated.iterations.count", i32 0} + [[intel::speculated_iterations(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; } template @@ -143,8 +156,8 @@ int main() { initiation_interval<6>(); max_concurrency<0>(); loop_coalesce<2>(); - max_interleaving<3>(); - speculated_iterations<4>(); + max_interleaving<3, 0>(); + speculated_iterations<4, 0>(); loop_count_control<12>(); }); return 0; diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index ce26f4d977780..af15a5799ac65 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -156,19 +156,19 @@ void goo() { // expected-error@+1 {{unknown argument to 'ivdep'; expected integer or array variable}} [[intel::ivdep("test123")]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'initiation_interval' attribute requires an integer constant}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} [[intel::initiation_interval("test123")]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'max_concurrency' attribute requires an integer constant}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} [[intel::max_concurrency("test123")]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'loop_coalesce' attribute requires an integer constant}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} [[intel::loop_coalesce("test123")]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'max_interleaving' attribute requires an integer constant}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} [[intel::max_interleaving("test123")]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'speculated_iterations' attribute requires an integer constant}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} [[intel::speculated_iterations("test123")]] for (int i = 0; i != 10; ++i) a[i] = 0; // expected-error@+1 {{unknown argument to 'ivdep'; expected integer or array variable}} @@ -207,8 +207,8 @@ void goo() { // expected-error@+1 {{'loop_count_avg' attribute requires a non-negative integral compile time constant expression}} [[intel::loop_count_avg(-1)]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'loop_count_avg' attribute requires an integer constant}} - [[intel::loop_count_avg("abc")]] for (int i = 0; i != 10; ++i) + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}} + [[intel::loop_count_avg("abc")]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -419,7 +419,7 @@ void ii_dependent() { a[i] = 0; } -template +template void max_concurrency_dependent() { int a[10]; // expected-error@+1 {{'max_concurrency' attribute requires a non-negative integral compile time constant expression}} @@ -430,6 +430,61 @@ void max_concurrency_dependent() { [[intel::max_concurrency(A)]] [[intel::max_concurrency(B)]] for (int i = 0; i != 10; ++i) a[i] = 0; + + // max_concurrency attribute accepts value 0. + [[intel::max_concurrency(D)]] for (int i = 0; i != 10; ++i) + a[i] = 0; +} + +template +void max_interleaving_dependent() { + int a[10]; + // expected-error@+1 {{'max_interleaving' attribute requires a non-negative integral compile time constant expression}} + [[intel::max_interleaving(C)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{duplicate Intel FPGA loop attribute 'max_interleaving'}} + [[intel::max_interleaving(A)]] + [[intel::max_interleaving(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // max_interleaving attribute accepts value 0. + [[intel::max_interleaving(D)]] for (int i = 0; i != 10; ++i) + a[i] = 0; +} + +template +void speculated_iterations_dependent() { + int a[10]; + // expected-error@+1 {{'speculated_iterations' attribute requires a non-negative integral compile time constant expression}} + [[intel::speculated_iterations(C)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{duplicate Intel FPGA loop attribute 'speculated_iterations'}} + [[intel::speculated_iterations(A)]] + [[intel::speculated_iterations(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // speculated_iterations attribute accepts value 0. + [[intel::speculated_iterations(D)]] for (int i = 0; i != 10; ++i) + a[i] = 0; +} + +template +void loop_coalesce_dependent() { + int a[10]; + // expected-error@+1 {{'loop_coalesce' attribute requires a positive integral compile time constant expression}} + [[intel::loop_coalesce(A)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{duplicate Intel FPGA loop attribute 'loop_coalesce'}} + [[intel::loop_coalesce]] + [[intel::loop_coalesce(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+1 {{'loop_coalesce' attribute requires a positive integral compile time constant expression}} + [[intel::loop_coalesce(C)]] for (int i = 0; i != 10; ++i) + a[i] = 0; } template @@ -468,6 +523,140 @@ void loop_count_control_dependent() { } +void check_max_concurrency_expression() { + int a[10]; + // Test that checks expression is not a constant expression. + // expected-note@+1{{declared here}} + int foo; + // expected-error@+2{{expression is not an integral constant expression}} + // expected-note@+1{{read of non-const variable 'foo' is not allowed in a constant expression}} + [[intel::max_concurrency(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // Test that checks expression is a constant expression. + constexpr int bar = 0; + [[intel::max_concurrency(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; +} + +void check_max_interleaving_expression() { + int a[10]; + // Test that checks expression is not a constant expression. + // expected-note@+1{{declared here}} + int foo; + // expected-error@+2{{expression is not an integral constant expression}} + // expected-note@+1{{read of non-const variable 'foo' is not allowed in a constant expression}} + [[intel::max_interleaving(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // Test that checks expression is a constant expression. + constexpr int bar = 0; + [[intel::max_interleaving(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; +} + +void check_initiation_interval_expression() { + int a[10]; + // Test that checks expression is not a constant expression. + // expected-note@+1{{declared here}} + int foo; + // expected-error@+2{{expression is not an integral constant expression}} + // expected-note@+1{{read of non-const variable 'foo' is not allowed in a constant expression}} + [[intel::initiation_interval(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // Test that checks expression is a constant expression. + constexpr int bar = 0; + [[intel::initiation_interval(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; +} + +void check_speculated_iterations_expression() { + int a[10]; + // Test that checks expression is not a constant expression. + // expected-note@+1{{declared here}} + int foo; + // expected-error@+2{{expression is not an integral constant expression}} + // expected-note@+1{{read of non-const variable 'foo' is not allowed in a constant expression}} + [[intel::speculated_iterations(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // Test that checks expression is a constant expression. + constexpr int bar = 0; + [[intel::speculated_iterations(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; +} + +void check_loop_coalesce_expression() { + int a[10]; + // Test that checks expression is not a constant expression. + // expected-note@+1{{declared here}} + int foo; + // expected-error@+2{{expression is not an integral constant expression}} + // expected-note@+1{{read of non-const variable 'foo' is not allowed in a constant expression}} + [[intel::loop_coalesce(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // Test that checks expression is a constant expression. + constexpr int bar = 0; + [[intel::loop_coalesce(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; +} + +void check_loop_count_expression() { + int a[10]; + + // Test that checks expression is not a constant expression. + int foo; // expected-note {{declared here}} + // expected-error@+2{{expression is not an integral constant expression}} + // expected-note@+1{{read of non-const variable 'foo' is not allowed in a constant expression}} + [[intel::loop_count_max(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // Test that checks expression is a constant expression. + constexpr int bar = 0; + [[intel::loop_count_max(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; +} + +// Test that checks wrong template instantiation and ensures that the type +// is checked properly when instantiating from the template definition. +struct S {}; +template +void check_loop_attr_template_instantiation() { + int a[10]; + + // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::initiation_interval(Ty{})]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::loop_coalesce(Ty{})]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::speculated_iterations(Ty{})]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::max_interleaving(Ty{})]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::max_concurrency(Ty{})]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}} + [[intel::loop_count_min(Ty{})]] for (int i = 0; i != 10; ++i) + a[i] = 0; +} + int main() { deviceQueue.submit([&](sycl::handler &h) { h.single_task([]() { @@ -483,12 +672,20 @@ int main() { //expected-note@-1 +{{in instantiation of function template specialization}} ii_dependent<2, 4, -1>(); //expected-note@-1 +{{in instantiation of function template specialization}} - max_concurrency_dependent<1, 4, -2>(); - //expected-note@-1 +{{in instantiation of function template specialization}} - - loop_count_control_dependent<3, 2, -1>(); - //expected-note@-1{{in instantiation of function template specialization 'loop_count_control_dependent<3, 2, -1>' requested here}} -}); + max_concurrency_dependent<1, 4, -2, 0>(); // expected-note{{in instantiation of function template specialization 'max_concurrency_dependent<1, 4, -2, 0>' requested here}} + max_interleaving_dependent<1, 4, -1, 0>(); // expected-note{{in instantiation of function template specialization 'max_interleaving_dependent<1, 4, -1, 0>' requested here}} + speculated_iterations_dependent<1, 8, -3, 0>(); // expected-note{{in instantiation of function template specialization 'speculated_iterations_dependent<1, 8, -3, 0>' requested here}} + loop_coalesce_dependent<-1, 4, 0>(); // expected-note{{in instantiation of function template specialization 'loop_coalesce_dependent<-1, 4, 0>' requested here}} + loop_count_control_dependent<3, 2, -1>(); // expected-note{{in instantiation of function template specialization 'loop_count_control_dependent<3, 2, -1>' requested here}} + check_max_concurrency_expression(); + check_max_interleaving_expression(); + check_speculated_iterations_expression(); + check_loop_coalesce_expression(); + check_initiation_interval_expression(); + check_loop_count_expression(); + check_loop_attr_template_instantiation(); //expected-note{{in instantiation of function template specialization 'check_loop_attr_template_instantiation' requested here}} + check_loop_attr_template_instantiation(); //expected-note{{in instantiation of function template specialization 'check_loop_attr_template_instantiation' requested here}} + }); }); return 0;