From f9d050dd9b0718efb1d6c308398e26d7a24c826c Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 05:45:48 -0700 Subject: [PATCH 01/24] Refactor of loop attributes Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 64 ++----- clang/lib/Sema/SemaStmtAttr.cpp | 201 ++++++++++++++++++--- clang/lib/Sema/SemaTemplateInstantiate.cpp | 18 +- 3 files changed, 193 insertions(+), 90 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index aa5a33a67875..096a246d771f 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2182,17 +2182,26 @@ 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 +13536,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/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 5f5e40df21ae..dfdc904e84b4 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -68,12 +68,153 @@ 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); + + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGAMaxInterleavingAttr(A, E); +} + +SYCLIntelFPGALoopCoalesceAttr * +Sema::BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, + Expr *E) { + if (!E) + return nullptr; + + 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) SYCLIntelFPGALoopCoalesceAttr(Context, CI, E); +} + +static Attr *handleSYCLIntelFPGALoopCoalesceAttr(Sema &S, Stmt *St, + const ParsedAttr &A) { + S.CheckDeprecatedSYCLAttributeSpelling(A); + + Expr *E = A.getNumArgs() ? 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); - return S.BuildSYCLIntelFPGALoopAttr( - A, A.getNumArgs() ? A.getArgAsExpr(0) : nullptr); + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGASpeculatedIterationsAttr(A, E); } static Attr *handleSYCLIntelFPGADisableLoopPipeliningAttr(Sema &S, Stmt *, @@ -270,27 +411,33 @@ 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 +878,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 d8d6203fe586..1958651de8ca 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1542,8 +1542,7 @@ TemplateInstantiator::TransformSYCLIntelFPGAInitiationIntervalAttr( Expr *TransformedExpr = getDerived().TransformExpr(II->getIntervalExpr()).get(); return getSema() - .BuildSYCLIntelFPGALoopAttr( - *II, TransformedExpr); + .BuildSYCLIntelFPGAInitiationIntervalAttr(*II, TransformedExpr); } const SYCLIntelFPGAMaxConcurrencyAttr * @@ -1551,24 +1550,21 @@ 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 * @@ -1576,8 +1572,7 @@ TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI) { Expr *TransformedExpr = getDerived().TransformExpr(SI->getNExpr()).get(); return getSema() - .BuildSYCLIntelFPGALoopAttr( - *SI, TransformedExpr); + .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( From 9f7cee7a376f4ecb3f6c3c8b9701a234ef5d8601 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 05:59:03 -0700 Subject: [PATCH 02/24] fix clang format problems Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 8 ++++++-- clang/lib/Sema/SemaStmtAttr.cpp | 3 +-- clang/lib/Sema/SemaTemplateInstantiate.cpp | 8 ++++---- 3 files changed, 11 insertions(+), 8 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 096a246d771f..8c136796a694 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2192,14 +2192,18 @@ class Sema final { SYCLIntelFPGAInitiationIntervalAttr * BuildSYCLIntelFPGAInitiationIntervalAttr(const AttributeCommonInfo &CI, - Expr *E); + Expr *E); + SYCLIntelFPGAMaxConcurrencyAttr * BuildSYCLIntelFPGAMaxConcurrencyAttr(const AttributeCommonInfo &CI, Expr *E); + SYCLIntelFPGAMaxInterleavingAttr * BuildSYCLIntelFPGAMaxInterleavingAttr(const AttributeCommonInfo &CI, Expr *E); + SYCLIntelFPGASpeculatedIterationsAttr * BuildSYCLIntelFPGASpeculatedIterationsAttr(const AttributeCommonInfo &CI, - Expr *E); + Expr *E); + SYCLIntelFPGALoopCoalesceAttr * BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E); diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index dfdc904e84b4..5f59a63e3c10 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -412,8 +412,7 @@ CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, } SYCLIntelFPGALoopCountAttr * -Sema::BuildSYCLIntelFPGALoopCountAttr(const AttributeCommonInfo &CI, - Expr *E) { +Sema::BuildSYCLIntelFPGALoopCountAttr(const AttributeCommonInfo &CI, Expr *E) { if (!E->isValueDependent()) { llvm::APSInt ArgVal; ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index 1958651de8ca..f0cee17d30ec 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1541,8 +1541,8 @@ TemplateInstantiator::TransformSYCLIntelFPGAInitiationIntervalAttr( const SYCLIntelFPGAInitiationIntervalAttr *II) { Expr *TransformedExpr = getDerived().TransformExpr(II->getIntervalExpr()).get(); - return getSema() - .BuildSYCLIntelFPGAInitiationIntervalAttr(*II, TransformedExpr); + return getSema().BuildSYCLIntelFPGAInitiationIntervalAttr(*II, + TransformedExpr); } const SYCLIntelFPGAMaxConcurrencyAttr * @@ -1571,8 +1571,8 @@ const SYCLIntelFPGASpeculatedIterationsAttr * TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI) { Expr *TransformedExpr = getDerived().TransformExpr(SI->getNExpr()).get(); - return getSema() - .BuildSYCLIntelFPGASpeculatedIterationsAttr(*SI, TransformedExpr); + return getSema().BuildSYCLIntelFPGASpeculatedIterationsAttr(*SI, + TransformedExpr); } const SYCLIntelFPGALoopCountAttr * From 275e53e0c12f37ee350b676ba318c4ae2e019974 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 06:10:26 -0700 Subject: [PATCH 03/24] fix clang format problems Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaStmtAttr.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 5f59a63e3c10..395e548aa6eb 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -147,7 +147,6 @@ Sema::BuildSYCLIntelFPGAMaxInterleavingAttr(const AttributeCommonInfo &CI, return new (Context) SYCLIntelFPGAMaxInterleavingAttr(Context, CI, E); } - static Attr *handleSYCLIntelFPGAMaxInterleavingAttr(Sema &S, Stmt *St, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); From b8d823de39fc4f72aa658292c2498d3de1cb7947 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 07:28:30 -0700 Subject: [PATCH 04/24] update codegen Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 3 ++- clang/lib/CodeGen/CGLoopInfo.cpp | 25 ++++++++++++------------ clang/lib/Sema/SemaStmtAttr.cpp | 4 +++- clang/test/SemaSYCL/intel-fpga-loops.cpp | 12 ++++++------ 4 files changed, 23 insertions(+), 21 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 8c136796a694..c7fe72c79eb1 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2205,7 +2205,8 @@ class Sema final { Expr *E); SYCLIntelFPGALoopCoalesceAttr * - BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E); + BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, + Expr *E = nullptr); bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc); diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 0486afbda2fe..e5de1cb0e3a1 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1032,17 +1032,17 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGAMaxConcurrency = dyn_cast(A)) { + const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); setSYCLMaxConcurrencyEnable(); - setSYCLMaxConcurrencyNThreads(IntelFPGAMaxConcurrency->getNThreadsExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + 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(); + unsigned int Count = ArgVal.getSExtValue(); const char *Var = IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" : IntelFPGALoopCountAvg->isMin() @@ -1065,19 +1065,18 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGAMaxInterleaving = dyn_cast(A)) { + const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); setSYCLMaxInterleavingEnable(); - setSYCLMaxInterleavingNInvocations(IntelFPGAMaxInterleaving->getNExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + setSYCLMaxInterleavingNInvocations(ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = dyn_cast(A)) { + const auto *CE = cast(IntelFPGASpeculatedIterations->getNExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); setSYCLSpeculatedIterationsEnable(); - setSYCLSpeculatedIterationsNIterations( - IntelFPGASpeculatedIterations->getNExpr() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + setSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); } if (isa(A)) diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 395e548aa6eb..a562a8e9befc 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -183,7 +183,9 @@ static Attr *handleSYCLIntelFPGALoopCoalesceAttr(Sema &S, Stmt *St, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - Expr *E = A.getNumArgs() ? A.getArgAsExpr(0) : nullptr; + Expr *E = A.isArgExpr(0) + ? A.getArgAsExpr(0) + : nullptr; return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); } diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index ce26f4d97778..32d7f2742995 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,7 +207,7 @@ 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}} + // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} [[intel::loop_count_avg("abc")]] for (int i = 0; i != 10; ++i) a[i] = 0; } From 24d09e3b8936d7bcd1fe49cb967fd419f70aca30 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 07:34:42 -0700 Subject: [PATCH 05/24] fix format errors Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/CodeGen/CGLoopInfo.cpp | 9 ++++++--- clang/lib/Sema/SemaStmtAttr.cpp | 4 +--- 3 files changed, 8 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c7fe72c79eb1..991a7e654edc 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2206,7 +2206,7 @@ class Sema final { SYCLIntelFPGALoopCoalesceAttr * BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, - Expr *E = nullptr); + Expr *E = nullptr); bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc); diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index e5de1cb0e3a1..8d04dc1f14dc 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1032,7 +1032,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGAMaxConcurrency = dyn_cast(A)) { - const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); + const auto *CE = + cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); setSYCLMaxConcurrencyEnable(); setSYCLMaxConcurrencyNThreads(ArgVal.getSExtValue()); @@ -1040,7 +1041,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopCountAvg = dyn_cast(A)) { - const auto *CE = cast(IntelFPGALoopCountAvg->getNTripCount()); + const auto *CE = + cast(IntelFPGALoopCountAvg->getNTripCount()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); unsigned int Count = ArgVal.getSExtValue(); const char *Var = IntelFPGALoopCountAvg->isMax() @@ -1073,7 +1075,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGASpeculatedIterations = dyn_cast(A)) { - const auto *CE = cast(IntelFPGASpeculatedIterations->getNExpr()); + const auto *CE = + cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); setSYCLSpeculatedIterationsEnable(); setSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index a562a8e9befc..a08843a7e78a 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -183,9 +183,7 @@ static Attr *handleSYCLIntelFPGALoopCoalesceAttr(Sema &S, Stmt *St, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - Expr *E = A.isArgExpr(0) - ? A.getArgAsExpr(0) - : nullptr; + Expr *E = A.isArgExpr(0) ? A.getArgAsExpr(0) : nullptr; return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); } From dd39f3a09b5016a8b5431af0eaea44f5bf185846 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 08:21:31 -0700 Subject: [PATCH 06/24] fix Lit test failures Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 4 ++-- clang/lib/Sema/SemaStmtAttr.cpp | 12 ++++++++---- clang/test/SemaSYCL/intel-fpga-loops.cpp | 4 ++-- 3 files changed, 12 insertions(+), 8 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 991a7e654edc..d6a6d4f3e182 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2205,8 +2205,8 @@ class Sema final { Expr *E); SYCLIntelFPGALoopCoalesceAttr * - BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, - Expr *E = nullptr); + BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E); + //Expr *E = nullptr); bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc); diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index a08843a7e78a..87e820c5ef85 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -158,8 +158,8 @@ static Attr *handleSYCLIntelFPGAMaxInterleavingAttr(Sema &S, Stmt *St, SYCLIntelFPGALoopCoalesceAttr * Sema::BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E) { - if (!E) - return nullptr; + //if (!E) + //return nullptr; if (!E->isValueDependent()) { llvm::APSInt ArgVal; @@ -183,8 +183,12 @@ 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); + //Expr *E = A.isArgExpr(0) ? A.getArgAsExpr(0) : nullptr; + if (A.isArgExpr(0)) { + Expr *E = A.getArgAsExpr(0); + return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); + } + return new (S.Context) SYCLIntelFPGALoopCoalesceAttr(S.Context, A); } SYCLIntelFPGASpeculatedIterationsAttr * diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 32d7f2742995..5d2ad7642314 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -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 {{integral constant expression must have integral or unscoped enumeration type, not 'const char [8]'}} - [[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; } From 2cca0ce3f0a30c69f5efe6760ce7c4eb45a66a39 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 13:13:29 -0700 Subject: [PATCH 07/24] update patch Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 9 +--- clang/lib/CodeGen/CGLoopInfo.cpp | 88 ++++++++++++++++---------------- clang/lib/CodeGen/CGLoopInfo.h | 57 +++++++-------------- clang/lib/Sema/SemaStmtAttr.cpp | 15 ++---- 4 files changed, 70 insertions(+), 99 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index d6a6d4f3e182..b90bf8ad9759 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2189,24 +2189,19 @@ class Sema final { SYCLIntelFPGALoopCountAttr * 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); - //Expr *E = nullptr); + BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, + Expr *E = nullptr); bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc); diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 8d04dc1f14dc..49f11b83d989 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -552,11 +552,10 @@ MDNode *LoopInfo::createMetadata( } // Setting max_concurrency attribute with number of threads - if (Attrs.SYCLMaxConcurrencyEnable) { - Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_concurrency.count"), + for (auto &MC : Attrs.SYCLMaxConcurrencyNThreads) { + Metadata *Vals[] = {MDString::get(Ctx, MC.first), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLMaxConcurrencyNThreads))}; + llvm::Type::getInt32Ty(Ctx), MC.second))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -582,11 +581,10 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLMaxInterleavingEnable) { - Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), + for (auto &MI : Attrs.SYCLMaxInterleavingNInvocations) { + Metadata *Vals[] = {MDString::get(Ctx, MI.first), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLMaxInterleavingNInvocations))}; + llvm::Type::getInt32Ty(Ctx), MI.second))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -596,12 +594,10 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLSpeculatedIterationsEnable) { - Metadata *Vals[] = { - MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLSpeculatedIterationsNIterations))}; + for (auto &SI : Attrs.SYCLSpeculatedIterationsNIterations) { + Metadata *Vals[] = {MDString::get(Ctx, SI.first), + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), SI.second))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -622,12 +618,9 @@ 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), + UnrollCount(0), UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), MustProgress(false) {} @@ -640,15 +633,12 @@ void LoopAttributes::clear() { GlobalSYCLIVDepInfo.reset(); ArraySYCLIVDepInfo.clear(); SYCLIInterval = 0; - SYCLMaxConcurrencyEnable = false; - SYCLMaxConcurrencyNThreads = 0; + SYCLMaxConcurrencyNThreads.clear(); SYCLLoopCoalesceEnable = false; SYCLLoopCoalesceNLevels = 0; SYCLLoopPipeliningDisable = false; - SYCLMaxInterleavingEnable = false; - SYCLMaxInterleavingNInvocations = 0; - SYCLSpeculatedIterationsEnable = false; - SYCLSpeculatedIterationsNIterations = 0; + SYCLMaxInterleavingNInvocations.clear();; + SYCLSpeculatedIterationsNIterations.clear(); SYCLIntelFPGAVariantCount.clear(); UnrollCount = 0; UnrollAndJamCount = 0; @@ -679,14 +669,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.empty() && Attrs.SYCLLoopCoalesceEnable == false && Attrs.SYCLLoopCoalesceNLevels == 0 && Attrs.SYCLLoopPipeliningDisable == false && - Attrs.SYCLMaxInterleavingEnable == false && - Attrs.SYCLMaxInterleavingNInvocations == 0 && - Attrs.SYCLSpeculatedIterationsEnable == false && - Attrs.SYCLSpeculatedIterationsNIterations == 0 && + Attrs.SYCLMaxInterleavingNInvocations.empty() && + Attrs.SYCLSpeculatedIterationsNIterations.empty() && Attrs.SYCLIntelFPGAVariantCount.empty() && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && @@ -1025,18 +1013,23 @@ 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)) { const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - setSYCLMaxConcurrencyEnable(); - setSYCLMaxConcurrencyNThreads(ArgVal.getSExtValue()); + unsigned int Value = ArgVal.getSExtValue(); + const char *Var = "llvm.loop.max_concurrency.count"; + if (Value == 0) + setSYCLMaxConcurrencyNThreads(Var, 0); + setSYCLMaxConcurrencyNThreads(Var, Value); } if (const auto *IntelFPGALoopCountAvg = @@ -1055,9 +1048,12 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopCoalesce = dyn_cast(A)) { - if (auto *LCE = IntelFPGALoopCoalesce->getNExpr()) - setSYCLLoopCoalesceNLevels( - LCE->getIntegerConstantExpr(Ctx)->getSExtValue()); + const auto *CE = + cast(IntelFPGALoopCoalesce->getNExpr()); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + unsigned int Value = ArgVal.getSExtValue(); + if (Value > 0) + setSYCLLoopCoalesceNLevels(Value); else setSYCLLoopCoalesceEnable(); } @@ -1069,8 +1065,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - setSYCLMaxInterleavingEnable(); - setSYCLMaxInterleavingNInvocations(ArgVal.getSExtValue()); + unsigned int Value = ArgVal.getSExtValue(); + const char *Var = "llvm.loop.max_interleaving.count"; + if (Value == 0) + setSYCLMaxInterleavingNInvocations(Var, 0); + setSYCLMaxInterleavingNInvocations(Var, Value); } if (const auto *IntelFPGASpeculatedIterations = @@ -1078,8 +1077,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - setSYCLSpeculatedIterationsEnable(); - setSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); + const char *Var = "llvm.loop.intel.speculated.iterations.count"; + unsigned int Value = ArgVal.getSExtValue(); + if (Value == 0) + setSYCLSpeculatedIterationsNIterations(Var, 0); + setSYCLSpeculatedIterationsNIterations(Var, Value); } if (isa(A)) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 317972a34ebe..9c1444978e50 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -111,11 +111,9 @@ 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; + /// Value for max_concurrency variant and metadata. + llvm::SmallVector, 2> + SYCLMaxConcurrencyNThreads; /// Value for count variant (min/max/avg) and count metadata. llvm::SmallVector, 2> @@ -130,17 +128,13 @@ 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; + /// Value for max_interleaving variant and metadata. + llvm::SmallVector, 2> + SYCLMaxInterleavingNInvocations; - /// Value for llvm.loop.intel.speculated.iterations.count metadata. - unsigned SYCLSpeculatedIterationsNIterations; + /// Value for speculated.iterations variant and metadata. + llvm::SmallVector, 2> + SYCLSpeculatedIterationsNIterations; /// llvm.unroll. unsigned UnrollCount; @@ -363,14 +357,9 @@ 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. - void setSYCLMaxConcurrencyNThreads(unsigned C) { - StagedAttrs.SYCLMaxConcurrencyNThreads = C; + /// Set variant and value of max_concurrency for the next loop pushed. + void setSYCLMaxConcurrencyNThreads(const char *Var, unsigned int Value) { + StagedAttrs.SYCLMaxConcurrencyNThreads.push_back({Var, Value}); } /// Set flag of loop_coalesce for the next loop pushed. @@ -388,24 +377,14 @@ 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 variant and value of max interleaved invocations for the next loop pushed. + void setSYCLMaxInterleavingNInvocations(const char *Var, unsigned int Value) { + StagedAttrs.SYCLMaxInterleavingNInvocations.push_back({Var, Value}); } - /// Set value of concurrent speculated iterations for the next loop pushed. - void setSYCLSpeculatedIterationsNIterations(unsigned C) { - StagedAttrs.SYCLSpeculatedIterationsNIterations = C; + /// Set variant and value of speculated iterations for the next loop pushed. + void setSYCLSpeculatedIterationsNIterations(const char *Var, unsigned int Value) { + StagedAttrs.SYCLSpeculatedIterationsNIterations.push_back({Var, Value}); } /// Set value of variant and loop count for the next loop pushed. diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 87e820c5ef85..0f777e9cbfde 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -158,10 +158,8 @@ static Attr *handleSYCLIntelFPGAMaxInterleavingAttr(Sema &S, Stmt *St, SYCLIntelFPGALoopCoalesceAttr * Sema::BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E) { - //if (!E) - //return nullptr; - - if (!E->isValueDependent()) { + + if (E && !E->isValueDependent()) { llvm::APSInt ArgVal; ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); if (Res.isInvalid()) @@ -183,12 +181,9 @@ static Attr *handleSYCLIntelFPGALoopCoalesceAttr(Sema &S, Stmt *St, const ParsedAttr &A) { S.CheckDeprecatedSYCLAttributeSpelling(A); - //Expr *E = A.isArgExpr(0) ? A.getArgAsExpr(0) : nullptr; - if (A.isArgExpr(0)) { - Expr *E = A.getArgAsExpr(0); - return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); - } - return new (S.Context) SYCLIntelFPGALoopCoalesceAttr(S.Context, A); + Expr *E = A.isArgExpr(0) ? A.getArgAsExpr(0) : nullptr; + return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); + } SYCLIntelFPGASpeculatedIterationsAttr * From 15844804855d012e9d268c8bea31b8cf2b83d0dd Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 13:18:02 -0700 Subject: [PATCH 08/24] fix format issues Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 49f11b83d989..05a56fb9fc65 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1068,7 +1068,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, unsigned int Value = ArgVal.getSExtValue(); const char *Var = "llvm.loop.max_interleaving.count"; if (Value == 0) - setSYCLMaxInterleavingNInvocations(Var, 0); + setSYCLMaxInterleavingNInvocations(Var, 0); setSYCLMaxInterleavingNInvocations(Var, Value); } From 9b77499b39259a7f1cf0ceec2f75d031db6fcdf6 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 16:03:34 -0700 Subject: [PATCH 09/24] Fix LIT test failures Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 31 ++++++++++--------------------- 1 file changed, 10 insertions(+), 21 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 05a56fb9fc65..25cca6bee339 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1025,11 +1025,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - unsigned int Value = ArgVal.getSExtValue(); const char *Var = "llvm.loop.max_concurrency.count"; - if (Value == 0) - setSYCLMaxConcurrencyNThreads(Var, 0); - setSYCLMaxConcurrencyNThreads(Var, Value); + setSYCLMaxConcurrencyNThreads(Var, ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCountAvg = @@ -1037,25 +1034,23 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGALoopCountAvg->getNTripCount()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - unsigned int Count = ArgVal.getSExtValue(); 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)) { - const auto *CE = - cast(IntelFPGALoopCoalesce->getNExpr()); - llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - unsigned int Value = ArgVal.getSExtValue(); - if (Value > 0) - setSYCLLoopCoalesceNLevels(Value); - else + if (auto *LCE = IntelFPGALoopCoalesce->getNExpr()) { + const auto *CE = cast(LCE); + llvm::APSInt ArgVal = CE->getResultAsAPSInt(); + setSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); + } else { setSYCLLoopCoalesceEnable(); + } } if (isa(A)) @@ -1065,11 +1060,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - unsigned int Value = ArgVal.getSExtValue(); const char *Var = "llvm.loop.max_interleaving.count"; - if (Value == 0) - setSYCLMaxInterleavingNInvocations(Var, 0); - setSYCLMaxInterleavingNInvocations(Var, Value); + setSYCLMaxInterleavingNInvocations(Var, ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = @@ -1078,10 +1070,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); const char *Var = "llvm.loop.intel.speculated.iterations.count"; - unsigned int Value = ArgVal.getSExtValue(); - if (Value == 0) - setSYCLSpeculatedIterationsNIterations(Var, 0); - setSYCLSpeculatedIterationsNIterations(Var, Value); + setSYCLSpeculatedIterationsNIterations(Var, ArgVal.getSExtValue()); } if (isa(A)) From f00efdabd5ff91e8d6e6726622f65f315e88236c Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 19 Jul 2021 16:13:20 -0700 Subject: [PATCH 10/24] fix format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 17 ++++++++--------- clang/lib/CodeGen/CGLoopInfo.h | 6 ++++-- clang/lib/Sema/SemaStmtAttr.cpp | 2 -- 3 files changed, 12 insertions(+), 13 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 25cca6bee339..6d776fc05fa1 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -620,10 +620,10 @@ LoopAttributes::LoopAttributes(bool IsParallel) VectorizeScalable(LoopAttributes::Unspecified), InterleaveCount(0), SYCLIInterval(0), SYCLLoopCoalesceEnable(false), SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false), - 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; @@ -637,7 +637,7 @@ void LoopAttributes::clear() { SYCLLoopCoalesceEnable = false; SYCLLoopCoalesceNLevels = 0; SYCLLoopPipeliningDisable = false; - SYCLMaxInterleavingNInvocations.clear();; + SYCLMaxInterleavingNInvocations.clear(); SYCLSpeculatedIterationsNIterations.clear(); SYCLIntelFPGAVariantCount.clear(); UnrollCount = 0; @@ -1014,8 +1014,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGAII = dyn_cast(A)) { - const auto *CE = - cast(IntelFPGAII->getIntervalExpr()); + const auto *CE = cast(IntelFPGAII->getIntervalExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); setSYCLIInterval(ArgVal.getSExtValue()); } @@ -1045,9 +1044,9 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopCoalesce = dyn_cast(A)) { if (auto *LCE = IntelFPGALoopCoalesce->getNExpr()) { - const auto *CE = cast(LCE); + const auto *CE = cast(LCE); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - setSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); + setSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); } else { setSYCLLoopCoalesceEnable(); } diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 9c1444978e50..483b1824fb58 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -377,13 +377,15 @@ class LoopInfoStack { StagedAttrs.SYCLLoopPipeliningDisable = true; } - /// Set variant and value of max interleaved invocations for the next loop pushed. + /// Set variant and value of max interleaved invocations for the next loop + /// pushed. void setSYCLMaxInterleavingNInvocations(const char *Var, unsigned int Value) { StagedAttrs.SYCLMaxInterleavingNInvocations.push_back({Var, Value}); } /// Set variant and value of speculated iterations for the next loop pushed. - void setSYCLSpeculatedIterationsNIterations(const char *Var, unsigned int Value) { + void setSYCLSpeculatedIterationsNIterations(const char *Var, + unsigned int Value) { StagedAttrs.SYCLSpeculatedIterationsNIterations.push_back({Var, Value}); } diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 0f777e9cbfde..623bd6c6b546 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -158,7 +158,6 @@ static Attr *handleSYCLIntelFPGAMaxInterleavingAttr(Sema &S, Stmt *St, SYCLIntelFPGALoopCoalesceAttr * Sema::BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E) { - if (E && !E->isValueDependent()) { llvm::APSInt ArgVal; ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal); @@ -183,7 +182,6 @@ static Attr *handleSYCLIntelFPGALoopCoalesceAttr(Sema &S, Stmt *St, Expr *E = A.isArgExpr(0) ? A.getArgAsExpr(0) : nullptr; return S.BuildSYCLIntelFPGALoopCoalesceAttr(A, E); - } SYCLIntelFPGASpeculatedIterationsAttr * From e15747b8c3ce27942fe2affa082b3dc19033ca95 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 08:39:35 -0700 Subject: [PATCH 11/24] address review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 30 +++++++++++++++--------------- clang/lib/CodeGen/CGLoopInfo.h | 28 ++++++++++++++-------------- 2 files changed, 29 insertions(+), 29 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 6d776fc05fa1..1264aacdeb1e 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -552,7 +552,7 @@ MDNode *LoopInfo::createMetadata( } // Setting max_concurrency attribute with number of threads - for (auto &MC : Attrs.SYCLMaxConcurrencyNThreads) { + for (const auto &MC : Attrs.SYCLMaxConcurrencyNThreads) { Metadata *Vals[] = {MDString::get(Ctx, MC.first), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), MC.second))}; @@ -581,7 +581,7 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - for (auto &MI : Attrs.SYCLMaxInterleavingNInvocations) { + for (const auto &MI : Attrs.SYCLMaxInterleavingNInvocations) { Metadata *Vals[] = {MDString::get(Ctx, MI.first), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), MI.second))}; @@ -594,14 +594,14 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - for (auto &SI : Attrs.SYCLSpeculatedIterationsNIterations) { + for (const auto &SI : Attrs.SYCLSpeculatedIterationsNIterations) { Metadata *Vals[] = {MDString::get(Ctx, SI.first), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), SI.second))}; 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))}; @@ -1016,7 +1016,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAII->getIntervalExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - setSYCLIInterval(ArgVal.getSExtValue()); + addSYCLIInterval(ArgVal.getSExtValue()); } if (const auto *IntelFPGAMaxConcurrency = @@ -1024,8 +1024,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - const char *Var = "llvm.loop.max_concurrency.count"; - setSYCLMaxConcurrencyNThreads(Var, ArgVal.getSExtValue()); + addSYCLMaxConcurrencyNThreads("llvm.loop.max_concurrency.count", + ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCountAvg = @@ -1038,17 +1038,17 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" : "llvm.loop.intel.loopcount_avg"; - setSYCLIntelFPGAVariantCount(Var, ArgVal.getSExtValue()); + addSYCLIntelFPGAVariantCount(Var, ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = dyn_cast(A)) { - if (auto *LCE = IntelFPGALoopCoalesce->getNExpr()) { + if (const auto *LCE = IntelFPGALoopCoalesce->getNExpr()) { const auto *CE = cast(LCE); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - setSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); + addSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); } else { - setSYCLLoopCoalesceEnable(); + addSYCLLoopCoalesceEnable(); } } @@ -1059,8 +1059,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - const char *Var = "llvm.loop.max_interleaving.count"; - setSYCLMaxInterleavingNInvocations(Var, ArgVal.getSExtValue()); + addSYCLMaxInterleavingNInvocations("llvm.loop.max_interleaving.count", + ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = @@ -1068,8 +1068,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - const char *Var = "llvm.loop.intel.speculated.iterations.count"; - setSYCLSpeculatedIterationsNIterations(Var, ArgVal.getSExtValue()); + addSYCLSpeculatedIterationsNIterations( + "llvm.loop.intel.speculated.iterations.count", ArgVal.getSExtValue()); } if (isa(A)) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 483b1824fb58..f0ab63cb2c8e 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -354,21 +354,21 @@ class LoopInfoStack { void addIVDepMetadata(const ValueDecl *Array, llvm::Instruction *GEP); - /// Set value of an initiation interval for the next loop pushed. - void setSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } + /// Add value of an initiation interval for the next loop pushed. + void addSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } - /// Set variant and value of max_concurrency for the next loop pushed. - void setSYCLMaxConcurrencyNThreads(const char *Var, unsigned int Value) { + /// Add variant and value of max_concurrency for the next loop pushed. + void addSYCLMaxConcurrencyNThreads(const char *Var, unsigned int Value) { StagedAttrs.SYCLMaxConcurrencyNThreads.push_back({Var, Value}); } - /// Set flag of loop_coalesce for the next loop pushed. - void setSYCLLoopCoalesceEnable() { + /// Add flag of loop_coalesce for the next loop pushed. + void addSYCLLoopCoalesceEnable() { StagedAttrs.SYCLLoopCoalesceEnable = true; } - /// Set value of coalesced levels for the next loop pushed. - void setSYCLLoopCoalesceNLevels(unsigned C) { + /// Add value of coalesced levels for the next loop pushed. + void addSYCLLoopCoalesceNLevels(unsigned C) { StagedAttrs.SYCLLoopCoalesceNLevels = C; } @@ -377,20 +377,20 @@ class LoopInfoStack { StagedAttrs.SYCLLoopPipeliningDisable = true; } - /// Set variant and value of max interleaved invocations for the next loop + /// Add variant and value of max interleaved invocations for the next loop /// pushed. - void setSYCLMaxInterleavingNInvocations(const char *Var, unsigned int Value) { + void addSYCLMaxInterleavingNInvocations(const char *Var, unsigned int Value) { StagedAttrs.SYCLMaxInterleavingNInvocations.push_back({Var, Value}); } - /// Set variant and value of speculated iterations for the next loop pushed. - void setSYCLSpeculatedIterationsNIterations(const char *Var, + /// Add variant and value of speculated iterations for the next loop pushed. + void addSYCLSpeculatedIterationsNIterations(const char *Var, unsigned int Value) { StagedAttrs.SYCLSpeculatedIterationsNIterations.push_back({Var, Value}); } - /// Set value of variant and loop count for the next loop pushed. - void setSYCLIntelFPGAVariantCount(const char *Var, unsigned int Count) { + /// Add value of variant and loop count for the next loop pushed. + void addSYCLIntelFPGAVariantCount(const char *Var, unsigned int Count) { StagedAttrs.SYCLIntelFPGAVariantCount.push_back({Var, Count}); } From 980677467b10407ae3867a6fd12b7261a7454df8 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 08:45:34 -0700 Subject: [PATCH 12/24] fix format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 1264aacdeb1e..993250347051 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1025,7 +1025,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); addSYCLMaxConcurrencyNThreads("llvm.loop.max_concurrency.count", - ArgVal.getSExtValue()); + ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCountAvg = @@ -1060,7 +1060,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); addSYCLMaxInterleavingNInvocations("llvm.loop.max_interleaving.count", - ArgVal.getSExtValue()); + ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = @@ -1069,7 +1069,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); addSYCLSpeculatedIterationsNIterations( - "llvm.loop.intel.speculated.iterations.count", ArgVal.getSExtValue()); + "llvm.loop.intel.speculated.iterations.count", ArgVal.getSExtValue()); } if (isa(A)) From 309ea84057fe9fe55e4fcca88e22d533e01e46da Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 09:59:01 -0700 Subject: [PATCH 13/24] add more template instantiation test coverage for the improvements in SemaStmtAttr.cpp Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/intel-fpga-loops.cpp | 231 ++++++++++++++++++++++- 1 file changed, 224 insertions(+), 7 deletions(-) diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 5d2ad7642314..afe634bae5c5 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -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,160 @@ 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; +} + +// Template instantiation test coverages for Intel FPGA loop attribute max_concurrency. +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; +} + +// Template instantiation test coverages for Intel FPGA loop attribute max_concurrency. +void check_loop_count_expression() { + int a[10]; + + // Test that checks expression is not a constant expression. + int foo; // expected-note 3{{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_avg(foo + 1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + // 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; + + // 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_min(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_avg(bar + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; + + constexpr int bar1 = 0; + [[intel::loop_count_max(bar1 + 2)]] for (int i = 0; i != 10; ++i) // OK + a[i] = 0; + + constexpr int bar2 = 0; + [[intel::loop_count_min(bar2 + 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 +692,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; From 800c59ae6d9ed84bc8fc055a340972cdf84977a5 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 10:30:13 -0700 Subject: [PATCH 14/24] add new codegen tests Signed-off-by: Soumi Manna --- clang/test/CodeGenSYCL/intel-fpga-loops.cpp | 21 +++++++++++++++++---- 1 file changed, 17 insertions(+), 4 deletions(-) diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index a0dfafdbcd92..06c0d80d7988 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; From 81f39ed91fcb2bed98586cb2d5f654d22c445c5c Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 11:55:47 -0700 Subject: [PATCH 15/24] Address @aaron's review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 25 +++++++++------------ clang/lib/CodeGen/CGLoopInfo.h | 38 ++++++++++++++------------------ 2 files changed, 28 insertions(+), 35 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 993250347051..5f37db924317 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -553,9 +553,9 @@ MDNode *LoopInfo::createMetadata( // Setting max_concurrency attribute with number of threads for (const auto &MC : Attrs.SYCLMaxConcurrencyNThreads) { - Metadata *Vals[] = {MDString::get(Ctx, MC.first), + Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_concurrency.count"), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), MC.second))}; + llvm::Type::getInt32Ty(Ctx), MC))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -582,9 +582,9 @@ MDNode *LoopInfo::createMetadata( } for (const auto &MI : Attrs.SYCLMaxInterleavingNInvocations) { - Metadata *Vals[] = {MDString::get(Ctx, MI.first), + Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), MI.second))}; + llvm::Type::getInt32Ty(Ctx), MI))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -595,9 +595,9 @@ MDNode *LoopInfo::createMetadata( } for (const auto &SI : Attrs.SYCLSpeculatedIterationsNIterations) { - Metadata *Vals[] = {MDString::get(Ctx, SI.first), + Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), SI.second))}; + llvm::Type::getInt32Ty(Ctx), SI))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -1016,7 +1016,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAII->getIntervalExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLIInterval(ArgVal.getSExtValue()); + setSYCLIInterval(ArgVal.getSExtValue()); } if (const auto *IntelFPGAMaxConcurrency = @@ -1024,8 +1024,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLMaxConcurrencyNThreads("llvm.loop.max_concurrency.count", - ArgVal.getSExtValue()); + addSYCLMaxConcurrencyNThreads(ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCountAvg = @@ -1048,7 +1047,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, llvm::APSInt ArgVal = CE->getResultAsAPSInt(); addSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); } else { - addSYCLLoopCoalesceEnable(); + setSYCLLoopCoalesceEnable(); } } @@ -1059,8 +1058,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLMaxInterleavingNInvocations("llvm.loop.max_interleaving.count", - ArgVal.getSExtValue()); + addSYCLMaxInterleavingNInvocations(ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = @@ -1068,8 +1066,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLSpeculatedIterationsNIterations( - "llvm.loop.intel.speculated.iterations.count", ArgVal.getSExtValue()); + addSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); } if (isa(A)) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index f0ab63cb2c8e..a54d30183516 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -112,8 +112,7 @@ struct LoopAttributes { unsigned SYCLIInterval; /// Value for max_concurrency variant and metadata. - llvm::SmallVector, 2> - SYCLMaxConcurrencyNThreads; + llvm::SmallVector SYCLMaxConcurrencyNThreads; /// Value for count variant (min/max/avg) and count metadata. llvm::SmallVector, 2> @@ -129,12 +128,10 @@ struct LoopAttributes { bool SYCLLoopPipeliningDisable; /// Value for max_interleaving variant and metadata. - llvm::SmallVector, 2> - SYCLMaxInterleavingNInvocations; + llvm::SmallVector SYCLMaxInterleavingNInvocations; /// Value for speculated.iterations variant and metadata. - llvm::SmallVector, 2> - SYCLSpeculatedIterationsNIterations; + llvm::SmallVector SYCLSpeculatedIterationsNIterations; /// llvm.unroll. unsigned UnrollCount; @@ -354,22 +351,22 @@ class LoopInfoStack { void addIVDepMetadata(const ValueDecl *Array, llvm::Instruction *GEP); - /// Add value of an initiation interval for the next loop pushed. - void addSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } + /// Set value of an initiation interval for the next loop pushed. + void setSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } - /// Add variant and value of max_concurrency for the next loop pushed. - void addSYCLMaxConcurrencyNThreads(const char *Var, unsigned int Value) { - StagedAttrs.SYCLMaxConcurrencyNThreads.push_back({Var, Value}); + /// Add value of max_concurrency for the next loop pushed. + void addSYCLMaxConcurrencyNThreads(unsigned int Value) { + StagedAttrs.SYCLMaxConcurrencyNThreads.push_back(Value); } - /// Add flag of loop_coalesce for the next loop pushed. - void addSYCLLoopCoalesceEnable() { + /// Set flag of loop_coalesce for the next loop pushed. + void setSYCLLoopCoalesceEnable() { StagedAttrs.SYCLLoopCoalesceEnable = true; } /// Add value of coalesced levels for the next loop pushed. - void addSYCLLoopCoalesceNLevels(unsigned C) { - StagedAttrs.SYCLLoopCoalesceNLevels = C; + void addSYCLLoopCoalesceNLevels(unsigned int Value) { + StagedAttrs.SYCLLoopCoalesceNLevels = Value; } /// Set flag of disable_loop_pipelining for the next loop pushed. @@ -379,14 +376,13 @@ class LoopInfoStack { /// Add variant and value of max interleaved invocations for the next loop /// pushed. - void addSYCLMaxInterleavingNInvocations(const char *Var, unsigned int Value) { - StagedAttrs.SYCLMaxInterleavingNInvocations.push_back({Var, Value}); + void addSYCLMaxInterleavingNInvocations(unsigned int Value) { + StagedAttrs.SYCLMaxInterleavingNInvocations.push_back(Value); } - /// Add variant and value of speculated iterations for the next loop pushed. - void addSYCLSpeculatedIterationsNIterations(const char *Var, - unsigned int Value) { - StagedAttrs.SYCLSpeculatedIterationsNIterations.push_back({Var, Value}); + /// Add value of speculated iterations for the next loop pushed. + void addSYCLSpeculatedIterationsNIterations(unsigned int Value) { + StagedAttrs.SYCLSpeculatedIterationsNIterations.push_back(Value); } /// Add value of variant and loop count for the next loop pushed. From 0aac37e8e0b1b085e632983f93f8180a3551ef70 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 13:42:47 -0700 Subject: [PATCH 16/24] address review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 37 ++++++++++++++++---------------- clang/lib/CodeGen/CGLoopInfo.h | 34 ++++++++++++++--------------- 2 files changed, 36 insertions(+), 35 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 5f37db924317..49fd2e96af6e 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -552,10 +552,10 @@ MDNode *LoopInfo::createMetadata( } // Setting max_concurrency attribute with number of threads - for (const auto &MC : Attrs.SYCLMaxConcurrencyNThreads) { + if (Attrs.SYCLMaxConcurrencyNThreads.hasValue()) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_concurrency.count"), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), MC))}; + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxConcurrencyNThreads))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -581,10 +581,10 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - for (const auto &MI : Attrs.SYCLMaxInterleavingNInvocations) { + if (Attrs.SYCLMaxInterleavingNInvocations.hasValue()) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), MI))}; + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxInterleavingNInvocations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -594,10 +594,11 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - for (const auto &SI : Attrs.SYCLSpeculatedIterationsNIterations) { - Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), SI))}; + if (Attrs.SYCLSpeculatedIterationsNIterations.hasValue()) { + Metadata *Vals[] = { + MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLSpeculatedIterationsNIterations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -633,12 +634,12 @@ void LoopAttributes::clear() { GlobalSYCLIVDepInfo.reset(); ArraySYCLIVDepInfo.clear(); SYCLIInterval = 0; - SYCLMaxConcurrencyNThreads.clear(); + SYCLMaxConcurrencyNThreads.reset(); SYCLLoopCoalesceEnable = false; SYCLLoopCoalesceNLevels = 0; SYCLLoopPipeliningDisable = false; - SYCLMaxInterleavingNInvocations.clear(); - SYCLSpeculatedIterationsNIterations.clear(); + SYCLMaxInterleavingNInvocations.reset(); + SYCLSpeculatedIterationsNIterations.reset(); SYCLIntelFPGAVariantCount.clear(); UnrollCount = 0; UnrollAndJamCount = 0; @@ -669,12 +670,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.SYCLMaxConcurrencyNThreads.empty() && + !Attrs.SYCLMaxConcurrencyNThreads.hasValue() && Attrs.SYCLLoopCoalesceEnable == false && Attrs.SYCLLoopCoalesceNLevels == 0 && Attrs.SYCLLoopPipeliningDisable == false && - Attrs.SYCLMaxInterleavingNInvocations.empty() && - Attrs.SYCLSpeculatedIterationsNIterations.empty() && + !Attrs.SYCLMaxInterleavingNInvocations.hasValue() && + !Attrs.SYCLSpeculatedIterationsNIterations.hasValue() && Attrs.SYCLIntelFPGAVariantCount.empty() && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && @@ -1037,7 +1038,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" : "llvm.loop.intel.loopcount_avg"; - addSYCLIntelFPGAVariantCount(Var, ArgVal.getSExtValue()); + setSYCLIntelFPGAVariantCount(Var, ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = @@ -1045,7 +1046,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *LCE = IntelFPGALoopCoalesce->getNExpr()) { const auto *CE = cast(LCE); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); + setSYCLLoopCoalesceNLevels(ArgVal.getSExtValue()); } else { setSYCLLoopCoalesceEnable(); } diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index a54d30183516..77cc6dc1d8de 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -111,8 +111,8 @@ struct LoopAttributes { /// Value for llvm.loop.ii.count metadata. unsigned SYCLIInterval; - /// Value for max_concurrency variant and metadata. - llvm::SmallVector SYCLMaxConcurrencyNThreads; + /// Value for llvm.loop.max_concurrency.count metadata. + llvm::Optional SYCLMaxConcurrencyNThreads; /// Value for count variant (min/max/avg) and count metadata. llvm::SmallVector, 2> @@ -127,11 +127,11 @@ struct LoopAttributes { /// Flag for llvm.loop.intel.pipelining.enable, i32 0 metadata. bool SYCLLoopPipeliningDisable; - /// Value for max_interleaving variant and metadata. - llvm::SmallVector SYCLMaxInterleavingNInvocations; + /// Value for llvm.loop.max_interleaving.count metadata. + llvm::Optional SYCLMaxInterleavingNInvocations; - /// Value for speculated.iterations variant and metadata. - llvm::SmallVector SYCLSpeculatedIterationsNIterations; + /// Value for llvm.loop.intel.speculated.iterations.count metadata. + llvm::Optional SYCLSpeculatedIterationsNIterations; /// llvm.unroll. unsigned UnrollCount; @@ -355,8 +355,8 @@ class LoopInfoStack { void setSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } /// Add value of max_concurrency for the next loop pushed. - void addSYCLMaxConcurrencyNThreads(unsigned int Value) { - StagedAttrs.SYCLMaxConcurrencyNThreads.push_back(Value); + void addSYCLMaxConcurrencyNThreads(unsigned C) { + StagedAttrs.SYCLMaxConcurrencyNThreads = C; } /// Set flag of loop_coalesce for the next loop pushed. @@ -364,9 +364,9 @@ class LoopInfoStack { StagedAttrs.SYCLLoopCoalesceEnable = true; } - /// Add value of coalesced levels for the next loop pushed. - void addSYCLLoopCoalesceNLevels(unsigned int Value) { - StagedAttrs.SYCLLoopCoalesceNLevels = Value; + /// Set value of coalesced levels for the next loop pushed. + void setSYCLLoopCoalesceNLevels(unsigned C) { + StagedAttrs.SYCLLoopCoalesceNLevels = C; } /// Set flag of disable_loop_pipelining for the next loop pushed. @@ -376,17 +376,17 @@ class LoopInfoStack { /// Add variant and value of max interleaved invocations for the next loop /// pushed. - void addSYCLMaxInterleavingNInvocations(unsigned int Value) { - StagedAttrs.SYCLMaxInterleavingNInvocations.push_back(Value); + void addSYCLMaxInterleavingNInvocations(unsigned C) { + StagedAttrs.SYCLMaxInterleavingNInvocations = C; } /// Add value of speculated iterations for the next loop pushed. - void addSYCLSpeculatedIterationsNIterations(unsigned int Value) { - StagedAttrs.SYCLSpeculatedIterationsNIterations.push_back(Value); + void addSYCLSpeculatedIterationsNIterations(unsigned C) { + StagedAttrs.SYCLSpeculatedIterationsNIterations = C; } - /// Add value of variant and loop count for the next loop pushed. - void addSYCLIntelFPGAVariantCount(const char *Var, unsigned int Count) { + /// Set value of variant and loop count for the next loop pushed. + void setSYCLIntelFPGAVariantCount(const char *Var, unsigned int Count) { StagedAttrs.SYCLIntelFPGAVariantCount.push_back({Var, Count}); } From 821269a00ca8f0af2b670f2a3bb9d9cf0ee0b8f7 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 13:50:42 -0700 Subject: [PATCH 17/24] fix format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 15 +++++++++------ clang/lib/CodeGen/CGLoopInfo.h | 5 ++--- 2 files changed, 11 insertions(+), 9 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 49fd2e96af6e..a25355ac0ced 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -553,9 +553,10 @@ MDNode *LoopInfo::createMetadata( // Setting max_concurrency attribute with number of threads if (Attrs.SYCLMaxConcurrencyNThreads.hasValue()) { - Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_concurrency.count"), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *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)); } @@ -583,8 +584,9 @@ MDNode *LoopInfo::createMetadata( if (Attrs.SYCLMaxInterleavingNInvocations.hasValue()) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxInterleavingNInvocations))}; + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), + *Attrs.SYCLMaxInterleavingNInvocations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -598,7 +600,8 @@ MDNode *LoopInfo::createMetadata( Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLSpeculatedIterationsNIterations))}; + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), + *Attrs.SYCLSpeculatedIterationsNIterations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 77cc6dc1d8de..b24ba83b5c06 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -364,7 +364,7 @@ class LoopInfoStack { StagedAttrs.SYCLLoopCoalesceEnable = true; } - /// Set value of coalesced levels for the next loop pushed. + /// Set value of coalesced levels for the next loop pushed. void setSYCLLoopCoalesceNLevels(unsigned C) { StagedAttrs.SYCLLoopCoalesceNLevels = C; } @@ -374,8 +374,7 @@ class LoopInfoStack { StagedAttrs.SYCLLoopPipeliningDisable = true; } - /// Add variant and value of max interleaved invocations for the next loop - /// pushed. + /// Add value of max interleaved invocations for the next loop pushed. void addSYCLMaxInterleavingNInvocations(unsigned C) { StagedAttrs.SYCLMaxInterleavingNInvocations = C; } From af627626e30da1cbf8c3a1f1c43f32506c036bf5 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 14:02:40 -0700 Subject: [PATCH 18/24] fix format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index a25355ac0ced..fa9e76bcc43a 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -584,7 +584,7 @@ MDNode *LoopInfo::createMetadata( if (Attrs.SYCLMaxInterleavingNInvocations.hasValue()) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), - ConstantAsMetadata::get(ConstantInt::get( + ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxInterleavingNInvocations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); @@ -600,7 +600,7 @@ MDNode *LoopInfo::createMetadata( Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLSpeculatedIterationsNIterations))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } From e73277a425191da2799e557c64469932fbe7e9f1 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 14:06:22 -0700 Subject: [PATCH 19/24] update codegen Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index fa9e76bcc43a..47785a8aa7ef 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -556,7 +556,7 @@ MDNode *LoopInfo::createMetadata( Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.max_concurrency.count"), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxConcurrencyNThreads))}; + llvm::Type::getInt32Ty(Ctx), Attrs.SYCLMaxConcurrencyNThreads))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -586,7 +586,7 @@ MDNode *LoopInfo::createMetadata( 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)); } @@ -601,7 +601,7 @@ MDNode *LoopInfo::createMetadata( 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)); } From eefabb8d4a2d416eed5398fd50ecf49dfddc581e Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 14:21:33 -0700 Subject: [PATCH 20/24] fix codegen errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 47785a8aa7ef..fa9e76bcc43a 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -556,7 +556,7 @@ MDNode *LoopInfo::createMetadata( Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.max_concurrency.count"), ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), Attrs.SYCLMaxConcurrencyNThreads))}; + llvm::Type::getInt32Ty(Ctx), *Attrs.SYCLMaxConcurrencyNThreads))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } @@ -586,7 +586,7 @@ MDNode *LoopInfo::createMetadata( 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)); } @@ -601,7 +601,7 @@ MDNode *LoopInfo::createMetadata( 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)); } From 0acb86bff9b132a0d1a8c09d4df7c77732758e91 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 22 Jul 2021 19:07:10 -0700 Subject: [PATCH 21/24] fix test and remove redundant comments Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/intel-fpga-loops.cpp | 24 ++---------------------- 1 file changed, 2 insertions(+), 22 deletions(-) diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index afe634bae5c5..af15a5799ac6 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -571,7 +571,6 @@ void check_initiation_interval_expression() { a[i] = 0; } -// Template instantiation test coverages for Intel FPGA loop attribute max_concurrency. void check_speculated_iterations_expression() { int a[10]; // Test that checks expression is not a constant expression. @@ -604,38 +603,19 @@ void check_loop_coalesce_expression() { a[i] = 0; } -// Template instantiation test coverages for Intel FPGA loop attribute max_concurrency. void check_loop_count_expression() { int a[10]; // Test that checks expression is not a constant expression. - int foo; // expected-note 3{{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_avg(foo + 1)]] for (int i = 0; i != 10; ++i) - a[i] = 0; - + 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; - // 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_min(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_avg(bar + 2)]] for (int i = 0; i != 10; ++i) // OK - a[i] = 0; - - constexpr int bar1 = 0; - [[intel::loop_count_max(bar1 + 2)]] for (int i = 0; i != 10; ++i) // OK - a[i] = 0; - - constexpr int bar2 = 0; - [[intel::loop_count_min(bar2 + 2)]] for (int i = 0; i != 10; ++i) // OK + [[intel::loop_count_max(bar + 2)]] for (int i = 0; i != 10; ++i) // OK a[i] = 0; } From 4a0bb25c86d8c3b434eb8818bb78467875f33527 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 23 Jul 2021 04:32:52 -0700 Subject: [PATCH 22/24] address review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 18 +++++++++--------- clang/lib/CodeGen/CGLoopInfo.h | 13 +++++++------ 2 files changed, 16 insertions(+), 15 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index fa9e76bcc43a..5083b2bd757a 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -552,7 +552,7 @@ MDNode *LoopInfo::createMetadata( } // Setting max_concurrency attribute with number of threads - if (Attrs.SYCLMaxConcurrencyNThreads.hasValue()) { + if (Attrs.SYCLMaxConcurrencyNThreads) { Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.max_concurrency.count"), ConstantAsMetadata::get(ConstantInt::get( @@ -582,7 +582,7 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLMaxInterleavingNInvocations.hasValue()) { + if (Attrs.SYCLMaxInterleavingNInvocations) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.max_interleaving.count"), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), @@ -596,7 +596,7 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLSpeculatedIterationsNIterations.hasValue()) { + if (Attrs.SYCLSpeculatedIterationsNIterations) { Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"), ConstantAsMetadata::get( @@ -673,12 +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.SYCLMaxConcurrencyNThreads.hasValue() && + !Attrs.SYCLMaxConcurrencyNThreads && Attrs.SYCLLoopCoalesceEnable == false && Attrs.SYCLLoopCoalesceNLevels == 0 && Attrs.SYCLLoopPipeliningDisable == false && - !Attrs.SYCLMaxInterleavingNInvocations.hasValue() && - !Attrs.SYCLSpeculatedIterationsNIterations.hasValue() && + !Attrs.SYCLMaxInterleavingNInvocations && + !Attrs.SYCLSpeculatedIterationsNIterations && Attrs.SYCLIntelFPGAVariantCount.empty() && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && @@ -1028,7 +1028,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGAMaxConcurrency->getNThreadsExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLMaxConcurrencyNThreads(ArgVal.getSExtValue()); + setSYCLMaxConcurrencyNThreads(ArgVal.getSExtValue()); } if (const auto *IntelFPGALoopCountAvg = @@ -1062,7 +1062,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { const auto *CE = cast(IntelFPGAMaxInterleaving->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLMaxInterleavingNInvocations(ArgVal.getSExtValue()); + setSYCLMaxInterleavingNInvocations(ArgVal.getSExtValue()); } if (const auto *IntelFPGASpeculatedIterations = @@ -1070,7 +1070,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const auto *CE = cast(IntelFPGASpeculatedIterations->getNExpr()); llvm::APSInt ArgVal = CE->getResultAsAPSInt(); - addSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); + setSYCLSpeculatedIterationsNIterations(ArgVal.getSExtValue()); } if (isa(A)) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index b24ba83b5c06..1d86537470a2 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -19,6 +19,7 @@ #include "llvm/IR/DebugLoc.h" #include "llvm/IR/Value.h" #include "llvm/Support/Compiler.h" +#include "llvm/ADT/Optional.h" namespace llvm { class BasicBlock; @@ -354,8 +355,8 @@ class LoopInfoStack { /// Set value of an initiation interval for the next loop pushed. void setSYCLIInterval(unsigned C) { StagedAttrs.SYCLIInterval = C; } - /// Add value of max_concurrency for the next loop pushed. - void addSYCLMaxConcurrencyNThreads(unsigned C) { + /// Set value of max_concurrency for the next loop pushed. + void setSYCLMaxConcurrencyNThreads(unsigned C) { StagedAttrs.SYCLMaxConcurrencyNThreads = C; } @@ -374,13 +375,13 @@ class LoopInfoStack { StagedAttrs.SYCLLoopPipeliningDisable = true; } - /// Add value of max interleaved invocations for the next loop pushed. - void addSYCLMaxInterleavingNInvocations(unsigned C) { + /// Set value of max interleaved invocations for the next loop pushed. + void setSYCLMaxInterleavingNInvocations(unsigned C) { StagedAttrs.SYCLMaxInterleavingNInvocations = C; } - /// Add value of speculated iterations for the next loop pushed. - void addSYCLSpeculatedIterationsNIterations(unsigned C) { + /// Set value of speculated iterations for the next loop pushed. + void setSYCLSpeculatedIterationsNIterations(unsigned C) { StagedAttrs.SYCLSpeculatedIterationsNIterations = C; } From cfa9a74858fbb923c4a761792f45b80ad589d48b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 23 Jul 2021 04:38:35 -0700 Subject: [PATCH 23/24] Fix format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 1d86537470a2..f48cfb248c3c 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -15,11 +15,11 @@ #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" #include "llvm/Support/Compiler.h" -#include "llvm/ADT/Optional.h" namespace llvm { class BasicBlock; From 4fc951f2f30edaf7b62be9b6edd4b2eea7b320b4 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 23 Jul 2021 10:03:08 -0700 Subject: [PATCH 24/24] remove default value Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index b90bf8ad9759..e4e7d7c66789 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2200,8 +2200,7 @@ class Sema final { BuildSYCLIntelFPGASpeculatedIterationsAttr(const AttributeCommonInfo &CI, Expr *E); SYCLIntelFPGALoopCoalesceAttr * - BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, - Expr *E = nullptr); + BuildSYCLIntelFPGALoopCoalesceAttr(const AttributeCommonInfo &CI, Expr *E); bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc);