From d73d5052ef156c18417ef5e8276a8fa5cbd62813 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 1 Nov 2020 19:25:35 -0800 Subject: [PATCH 1/7] [SYCL] Add support for new FPGA loop attribute nofusion This patch adds support a new loop attribute for FPGA, intel::nofusion. This attribute should be passed to the FPGA backend, and ignored by the emulator. The attribute indicates that the annotated loop should not be fused with any adjacent loop. Note: this does not include a corresponding [[intel::fusion]] attribute, because a different mechanism (loop_fuse) will be built for FPGA. Syntax: [[intel::nofusion]] The LLVM IR representation should be similar to the representation used for #pragma nofusion. The llvm.loop metadata should specify llvm.loop.fusion.disable. Signed-off-by: Soumi Manna --- clang/include/clang/Basic/Attr.td | 12 ++++ clang/include/clang/Basic/AttrDocs.td | 9 +++ clang/lib/CodeGen/CGLoopInfo.cpp | 23 ++++++- clang/lib/CodeGen/CGLoopInfo.h | 7 ++ clang/lib/Parse/ParseStmt.cpp | 3 +- clang/lib/Sema/SemaStmtAttr.cpp | 17 +++++ .../test/CodeGenSYCL/intel-fpga-nofusion.cpp | 69 +++++++++++++++++++ clang/test/SemaSYCL/intel-fpga-loops.cpp | 18 +++++ clang/test/SemaSYCL/intel-fpga-nofusion.cpp | 44 ++++++++++++ 9 files changed, 199 insertions(+), 3 deletions(-) create mode 100644 clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp create mode 100644 clang/test/SemaSYCL/intel-fpga-nofusion.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index fb903290706db..993115e737d0b 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/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 : SubsetSubjecthasLocalStorage() && S->getKind() != Decl::ImplicitParam && diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index b94d60a251cb7..3bda54eb51ee9 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/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"; diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index e670c947e03dc..61730ffa57cf2 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/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, @@ -656,6 +664,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && + Attrs.SYCLNofusionEnable == false && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -970,6 +979,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(Attr); @@ -986,10 +997,13 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, dyn_cast(Attr); const SYCLIntelFPGASpeculatedIterationsAttr *IntelFPGASpeculatedIterations = dyn_cast(Attr); + const SYCLIntelFPGANofusionAttr *IntelFPGANofusion = + dyn_cast(Attr); if (!IntelFPGAIVDep && !IntelFPGAII && !IntelFPGAMaxConcurrency && !IntelFPGALoopCoalesce && !IntelFPGADisableLoopPipelining && - !IntelFPGAMaxInterleaving && !IntelFPGASpeculatedIterations) + !IntelFPGAMaxInterleaving && !IntelFPGASpeculatedIterations && + !IntelFPGANofusion) continue; if (IntelFPGAIVDep) @@ -1034,6 +1048,11 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, ->getIntegerConstantExpr(Ctx) ->getSExtValue()); } + + if (IntelFPGANofusion) { + setSYCLNofusionEnable(); + } + } if (CGOpts.OptimizationLevel > 0) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index f43db7a7bb63f..9ee77d52325d3 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/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,10 @@ 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(); } diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp index dc946bd94bbba..cdef0450cd7ae 100644 --- a/clang/lib/Parse/ParseStmt.cpp +++ b/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); diff --git a/clang/lib/Sema/SemaStmtAttr.cpp b/clang/lib/Sema/SemaStmtAttr.cpp index cff377c7aed43..5c72b57dc22b5 100644 --- a/clang/lib/Sema/SemaStmtAttr.cpp +++ b/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(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 diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp new file mode 100644 index 0000000000000..1cbba833cf233 --- /dev/null +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -0,0 +1,69 @@ +// 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]+]] + +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; + } + +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { + kernelFunc(); +} + +int main() { + kernel_single_task([]() { + 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]]} diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index 92c5434464838..f23334ea7d3f6 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/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 diff --git a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp new file mode 100644 index 0000000000000..f284efde96554 --- /dev/null +++ b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp @@ -0,0 +1,44 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify %s +// expected-no-diagnostics + +void foo() { + 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]] + do { + a1[i] += 4; + } while (i < 10); + + // CHECK: AttributedStmt + // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}} + [[intel::nofusion]] + for (int i = 0; i < 10; ++i) { + for (int j = 0; j < 10; ++j) { + a1[i] += a1[j]; + } + } + +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { + kernelFunc(); +} + +int main() { + kernel_single_task([]() { + foo(); + }); + return 0; +} + From 48d3465e6ec75d4f55027fe7d028f114e70ea30b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 1 Nov 2020 19:52:39 -0800 Subject: [PATCH 2/7] fix clang-format issues Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 9 ++++--- clang/lib/CodeGen/CGLoopInfo.h | 5 ++-- .../test/CodeGenSYCL/intel-fpga-nofusion.cpp | 24 +++++++------------ clang/test/SemaSYCL/intel-fpga-nofusion.cpp | 16 +++++-------- 4 files changed, 21 insertions(+), 33 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 61730ffa57cf2..dcb61d4692816 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -638,7 +638,7 @@ void LoopAttributes::clear() { DistributeEnable = LoopAttributes::Unspecified; PipelineDisabled = false; PipelineInitiationInterval = 0; - SYCLNofusionEnable = false; + SYCLNofusionEnable = false; } LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, @@ -664,7 +664,6 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.SYCLMaxInterleavingNInvocations == 0 && Attrs.SYCLSpeculatedIterationsEnable == false && Attrs.SYCLSpeculatedIterationsNIterations == 0 && - Attrs.SYCLNofusionEnable == false && Attrs.UnrollCount == 0 && Attrs.UnrollAndJamCount == 0 && !Attrs.PipelineDisabled && Attrs.PipelineInitiationInterval == 0 && Attrs.VectorizePredicateEnable == LoopAttributes::Unspecified && @@ -672,6 +671,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.UnrollEnable == LoopAttributes::Unspecified && Attrs.UnrollAndJamEnable == LoopAttributes::Unspecified && Attrs.DistributeEnable == LoopAttributes::Unspecified && !StartLoc && + Attrs.SYCLNofusionEnable == false && !EndLoc) return; @@ -998,12 +998,12 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, const SYCLIntelFPGASpeculatedIterationsAttr *IntelFPGASpeculatedIterations = dyn_cast(Attr); const SYCLIntelFPGANofusionAttr *IntelFPGANofusion = - dyn_cast(Attr); + dyn_cast(Attr); if (!IntelFPGAIVDep && !IntelFPGAII && !IntelFPGAMaxConcurrency && !IntelFPGALoopCoalesce && !IntelFPGADisableLoopPipelining && !IntelFPGAMaxInterleaving && !IntelFPGASpeculatedIterations && - !IntelFPGANofusion) + !IntelFPGANofusion) continue; if (IntelFPGAIVDep) @@ -1052,7 +1052,6 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, if (IntelFPGANofusion) { setSYCLNofusionEnable(); } - } if (CGOpts.OptimizationLevel > 0) diff --git a/clang/lib/CodeGen/CGLoopInfo.h b/clang/lib/CodeGen/CGLoopInfo.h index 9ee77d52325d3..5648a1a9d97fa 100644 --- a/clang/lib/CodeGen/CGLoopInfo.h +++ b/clang/lib/CodeGen/CGLoopInfo.h @@ -409,9 +409,8 @@ class LoopInfoStack { } /// Set flag of nofusion for the next loop pushed. - void setSYCLNofusionEnable() { - StagedAttrs.SYCLNofusionEnable = true; - } + void setSYCLNofusionEnable() { StagedAttrs.SYCLNofusionEnable = true; } + private: /// Returns true if there is LoopInfo on the stack. bool hasInfo() const { return !Active.empty(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp index 1cbba833cf233..815a90bb16f11 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -11,41 +11,35 @@ void nofusion() { int a[10]; int i = 0; - [[intel::nofusion]] - while (i < 10) { + [[intel::nofusion]] while (i < 10) { a[i] += 7; } for (int i = 0; i < 10; ++i) { - [[intel::nofusion]] - for (int j = 0; j < 10; ++j) { + [[intel::nofusion]] for (int j = 0; j < 10; ++j) { a[i] += a[j]; } } - [[intel::nofusion]] - do { + [[intel::nofusion]] do { a[i] += 4; - } while (i < 10); + } + while (i < 10); - [[intel::nofusion]] - for (int i = 0; i < 10; ++i) { + [[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) { + int k = 0; + [[intel::nofusion]] for (auto k: a) { k += 2; } - [[intel::nofusion]] - for (int i = 0; i < 10; ++i) { + [[intel::nofusion]] for (int i = 0; i < 10; ++i) { a[i] += 3; } - } template diff --git a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp index f284efde96554..72126941d9a72 100644 --- a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp @@ -6,28 +6,25 @@ void foo() { // CHECK: AttributedStmt // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}} - [[intel::nofusion]] - for (int p = 0; p < 10; ++p) { + [[intel::nofusion]] for (int p = 0; p < 10; ++p) { a1[p] = a2[p] = 0; } // CHECK: AttributedStmt // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}} - int i=0; - [[intel::nofusion]] - do { + int i = 0; + [[intel::nofusion]] do { a1[i] += 4; - } while (i < 10); + } + while (i < 10); // CHECK: AttributedStmt // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}} - [[intel::nofusion]] - for (int i = 0; i < 10; ++i) { + [[intel::nofusion]] for (int i = 0; i < 10; ++i) { for (int j = 0; j < 10; ++j) { a1[i] += a1[j]; } } - } template @@ -41,4 +38,3 @@ int main() { }); return 0; } - From cf154405e8d50b615340ca49708c678fcfe4adb4 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 1 Nov 2020 20:02:58 -0800 Subject: [PATCH 3/7] fix clang-format issues Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 3 +-- clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp | 9 +++++---- clang/test/SemaSYCL/intel-fpga-nofusion.cpp | 13 ++++++------- 3 files changed, 12 insertions(+), 13 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index dcb61d4692816..0cb831c8ce97b 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -671,8 +671,7 @@ LoopInfo::LoopInfo(BasicBlock *Header, const LoopAttributes &Attrs, Attrs.UnrollEnable == LoopAttributes::Unspecified && Attrs.UnrollAndJamEnable == LoopAttributes::Unspecified && Attrs.DistributeEnable == LoopAttributes::Unspecified && !StartLoc && - Attrs.SYCLNofusionEnable == false && - !EndLoc) + Attrs.SYCLNofusionEnable == false && !EndLoc) return; TempLoopID = MDNode::getTemporary(Header->getContext(), None); diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp index 815a90bb16f11..65d5ada2c69cd 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -23,17 +23,18 @@ void nofusion() { [[intel::nofusion]] do { a[i] += 4; - } - while (i < 10); + } + 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) { + [[intel::nofusion]] for (auto k : a) { k += 2; } diff --git a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp index 72126941d9a72..26f6868f78198 100644 --- a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp @@ -12,16 +12,15 @@ void foo() { // CHECK: AttributedStmt // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}} - int i = 0; - [[intel::nofusion]] do { - a1[i] += 4; - } - while (i < 10); + int i = 0; + [[intel::nofusion]] while (i < 10) { + a1[i] += 3; + } // CHECK: AttributedStmt // CHECK-NEXT: SYCLIntelFPGANofusionAttr {{.*}} - [[intel::nofusion]] for (int i = 0; i < 10; ++i) { - for (int j = 0; j < 10; ++j) { + for (int i = 0; i < 10; ++i) { + [[intel::nofusion]] for (int j = 0; j < 10; ++j) { a1[i] += a1[j]; } } From 824625059f5ce777baae2d1c1520cec92b15986d Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sun, 1 Nov 2020 20:06:16 -0800 Subject: [PATCH 4/7] fix clang-format issues Signed-off-by: Soumi Manna --- clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp index 65d5ada2c69cd..5edcc42e7e3b4 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -25,7 +25,7 @@ void nofusion() { a[i] += 4; } while (i < 10) - ; + ; [[intel::nofusion]] for (int i = 0; i < 10; ++i) { for (int j = 0; j < 10; ++j) { From d5137482599e241b2ae58e6fc094ecbd70d7bb06 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 2 Nov 2020 07:03:22 -0800 Subject: [PATCH 5/7] Address review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CGLoopInfo.cpp | 3 +-- clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp | 9 +++------ clang/test/SemaSYCL/intel-fpga-nofusion.cpp | 15 ++++++--------- 3 files changed, 10 insertions(+), 17 deletions(-) diff --git a/clang/lib/CodeGen/CGLoopInfo.cpp b/clang/lib/CodeGen/CGLoopInfo.cpp index 0cb831c8ce97b..fc9c19a1aab45 100644 --- a/clang/lib/CodeGen/CGLoopInfo.cpp +++ b/clang/lib/CodeGen/CGLoopInfo.cpp @@ -1048,9 +1048,8 @@ void LoopInfoStack::push(BasicBlock *Header, clang::ASTContext &Ctx, ->getSExtValue()); } - if (IntelFPGANofusion) { + if (IntelFPGANofusion) setSYCLNofusionEnable(); - } } if (CGOpts.OptimizationLevel > 0) diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp index 5edcc42e7e3b4..ddf4dbdeb643e 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -7,6 +7,8 @@ // 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]; @@ -43,13 +45,8 @@ void nofusion() { } } -template -__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { - kernelFunc(); -} - int main() { - kernel_single_task([]() { + cl::sycl::kernel_single_task([]() { nofusion(); }); return 0; diff --git a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp index 26f6868f78198..ea2f45fe93c88 100644 --- a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp @@ -1,7 +1,9 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify %s | FileCheck %s // expected-no-diagnostics -void foo() { +#include "Inputs/sycl.hpp" + +void nofusion() { int a1[10], a2[10]; // CHECK: AttributedStmt @@ -26,14 +28,9 @@ void foo() { } } -template -__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { - kernelFunc(); -} - int main() { - kernel_single_task([]() { - foo(); + cl::sycl::kernel_single_task([]() { + nofusion(); }); return 0; } From 822be0ecb3d985fcdbfc004b5f6b25d061834555 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 2 Nov 2020 20:09:52 -0800 Subject: [PATCH 6/7] update tests based on review comments Signed-off-by: Soumi Manna --- .../test/CodeGenSYCL/intel-fpga-nofusion.cpp | 50 ++++++++++--------- clang/test/SemaSYCL/intel-fpga-nofusion.cpp | 11 ++-- 2 files changed, 34 insertions(+), 27 deletions(-) diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp index ddf4dbdeb643e..1f532c25eb384 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -1,53 +1,57 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl -fsycl-is-device -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -fsycl -fsycl-is-device -internal-isystem %S/Inputs -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 "sycl.hpp" -#include "Inputs/sycl.hpp" +using namespace cl::sycl; +queue q; 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]; - } + // CHECK: br label {{.*}}, !llvm.loop ![[MD_NF_1:.*]] + a[i] += 2; } [[intel::nofusion]] do { - a[i] += 4; + // CHECK: br i1 %{{.*}}, !llvm.loop ![[MD_NF_2:.*]] + a[i] += 3; } while (i < 10) ; [[intel::nofusion]] for (int i = 0; i < 10; ++i) { + // CHECK: br label %{{.*}}, !llvm.loop ![[MD_NF_3:.*]] for (int j = 0; j < 10; ++j) { + // CHECK-NOT: br label %{{.*}}, !llvm.loop !{{.*}} a[i] += a[j]; } } - int k = 0; - [[intel::nofusion]] for (auto k : a) { - k += 2; + int k; + [[intel::nofusion]] for (auto k: a) { + // CHECK: br label %{{.*}}, !llvm.loop ![[MD_NF_5:.*]] + k += 4; } [[intel::nofusion]] for (int i = 0; i < 10; ++i) { - a[i] += 3; + // CHECK: br label %{{.*}}, !llvm.loop ![[MD_NF_6:.*]] + a[i] += 5; + } + + for (int i = 0; i < 10; ++i) { + // CHECK-NOT: br label %{{.*}}, !llvm.loop !{{.*}} + [[intel::nofusion]] for (int j = 0; j < 10; ++j) { + // CHECK: br label %{{.*}}, !llvm.loop ![[MD_NF_8:.*]] + a[i] += a[j]; + } } } int main() { - cl::sycl::kernel_single_task([]() { - nofusion(); + q.submit([&](handler &h) { + h.single_task([]() { nofusion(); }); }); return 0; } @@ -56,6 +60,6 @@ int main() { // 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]]} +// CHECK: ![[MD_NF_8]] = distinct !{![[MD_NF_8]], ![[MD_Nofusion]]} diff --git a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp index ea2f45fe93c88..bea1362ba2835 100644 --- a/clang/test/SemaSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/SemaSYCL/intel-fpga-nofusion.cpp @@ -1,7 +1,10 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify %s | FileCheck %s // expected-no-diagnostics -#include "Inputs/sycl.hpp" +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; void nofusion() { int a1[10], a2[10]; @@ -29,8 +32,8 @@ void nofusion() { } int main() { - cl::sycl::kernel_single_task([]() { - nofusion(); + q.submit([&](handler &h) { + h.single_task([]() { nofusion(); }); }); return 0; } From 14e71a196a519d0ede79ac28602e5552df029632 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 2 Nov 2020 20:14:12 -0800 Subject: [PATCH 7/7] fix clang-fromat issues Signed-off-by: Soumi Manna --- clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp index 1f532c25eb384..7d045ced14477 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-nofusion.cpp @@ -30,7 +30,7 @@ void nofusion() { } int k; - [[intel::nofusion]] for (auto k: a) { + [[intel::nofusion]] for (auto k : a) { // CHECK: br label %{{.*}}, !llvm.loop ![[MD_NF_5:.*]] k += 4; }