[SYCL][FPGA] Add support for loop count attribute.

clang/include/clang/Basic/Attr.td

@@ -1953,6 +1953,17 @@ def SYCLIntelFPGADisableLoopPipelining : DeclOrStmtAttr {
 def : MutualExclusions<[SYCLIntelFPGAInitiationInterval,
                         SYCLIntelFPGADisableLoopPipelining]>;
 
+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<"NTripCount">];
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let HasCustomTypeTransform = 1;
+  let Documentation = [SYCLIntelFPGALoopControlAvgAttrDocs];
+}
+
 def SYCLIntelFPGAMaxInterleaving : StmtAttr {
   let Spellings = [CXX11<"intelfpga","max_interleaving">,
                    CXX11<"intel","max_interleaving">];

clang/include/clang/Basic/AttrDocs.td

@@ -2950,6 +2950,29 @@ max_concurrency, initiation_interval, or ivdep.
   }];
 }
 
+def SYCLIntelFPGALoopControlAvgAttrDocs : 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<int N>
+  void bar() {
+    [[intel::loop_control_avg(N)]] for(;;) { }
+  }
+
+  }];
+}
+
 def SYCLIntelFPGAMaxInterleavingAttrDocs : Documentation {
   let Category = DocCatVariable;
   let Heading = "intel::max_interleaving";

clang/include/clang/Sema/Sema.h

@@ -13419,7 +13419,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;

clang/lib/CodeGen/CGLoopInfo.cpp

@@ -605,6 +605,13 @@ MDNode *LoopInfo::createMetadata(
     LoopProperties.push_back(MDNode::get(Ctx, Vals));
   }
 
+  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(),
                         AdditionalLoopProperties.end());
   return createFullUnrollMetadata(Attrs, LoopProperties, HasUserTransforms);
@@ -621,7 +628,8 @@ LoopAttributes::LoopAttributes(bool IsParallel)
       SYCLLoopCoalesceNLevels(0), SYCLLoopPipeliningDisable(false),
       SYCLMaxInterleavingEnable(false), SYCLMaxInterleavingNInvocations(0),
       SYCLSpeculatedIterationsEnable(false),
-      SYCLSpeculatedIterationsNIterations(0), UnrollCount(0),
+      SYCLSpeculatedIterationsNIterations(0),
+      SYCLIntelFPGALoopControlAverageEnable(false), UnrollCount(0),
       UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified),
       PipelineDisabled(false), PipelineInitiationInterval(0),
       SYCLNofusionEnable(false), MustProgress(false) {}
@@ -643,6 +651,7 @@ void LoopAttributes::clear() {
   SYCLMaxInterleavingNInvocations = 0;
   SYCLSpeculatedIterationsEnable = false;
   SYCLSpeculatedIterationsNIterations = 0;
+  SYCLIntelFPGALoopControlAverageEnable = false;
   UnrollCount = 0;
   UnrollAndJamCount = 0;
   VectorizeEnable = LoopAttributes::Unspecified;
@@ -680,8 +689,10 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs,
       Attrs.SYCLMaxInterleavingNInvocations == 0 &&
       Attrs.SYCLSpeculatedIterationsEnable == false &&
       Attrs.SYCLSpeculatedIterationsNIterations == 0 &&
-      Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 &&
-      !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 &&
+      Attrs.SYCLIntelFPGALoopControlAverageEnable == 0 &&
+      Attrs.UnrollCount == 0 &&
+      Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled &&
+      Attrs.PipelineInitiationInterval == 0 &&
       Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified &&
       Attrs.VectorizeEnable == LoopAttributes::Unspecified &&
       Attrs.UnrollEnable == LoopAttributes::Unspecified &&
@@ -1030,6 +1041,14 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
                                         ->getSExtValue());
     }
 
+    if (const auto *IntelFPGALoopControlAvg =
+            dyn_cast<SYCLIntelFPGALoopControlAvgAttr>(A)) {
+      setSYCLIntelFPGALoopControlAvgEnable();
+      setSYCLIntelFPGALoopControlAverage(IntelFPGALoopControlAvg->getNTripCount()
+              ->getIntegerConstantExpr(Ctx)
+              ->getSExtValue());
+    }
+
     if (const auto *IntelFPGALoopCoalesce =
             dyn_cast<SYCLIntelFPGALoopCoalesceAttr>(A)) {
       if (auto *LCE = IntelFPGALoopCoalesce->getNExpr())

clang/lib/CodeGen/CGLoopInfo.h

@@ -117,6 +117,12 @@ struct LoopAttributes {
   /// 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;
 
@@ -404,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; }
 

clang/lib/Sema/SemaStmtAttr.cpp

@@ -572,6 +572,8 @@ static void CheckForIncompatibleSYCLLoopAttributes(
                                                                          Attrs);
   CheckForDuplicationSYCLLoopAttribute<SYCLIntelFPGASpeculatedIterationsAttr>(
       S, Attrs);
+  CheckForDuplicationSYCLLoopAttribute<SYCLIntelFPGALoopControlAvgAttr>(S,
+                                                                        Attrs);
   CheckForDuplicationSYCLLoopAttribute<LoopUnrollHintAttr>(S, Attrs, false);
   CheckMutualExclusionSYCLLoopAttribute<SYCLIntelFPGADisableLoopPipeliningAttr,
                                         SYCLIntelFPGAMaxInterleavingAttr>(
@@ -716,6 +718,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
   case ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations:
     return handleIntelFPGALoopAttr<SYCLIntelFPGASpeculatedIterationsAttr>(S, St,
                                                                           A);
+  case ParsedAttr::AT_SYCLIntelFPGALoopControlAvg:
+    return handleIntelFPGALoopAttr<SYCLIntelFPGALoopControlAvgAttr>(S, St, A);
   case ParsedAttr::AT_OpenCLUnrollHint:
   case ParsedAttr::AT_LoopUnrollHint:
     return handleLoopUnrollHint(S, St, A, Range);

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,15 @@ TemplateInstantiator::TransformSYCLIntelFPGASpeculatedIterationsAttr(
           *SI, TransformedExpr);
 }
 
+const SYCLIntelFPGALoopControlAvgAttr *
+TemplateInstantiator::TransformSYCLIntelFPGALoopControlAvgAttr(
+    const SYCLIntelFPGALoopControlAvgAttr *LCA) {
+  Expr *TransformedExpr =
+      getDerived().TransformExpr(LCA->getNTripCount()).get();
+  return getSema().BuildSYCLIntelFPGALoopAttr<SYCLIntelFPGALoopControlAvgAttr>(
+      *LCA, TransformedExpr);
+}
+
 const LoopUnrollHintAttr *TemplateInstantiator::TransformLoopUnrollHintAttr(
     const LoopUnrollHintAttr *LU) {
   Expr *TransformedExpr =

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 <int A>
+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 <typename name, typename Func>
 __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;
 }

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<int A, int B, int C>
+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<class kernel_function>([]() {
@@ -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;