From 488dfb9d9e39b5e74c77e9570d39b3445e3c2373 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 29 Mar 2021 08:02:23 -0400 Subject: [PATCH 01/28] [SYCL] Implementation of loop attribute control_avg. --- clang/include/clang/Basic/Attr.td | 11 +++++++++++ clang/include/clang/Basic/AttrDocs.td | 23 +++++++++++++++++++++++ 2 files changed, 34 insertions(+) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 600f94f19ab18..680c71ea0641e 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1904,6 +1904,17 @@ def SYCLIntelFPGADisableLoopPipelining : StmtAttr { let Documentation = [SYCLIntelFPGADisableLoopPipeliningAttrDocs]; } +def SYCLIntelFPGALoopControlAvg : StmtAttr { + let Spellings = [CXX11<"intelfpga","loop_control_avg">, + CXX11<"intel","loop_control_avg">]; + let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt], + ErrorDiag, "'for', 'while', and 'do' statements">; + let Args = [ExprArgument<"N">]; + let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; + let HasCustomTypeTransform = 1; + let Documentation = [SYCLIntelFPGALoopControlAvgDocs]; +} + def SYCLIntelFPGAMaxInterleaving : StmtAttr { let Spellings = [CXX11<"intelfpga","max_interleaving">, CXX11<"intel","max_interleaving">]; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 8a3b2e207b58d..b32a6aad598f7 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2877,6 +2877,29 @@ or ivdep. }]; } +def SYCLIntelFPGADisableLoopPipeliningAttrDocs : Documentation { + let Category = DocCatVariable; + let Heading = "intel::loo_control_avg"; + let Content = [{ +This attribute applies to a loop. Its purpose is to pass to the optimizer the +average trip count of a loop. This will allow the optimizer to figure out if +the optimization is beneficial. + +.. code-block:: c++ + + void foo() { + int a[10]; + [[intel::loop_control_avg(40)] for (int i = 0; i < 10; ++i) a[i] = 0; + } + + template + void bar() { + [[intel::loop_control_avg(N)]] for(;;) { } + } + + }]; +} + def SYCLIntelFPGAMaxInterleavingAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "intel::max_interleaving"; From 0324f5fafe3ff2be7fa2965520e070a32355046c Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 29 Mar 2021 10:03:05 -0400 Subject: [PATCH 02/28] Fixing Attr name Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Basic/AttrDocs.td | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 680c71ea0641e..6f3adf377b1ff 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1912,7 +1912,7 @@ def SYCLIntelFPGALoopControlAvg : StmtAttr { let Args = [ExprArgument<"N">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let HasCustomTypeTransform = 1; - let Documentation = [SYCLIntelFPGALoopControlAvgDocs]; + let Documentation = [SYCLIntelFPGALoopControlAvgAttrDocs]; } def SYCLIntelFPGAMaxInterleaving : StmtAttr { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index b32a6aad598f7..e79f5741d90c5 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2877,7 +2877,7 @@ or ivdep. }]; } -def SYCLIntelFPGADisableLoopPipeliningAttrDocs : Documentation { +def SYCLIntelFPGALoopControlAvgAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "intel::loo_control_avg"; let Content = [{ From a3a861373fcd9eb31a3acc6d259f2e254c61a382 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 9 Apr 2021 12:14:21 -0400 Subject: [PATCH 03/28] Complete implementation Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Sema/Sema.h | 4 +++- clang/lib/CodeGen/CGLoopInfo.cpp | 10 ++++++++++ clang/lib/CodeGen/CGLoopInfo.h | 10 ++++++++++ clang/lib/Sema/SemaStmtAttr.cpp | 6 +++++- clang/lib/Sema/SemaTemplateInstantiate.cpp | 11 +++++++++++ 6 files changed, 40 insertions(+), 3 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 6f3adf377b1ff..8ba9139d8c70b 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1909,7 +1909,7 @@ def SYCLIntelFPGALoopControlAvg : StmtAttr { CXX11<"intel","loop_control_avg">]; let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt], ErrorDiag, "'for', 'while', and 'do' statements">; - let Args = [ExprArgument<"N">]; + let Args = [ExprArgument<"NTripCount">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let HasCustomTypeTransform = 1; let Documentation = [SYCLIntelFPGALoopControlAvgAttrDocs]; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 54cc16faf56e4..90f1710d8f57e 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13253,7 +13253,9 @@ FPGALoopAttrT *Sema::BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving || A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations) { + ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations || + A.getParsedKind() == + ParsedAttr::AT_SYCLIntelFPGALoopControlAvg) { if (Val < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << A.getAttrName() << /* non-negative */ 1; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index fb3fd46ca293a..25c3d536a4acd 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1030,6 +1030,16 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, ->getSExtValue()); } + if (const auto *IntelFPGALoopControlAvg = + dyn_cast( + A)) { + setSYCLIntelFPGALoopControlAvgEnable(); + setSYCLIntelFPGALoopControlAvgNTripCount( + IntelFPGALoopControlAvg->getNTripCount() + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); + } + if (const auto *IntelFPGALoopCoalesce = dyn_cast(A)) { if (auto *LCE = IntelFPGALoopCoalesce->getNExpr()) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 2aa04eb5974a2..e5d728d7c6785 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -114,6 +114,8 @@ struct LoopAttributes { /// Flag for llvm.loop.max_concurrency.count metadata. bool SYCLMaxConcurrencyEnable; + bool SYCLIntelFPGALoopControlAvgEnable; + /// Value for llvm.loop.max_concurrency.count metadata. unsigned SYCLMaxConcurrencyNThreads; @@ -364,6 +366,14 @@ class LoopInfoStack { StagedAttrs.SYCLMaxConcurrencyEnable = true; } + void setSYCLIntelFPGALoopControlAvgEnable() { + StagedAttrs.SYCLIntelFPGALoopControlAvgEnable = true; + } + + void setSYCLIntelFPGALoopControlAvgNTripCount(unsigned C) { + StagedAttrs.SYCLIntelFPGALoopControlAvgEnable = C; + } + /// Set value of threads for the next loop pushed. void setSYCLMaxConcurrencyNThreads(unsigned C) { StagedAttrs.SYCLMaxConcurrencyNThreads = C; diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index e4e679099cadb..973038f67bc07 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -96,7 +96,7 @@ static Attr *handleIntelFPGALoopAttr(Sema &S, Stmt *St, const ParsedAttr &A) { if (A.getKind() == ParsedAttr::AT_SYCLIntelFPGAInitiationInterval || A.getKind() == ParsedAttr::AT_SYCLIntelFPGAMaxConcurrency || A.getKind() == ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving || - A.getKind() == ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations) { + A.getKind() == ParsedAttr::AT_SYCLIntelFPGALoopControlAvg) { S.Diag(A.getLoc(), diag::warn_attribute_too_few_arguments) << A << 1; return nullptr; } @@ -675,6 +675,8 @@ static void CheckForIncompatibleSYCLLoopAttributes( Attrs); CheckForDuplicationSYCLLoopAttribute( S, Attrs); + CheckForDuplicationSYCLLoopAttribute( + S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs, false); CheckMutualExclusionSYCLLoopAttribute( @@ -827,6 +829,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A, case ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations: return handleIntelFPGALoopAttr(S, St, A); + case ParsedAttr::AT_SYCLIntelFPGALoopControlAvg: + return handleIntelFPGALoopAttr(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 e59200cfdd657..4561add53fd49 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1093,6 +1093,9 @@ namespace { const SYCLIntelFPGASpeculatedIterationsAttr * TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI); + const SYCLIntelFPGALoopControlAvgAttr * + TransformSYCLIntelFPGALoopControlAvgAttr( + const SYCLIntelFPGALoopControlAvgAttr *SI); ExprResult TransformPredefinedExpr(PredefinedExpr *E); ExprResult TransformDeclRefExpr(DeclRefExpr *E); @@ -1618,6 +1621,14 @@ TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( *SI, TransformedExpr); } +const SYCLIntelFPGALoopControlAvgAttr * +TemplateInstantiator::TransformSYCLIntelFPGALoopControlAvgAttr( + const SYCLIntelFPGALoopControlAvgAttr *LCA) { + Expr *TransformedExpr = getDerived().TransformExpr(LCA->getNTripCount()).get(); + return getSema().BuildSYCLIntelFPGALoopAttr( + *LCA, TransformedExpr); +} + const LoopUnrollHintAttr *TemplateInstantiator::TransformLoopUnrollHintAttr( const LoopUnrollHintAttr *LU) { Expr *TransformedExpr = From 48d6df05c5b97d7a5e5552b50e69daa228b438db Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 9 Apr 2021 17:21:22 -0400 Subject: [PATCH 04/28] Adding test cases Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 15 ++++++-- clang/lib/CodeGen/CGLoopInfo.h | 23 +++++++----- clang/test/SemaSYCL/intel-fpga-loops.cpp | 45 +++++++++++++++++++++++- 3 files changed, 71 insertions(+), 12 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 25c3d536a4acd..b8bb312b983ac 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -605,6 +605,14 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } + if (Attrs.SYCLIntelFPGALoopControlAvgEnable) { + Metadata *Vals[] = { + MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), + Attrs.SYCLIntelFPGALoopCountAverage))}; + LoopProperties.push_back(MDNode::get(Ctx, Vals)); + } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), AdditionalLoopProperties.end()); return createFullUnrollMetadata(Attrs, LoopProperties, HasUserTransforms); @@ -621,7 +629,8 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false), SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), - SYCLSpeculatedIterationsNIterations(0), UnrollCount(0), + SYCLSpeculatedIterationsNIterations(0), + SYCLIntelFPGALoopCountAverage(0), UnrollCount(0), UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), MustProgress(false) {} @@ -643,6 +652,7 @@ void LoopAttributes::clear() { SYCLMaxInterleavingNInvocations = 0; SYCLSpeculatedIterationsEnable = false; SYCLSpeculatedIterationsNIterations = 0; + SYCLIntelFPGALoopCountAverage = 0; UnrollCount = 0; UnrollAndJamCount = 0; VectorizeEnable = LoopAttributes::Unspecified; @@ -680,6 +690,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && + Attrs.SYCLIntelFPGALoopCountAverage == 0 && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -1034,7 +1045,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast( A)) { setSYCLIntelFPGALoopControlAvgEnable(); - setSYCLIntelFPGALoopControlAvgNTripCount( + setSYCLIntelFPGALoopCountAverage( IntelFPGALoopControlAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index e5d728d7c6785..da64937d114f4 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -140,7 +140,10 @@ struct LoopAttributes { /// Value for llvm.loop.intel.speculated.iterations.count metadata. unsigned SYCLSpeculatedIterationsNIterations; - /// llvm.unroll. + /// Value for llvm.loop.intel.loopcount_average metadata. + unsigned SYCLIntelFPGALoopCountAverage; + + /// llvm.unroll. unsigned UnrollCount; /// llvm.unroll. @@ -366,19 +369,21 @@ class LoopInfoStack { StagedAttrs.SYCLMaxConcurrencyEnable = true; } - void setSYCLIntelFPGALoopControlAvgEnable() { - StagedAttrs.SYCLIntelFPGALoopControlAvgEnable = true; - } - - void setSYCLIntelFPGALoopControlAvgNTripCount(unsigned C) { - StagedAttrs.SYCLIntelFPGALoopControlAvgEnable = C; - } - /// Set value of threads for the next loop pushed. void setSYCLMaxConcurrencyNThreads(unsigned C) { StagedAttrs.SYCLMaxConcurrencyNThreads = C; } + /// Set flag of loop_control_avg for the next loop pushed. + void setSYCLIntelFPGALoopControlAvgEnable() { + StagedAttrs.SYCLIntelFPGALoopCountAverage = true; + } + /// Set value of loop control count average for the next + /// loop pushed. + void setSYCLIntelFPGALoopCountAverage(unsigned C) { + StagedAttrs.SYCLIntelFPGALoopCountAverage = C; + } + /// Set flag of loop_coalesce for the next loop pushed. void setSYCLLoopCoalesceEnable() { StagedAttrs.SYCLLoopCoalesceEnable = true; diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 082899e84ab28..7b4da9d643e25 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -22,6 +22,8 @@ void foo() { [[intel::speculated_iterations(6)]] int j[10]; // expected-error@+1 {{'nofusion' attribute cannot be applied to a declaration}} [[intel::nofusion]] int k[10]; + // expected-error@+1{{'loop_control_avg' attribute cannot be applied to a declaration}} + [[intel::loop_control_avg(6)]] int p[10]; } // Test for deprecated spelling of Intel FPGA loop attributes @@ -66,6 +68,11 @@ void foo_deprecated() { // expected-note@+1 {{did you mean to use 'intel::speculated_iterations' instead?}} [[intelfpga::speculated_iterations(6)]] for (int i = 0; i != 10; ++i) a[i] = 0; + + // expected-warning@+2 {{attribute 'intelfpga::loop_control_avg' is deprecated}} + // expected-note@+1 {{did you mean to use 'intel::loop_control_avg' instead?}} + [[intelfpga::loop_control_avg(6)]] for (int i = 0; i != 10; ++i) + a[i] = 0; } // Test for incorrect number of arguments for Intel FPGA loop attributes @@ -119,6 +126,9 @@ void boo() { // expected-error@+1 {{'nofusion' attribute takes no arguments}} [[intel::nofusion(0)]] for (int i = 0; i != 10; ++i) a[i] = 0; + // expected-error@+1 {{'loop_control_avg' attribute takes one argument}} + [[intel::loop_control_avg(3, 6)]] for (int i = 0; i != 10; ++i) + a[i] = 0; } // Test for incorrect argument value for Intel FPGA loop attributes @@ -196,6 +206,15 @@ void goo() { // no diagnostics are expected [[intel::nofusion]] for (int i = 0; i != 10; ++i) a[i] = 0; + + [[intel::loop_control_avg(0)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + // expected-error@+1 {{'loop_control_avg' attribute requires a non-negative integral compile time constant expression}} + [[intel::loop_control_avg(-1)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + // expected-error@+1 {{'loop_control_avg' attribute requires an integer constant}} + [[intel::loop_control_avg("abc")]] for (int i = 0; i != 10; ++i) + a[i] = 0; } // Test for Intel FPGA loop attributes duplication @@ -304,6 +323,11 @@ void zoo() { // expected-error@+1 {{duplicate Intel FPGA loop attribute 'nofusion'}} [[intel::nofusion]] for (int i = 0; i != 10; ++i) a[i] = 0; + + [[intel::loop_control_avg(2)]] + // expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_control_avg'}} + [[intel::loop_control_avg(2)]] for (int i = 0; i != 10; ++i) + a[i] = 0; } // Test for Intel FPGA loop attributes compatibility @@ -397,6 +421,22 @@ void max_concurrency_dependent() { a[i] = 0; } +template +void loop_control_avg_dependent() { + int a[10]; + + //expected-error@+1{{'loop_control_avg' attribute requires a non-negative integral compile time constant expression}} + [[intel::loop_control_avg(C)]] + for (int i = 0; i != 10; ++i) + a[i] = 0; + + [[intel::loop_control_avg(A)]] + //expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_control_avg'}} + [[intel::loop_control_avg(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + +} + int main() { deviceQueue.submit([&](sycl::handler &h) { h.single_task([]() { @@ -414,7 +454,10 @@ int main() { //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_control_avg_dependent<3, 2, -1>(); + //expected-note@-1{{in instantiation of function template specialization 'loop_control_avg_dependent<3, 2, -1>' requested here}} +}); }); return 0; From 44b86814a41d5f7c46d0d7d5fe4e8004f729408b Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 08:39:33 -0400 Subject: [PATCH 05/28] Indent Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 40 ++++++++++------------ clang/lib/CodeGen/CGLoopInfo.h | 4 +-- clang/lib/Sema/SemaStmtAttr.cpp | 4 +-- clang/lib/Sema/SemaTemplateInstantiate.cpp | 5 +-- 4 files changed, 26 insertions(+), 27 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index b8bb312b983ac..97c6aa01d780e 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -606,13 +606,13 @@ MDNode *LoopInfo::createMetadata( } if (Attrs.SYCLIntelFPGALoopControlAvgEnable) { - Metadata *Vals[] = { - MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLIntelFPGALoopCountAverage))}; - LoopProperties.push_back(MDNode::get(Ctx, Vals)); - } + Metadata *Vals[] = { + MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), + Attrs.SYCLIntelFPGALoopCountAverage))}; + LoopProperties.push_back(MDNode::get(Ctx, Vals)); + } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), AdditionalLoopProperties.end()); return createFullUnrollMetadata(Attrs, LoopProperties, HasUserTransforms); @@ -629,11 +629,11 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false), SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), - SYCLSpeculatedIterationsNIterations(0), - SYCLIntelFPGALoopCountAverage(0), UnrollCount(0), - UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), - PipelineDisabled(false), PipelineInitiationInterval(0), - SYCLNofusionEnable(false), MustProgress(false) {} + SYCLSpeculatedIterationsNIterations(0), SYCLIntelFPGALoopCountAverage(0), + UnrollCount(0), UnrollAndJamCount(0), + DistributeEnable(LoopAttributes::Unspecified), + PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), + MustProgress(false) {} void LoopAttributes::clear() { IsParallel = false; @@ -690,9 +690,9 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && - Attrs.SYCLIntelFPGALoopCountAverage == 0 && - Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && - !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && + Attrs.SYCLIntelFPGALoopCountAverage == 0 && Attrs.UnrollCount == 0 && + Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && + Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && Attrs.VectorizeEnable == LoopAttributes::Unspecified && Attrs.UnrollEnable == LoopAttributes::Unspecified && @@ -1042,13 +1042,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, } if (const auto *IntelFPGALoopControlAvg = - dyn_cast( - A)) { + dyn_cast(A)) { setSYCLIntelFPGALoopControlAvgEnable(); - setSYCLIntelFPGALoopCountAverage( - IntelFPGALoopControlAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + setSYCLIntelFPGALoopCountAverage(IntelFPGALoopControlAvg->getNTripCount() + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index da64937d114f4..3d224baf90daf 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -140,10 +140,10 @@ struct LoopAttributes { /// Value for llvm.loop.intel.speculated.iterations.count metadata. unsigned SYCLSpeculatedIterationsNIterations; - /// Value for llvm.loop.intel.loopcount_average metadata. + /// Value for llvm.loop.intel.loopcount_average metadata. unsigned SYCLIntelFPGALoopCountAverage; - /// llvm.unroll. + /// llvm.unroll. unsigned UnrollCount; /// llvm.unroll. diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 615c8b7924087..c6f946d94b7fc 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -572,8 +572,8 @@ static void CheckForIncompatibleSYCLLoopAttributes( Attrs); CheckForDuplicationSYCLLoopAttribute( S, Attrs); - CheckForDuplicationSYCLLoopAttribute( - S, Attrs); + CheckForDuplicationSYCLLoopAttribute(S, + Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs, false); CheckMutualExclusionSYCLLoopAttribute( diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index 2f9c99804d6c6..ce5db06269ac2 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1624,9 +1624,10 @@ TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGALoopControlAvgAttr * TemplateInstantiator::TransformSYCLIntelFPGALoopControlAvgAttr( const SYCLIntelFPGALoopControlAvgAttr *LCA) { - Expr *TransformedExpr = getDerived().TransformExpr(LCA->getNTripCount()).get(); + Expr *TransformedExpr = + getDerived().TransformExpr(LCA->getNTripCount()).get(); return getSema().BuildSYCLIntelFPGALoopAttr( - *LCA, TransformedExpr); + *LCA, TransformedExpr); } const LoopUnrollHintAttr *TemplateInstantiator::TransformLoopUnrollHintAttr( From 7622568b3e45cba51243f1457f4777cf10407bc1 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 08:50:36 -0400 Subject: [PATCH 06/28] Indent Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 11 +++++------ clang/lib/Sema/SemaTemplateInstantiate.cpp | 2 +- 2 files changed, 6 insertions(+), 7 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 97c6aa01d780e..08b4458a05400 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -608,11 +608,10 @@ MDNode *LoopInfo::createMetadata( if (Attrs.SYCLIntelFPGALoopControlAvgEnable) { Metadata *Vals[] = { MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLIntelFPGALoopCountAverage))}; + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), Attrs.SYCLIntelFPGALoopCountAverage))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); - } + } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), AdditionalLoopProperties.end()); return createFullUnrollMetadata(Attrs, LoopProperties, HasUserTransforms); @@ -631,8 +630,8 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLSpeculatedIterationsEnable(false), SYCLSpeculatedIterationsNIterations(0), SYCLIntelFPGALoopCountAverage(0), UnrollCount(0), UnrollAndJamCount(0), - DistributeEnable(LoopAttributes::Unspecified), - PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), + DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), + PipelineInitiationInterval(0), SYCLNofusionEnable(false), MustProgress(false) {} void LoopAttributes::clear() { diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index ce5db06269ac2..04fbe31eeb455 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1627,7 +1627,7 @@ TemplateInstantiator::TransformSYCLIntelFPGALoopControlAvgAttr( Expr *TransformedExpr = getDerived().TransformExpr(LCA->getNTripCount()).get(); return getSema().BuildSYCLIntelFPGALoopAttr( - *LCA, TransformedExpr); + *LCA, TransformedExpr); } const LoopUnrollHintAttr *TemplateInstantiator::TransformLoopUnrollHintAttr( From c3d8c59bda9ab16a059e34b0f717c94eddcefa26 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 09:09:18 -0400 Subject: [PATCH 07/28] Indent Signed-off-by: Zahira Ammarguellat --- 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 3d224baf90daf..143411f65cc43 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -142,7 +142,7 @@ struct LoopAttributes { /// Value for llvm.loop.intel.loopcount_average metadata. unsigned SYCLIntelFPGALoopCountAverage; - + /// llvm.unroll. unsigned UnrollCount; From 4dba2ea02719d2b92d621db53a902c0c1e6946fb Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 11:22:54 -0400 Subject: [PATCH 08/28] CodeGen impl Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 31 ++++++++++---------- clang/lib/CodeGen/CGLoopInfo.h | 32 +++++++++++---------- clang/test/CodeGenSYCL/intel-fpga-loops.cpp | 16 +++++++++++ 3 files changed, 49 insertions(+), 30 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 08b4458a05400..6fe4c6b5078a0 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -605,11 +605,11 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLIntelFPGALoopControlAvgEnable) { - Metadata *Vals[] = { - MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), Attrs.SYCLIntelFPGALoopCountAverage))}; + if (Attrs.SYCLIntelFPGALoopControlAverageEnable) { + Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), + Attrs.SYCLIntelFPGALoopControlAverage))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), @@ -628,11 +628,11 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false), SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), - SYCLSpeculatedIterationsNIterations(0), SYCLIntelFPGALoopCountAverage(0), - UnrollCount(0), UnrollAndJamCount(0), - DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), - PipelineInitiationInterval(0), SYCLNofusionEnable(false), - MustProgress(false) {} + SYCLSpeculatedIterationsNIterations(0), + SYCLIntelFPGALoopControlAverageEnable(false), UnrollCount(0), + UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), + PipelineDisabled(false), PipelineInitiationInterval(0), + SYCLNofusionEnable(false), MustProgress(false) {} void LoopAttributes::clear() { IsParallel = false; @@ -651,7 +651,7 @@ void LoopAttributes::clear() { SYCLMaxInterleavingNInvocations = 0; SYCLSpeculatedIterationsEnable = false; SYCLSpeculatedIterationsNIterations = 0; - SYCLIntelFPGALoopCountAverage = 0; + SYCLIntelFPGALoopControlAverageEnable = false; UnrollCount = 0; UnrollAndJamCount = 0; VectorizeEnable = LoopAttributes::Unspecified; @@ -689,7 +689,8 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && - Attrs.SYCLIntelFPGALoopCountAverage == 0 && Attrs.UnrollCount == 0 && + Attrs.SYCLIntelFPGALoopControlAverageEnable == 0 && + Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -1043,9 +1044,9 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopControlAvg = dyn_cast(A)) { setSYCLIntelFPGALoopControlAvgEnable(); - setSYCLIntelFPGALoopCountAverage(IntelFPGALoopControlAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + setSYCLIntelFPGALoopControlAverage(IntelFPGALoopControlAvg->getNTripCount() + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 143411f65cc43..02d8df21c4a33 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -114,11 +114,15 @@ struct LoopAttributes { /// Flag for llvm.loop.max_concurrency.count metadata. bool SYCLMaxConcurrencyEnable; - bool SYCLIntelFPGALoopControlAvgEnable; - /// Value for llvm.loop.max_concurrency.count metadata. unsigned SYCLMaxConcurrencyNThreads; + /// Flag for llvm.loop.intel.loopcount_average metadata. + bool SYCLIntelFPGALoopControlAverageEnable; + + /// Value for llvm.loop.intel.loopcount_average metadata. + unsigned SYCLIntelFPGALoopControlAverage; + /// Flag for llvm.loop.coalesce metadata. bool SYCLLoopCoalesceEnable; @@ -140,9 +144,6 @@ struct LoopAttributes { /// Value for llvm.loop.intel.speculated.iterations.count metadata. unsigned SYCLSpeculatedIterationsNIterations; - /// Value for llvm.loop.intel.loopcount_average metadata. - unsigned SYCLIntelFPGALoopCountAverage; - /// llvm.unroll. unsigned UnrollCount; @@ -374,16 +375,6 @@ class LoopInfoStack { StagedAttrs.SYCLMaxConcurrencyNThreads = C; } - /// Set flag of loop_control_avg for the next loop pushed. - void setSYCLIntelFPGALoopControlAvgEnable() { - StagedAttrs.SYCLIntelFPGALoopCountAverage = true; - } - /// Set value of loop control count average for the next - /// loop pushed. - void setSYCLIntelFPGALoopCountAverage(unsigned C) { - StagedAttrs.SYCLIntelFPGALoopCountAverage = C; - } - /// Set flag of loop_coalesce for the next loop pushed. void setSYCLLoopCoalesceEnable() { StagedAttrs.SYCLLoopCoalesceEnable = true; @@ -419,6 +410,17 @@ class LoopInfoStack { StagedAttrs.SYCLSpeculatedIterationsNIterations = C; } + /// Set flag of loop_control_avg for the next loop pushed. + void setSYCLIntelFPGALoopControlAvgEnable() { + StagedAttrs.SYCLIntelFPGALoopControlAverageEnable = true; + } + + /// Set value of loop control average for the next loop pushed. + void setSYCLIntelFPGALoopControlAverage(unsigned C) { + StagedAttrs.SYCLIntelFPGALoopControlAverage = C; + } + + /// Set the unroll count for the next loop pushed. void setUnrollCount(unsigned C) { StagedAttrs.UnrollCount = C; } diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index 8fb40e31ada3d..92347eb53c67b 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -14,6 +14,8 @@ // CHECK: br label %for.cond2, !llvm.loop ![[MD_MI_2:[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.cond, !llvm.loop ![[MD_LCA:[0-9]+]] +// CHECK: br label %for.cond2, !llvm.loop ![[MD_LCA_1:[0-9]+]] void disable_loop_pipelining() { int a[10]; @@ -109,6 +111,19 @@ void speculated_iterations() { a[i] = 0; } +template +void loop_control_avg() { + int a[10]; + // CHECK: ![[MD_LCA]] = distinct !{![[MD_LCA]], ![[MP:[0-9]+]], ![[MD_lca:[0-9]+]]} + // CHECK-NEXT: ![[MD_lca]] = !{!"llvm.loop.intel.loopcount_average", i32 12} + [[intel::loop_control_avg(A)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + // CHECK: ![[MD_LCA_1]] = distinct !{![[MD_LCA_1]], ![[MP:[0-9]+]], ![[MD_lca_1:[0-9]+]]} + // CHECK-NEXT: ![[MD_lca_1]] = !{!"llvm.loop.intel.loopcount_average", i32 4} + [[intel::loop_control_avg(4)]] for (int i = 0; i != 10; ++i) + a[i] = 0; +} + template __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); @@ -123,6 +138,7 @@ int main() { loop_coalesce<2>(); max_interleaving<3>(); speculated_iterations<4>(); + loop_control_avg<12>(); }); return 0; } From 50a06c78811266b974c0e197aefb0339dc718e04 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 12:55:17 -0400 Subject: [PATCH 09/28] After review comments Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 3 +-- clang/test/SemaSYCL/intel-fpga-loops.cpp | 10 +++++----- 2 files changed, 6 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 610f7c5c8761a..2036b96391509 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1954,8 +1954,7 @@ def : MutualExclusions<[SYCLIntelFPGAInitiationInterval, SYCLIntelFPGADisableLoopPipelining]>; def SYCLIntelFPGALoopControlAvg : StmtAttr { - let Spellings = [CXX11<"intelfpga","loop_control_avg">, - CXX11<"intel","loop_control_avg">]; + let Spellings = [CXX11<"intel","loop_control_avg">]; let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt], ErrorDiag, "'for', 'while', and 'do' statements">; let Args = [ExprArgument<"NTripCount">]; diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 7b4da9d643e25..b8c1e750cf6a1 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -68,11 +68,6 @@ void foo_deprecated() { // expected-note@+1 {{did you mean to use 'intel::speculated_iterations' instead?}} [[intelfpga::speculated_iterations(6)]] for (int i = 0; i != 10; ++i) a[i] = 0; - - // expected-warning@+2 {{attribute 'intelfpga::loop_control_avg' is deprecated}} - // expected-note@+1 {{did you mean to use 'intel::loop_control_avg' instead?}} - [[intelfpga::loop_control_avg(6)]] for (int i = 0; i != 10; ++i) - a[i] = 0; } // Test for incorrect number of arguments for Intel FPGA loop attributes @@ -361,6 +356,11 @@ void loop_attrs_compatibility() { [[intel::disable_loop_pipelining]] [[intel::nofusion]] for (int i = 0; i != 10; ++i) a[i] = 0; + // no diagnostics are expected + [[intel::disable_loop_pipelining]] + [[intel::loop_control_avg(8)]] + for (int i = 0; i != 10; ++i) + a[i] = 0; } template From 1bb84bcd9bc81e23cead102b6fb57137133bda9b Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 13:26:53 -0400 Subject: [PATCH 10/28] Changing name Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 6 +-- clang/include/clang/Basic/AttrDocs.td | 6 +-- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/CodeGen/CGLoopInfo.cpp | 12 +++--- clang/lib/CodeGen/CGLoopInfo.h | 6 +-- clang/lib/Sema/SemaStmtAttr.cpp | 6 +-- clang/lib/Sema/SemaTemplateInstantiate.cpp | 14 +++---- clang/test/CodeGenSYCL/intel-fpga-loops.cpp | 8 ++-- clang/test/SemaSYCL/intel-fpga-loops.cpp | 42 ++++++++++----------- 9 files changed, 51 insertions(+), 51 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 2036b96391509..84941de9ebdaf 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1953,14 +1953,14 @@ def SYCLIntelFPGADisableLoopPipelining : DeclOrStmtAttr { def : MutualExclusions<[SYCLIntelFPGAInitiationInterval, SYCLIntelFPGADisableLoopPipelining]>; -def SYCLIntelFPGALoopControlAvg : StmtAttr { - let Spellings = [CXX11<"intel","loop_control_avg">]; +def SYCLIntelFPGALoopCountAvg : StmtAttr { + let Spellings = [CXX11<"intel","loop_count_avg">]; let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt], ErrorDiag, "'for', 'while', and 'do' statements">; let Args = [ExprArgument<"NTripCount">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let HasCustomTypeTransform = 1; - let Documentation = [SYCLIntelFPGALoopControlAvgAttrDocs]; + let Documentation = [SYCLIntelFPGALoopCountAvgAttrDocs]; } def SYCLIntelFPGAMaxInterleaving : StmtAttr { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index c1f52af4d6268..855ccae167969 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2950,7 +2950,7 @@ max_concurrency, initiation_interval, or ivdep. }]; } -def SYCLIntelFPGALoopControlAvgAttrDocs : Documentation { +def SYCLIntelFPGALoopCountAvgAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "intel::loo_control_avg"; let Content = [{ @@ -2962,12 +2962,12 @@ the optimization is beneficial. void foo() { int a[10]; - [[intel::loop_control_avg(40)] for (int i = 0; i < 10; ++i) a[i] = 0; + [[intel::loop_count_avg(40)] for (int i = 0; i < 10; ++i) a[i] = 0; } template void bar() { - [[intel::loop_control_avg(N)]] for(;;) { } + [[intel::loop_count_avg(N)]] for(;;) { } } }]; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index d8a382d08865e..0dd975f3b0819 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13421,7 +13421,7 @@ FPGALoopAttrT *Sema::BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations || A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGALoopControlAvg) { + ParsedAttr::AT_SYCLIntelFPGALoopCountAvg) { if (Val < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << A.getAttrName() << /* non-negative */ 1; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 6fe4c6b5078a0..fc6decd7e063e 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -605,7 +605,7 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLIntelFPGALoopControlAverageEnable) { + if (Attrs.SYCLIntelFPGALoopCountAverageEnable) { Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), ConstantAsMetadata::get(ConstantInt::get( llvm::Type::getInt32Ty(Ctx), @@ -629,7 +629,7 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), SYCLSpeculatedIterationsNIterations(0), - SYCLIntelFPGALoopControlAverageEnable(false), UnrollCount(0), + SYCLIntelFPGALoopCountAverageEnable(false), UnrollCount(0), UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), MustProgress(false) {} @@ -651,7 +651,7 @@ void LoopAttributes::clear() { SYCLMaxInterleavingNInvocations = 0; SYCLSpeculatedIterationsEnable = false; SYCLSpeculatedIterationsNIterations = 0; - SYCLIntelFPGALoopControlAverageEnable = false; + SYCLIntelFPGALoopCountAverageEnable = false; UnrollCount = 0; UnrollAndJamCount = 0; VectorizeEnable = LoopAttributes::Unspecified; @@ -689,7 +689,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && - Attrs.SYCLIntelFPGALoopControlAverageEnable == 0 && + Attrs.SYCLIntelFPGALoopCountAverageEnable == 0 && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && @@ -1042,8 +1042,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, } if (const auto *IntelFPGALoopControlAvg = - dyn_cast(A)) { - setSYCLIntelFPGALoopControlAvgEnable(); + dyn_cast(A)) { + setSYCLIntelFPGALoopCountAvgEnable(); setSYCLIntelFPGALoopControlAverage(IntelFPGALoopControlAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 02d8df21c4a33..8c074d4aee7c7 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -118,7 +118,7 @@ struct LoopAttributes { unsigned SYCLMaxConcurrencyNThreads; /// Flag for llvm.loop.intel.loopcount_average metadata. - bool SYCLIntelFPGALoopControlAverageEnable; + bool SYCLIntelFPGALoopCountAverageEnable; /// Value for llvm.loop.intel.loopcount_average metadata. unsigned SYCLIntelFPGALoopControlAverage; @@ -411,8 +411,8 @@ class LoopInfoStack { } /// Set flag of loop_control_avg for the next loop pushed. - void setSYCLIntelFPGALoopControlAvgEnable() { - StagedAttrs.SYCLIntelFPGALoopControlAverageEnable = true; + void setSYCLIntelFPGALoopCountAvgEnable() { + StagedAttrs.SYCLIntelFPGALoopCountAverageEnable = true; } /// Set value of loop control average for the next loop pushed. diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index c6f946d94b7fc..5a64be3f250af 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -572,7 +572,7 @@ static void CheckForIncompatibleSYCLLoopAttributes( Attrs); CheckForDuplicationSYCLLoopAttribute( S, Attrs); - CheckForDuplicationSYCLLoopAttribute(S, + CheckForDuplicationSYCLLoopAttribute(S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs, false); CheckMutualExclusionSYCLLoopAttribute(S, St, A); - case ParsedAttr::AT_SYCLIntelFPGALoopControlAvg: - return handleIntelFPGALoopAttr(S, St, A); + case ParsedAttr::AT_SYCLIntelFPGALoopCountAvg: + return handleIntelFPGALoopAttr(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 04fbe31eeb455..545329e6fb968 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1093,9 +1093,9 @@ namespace { const SYCLIntelFPGASpeculatedIterationsAttr * TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI); - const SYCLIntelFPGALoopControlAvgAttr * - TransformSYCLIntelFPGALoopControlAvgAttr( - const SYCLIntelFPGALoopControlAvgAttr *SI); + const SYCLIntelFPGALoopCountAvgAttr * + TransformSYCLIntelFPGALoopCountAvgAttr( + const SYCLIntelFPGALoopCountAvgAttr *SI); ExprResult TransformPredefinedExpr(PredefinedExpr *E); ExprResult TransformDeclRefExpr(DeclRefExpr *E); @@ -1621,12 +1621,12 @@ TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( *SI, TransformedExpr); } -const SYCLIntelFPGALoopControlAvgAttr * -TemplateInstantiator::TransformSYCLIntelFPGALoopControlAvgAttr( - const SYCLIntelFPGALoopControlAvgAttr *LCA) { +const SYCLIntelFPGALoopCountAvgAttr * +TemplateInstantiator::TransformSYCLIntelFPGALoopCountAvgAttr( + const SYCLIntelFPGALoopCountAvgAttr *LCA) { Expr *TransformedExpr = getDerived().TransformExpr(LCA->getNTripCount()).get(); - return getSema().BuildSYCLIntelFPGALoopAttr( + return getSema().BuildSYCLIntelFPGALoopAttr( *LCA, TransformedExpr); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index 92347eb53c67b..cfbf6d09436e3 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -112,15 +112,15 @@ void speculated_iterations() { } template -void loop_control_avg() { +void loop_count_avg() { int a[10]; // CHECK: ![[MD_LCA]] = distinct !{![[MD_LCA]], ![[MP:[0-9]+]], ![[MD_lca:[0-9]+]]} // CHECK-NEXT: ![[MD_lca]] = !{!"llvm.loop.intel.loopcount_average", i32 12} - [[intel::loop_control_avg(A)]] for (int i = 0; i != 10; ++i) + [[intel::loop_count_avg(A)]] for (int i = 0; i != 10; ++i) a[i] = 0; // CHECK: ![[MD_LCA_1]] = distinct !{![[MD_LCA_1]], ![[MP:[0-9]+]], ![[MD_lca_1:[0-9]+]]} // CHECK-NEXT: ![[MD_lca_1]] = !{!"llvm.loop.intel.loopcount_average", i32 4} - [[intel::loop_control_avg(4)]] for (int i = 0; i != 10; ++i) + [[intel::loop_count_avg(4)]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -138,7 +138,7 @@ int main() { loop_coalesce<2>(); max_interleaving<3>(); speculated_iterations<4>(); - loop_control_avg<12>(); + loop_count_avg<12>(); }); return 0; } diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index b8c1e750cf6a1..67644a5157a7f 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -22,8 +22,8 @@ void foo() { [[intel::speculated_iterations(6)]] int j[10]; // expected-error@+1 {{'nofusion' attribute cannot be applied to a declaration}} [[intel::nofusion]] int k[10]; - // expected-error@+1{{'loop_control_avg' attribute cannot be applied to a declaration}} - [[intel::loop_control_avg(6)]] int p[10]; + // expected-error@+1{{'loop_count_avg' attribute cannot be applied to a declaration}} + [[intel::loop_count_avg(6)]] int p[10]; } // Test for deprecated spelling of Intel FPGA loop attributes @@ -121,8 +121,8 @@ void boo() { // expected-error@+1 {{'nofusion' attribute takes no arguments}} [[intel::nofusion(0)]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'loop_control_avg' attribute takes one argument}} - [[intel::loop_control_avg(3, 6)]] for (int i = 0; i != 10; ++i) + // expected-error@+1 {{'loop_count_avg' attribute takes one argument}} + [[intel::loop_count_avg(3, 6)]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -202,13 +202,13 @@ void goo() { [[intel::nofusion]] for (int i = 0; i != 10; ++i) a[i] = 0; - [[intel::loop_control_avg(0)]] for (int i = 0; i != 10; ++i) + [[intel::loop_count_avg(0)]] for (int i = 0; i != 10; ++i) a[i] = 0; - // expected-error@+1 {{'loop_control_avg' attribute requires a non-negative integral compile time constant expression}} - [[intel::loop_control_avg(-1)]] for (int i = 0; i != 10; ++i) + // 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_control_avg' attribute requires an integer constant}} - [[intel::loop_control_avg("abc")]] for (int i = 0; i != 10; ++i) + // expected-error@+1 {{'loop_count_avg' attribute requires an integer constant}} + [[intel::loop_count_avg("abc")]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -319,9 +319,9 @@ void zoo() { [[intel::nofusion]] for (int i = 0; i != 10; ++i) a[i] = 0; - [[intel::loop_control_avg(2)]] - // expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_control_avg'}} - [[intel::loop_control_avg(2)]] for (int i = 0; i != 10; ++i) + [[intel::loop_count_avg(2)]] + // expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_count_avg'}} + [[intel::loop_count_avg(2)]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -358,7 +358,7 @@ void loop_attrs_compatibility() { a[i] = 0; // no diagnostics are expected [[intel::disable_loop_pipelining]] - [[intel::loop_control_avg(8)]] + [[intel::loop_count_avg(8)]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -422,17 +422,17 @@ void max_concurrency_dependent() { } template -void loop_control_avg_dependent() { +void loop_count_avg_dependent() { int a[10]; - //expected-error@+1{{'loop_control_avg' attribute requires a non-negative integral compile time constant expression}} - [[intel::loop_control_avg(C)]] + //expected-error@+1{{'loop_count_avg' attribute requires a non-negative integral compile time constant expression}} + [[intel::loop_count_avg(C)]] for (int i = 0; i != 10; ++i) a[i] = 0; - [[intel::loop_control_avg(A)]] - //expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_control_avg'}} - [[intel::loop_control_avg(B)]] for (int i = 0; i != 10; ++i) + [[intel::loop_count_avg(A)]] + //expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_count_avg'}} + [[intel::loop_count_avg(B)]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -455,8 +455,8 @@ int main() { max_concurrency_dependent<1, 4, -2>(); //expected-note@-1 +{{in instantiation of function template specialization}} - loop_control_avg_dependent<3, 2, -1>(); - //expected-note@-1{{in instantiation of function template specialization 'loop_control_avg_dependent<3, 2, -1>' requested here}} + loop_count_avg_dependent<3, 2, -1>(); + //expected-note@-1{{in instantiation of function template specialization 'loop_count_avg_dependent<3, 2, -1>' requested here}} }); }); From 40a32799b19cd0982511049db66b312014d977bb Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 12 Apr 2021 14:16:54 -0400 Subject: [PATCH 11/28] Format Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Sema/Sema.h | 3 +-- clang/lib/CodeGen/CGLoopInfo.cpp | 9 ++++----- clang/lib/CodeGen/CGLoopInfo.h | 1 - clang/lib/Sema/SemaStmtAttr.cpp | 3 +-- clang/lib/Sema/SemaTemplateInstantiate.cpp | 3 +-- 5 files changed, 7 insertions(+), 12 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 0dd975f3b0819..aa4b361224be8 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13420,8 +13420,7 @@ FPGALoopAttrT *Sema::BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving || A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations || - A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGALoopCountAvg) { + A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGALoopCountAvg) { if (Val < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << A.getAttrName() << /* non-negative */ 1; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index fc6decd7e063e..dc2a145f687ce 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -690,9 +690,8 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && Attrs.SYCLIntelFPGALoopCountAverageEnable == 0 && - Attrs.UnrollCount == 0 && - Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && - Attrs.PipelineInitiationInterval == 0 && + Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && + !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && Attrs.VectorizeEnable == LoopAttributes::Unspecified && Attrs.UnrollEnable == LoopAttributes::Unspecified && @@ -1041,10 +1040,10 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, ->getSExtValue()); } - if (const auto *IntelFPGALoopControlAvg = + if (const auto *IntelFPGALoopCountAvg = dyn_cast(A)) { setSYCLIntelFPGALoopCountAvgEnable(); - setSYCLIntelFPGALoopControlAverage(IntelFPGALoopControlAvg->getNTripCount() + setSYCLIntelFPGALoopControlAverage(IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); } diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 8c074d4aee7c7..b9cbd2a19c57c 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -419,7 +419,6 @@ class LoopInfoStack { void setSYCLIntelFPGALoopControlAverage(unsigned C) { StagedAttrs.SYCLIntelFPGALoopControlAverage = C; } - /// Set the unroll count for the next loop pushed. void setUnrollCount(unsigned C) { StagedAttrs.UnrollCount = C; } diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 5a64be3f250af..3c4566b673405 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -572,8 +572,7 @@ static void CheckForIncompatibleSYCLLoopAttributes( Attrs); CheckForDuplicationSYCLLoopAttribute( S, Attrs); - CheckForDuplicationSYCLLoopAttribute(S, - Attrs); + CheckForDuplicationSYCLLoopAttribute(S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs, false); CheckMutualExclusionSYCLLoopAttribute( diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index 545329e6fb968..ded1969969e6f 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1093,8 +1093,7 @@ namespace { const SYCLIntelFPGASpeculatedIterationsAttr * TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI); - const SYCLIntelFPGALoopCountAvgAttr * - TransformSYCLIntelFPGALoopCountAvgAttr( + const SYCLIntelFPGALoopCountAvgAttr *TransformSYCLIntelFPGALoopCountAvgAttr( const SYCLIntelFPGALoopCountAvgAttr *SI); ExprResult TransformPredefinedExpr(PredefinedExpr *E); From 6ab3b07d01c28a7a54dbf9ec6caaa36df617314a Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 13 Apr 2021 16:28:41 -0400 Subject: [PATCH 12/28] After review Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 3 +-- clang/include/clang/Basic/AttrDocs.td | 9 +++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 5472c6d897e8f..01bee833dcd64 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1958,12 +1958,11 @@ def : MutualExclusions<[SYCLIntelFPGAMaxConcurrency, SYCLIntelFPGADisableLoopPipelining]>; def SYCLIntelFPGALoopCountAvg : StmtAttr { - let Spellings = [CXX11<"intel","loop_count_avg">]; + let Spellings = [CXX11<"intel", "loop_count_avg">]; let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt], ErrorDiag, "'for', 'while', and 'do' statements">; let Args = [ExprArgument<"NTripCount">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let HasCustomTypeTransform = 1; let Documentation = [SYCLIntelFPGALoopCountAvgAttrDocs]; } diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 855ccae167969..86aa1fd326faf 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2952,11 +2952,12 @@ max_concurrency, initiation_interval, or ivdep. def SYCLIntelFPGALoopCountAvgAttrDocs : Documentation { let Category = DocCatVariable; - let Heading = "intel::loo_control_avg"; + let Heading = "intel::loop_count_avg"; let Content = [{ -This attribute applies to a loop. Its purpose is to pass to the optimizer the -average trip count of a loop. This will allow the optimizer to figure out if -the optimization is beneficial. +The "loop_count" pragma specifies the minimum, maximum, or average number of +iterations for a for loop. In addition, a list of commonly occurring values +can be specified to help the compiler generate multiple versions and perform +complete unrolling. .. code-block:: c++ From fd34d550ae489af54985c171dbdcdf95dd47173e Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 14 Apr 2021 08:20:46 -0400 Subject: [PATCH 13/28] Formatting Signed-off-by: Zahira Ammarguellat --- 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 dc2a145f687ce..c4a4d8dc0a276 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1044,8 +1044,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { setSYCLIntelFPGALoopCountAvgEnable(); setSYCLIntelFPGALoopControlAverage(IntelFPGALoopCountAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = From d812fa443883fe9aa4631cb0f3f4eb06fe5ab89d Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 14 Apr 2021 09:38:45 -0400 Subject: [PATCH 14/28] Formatting Signed-off-by: Zahira Ammarguellat --- 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 c4a4d8dc0a276..e9d516153d459 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1044,8 +1044,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { setSYCLIntelFPGALoopCountAvgEnable(); setSYCLIntelFPGALoopControlAverage(IntelFPGALoopCountAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); } if (const auto *IntelFPGALoopCoalesce = From fd748424d1e4f3c12ce1b4b02b2382c41bcbd594 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 14 Apr 2021 16:45:26 -0400 Subject: [PATCH 15/28] Looks like the HasCustomTypeTransform is needed Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index e9b6b0da52eb4..03c4b60939ffc 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1963,6 +1963,7 @@ def SYCLIntelFPGALoopCountAvg : StmtAttr { ErrorDiag, "'for', 'while', and 'do' statements">; let Args = [ExprArgument<"NTripCount">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; + let HasCustomTypeTransform = 1; let Documentation = [SYCLIntelFPGALoopCountAvgAttrDocs]; } From 73f92a4a55e4faa40b4da113ef2220ba7d19afb4 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 16 Apr 2021 18:10:57 -0400 Subject: [PATCH 16/28] Added min/max Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 19 +++++++-- clang/include/clang/Basic/AttrDocs.td | 11 ++++-- clang/include/clang/Sema/Sema.h | 6 ++- clang/lib/CodeGen/CGLoopInfo.cpp | 44 +++++++++++++-------- clang/lib/CodeGen/CGLoopInfo.h | 19 ++++++--- clang/lib/Sema/SemaStmtAttr.cpp | 30 ++++++++++++-- clang/lib/Sema/SemaTemplateInstantiate.cpp | 12 +++--- clang/test/CodeGenSYCL/intel-fpga-loops.cpp | 8 ++-- clang/test/SemaSYCL/intel-fpga-loops.cpp | 6 +-- 9 files changed, 109 insertions(+), 46 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 03c4b60939ffc..8659c19688c12 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1957,14 +1957,27 @@ def : MutualExclusions<[SYCLIntelFPGAIVDep, def : MutualExclusions<[SYCLIntelFPGAMaxConcurrency, SYCLIntelFPGADisableLoopPipelining]>; -def SYCLIntelFPGALoopCountAvg : StmtAttr { - let Spellings = [CXX11<"intel", "loop_count_avg">]; +def SYCLIntelFPGALoopCount : StmtAttr { + let Spellings = [CXX11<"intel", "loop_count_min">, + CXX11<"intel", "loop_count_max">, + CXX11<"intel", "loop_count_avg">]; let Subjects = SubjectList<[ForStmt, CXXForRangeStmt, WhileStmt, DoStmt], ErrorDiag, "'for', 'while', and 'do' statements">; + let Accessors = [Accessor<"isMin", [CXX11<"intel", "loop_count_min">]>, + Accessor<"isMax", [CXX11<"intel", "loop_count_max">]>, + Accessor<"isAvg", [CXX11<"intel", "loop_count_avg">]>]; let Args = [ExprArgument<"NTripCount">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; let HasCustomTypeTransform = 1; - let Documentation = [SYCLIntelFPGALoopCountAvgAttrDocs]; + let AdditionalMembers = [{ + enum CountKind { loop_count_min, loop_count_max, loop_count_avg }; + CountKind getCountKind() const { + return isMin() ? loop_count_min : + isMax() ? loop_count_max : + loop_count_avg; + } + }]; + let Documentation = [SYCLIntelFPGALoopCountAttrDocs]; } def : MutualExclusions<[SYCLIntelFPGAMaxConcurrency, diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 314a0a8f44bd3..30fb76b0e306c 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2951,9 +2951,9 @@ or ``ivdep``. }]; } -def SYCLIntelFPGALoopCountAvgAttrDocs : Documentation { +def SYCLIntelFPGALoopCountAttrDocs : Documentation { let Category = DocCatVariable; - let Heading = "intel::loop_count_avg"; + let Heading = "intel::loop_count_min, intel::loop_count_max, intel::loop_count_avg"; let Content = [{ The "loop_count" pragma specifies the minimum, maximum, or average number of iterations for a for loop. In addition, a list of commonly occurring values @@ -2964,7 +2964,12 @@ complete unrolling. void foo() { int a[10]; - [[intel::loop_count_avg(40)] for (int i = 0; i < 10; ++i) a[i] = 0; + [[intel::loop_count_min(40)] for (int i = 0; i < 10; ++i) a[i] = 0; + } + + void foo() { + int a[10]; + [[intel::loop_count_max(40)] for (int i = 0; i < 10; ++i) a[i] = 0; } template diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index aa4b361224be8..d992bd1186766 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -2121,6 +2121,9 @@ class Sema final { OpenCLUnrollHintAttr * BuildOpenCLLoopUnrollHintAttr(const AttributeCommonInfo &A, Expr *E); + SYCLIntelFPGALoopCountAttr * + BuildSYCLIntelFPGALoopCount(const AttributeCommonInfo &CI, Expr *E); + bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc); bool CheckFunctionReturnType(QualType T, SourceLocation Loc); @@ -13420,7 +13423,8 @@ FPGALoopAttrT *Sema::BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving || A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations || - A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGALoopCountAvg) { + A.getParsedKind() == + ParsedAttr::AT_SYCLIntelFPGALoopCount) { if (Val < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << A.getAttrName() << /* non-negative */ 1; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index e9d516153d459..919f6e57a2a13 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -553,10 +553,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"), - 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)); } @@ -605,11 +605,11 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLIntelFPGALoopCountAverageEnable) { - Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.intel.loopcount_average"), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLIntelFPGALoopControlAverage))}; + if (Attrs.SYCLIntelFPGALoopCountEnable) { + Metadata *Vals[] = { + MDString::get(Ctx, Attrs.SYCLIntelFPGALoopCountVariation), + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), Attrs.SYCLIntelFPGALoopCount))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), @@ -629,7 +629,8 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), SYCLSpeculatedIterationsNIterations(0), - SYCLIntelFPGALoopCountAverageEnable(false), UnrollCount(0), + SYCLIntelFPGALoopCountEnable(false), + SYCLIntelFPGALoopCountVariation(nullptr), UnrollCount(0), UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), MustProgress(false) {} @@ -651,7 +652,8 @@ void LoopAttributes::clear() { SYCLMaxInterleavingNInvocations = 0; SYCLSpeculatedIterationsEnable = false; SYCLSpeculatedIterationsNIterations = 0; - SYCLIntelFPGALoopCountAverageEnable = false; + SYCLIntelFPGALoopCountEnable = false; + SYCLIntelFPGALoopCountVariation = nullptr; UnrollCount = 0; UnrollAndJamCount = 0; VectorizeEnable = LoopAttributes::Unspecified; @@ -689,7 +691,8 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && - Attrs.SYCLIntelFPGALoopCountAverageEnable == 0 && + Attrs.SYCLIntelFPGALoopCountEnable == 0 && + Attrs.SYCLIntelFPGALoopCountVariation == nullptr && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -1041,11 +1044,18 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, } if (const auto *IntelFPGALoopCountAvg = - dyn_cast(A)) { - setSYCLIntelFPGALoopCountAvgEnable(); - setSYCLIntelFPGALoopControlAverage(IntelFPGALoopCountAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); + dyn_cast(A)) { + setSYCLIntelFPGALoopCountEnable(); + setSYCLIntelFPGALoopCount(IntelFPGALoopCountAvg->getNTripCount() + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); + SYCLIntelFPGALoopCountAttr::CountKind K = + IntelFPGALoopCountAvg->getCountKind(); + const char *var = + IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" + : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" + : "llvm.loop.intel.loopcount_avg"; + setSYCLIntelFPGALoopCountVariation(var); } if (const auto *IntelFPGALoopCoalesce = diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index b9cbd2a19c57c..6920874ffea1b 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -118,10 +118,12 @@ struct LoopAttributes { unsigned SYCLMaxConcurrencyNThreads; /// Flag for llvm.loop.intel.loopcount_average metadata. - bool SYCLIntelFPGALoopCountAverageEnable; + bool SYCLIntelFPGALoopCountEnable; /// Value for llvm.loop.intel.loopcount_average metadata. - unsigned SYCLIntelFPGALoopControlAverage; + unsigned SYCLIntelFPGALoopCount; + + const char *SYCLIntelFPGALoopCountVariation; /// Flag for llvm.loop.coalesce metadata. bool SYCLLoopCoalesceEnable; @@ -411,13 +413,18 @@ class LoopInfoStack { } /// Set flag of loop_control_avg for the next loop pushed. - void setSYCLIntelFPGALoopCountAvgEnable() { - StagedAttrs.SYCLIntelFPGALoopCountAverageEnable = true; + void setSYCLIntelFPGALoopCountEnable() { + StagedAttrs.SYCLIntelFPGALoopCountEnable = true; + } + + /// Set value of loop control average for the next loop pushed. + void setSYCLIntelFPGALoopCount(unsigned C) { + StagedAttrs.SYCLIntelFPGALoopCount = C; } /// Set value of loop control average for the next loop pushed. - void setSYCLIntelFPGALoopControlAverage(unsigned C) { - StagedAttrs.SYCLIntelFPGALoopControlAverage = C; + void setSYCLIntelFPGALoopCountVariation(const char *var) { + StagedAttrs.SYCLIntelFPGALoopCountVariation = var; } /// Set the unroll count for the next loop pushed. diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 0e8103e4377f1..fcf2f159d5f46 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -243,6 +243,30 @@ static Attr *handleIntelFPGAIVDepAttr(Sema &S, Stmt *St, const ParsedAttr &A) { NumArgs == 2 ? A.getArgAsExpr(1) : nullptr); } +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.getAttrName() << AANT_ArgumentIntegerConstant + << E->getSourceRange(); + return nullptr; + } + + if (ArgVal->getSExtValue() < 0) { + S.Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) + << A.getAttrName() << /* non-negative */ 1; + return nullptr; + } + } + return new (S.Context) + SYCLIntelFPGALoopCountAttr(S.Context, A, A.getArgAsExpr(0)); +} + static Attr *handleIntelFPGANofusionAttr(Sema &S, Stmt *St, const ParsedAttr &A) { return new (S.Context) SYCLIntelFPGANofusionAttr(S.Context, A); @@ -558,7 +582,7 @@ static void CheckForIncompatibleSYCLLoopAttributes( Attrs); CheckForDuplicationSYCLLoopAttribute( S, Attrs); - CheckForDuplicationSYCLLoopAttribute(S, Attrs); + CheckForDuplicationSYCLLoopAttribute(S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs, false); CheckRedundantSYCLIntelFPGAIVDepAttrs(S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs); @@ -688,8 +712,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A, case ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations: return handleIntelFPGALoopAttr(S, St, A); - case ParsedAttr::AT_SYCLIntelFPGALoopCountAvg: - return handleIntelFPGALoopAttr(S, St, A); + case ParsedAttr::AT_SYCLIntelFPGALoopCount: + return handleIntelFPGALoopCountAttr(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 ded1969969e6f..522e4c532d3df 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1093,8 +1093,8 @@ namespace { const SYCLIntelFPGASpeculatedIterationsAttr * TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI); - const SYCLIntelFPGALoopCountAvgAttr *TransformSYCLIntelFPGALoopCountAvgAttr( - const SYCLIntelFPGALoopCountAvgAttr *SI); + const SYCLIntelFPGALoopCountAttr *TransformSYCLIntelFPGALoopCountAttr( + const SYCLIntelFPGALoopCountAttr *SI); ExprResult TransformPredefinedExpr(PredefinedExpr *E); ExprResult TransformDeclRefExpr(DeclRefExpr *E); @@ -1620,12 +1620,12 @@ TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr( *SI, TransformedExpr); } -const SYCLIntelFPGALoopCountAvgAttr * -TemplateInstantiator::TransformSYCLIntelFPGALoopCountAvgAttr( - const SYCLIntelFPGALoopCountAvgAttr *LCA) { +const SYCLIntelFPGALoopCountAttr * +TemplateInstantiator::TransformSYCLIntelFPGALoopCountAttr( + const SYCLIntelFPGALoopCountAttr *LCA) { Expr *TransformedExpr = getDerived().TransformExpr(LCA->getNTripCount()).get(); - return getSema().BuildSYCLIntelFPGALoopAttr( + return getSema().BuildSYCLIntelFPGALoopAttr( *LCA, TransformedExpr); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index cfbf6d09436e3..c05af5105c582 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -112,14 +112,14 @@ void speculated_iterations() { } template -void loop_count_avg() { +void loop_count_control() { int a[10]; // CHECK: ![[MD_LCA]] = distinct !{![[MD_LCA]], ![[MP:[0-9]+]], ![[MD_lca:[0-9]+]]} - // CHECK-NEXT: ![[MD_lca]] = !{!"llvm.loop.intel.loopcount_average", i32 12} + // CHECK-NEXT: ![[MD_lca]] = !{!"llvm.loop.intel.loopcount_avg", i32 12} [[intel::loop_count_avg(A)]] for (int i = 0; i != 10; ++i) a[i] = 0; // CHECK: ![[MD_LCA_1]] = distinct !{![[MD_LCA_1]], ![[MP:[0-9]+]], ![[MD_lca_1:[0-9]+]]} - // CHECK-NEXT: ![[MD_lca_1]] = !{!"llvm.loop.intel.loopcount_average", i32 4} + // CHECK-NEXT: ![[MD_lca_1]] = !{!"llvm.loop.intel.loopcount_avg", i32 4} [[intel::loop_count_avg(4)]] for (int i = 0; i != 10; ++i) a[i] = 0; } @@ -138,7 +138,7 @@ int main() { loop_coalesce<2>(); max_interleaving<3>(); speculated_iterations<4>(); - loop_count_avg<12>(); + loop_count_control<12>(); }); return 0; } diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 97d6a3819495c..6668d75941b31 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -427,7 +427,7 @@ void max_concurrency_dependent() { } template -void loop_count_avg_dependent() { +void loop_count_control_dependent() { int a[10]; //expected-error@+1{{'loop_count_avg' attribute requires a non-negative integral compile time constant expression}} @@ -460,8 +460,8 @@ int main() { max_concurrency_dependent<1, 4, -2>(); //expected-note@-1 +{{in instantiation of function template specialization}} - loop_count_avg_dependent<3, 2, -1>(); - //expected-note@-1{{in instantiation of function template specialization 'loop_count_avg_dependent<3, 2, -1>' requested here}} + 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}} }); }); From 80f372a8534659d5052f85c30ea8e14e4b7933a3 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 19 Apr 2021 08:51:25 -0400 Subject: [PATCH 17/28] Added min/max Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 919f6e57a2a13..882ca7890f6d7 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1049,8 +1049,6 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, setSYCLIntelFPGALoopCount(IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); - SYCLIntelFPGALoopCountAttr::CountKind K = - IntelFPGALoopCountAvg->getCountKind(); const char *var = IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" From d2884f0193b9750e72b53759b5c55ddec6c4935a Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 19 Apr 2021 08:59:40 -0400 Subject: [PATCH 18/28] Indentation Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Sema/Sema.h | 3 +-- clang/lib/CodeGen/CGLoopInfo.cpp | 9 +++++---- clang/lib/Sema/SemaStmtAttr.cpp | 2 +- clang/lib/Sema/SemaTemplateInstantiate.cpp | 4 ++-- 4 files changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index aabea3a266b47..9a1afd1309c63 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13425,8 +13425,7 @@ FPGALoopAttrT *Sema::BuildSYCLIntelFPGALoopAttr(const AttributeCommonInfo &A, ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving || A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations || - A.getParsedKind() == - ParsedAttr::AT_SYCLIntelFPGALoopCount) { + A.getParsedKind() == ParsedAttr::AT_SYCLIntelFPGALoopCount) { if (Val < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << A.getAttrName() << /* non-negative */ 1; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 882ca7890f6d7..f29832b8dd09d 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1049,10 +1049,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, setSYCLIntelFPGALoopCount(IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); - const char *var = - IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" - : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" - : "llvm.loop.intel.loopcount_avg"; + const char *var = IntelFPGALoopCountAvg->isMax() + ? "llvm.loop.intel.loopcount_max" + : IntelFPGALoopCountAvg->isMin() + ? "llvm.loop.intel.loopcount_min" + : "llvm.loop.intel.loopcount_avg"; setSYCLIntelFPGALoopCountVariation(var); } diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index fcf2f159d5f46..7292e01dd3667 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -244,7 +244,7 @@ static Attr *handleIntelFPGAIVDepAttr(Sema &S, Stmt *St, const ParsedAttr &A) { } static SYCLIntelFPGALoopCountAttr * - handleIntelFPGALoopCountAttr(Sema &S, Stmt *St, const ParsedAttr &A) { +handleIntelFPGALoopCountAttr(Sema &S, Stmt *St, const ParsedAttr &A) { Expr *E = A.getArgAsExpr(0); if (E && !E->isInstantiationDependent()) { Optional ArgVal = diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index 522e4c532d3df..ba407993e90e9 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1093,8 +1093,8 @@ namespace { const SYCLIntelFPGASpeculatedIterationsAttr * TransformSYCLIntelFPGASpeculatedIterationsAttr( const SYCLIntelFPGASpeculatedIterationsAttr *SI); - const SYCLIntelFPGALoopCountAttr *TransformSYCLIntelFPGALoopCountAttr( - const SYCLIntelFPGALoopCountAttr *SI); + const SYCLIntelFPGALoopCountAttr * + TransformSYCLIntelFPGALoopCountAttr(const SYCLIntelFPGALoopCountAttr *SI); ExprResult TransformPredefinedExpr(PredefinedExpr *E); ExprResult TransformDeclRefExpr(DeclRefExpr *E); From 5ab429b506894b53a6be9479146905307afe8574 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 13:35:10 -0400 Subject: [PATCH 19/28] Added all 3 variants Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/AttrDocs.td | 12 +++++-- clang/lib/CodeGen/CGLoopInfo.cpp | 29 +++++++++------- clang/lib/CodeGen/CGLoopInfo.h | 24 +++++++------ clang/lib/Sema/SemaStmtAttr.cpp | 38 ++++++++++++++++++++- clang/test/CodeGenSYCL/intel-fpga-loops.cpp | 17 ++++++--- 5 files changed, 89 insertions(+), 31 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index b0675e73a15b7..a98defb17f4c5 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2955,7 +2955,7 @@ def SYCLIntelFPGALoopCountAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "intel::loop_count_min, intel::loop_count_max, intel::loop_count_avg"; let Content = [{ -The "loop_count" pragma specifies the minimum, maximum, or average number of +The "loop_count" attribute specifies the minimum, maximum, or average number of iterations for a for loop. In addition, a list of commonly occurring values can be specified to help the compiler generate multiple versions and perform complete unrolling. @@ -2967,11 +2967,19 @@ complete unrolling. [[intel::loop_count_min(40)] for (int i = 0; i < 10; ++i) a[i] = 0; } - void foo() { + void zoo() { int a[10]; [[intel::loop_count_max(40)] for (int i = 0; i < 10; ++i) a[i] = 0; } + void goo() { + int a[10]; + [[intel::loop_count_min(10)] + [[intel::loop_count_max(40)] + [[intel::loop_count_avg(15)] + for (int i = 0; i < 10; ++i) a[i] = 0; + } + template void bar() { [[intel::loop_count_avg(N)]] for(;;) { } diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index f29832b8dd09d..b5c957f48dc48 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -606,11 +606,14 @@ MDNode *LoopInfo::createMetadata( } if (Attrs.SYCLIntelFPGALoopCountEnable) { + for (int i = 0; i < Attrs.SYCLIntelFPGALoopCountVariant.size(); i++) { Metadata *Vals[] = { - MDString::get(Ctx, Attrs.SYCLIntelFPGALoopCountVariation), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), Attrs.SYCLIntelFPGALoopCount))}; + MDString::get(Ctx, Attrs.SYCLIntelFPGALoopCountVariant[i]), + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), + Attrs.SYCLIntelFPGALoopCountValue[i]))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); + } } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), AdditionalLoopProperties.end()); @@ -629,11 +632,11 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), SYCLSpeculatedIterationsNIterations(0), - SYCLIntelFPGALoopCountEnable(false), - SYCLIntelFPGALoopCountVariation(nullptr), UnrollCount(0), - UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), - PipelineDisabled(false), PipelineInitiationInterval(0), - SYCLNofusionEnable(false), MustProgress(false) {} + SYCLIntelFPGALoopCountEnable(false), SYCLIntelFPGALoopCountValue(0), + SYCLIntelFPGALoopCountVariant(0), UnrollCount(0), UnrollAndJamCount(0), + DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), + PipelineInitiationInterval(0), SYCLNofusionEnable(false), + MustProgress(false) {} void LoopAttributes::clear() { IsParallel = false; @@ -653,7 +656,8 @@ void LoopAttributes::clear() { SYCLSpeculatedIterationsEnable = false; SYCLSpeculatedIterationsNIterations = 0; SYCLIntelFPGALoopCountEnable = false; - SYCLIntelFPGALoopCountVariation = nullptr; + SYCLIntelFPGALoopCountVariant.clear(); + SYCLIntelFPGALoopCountValue.clear(); UnrollCount = 0; UnrollAndJamCount = 0; VectorizeEnable = LoopAttributes::Unspecified; @@ -692,7 +696,8 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && Attrs.SYCLIntelFPGALoopCountEnable == 0 && - Attrs.SYCLIntelFPGALoopCountVariation == nullptr && + Attrs.SYCLIntelFPGALoopCountVariant.empty() && + Attrs.SYCLIntelFPGALoopCountValue.empty() && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -1046,7 +1051,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopCountAvg = dyn_cast(A)) { setSYCLIntelFPGALoopCountEnable(); - setSYCLIntelFPGALoopCount(IntelFPGALoopCountAvg->getNTripCount() + setSYCLIntelFPGALoopCountValue(IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); const char *var = IntelFPGALoopCountAvg->isMax() @@ -1054,7 +1059,7 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" : "llvm.loop.intel.loopcount_avg"; - setSYCLIntelFPGALoopCountVariation(var); + setSYCLIntelFPGALoopCountVariant(var); } if (const auto *IntelFPGALoopCoalesce = diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 6920874ffea1b..c3414441dc9a3 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/DenseMap.h" namespace llvm { class BasicBlock; @@ -117,13 +118,14 @@ struct LoopAttributes { /// Value for llvm.loop.max_concurrency.count metadata. unsigned SYCLMaxConcurrencyNThreads; - /// Flag for llvm.loop.intel.loopcount_average metadata. + /// Flag for llvm.loop.intel.loopcount metadata. bool SYCLIntelFPGALoopCountEnable; - /// Value for llvm.loop.intel.loopcount_average metadata. - unsigned SYCLIntelFPGALoopCount; + /// Value for llvm.loop.intel.loopcount value metadata. + llvm::SmallVector SYCLIntelFPGALoopCountValue; - const char *SYCLIntelFPGALoopCountVariation; + /// Value for llvm.loop.intel.loopcount variant(min/max/avg) metadata. + llvm::SmallVector SYCLIntelFPGALoopCountVariant; /// Flag for llvm.loop.coalesce metadata. bool SYCLLoopCoalesceEnable; @@ -412,19 +414,19 @@ class LoopInfoStack { StagedAttrs.SYCLSpeculatedIterationsNIterations = C; } - /// Set flag of loop_control_avg for the next loop pushed. + /// Set flag of loopcount for the next loop pushed. void setSYCLIntelFPGALoopCountEnable() { StagedAttrs.SYCLIntelFPGALoopCountEnable = true; } - /// Set value of loop control average for the next loop pushed. - void setSYCLIntelFPGALoopCount(unsigned C) { - StagedAttrs.SYCLIntelFPGALoopCount = C; + /// Set value of loopcount value for the next loop pushed. + void setSYCLIntelFPGALoopCountValue(unsigned C) { + StagedAttrs.SYCLIntelFPGALoopCountValue.push_back(C); } - /// Set value of loop control average for the next loop pushed. - void setSYCLIntelFPGALoopCountVariation(const char *var) { - StagedAttrs.SYCLIntelFPGALoopCountVariation = var; + /// Set value of loopcount variant for the next loop pushed. + void setSYCLIntelFPGALoopCountVariant(const char *var) { + StagedAttrs.SYCLIntelFPGALoopCountVariant.push_back(var); } /// Set the unroll count for the next loop pushed. diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 7292e01dd3667..f433765f234b6 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -243,6 +243,42 @@ static Attr *handleIntelFPGAIVDepAttr(Sema &S, Stmt *St, const ParsedAttr &A) { NumArgs == 2 ? A.getArgAsExpr(1) : nullptr); } +static void +CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, + ArrayRef Attrs) { + // Create a list of SYCLIntelFPGALoopCount attributes only. + SmallVector OnlyLoopCountAttrs; + llvm::transform( + Attrs, std::back_inserter(OnlyLoopCountAttrs), [](const Attr *A) { + return dyn_cast_or_null(A); + }); + OnlyLoopCountAttrs.erase( + std::remove(OnlyLoopCountAttrs.begin(), OnlyLoopCountAttrs.end(), + static_cast(nullptr)), + OnlyLoopCountAttrs.end()); + if (OnlyLoopCountAttrs.empty()) + return; + + unsigned int MinCount, MaxCount, AvgCount = 0; + for (const auto *A : OnlyLoopCountAttrs) { + const SYCLIntelFPGALoopCountAttr *At = + dyn_cast(A); + switch (At->getCountKind()) { + case SYCLIntelFPGALoopCountAttr::CountKind::loop_count_min: + MinCount++; + break; + case SYCLIntelFPGALoopCountAttr::CountKind::loop_count_max: + MaxCount++; + break; + case SYCLIntelFPGALoopCountAttr::CountKind::loop_count_avg: + AvgCount++; + break; + } + if (MinCount > 1 || MaxCount > 1 || AvgCount >1 ) + S.Diag(A->getLocation(), diag::err_sycl_loop_attr_duplication) << 1 << A; + } +} + static SYCLIntelFPGALoopCountAttr * handleIntelFPGALoopCountAttr(Sema &S, Stmt *St, const ParsedAttr &A) { Expr *E = A.getArgAsExpr(0); @@ -582,7 +618,7 @@ static void CheckForIncompatibleSYCLLoopAttributes( Attrs); CheckForDuplicationSYCLLoopAttribute( S, Attrs); - CheckForDuplicationSYCLLoopAttribute(S, Attrs); + CheckForDuplicateSYCLIntelLoopCountAttrs(S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs, false); CheckRedundantSYCLIntelFPGAIVDepAttrs(S, Attrs); CheckForDuplicationSYCLLoopAttribute(S, Attrs); diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index c05af5105c582..a0dfafdbcd92f 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -16,6 +16,7 @@ // CHECK: br label %for.cond2, !llvm.loop ![[MD_SI_2:[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]+]] void disable_loop_pipelining() { int a[10]; @@ -114,13 +115,19 @@ void speculated_iterations() { template void loop_count_control() { int a[10]; - // CHECK: ![[MD_LCA]] = distinct !{![[MD_LCA]], ![[MP:[0-9]+]], ![[MD_lca:[0-9]+]]} - // CHECK-NEXT: ![[MD_lca]] = !{!"llvm.loop.intel.loopcount_avg", i32 12} + // CHECK: ![[MD_LCA]] = distinct !{![[MD_LCA]], ![[MP:[0-9]+]], ![[MD_loop_count_avg:[0-9]+]]} + // CHECK-NEXT: ![[MD_loop_count_avg]] = !{!"llvm.loop.intel.loopcount_avg", i32 12} [[intel::loop_count_avg(A)]] for (int i = 0; i != 10; ++i) a[i] = 0; - // CHECK: ![[MD_LCA_1]] = distinct !{![[MD_LCA_1]], ![[MP:[0-9]+]], ![[MD_lca_1:[0-9]+]]} - // CHECK-NEXT: ![[MD_lca_1]] = !{!"llvm.loop.intel.loopcount_avg", i32 4} - [[intel::loop_count_avg(4)]] for (int i = 0; i != 10; ++i) + // CHECK: ![[MD_LCA_1]] = distinct !{![[MD_LCA_1]], ![[MP:[0-9]+]], ![[MD_loop_count_max:[0-9]+]]} + // CHECK-NEXT: ![[MD_loop_count_max]] = !{!"llvm.loop.intel.loopcount_max", i32 4} + [[intel::loop_count_max(4)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + // CHECK: ![[MD_LCA_2]] = distinct !{![[MD_LCA_2]], ![[MP:[0-9]+]], ![[MD_loop_count_min:[0-9]+]], ![[MD_loop_count_max_1:[0-9]+]], ![[MD_loop_count_avg_1:[0-9]+]]} + // CHECK: ![[MD_loop_count_min]] = !{!"llvm.loop.intel.loopcount_min", i32 4} + // CHECK: ![[MD_loop_count_max_1]] = !{!"llvm.loop.intel.loopcount_max", i32 40} + // CHECK-NEXT: ![[MD_loop_count_avg_1]] = !{!"llvm.loop.intel.loopcount_avg", i32 21} + [[intel::loop_count_min(4)]] [[intel::loop_count_max(40)]] [[intel::loop_count_avg(21)]] for (int i = 0; i != 10; ++i) a[i] = 0; } From 86ae532007738c2dcf971a83d09080a58e11cb85 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 15:57:41 -0400 Subject: [PATCH 20/28] Added all 3 variants Signed-off-by: Zahira Ammarguellat --- clang/lib/Sema/SemaStmtAttr.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index c7897c7cbe804..d9ca50ef1eff8 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -259,7 +259,9 @@ CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, if (OnlyLoopCountAttrs.empty()) return; - unsigned int MinCount, MaxCount, AvgCount = 0; + unsigned int MinCount = 0; + unsigned int MaxCount = 0; + unsigned int AvgCount = 0; for (const auto *A : OnlyLoopCountAttrs) { const SYCLIntelFPGALoopCountAttr *At = dyn_cast(A); From 3044f8818acb66156e6399e179e8ad683fc6e9a1 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 16:20:23 -0400 Subject: [PATCH 21/28] Indentation Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 31 +++++++++++++++---------------- clang/lib/CodeGen/CGLoopInfo.h | 2 +- clang/lib/Sema/SemaStmtAttr.cpp | 24 ++++++++++++------------ 3 files changed, 28 insertions(+), 29 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index b5c957f48dc48..d8527d05b8fce 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -607,12 +607,12 @@ MDNode *LoopInfo::createMetadata( if (Attrs.SYCLIntelFPGALoopCountEnable) { for (int i = 0; i < Attrs.SYCLIntelFPGALoopCountVariant.size(); i++) { - Metadata *Vals[] = { - MDString::get(Ctx, Attrs.SYCLIntelFPGALoopCountVariant[i]), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLIntelFPGALoopCountValue[i]))}; - LoopProperties.push_back(MDNode::get(Ctx, Vals)); + Metadata *Vals[] = { + MDString::get(Ctx, Attrs.SYCLIntelFPGALoopCountVariant[i]), + ConstantAsMetadata::get( + ConstantInt::get(llvm::Type::getInt32Ty(Ctx), + Attrs.SYCLIntelFPGALoopCountValue[i]))}; + LoopProperties.push_back(MDNode::get(Ctx, Vals)); } } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), @@ -697,9 +697,9 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLSpeculatedIterationsNIterations == 0 && Attrs.SYCLIntelFPGALoopCountEnable == 0 && Attrs.SYCLIntelFPGALoopCountVariant.empty() && - Attrs.SYCLIntelFPGALoopCountValue.empty() && - Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && - !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && + Attrs.SYCLIntelFPGALoopCountValue.empty() && Attrs.UnrollCount == 0 && + Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && + Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && Attrs.VectorizeEnable == LoopAttributes::Unspecified && Attrs.UnrollEnable == LoopAttributes::Unspecified && @@ -1052,13 +1052,12 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(A)) { setSYCLIntelFPGALoopCountEnable(); setSYCLIntelFPGALoopCountValue(IntelFPGALoopCountAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); - const char *var = IntelFPGALoopCountAvg->isMax() - ? "llvm.loop.intel.loopcount_max" - : IntelFPGALoopCountAvg->isMin() - ? "llvm.loop.intel.loopcount_min" - : "llvm.loop.intel.loopcount_avg"; + ->getIntegerConstantExpr(Ctx) + ->getSExtValue()); + const char *var = + IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" + : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" + : "llvm.loop.intel.loopcount_avg"; setSYCLIntelFPGALoopCountVariant(var); } diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index c3414441dc9a3..af0449826c561 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/DenseMap.h" #include "llvm/ADT/SmallVector.h" #include "llvm/IR/DebugLoc.h" #include "llvm/IR/Value.h" #include "llvm/Support/Compiler.h" -#include "llvm/ADT/DenseMap.h" namespace llvm { class BasicBlock; diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index d9ca50ef1eff8..38fc94aa93cc5 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -246,18 +246,18 @@ static Attr *handleIntelFPGAIVDepAttr(Sema &S, Stmt *St, const ParsedAttr &A) { static void CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, ArrayRef Attrs) { - // Create a list of SYCLIntelFPGALoopCount attributes only. - SmallVector OnlyLoopCountAttrs; - llvm::transform( - Attrs, std::back_inserter(OnlyLoopCountAttrs), [](const Attr *A) { - return dyn_cast_or_null(A); - }); - OnlyLoopCountAttrs.erase( + // Create a list of SYCLIntelFPGALoopCount attributes only. + SmallVector OnlyLoopCountAttrs; + llvm::transform( + Attrs, std::back_inserter(OnlyLoopCountAttrs), [](const Attr *A) { + return dyn_cast_or_null(A); + }); + OnlyLoopCountAttrs.erase( std::remove(OnlyLoopCountAttrs.begin(), OnlyLoopCountAttrs.end(), - static_cast(nullptr)), + static_cast(nullptr)), OnlyLoopCountAttrs.end()); - if (OnlyLoopCountAttrs.empty()) - return; + if (OnlyLoopCountAttrs.empty()) + return; unsigned int MinCount = 0; unsigned int MaxCount = 0; @@ -276,8 +276,8 @@ CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, AvgCount++; break; } - if (MinCount > 1 || MaxCount > 1 || AvgCount >1 ) - S.Diag(A->getLocation(), diag::err_sycl_loop_attr_duplication) << 1 << A; + if (MinCount > 1 || MaxCount > 1 || AvgCount > 1) + S.Diag(A->getLocation(), diag::err_sycl_loop_attr_duplication) << 1 << A; } } From 453b654d9c7602e82cf55c550cf3e1262c1d3a45 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 16:43:53 -0400 Subject: [PATCH 22/28] Indentation Signed-off-by: Zahira Ammarguellat --- clang/test/SemaSYCL/intel-fpga-loops.cpp | 26 ++++++++++++++++++++++++ 1 file changed, 26 insertions(+) diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 6668d75941b31..ce26f4d977780 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -366,6 +366,12 @@ void loop_attrs_compatibility() { [[intel::loop_count_avg(8)]] for (int i = 0; i != 10; ++i) a[i] = 0; + [[intel::loop_count_min(8)]] + for (int i = 0; i != 10; ++i) + a[i] = 0; + [[intel::loop_count_max(8)]] + for (int i = 0; i != 10; ++i) + a[i] = 0; } template @@ -435,11 +441,31 @@ void loop_count_control_dependent() { for (int i = 0; i != 10; ++i) a[i] = 0; + //expected-error@+1{{'loop_count_min' attribute requires a non-negative integral compile time constant expression}} + [[intel::loop_count_min(C)]] + for (int i = 0; i != 10; ++i) + a[i] = 0; + + //expected-error@+1{{'loop_count_max' attribute requires a non-negative integral compile time constant expression}} + [[intel::loop_count_max(C)]] + for (int i = 0; i != 10; ++i) + a[i] = 0; + [[intel::loop_count_avg(A)]] //expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_count_avg'}} [[intel::loop_count_avg(B)]] for (int i = 0; i != 10; ++i) a[i] = 0; + [[intel::loop_count_min(A)]] + //expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_count_min'}} + [[intel::loop_count_min(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + + [[intel::loop_count_max(A)]] + //expected-error@+1{{duplicate Intel FPGA loop attribute 'loop_count_max'}} + [[intel::loop_count_max(B)]] for (int i = 0; i != 10; ++i) + a[i] = 0; + } int main() { From 9d46fc03ef3f11539155e82e18d89eb92de14d01 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 21 Apr 2021 16:49:45 -0400 Subject: [PATCH 23/28] Indentation Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index d8527d05b8fce..8e762686062f9 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1054,10 +1054,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, setSYCLIntelFPGALoopCountValue(IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue()); - const char *var = - IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" - : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" - : "llvm.loop.intel.loopcount_avg"; + const char *var = IntelFPGALoopCountAvg->isMax() + ? "llvm.loop.intel.loopcount_max" + : IntelFPGALoopCountAvg->isMin() + ? "llvm.loop.intel.loopcount_min" + : "llvm.loop.intel.loopcount_avg"; setSYCLIntelFPGALoopCountVariant(var); } From 1a2063526da1b1c1888c7690a89bbb9dcfc856c2 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 27 Apr 2021 12:26:47 -0400 Subject: [PATCH 24/28] Fixes after review Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/Attr.td | 10 +------ clang/include/clang/Basic/AttrDocs.td | 21 ++++++------- clang/lib/CodeGen/CGLoopInfo.cpp | 43 +++++++++++---------------- clang/lib/CodeGen/CGLoopInfo.h | 27 ++++------------- clang/lib/Sema/SemaStmtAttr.cpp | 20 +++---------- 5 files changed, 39 insertions(+), 82 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 9a03711969069..d5ab575f50417 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1986,15 +1986,7 @@ def SYCLIntelFPGALoopCount : StmtAttr { Accessor<"isAvg", [CXX11<"intel", "loop_count_avg">]>]; let Args = [ExprArgument<"NTripCount">]; let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost]; - let HasCustomTypeTransform = 1; - let AdditionalMembers = [{ - enum CountKind { loop_count_min, loop_count_max, loop_count_avg }; - CountKind getCountKind() const { - return isMin() ? loop_count_min : - isMax() ? loop_count_max : - loop_count_avg; - } - }]; + let IsStmtDependent = 1; let Documentation = [SYCLIntelFPGALoopCountAttrDocs]; } diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index a98defb17f4c5..16cd62b0dcf43 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2955,29 +2955,30 @@ def SYCLIntelFPGALoopCountAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "intel::loop_count_min, intel::loop_count_max, intel::loop_count_avg"; let Content = [{ -The "loop_count" attribute specifies the minimum, maximum, or average number of -iterations for a for loop. In addition, a list of commonly occurring values -can be specified to help the compiler generate multiple versions and perform -complete unrolling. +The loop count attribute specifies the minimum, maximum, or average number of +iterations for a ``for`` loop. These are hints that the user specify that can be +used by some of the loop optimization to make decisions such as if the loop +should be unrolled. It is a way for the user to provide some information without +using PGO. .. code-block:: c++ void foo() { int a[10]; - [[intel::loop_count_min(40)] for (int i = 0; i < 10; ++i) a[i] = 0; + [[intel::loop_count_min(4)] for (int i = 0; i < n; ++i) a[i] = 0; } void zoo() { int a[10]; - [[intel::loop_count_max(40)] for (int i = 0; i < 10; ++i) a[i] = 0; + [[intel::loop_count_max(10)] for (int i = 0; i < m; ++i) a[i] = 0; } void goo() { int a[10]; - [[intel::loop_count_min(10)] - [[intel::loop_count_max(40)] - [[intel::loop_count_avg(15)] - for (int i = 0; i < 10; ++i) a[i] = 0; + [[intel::loop_count_min(3)] + [[intel::loop_count_max(10)] + [[intel::loop_count_avg(5)] + for (int i = start; i < end; ++i) a[i] = 0; } template diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 8e762686062f9..f9280171055b9 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -605,13 +605,11 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLIntelFPGALoopCountEnable) { - for (int i = 0; i < Attrs.SYCLIntelFPGALoopCountVariant.size(); i++) { - Metadata *Vals[] = { - MDString::get(Ctx, Attrs.SYCLIntelFPGALoopCountVariant[i]), - ConstantAsMetadata::get( - ConstantInt::get(llvm::Type::getInt32Ty(Ctx), - Attrs.SYCLIntelFPGALoopCountValue[i]))}; + if (Attrs.SYCLIntelFPGAVariantCount.size() > 0) { + for (auto &VC : Attrs.SYCLIntelFPGAVariantCount) { + Metadata *Vals[] = {MDString::get(Ctx, VC.first), + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), VC.second))}; LoopProperties.push_back(MDNode::get(Ctx, Vals)); } } @@ -631,9 +629,8 @@ LoopAttributes::LoopAttributes(bool IsParallel) SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false), SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0), SYCLSpeculatedIterationsEnable(false), - SYCLSpeculatedIterationsNIterations(0), - SYCLIntelFPGALoopCountEnable(false), SYCLIntelFPGALoopCountValue(0), - SYCLIntelFPGALoopCountVariant(0), UnrollCount(0), UnrollAndJamCount(0), + SYCLSpeculatedIterationsNIterations(0), SYCLIntelFPGAVariantCount(false), + UnrollCount(0), UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false), PipelineInitiationInterval(0), SYCLNofusionEnable(false), MustProgress(false) {} @@ -655,9 +652,7 @@ void LoopAttributes::clear() { SYCLMaxInterleavingNInvocations = 0; SYCLSpeculatedIterationsEnable = false; SYCLSpeculatedIterationsNIterations = 0; - SYCLIntelFPGALoopCountEnable = false; - SYCLIntelFPGALoopCountVariant.clear(); - SYCLIntelFPGALoopCountValue.clear(); + SYCLIntelFPGAVariantCount.clear(); UnrollCount = 0; UnrollAndJamCount = 0; VectorizeEnable = LoopAttributes::Unspecified; @@ -695,9 +690,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && - Attrs.SYCLIntelFPGALoopCountEnable == 0 && - Attrs.SYCLIntelFPGALoopCountVariant.empty() && - Attrs.SYCLIntelFPGALoopCountValue.empty() && Attrs.UnrollCount == 0 && + Attrs.SYCLIntelFPGAVariantCount.empty() && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -1050,16 +1043,14 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopCountAvg = dyn_cast(A)) { - setSYCLIntelFPGALoopCountEnable(); - setSYCLIntelFPGALoopCountValue(IntelFPGALoopCountAvg->getNTripCount() - ->getIntegerConstantExpr(Ctx) - ->getSExtValue()); - const char *var = IntelFPGALoopCountAvg->isMax() - ? "llvm.loop.intel.loopcount_max" - : IntelFPGALoopCountAvg->isMin() - ? "llvm.loop.intel.loopcount_min" - : "llvm.loop.intel.loopcount_avg"; - setSYCLIntelFPGALoopCountVariant(var); + unsigned int count = IntelFPGALoopCountAvg->getNTripCount() + ->getIntegerConstantExpr(Ctx) + ->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); } if (const auto *IntelFPGALoopCoalesce = diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index af0449826c561..6fbc6dcc9c6cf 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -118,14 +118,9 @@ struct LoopAttributes { /// Value for llvm.loop.max_concurrency.count metadata. unsigned SYCLMaxConcurrencyNThreads; - /// Flag for llvm.loop.intel.loopcount metadata. - bool SYCLIntelFPGALoopCountEnable; - - /// Value for llvm.loop.intel.loopcount value metadata. - llvm::SmallVector SYCLIntelFPGALoopCountValue; - - /// Value for llvm.loop.intel.loopcount variant(min/max/avg) metadata. - llvm::SmallVector SYCLIntelFPGALoopCountVariant; + /// Value for count variant (min/max/avg) and count metadata. + llvm::SmallVector, 2> + SYCLIntelFPGAVariantCount; /// Flag for llvm.loop.coalesce metadata. bool SYCLLoopCoalesceEnable; @@ -414,19 +409,9 @@ class LoopInfoStack { StagedAttrs.SYCLSpeculatedIterationsNIterations = C; } - /// Set flag of loopcount for the next loop pushed. - void setSYCLIntelFPGALoopCountEnable() { - StagedAttrs.SYCLIntelFPGALoopCountEnable = true; - } - - /// Set value of loopcount value for the next loop pushed. - void setSYCLIntelFPGALoopCountValue(unsigned C) { - StagedAttrs.SYCLIntelFPGALoopCountValue.push_back(C); - } - - /// Set value of loopcount variant for the next loop pushed. - void setSYCLIntelFPGALoopCountVariant(const char *var) { - StagedAttrs.SYCLIntelFPGALoopCountVariant.push_back(var); + /// Set value of variant and loop count for the next loop pushed. + void setSYCLIntelFPGAVariantCount(const char *var, unsigned int C) { + StagedAttrs.SYCLIntelFPGAVariantCount.push_back({var, C}); } /// Set the unroll count for the next loop pushed. diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index 38fc94aa93cc5..fad4c2533c638 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/clang/lib/Sema/SemaStmtAttr.cpp @@ -263,19 +263,8 @@ CheckForDuplicateSYCLIntelLoopCountAttrs(Sema &S, unsigned int MaxCount = 0; unsigned int AvgCount = 0; for (const auto *A : OnlyLoopCountAttrs) { - const SYCLIntelFPGALoopCountAttr *At = - dyn_cast(A); - switch (At->getCountKind()) { - case SYCLIntelFPGALoopCountAttr::CountKind::loop_count_min: - MinCount++; - break; - case SYCLIntelFPGALoopCountAttr::CountKind::loop_count_max: - MaxCount++; - break; - case SYCLIntelFPGALoopCountAttr::CountKind::loop_count_avg: - AvgCount++; - break; - } + const auto *At = dyn_cast(A); + At->isMin() ? MinCount++ : At->isMax() ? MaxCount++ : AvgCount++; if (MinCount > 1 || MaxCount > 1 || AvgCount > 1) S.Diag(A->getLocation(), diag::err_sycl_loop_attr_duplication) << 1 << A; } @@ -290,14 +279,13 @@ handleIntelFPGALoopCountAttr(Sema &S, Stmt *St, const ParsedAttr &A) { if (!ArgVal) { S.Diag(E->getExprLoc(), diag::err_attribute_argument_type) - << A.getAttrName() << AANT_ArgumentIntegerConstant - << E->getSourceRange(); + << A << AANT_ArgumentIntegerConstant << E->getSourceRange(); return nullptr; } if (ArgVal->getSExtValue() < 0) { S.Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) - << A.getAttrName() << /* non-negative */ 1; + << A << /* non-negative */ 1; return nullptr; } } From c10bfafade526cc234ee73d520c4d39584c6cbb2 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 27 Apr 2021 13:49:44 -0400 Subject: [PATCH 25/28] Identation Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index f9280171055b9..18c3a3919e9ca 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1046,10 +1046,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, unsigned int count = IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue(); - const char *var = - IntelFPGALoopCountAvg->isMax() ? "llvm.loop.intel.loopcount_max" - : IntelFPGALoopCountAvg->isMin() ? "llvm.loop.intel.loopcount_min" - : "llvm.loop.intel.loopcount_avg"; + 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); } From cefac8a2fe28d620908a2bf8636f4feec3183037 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 27 Apr 2021 15:20:36 -0400 Subject: [PATCH 26/28] Fixes after review Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/AttrDocs.td | 21 +++++++++------------ clang/lib/CodeGen/CGLoopInfo.cpp | 18 ++++++++---------- clang/lib/CodeGen/CGLoopInfo.h | 5 ++--- 3 files changed, 19 insertions(+), 25 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 16cd62b0dcf43..2b0936d249b71 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2963,22 +2963,19 @@ using PGO. .. code-block:: c++ - void foo() { - int a[10]; - [[intel::loop_count_min(4)] for (int i = 0; i < n; ++i) a[i] = 0; + void foo(int *array, size_t n) { + [[intel::loop_count_min(4)]] for (int i = 0; i < n; ++i) array[i] = 0; } - void zoo() { - int a[10]; - [[intel::loop_count_max(10)] for (int i = 0; i < m; ++i) a[i] = 0; + void zoo(int *array, size_t n) { + [[intel::loop_count_max(10)]] for (int i = 0; i < n; ++i) array[i] = 0; } - void goo() { - int a[10]; - [[intel::loop_count_min(3)] - [[intel::loop_count_max(10)] - [[intel::loop_count_avg(5)] - for (int i = start; i < end; ++i) a[i] = 0; + void goo(int *array, size_t n) { + [[intel::loop_count_min(3)]] + [[intel::loop_count_max(10)]] + [[intel::loop_count_avg(5)]] + for (int i = 0; i < n; ++i) array[i] = 0; } template diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 18c3a3919e9ca..daa97012274d9 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -605,13 +605,11 @@ MDNode *LoopInfo::createMetadata( LoopProperties.push_back(MDNode::get(Ctx, Vals)); } - if (Attrs.SYCLIntelFPGAVariantCount.size() > 0) { - for (auto &VC : Attrs.SYCLIntelFPGAVariantCount) { - Metadata *Vals[] = {MDString::get(Ctx, VC.first), - ConstantAsMetadata::get(ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), VC.second))}; - LoopProperties.push_back(MDNode::get(Ctx, Vals)); - } + for (auto &VC : Attrs.SYCLIntelFPGAVariantCount) { + Metadata *Vals[] = {MDString::get(Ctx, VC.first), + ConstantAsMetadata::get(ConstantInt::get( + llvm::Type::getInt32Ty(Ctx), VC.second))}; + LoopProperties.push_back(MDNode::get(Ctx, Vals)); } LoopProperties.insert(LoopProperties.end(), AdditionalLoopProperties.begin(), AdditionalLoopProperties.end()); @@ -1043,15 +1041,15 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (const auto *IntelFPGALoopCountAvg = dyn_cast(A)) { - unsigned int count = IntelFPGALoopCountAvg->getNTripCount() + unsigned int Count = IntelFPGALoopCountAvg->getNTripCount() ->getIntegerConstantExpr(Ctx) ->getSExtValue(); - const char *var = IntelFPGALoopCountAvg->isMax() + 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, Count); } if (const auto *IntelFPGALoopCoalesce = diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 6fbc6dcc9c6cf..317972a34ebeb 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -15,7 +15,6 @@ #define LLVM_CLANG_LIB_CODEGEN_CGLOOPINFO_H #include "llvm/ADT/ArrayRef.h" -#include "llvm/ADT/DenseMap.h" #include "llvm/ADT/SmallVector.h" #include "llvm/IR/DebugLoc.h" #include "llvm/IR/Value.h" @@ -410,8 +409,8 @@ class LoopInfoStack { } /// Set value of variant and loop count for the next loop pushed. - void setSYCLIntelFPGAVariantCount(const char *var, unsigned int C) { - StagedAttrs.SYCLIntelFPGAVariantCount.push_back({var, C}); + void setSYCLIntelFPGAVariantCount(const char *Var, unsigned int Count) { + StagedAttrs.SYCLIntelFPGAVariantCount.push_back({Var, Count}); } /// Set the unroll count for the next loop pushed. From 6ddf9cef576fe45a6080f94cc73a29d0e6173483 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 27 Apr 2021 15:31:18 -0400 Subject: [PATCH 27/28] Fixes after review Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/AttrDocs.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2b0936d249b71..c1c5a97ef4a85 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2955,7 +2955,7 @@ def SYCLIntelFPGALoopCountAttrDocs : Documentation { let Category = DocCatVariable; let Heading = "intel::loop_count_min, intel::loop_count_max, intel::loop_count_avg"; let Content = [{ -The loop count attribute specifies the minimum, maximum, or average number of +The loop count attributes specify the minimum, maximum, or average number of iterations for a ``for`` loop. These are hints that the user specify that can be used by some of the loop optimization to make decisions such as if the loop should be unrolled. It is a way for the user to provide some information without From d0ddcdd903101499c9a463c70f1209db96d32048 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 27 Apr 2021 15:46:27 -0400 Subject: [PATCH 28/28] Fixes after review Signed-off-by: Zahira Ammarguellat --- clang/lib/CodeGen/CGLoopInfo.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index daa97012274d9..ea63a29868263 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -553,10 +553,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"), - 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)); }