[SYCL] Add support for new FPGA loop attribute nofusion

clang/include/clang/Basic/Attr.td

@@ -1849,6 +1849,18 @@ def SYCLIntelFPGASpeculatedIterations : Attr {
   let Documentation = [SYCLIntelFPGASpeculatedIterationsAttrDocs];
 }
 
+def SYCLIntelFPGANofusion : Attr {
+  let Spellings = [CXX11<"intel","nofusion">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let HasCustomTypeTransform = 1;
+  let AdditionalMembers = [{
+    static const char *getName() {
+      return "nofusion";
+    }
+  }];
+  let Documentation = [SYCLIntelFPGANofusionAttrDocs];
+}
+
 def IntelFPGALocalNonConstVar : SubsetSubject<Var,
                               [{S->hasLocalStorage() &&
                                 S->getKind() != Decl::ImplicitParam &&

clang/include/clang/Basic/AttrDocs.td

@@ -2418,6 +2418,15 @@ used on the same loop in conjunction with disable_loop_pipelining.
   }];
 }
 
+def SYCLIntelFPGANofusionAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "intel::nofusion";
+  let Content = [{
+This attribute applies to a loop. Indicates that the annotated
+loop should not be fused with any adjacent loop.
+  }];
+}
+
 def SYCLDeviceIndirectlyCallableDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "intel::device_indirectly_callable";

clang/lib/CodeGen/CGLoopInfo.cpp

@@ -578,6 +578,12 @@ MDNode *LoopInfo::createMetadata(
     LoopProperties.push_back(MDNode::get(Ctx, Vals));
   }
 
+  // nofusion attribute corresponds to 'llvm.loop.fusion.disable' metadata
+  if (Attrs.SYCLNofusionEnable) {
+    Metadata *Vals[] = {MDString::get(Ctx, "llvm.loop.fusion.disable")};
+    LoopProperties.push_back(MDNode::get(Ctx, Vals));
+  }
+
   if (Attrs.SYCLSpeculatedIterationsEnable) {
     Metadata *Vals[] = {
         MDString::get(Ctx, "llvm.loop.intel.speculated.iterations.count"),
@@ -604,7 +610,8 @@ LoopAttributes::LoopAttributes(bool IsParallel)
       SYCLSpeculatedIterationsEnable(false),
       SYCLSpeculatedIterationsNIterations(0), UnrollCount(0),
       UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified),
-      PipelineDisabled(false), PipelineInitiationInterval(0) {}
+      PipelineDisabled(false), PipelineInitiationInterval(0),
+      SYCLNofusionEnable(false) {}
 
 void LoopAttributes::clear() {
   IsParallel = false;
@@ -631,6 +638,7 @@ void LoopAttributes::clear() {
   DistributeEnable = LoopAttributes::Unspecified;
   PipelineDisabled = false;
   PipelineInitiationInterval = 0;
+  SYCLNofusionEnable = false;
 }
 
 LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs,
@@ -663,7 +671,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs,
       Attrs.UnrollEnable == LoopAttributes::Unspecified &&
       Attrs.UnrollAndJamEnable == LoopAttributes::Unspecified &&
       Attrs.DistributeEnable == LoopAttributes::Unspecified && !StartLoc &&
-      !EndLoc)
+      Attrs.SYCLNofusionEnable == false && !EndLoc)
     return;
 
   TempLoopID = MDNode::getTemporary(Header->getContext(), None);
@@ -970,6 +978,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
   // For attribute speculated_iterations:
   // n - 'llvm.loop.intel.speculated.iterations.count, i32 n' metadata will be
   // emitted
+  // For attribute nofusion:
+  // 'llvm.loop.fusion.disable' metadata will be emitted
   for (const auto *Attr : Attrs) {
     const SYCLIntelFPGAIVDepAttr *IntelFPGAIVDep =
       dyn_cast<SYCLIntelFPGAIVDepAttr>(Attr);
@@ -986,10 +996,13 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
         dyn_cast<SYCLIntelFPGAMaxInterleavingAttr>(Attr);
     const SYCLIntelFPGASpeculatedIterationsAttr *IntelFPGASpeculatedIterations =
         dyn_cast<SYCLIntelFPGASpeculatedIterationsAttr>(Attr);
+    const SYCLIntelFPGANofusionAttr *IntelFPGANofusion =
+        dyn_cast<SYCLIntelFPGANofusionAttr>(Attr);
 
     if (!IntelFPGAIVDep && !IntelFPGAII && !IntelFPGAMaxConcurrency &&
         !IntelFPGALoopCoalesce && !IntelFPGADisableLoopPipelining &&
-        !IntelFPGAMaxInterleaving && !IntelFPGASpeculatedIterations)
+        !IntelFPGAMaxInterleaving && !IntelFPGASpeculatedIterations &&
+        !IntelFPGANofusion)
       continue;
 
     if (IntelFPGAIVDep)
@@ -1034,6 +1047,9 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx,
               ->getIntegerConstantExpr(Ctx)
               ->getSExtValue());
     }
+
+    if (IntelFPGANofusion)
+      setSYCLNofusionEnable();
   }
 
   if (CGOpts.OptimizationLevel > 0)

clang/lib/CodeGen/CGLoopInfo.h

@@ -149,6 +149,9 @@ struct LoopAttributes {
 
   /// Value for llvm.loop.pipeline.iicount metadata.
   unsigned PipelineInitiationInterval;
+
+  /// Flag for llvm.loop.fusion.disable metatdata.
+  bool SYCLNofusionEnable;
 };
 
 /// Information used when generating a structured loop.
@@ -405,6 +408,9 @@ class LoopInfoStack {
     StagedAttrs.PipelineInitiationInterval = C;
   }
 
+  /// Set flag of nofusion for the next loop pushed.
+  void setSYCLNofusionEnable() { StagedAttrs.SYCLNofusionEnable = true; }
+
 private:
   /// Returns true if there is LoopInfo on the stack.
   bool hasInfo() const { return !Active.empty(); }

clang/lib/Parse/ParseStmt.cpp

@@ -2580,7 +2580,8 @@ bool Parser::ParseSYCLLoopAttributes(ParsedAttributes &Attrs) {
       Attrs.begin()->getKind() != ParsedAttr::AT_SYCLIntelFPGAMaxInterleaving &&
       Attrs.begin()->getKind() !=
           ParsedAttr::AT_SYCLIntelFPGASpeculatedIterations &&
-      Attrs.begin()->getKind() != ParsedAttr::AT_LoopUnrollHint)
+      Attrs.begin()->getKind() != ParsedAttr::AT_LoopUnrollHint &&
+      Attrs.begin()->getKind() != ParsedAttr::AT_SYCLIntelFPGANofusion)
     return true;
 
   bool IsIntelFPGAAttribute = (Attrs.begin()->getKind() != ParsedAttr::AT_LoopUnrollHint);

clang/lib/Sema/SemaStmtAttr.cpp

@@ -317,6 +317,19 @@ static Attr *handleIntelFPGAIVDepAttr(Sema &S, const ParsedAttr &A) {
       NumArgs == 2 ? A.getArgAsExpr(1) : nullptr);
 }
 
+static Attr *handleIntelFPGANofusionAttr(Sema &S, const ParsedAttr &A) {
+  if (S.LangOpts.SYCLIsHost)
+    return nullptr;
+
+  unsigned NumArgs = A.getNumArgs();
+  if (NumArgs > 0) {
+    S.Diag(A.getLoc(), diag::warn_attribute_too_many_arguments) << A << 0;
+    return nullptr;
+  }
+
+  return new (S.Context) SYCLIntelFPGANofusionAttr(S.Context, A);
+}
+
 static Attr *handleLoopHintAttr(Sema &S, Stmt *St, const ParsedAttr &A,
                                 SourceRange) {
   IdentifierLoc *PragmaNameLoc = A.getArgAsIdent(0);
@@ -675,6 +688,8 @@ static void CheckForIncompatibleSYCLLoopAttributes(
       S, Attrs, Range);
 
   CheckRedundantSYCLIntelFPGAIVDepAttrs(S, Attrs);
+  CheckForDuplicationSYCLLoopAttribute<SYCLIntelFPGANofusionAttr>(S, Attrs,
+                                                                  Range);
 }
 
 void CheckForIncompatibleUnrollHintAttributes(
@@ -803,6 +818,8 @@ static Attr *ProcessStmtAttribute(Sema &S, Stmt *St, const ParsedAttr &A,
     return handleLikely(S, St, A, Range);
   case ParsedAttr::AT_Unlikely:
     return handleUnlikely(S, St, A, Range);
+  case ParsedAttr::AT_SYCLIntelFPGANofusion:
+    return handleIntelFPGANofusionAttr(S, A);
   default:
     // if we're here, then we parsed a known attribute, but didn't recognize
     // it as a statement attribute => it is declaration attribute

clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp

@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
+
+// CHECK: br label %while.cond, !llvm.loop ![[MD_NF_1:[0-9]+]]
+// CHECK: br label %for.cond3, !llvm.loop ![[MD_NF_2:[0-9]+]]
+// CHECK: i1 %cmp18, label %do.body, label %do.end, !llvm.loop ![[MD_NF_3:[0-9]+]]
+// CHECK: br label %for.cond20, !llvm.loop ![[MD_NF_4:[0-9]+]]
+// CHECK: br label %for.cond41, !llvm.loop ![[MD_NF_5:[0-9]+]]
+// CHECK: br label %for.cond50, !llvm.loop ![[MD_NF_6:[0-9]+]]
+
+#include "Inputs/sycl.hpp"
+
+void nofusion() {
+  int a[10];
+
+  int i = 0;
+  [[intel::nofusion]] while (i < 10) {
+    a[i] += 7;
+  }
+
+  for (int i = 0; i < 10; ++i) {
+    [[intel::nofusion]] for (int j = 0; j < 10; ++j) {
+      a[i] += a[j];
+    }
+  }
+
+  [[intel::nofusion]] do {
+    a[i] += 4;
+  }
+  while (i < 10)
+    ;
+
+  [[intel::nofusion]] for (int i = 0; i < 10; ++i) {
+    for (int j = 0; j < 10; ++j) {
+      a[i] += a[j];
+    }
+  }
+
+  int k = 0;
+  [[intel::nofusion]] for (auto k : a) {
+    k += 2;
+  }
+
+  [[intel::nofusion]] for (int i = 0; i < 10; ++i) {
+    a[i] += 3;
+  }
+}
+
+int main() {
+  cl::sycl::kernel_single_task<class kernel_function>([]() {
+    nofusion();
+  });
+  return 0;
+}
+
+// CHECK: ![[MD_NF_1]] = distinct !{![[MD_NF_1]], ![[MD_Nofusion:[0-9]+]]}
+// CHECK: ![[MD_Nofusion]] = !{!"llvm.loop.fusion.disable"}
+// CHECK: ![[MD_NF_2]] = distinct !{![[MD_NF_2]], ![[MD_Nofusion]]}
+// CHECK: ![[MD_NF_3]] = distinct !{![[MD_NF_3]], ![[MD_Nofusion]]}
+// CHECK: ![[MD_NF_4]] = distinct !{![[MD_NF_4]], ![[MD_Nofusion]]}
+// CHECK: ![[MD_NF_5]] = distinct !{![[MD_NF_5]], ![[MD_Nofusion]]}
+// CHECK: ![[MD_NF_6]] = distinct !{![[MD_NF_6]], ![[MD_Nofusion]]}

clang/test/SemaSYCL/intel-fpga-loops.cpp

@@ -25,6 +25,8 @@ void foo() {
   [[intel::max_interleaving(4)]] int i[10];
   // expected-error@+1 {{intelfpga loop attributes must be applied to for, while, or do statements}}
   [[intel::speculated_iterations(6)]] int j[10];
+  // expected-error@+1 {{intelfpga loop attributes must be applied to for, while, or do statements}}
+  [[intel::nofusion]] int k[10];
 }
 
 // Test for deprecated spelling of Intel FPGA loop attributes
@@ -114,6 +116,9 @@ void boo() {
   // expected-warning@+1 {{'speculated_iterations' attribute takes no more than 1 argument - attribute ignored}}
   [[intel::speculated_iterations(1, 2)]] for (int i = 0; i != 10; ++i)
       a[i] = 0;
+  // expected-warning@+1 {{'nofusion' attribute takes no more than 0 arguments - attribute ignored}}
+  [[intel::nofusion(0)]] for (int i = 0; i != 10; ++i)
+      a[i] = 0;
 }
 
 // Test for incorrect argument value for Intel FPGA loop attributes
@@ -187,6 +192,10 @@ void goo() {
   // no diagnostics are expected
   [[intel::ivdep(2, s.ptr)]] for (int i = 0; i != 10; ++i)
       s.ptr[i] = 0;
+
+  // no diagnostics are expected
+  [[intel::nofusion]] for (int i = 0; i != 10; ++i)
+      a[i] = 0;
 }
 
 // Test for Intel FPGA loop attributes duplication
@@ -290,6 +299,11 @@ void zoo() {
   // expected-note@+1 {{previous attribute is here}}
   [[intel::ivdep(a, 3)]] for (int i = 0; i != 10; ++i)
       a[i] = 0;
+
+  [[intel::nofusion]]
+  // expected-error@-1 {{duplicate Intel FPGA loop attribute 'nofusion'}}
+  [[intel::nofusion]] for (int i = 0; i != 10; ++i)
+      a[i] = 0;
 }
 
 // Test for Intel FPGA loop attributes compatibility
@@ -319,6 +333,10 @@ void loop_attrs_compatibility() {
   [[intel::disable_loop_pipelining]]
   [[intel::ivdep]] for (int i = 0; i != 10; ++i)
       a[i] = 0;
+  // no diagnostics are expected
+  [[intel::disable_loop_pipelining]]
+  [[intel::nofusion]] for (int i = 0; i != 10; ++i)
+      a[i] = 0;
 }
 
 template<int A, int B, int C>

clang/test/SemaSYCL/intel-fpga-nofusion.cpp

@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify %s | FileCheck %s
+// expected-no-diagnostics
+
+#include "Inputs/sycl.hpp"
+
+void nofusion() {
+  int a1[10], a2[10];
+
+  // CHECK: AttributedStmt
+  // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}}
+  [[intel::nofusion]] for (int p = 0; p < 10; ++p) {
+    a1[p] = a2[p] = 0;
+  }
+
+  // CHECK: AttributedStmt
+  // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}}
+  int i = 0;
+  [[intel::nofusion]] while (i < 10) {
+    a1[i] += 3;
+  }
+
+  // CHECK: AttributedStmt
+  // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}}
+  for (int i = 0; i < 10; ++i) {
+    [[intel::nofusion]] for (int j = 0; j < 10; ++j) {
+      a1[i] += a1[j];
+    }
+  }
+}
+
+int main() {
+  cl::sycl::kernel_single_task<class kernel_function>([]() {
+    nofusion();
+  });
+  return 0;
+}