Merge branch 'sycl' into clangusm

Author: bader

README.md

@@ -7,17 +7,17 @@ Home for Intel LLVM-based projects:
  - SYCL* Compiler and Runtimes - compiler and runtime libraries for SYCL ([https://www.khronos.org/sycl/](https://www.khronos.org/sycl/)). See **sycl** branch.
 
 ## License
-See [LICENSE.txt](https://github.com/intel/llvm/blob/intel/llvm/LICENSE.TXT) for details.
+See [LICENSE.txt](sycl/LICENSE.TXT) for details.
 
 
 ## Contributing
-See [CONTRIBUTING.md](https://github.com/intel/llvm/blob/intel/CONTRIBUTING.md) for details.
+See [CONTRIBUTING.md](CONTRIBUTING.md) for details.
 
 ## Sub-projects Documentation
- - SYCL Compiler and Runtimes - See [GetStartedWithSYCLCompiler.md](https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedWithSYCLCompiler.md)
+ - SYCL Compiler and Runtimes - See [GetStartedWithSYCLCompiler.md](sycl/doc/GetStartedWithSYCLCompiler.md)
 
 *Other names and brands may be claimed as the property of others.
 
 ## SYCL Extension Proposal Documents
 
-See [sycl/doc/extensions](https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions)
+See [sycl/doc/extensions](sycl/doc/extensions)

clang/include/clang/Basic/Attr.td

@@ -1560,6 +1560,41 @@ def IntelFPGAMaxPrivateCopies : InheritableAttr {
   }];
 }
 
+// Two string arguments.
+def IntelFPGAMerge : Attr {
+  let Spellings = [CXX11<"intelfpga","merge">];
+  let Args = [StringArgument<"Name">, StringArgument<"Direction">];
+  let Subjects = SubjectList<[IntelFPGAConstVar, IntelFPGALocalOrStaticVar,
+                              Field], ErrorDiag>;
+  let LangOpts = [SYCL];
+  let Documentation = [IntelFPGAMergeAttrDocs];
+}
+
+def IntelFPGAMaxReplicates : Attr {
+  let Spellings = [CXX11<"intelfpga","max_replicates">];
+  let Args = [ExprArgument<"Value">];
+  let Subjects = SubjectList<[IntelFPGAConstVar, IntelFPGALocalStaticSlaveMemVar,
+                              Field], ErrorDiag>;
+  let LangOpts = [SYCL];
+  let Documentation = [IntelFPGAMaxReplicatesAttrDocs];
+  let AdditionalMembers = [{
+    static unsigned getMinValue() {
+      return 1;
+    }
+    static unsigned getMaxValue() {
+      return 1024*1024;
+    }
+  }];
+}
+
+def IntelFPGASimpleDualPort : Attr {
+  let Spellings = [CXX11<"intelfpga","simple_dual_port">];
+  let Subjects = SubjectList<[IntelFPGAConstVar, IntelFPGALocalStaticSlaveMemVar,
+                              Field], ErrorDiag>;
+  let LangOpts = [SYCL];
+  let Documentation = [IntelFPGASimpleDualPortAttrDocs];
+}
+
 def Naked : InheritableAttr {
   let Spellings = [GCC<"naked">, Declspec<"naked">];
   let Subjects = SubjectList<[Function]>;

clang/include/clang/Basic/AttrDocs.td

@@ -1714,6 +1714,40 @@ threads or loop iterations.
   }];
 }
 
+def IntelFPGAMergeAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "merge (IntelFPGA)";
+  let Content = [{
+This attribute may be attached to a variable or struct member declaration and
+instructs the backend to merge the memories used to implement any variable or
+struct members that are annotated with this attribute and the same first
+argument. The second argument indicates if the memories should be merged in a
+depth-wise or width-wise manner.
+  }];
+}
+
+def IntelFPGAMaxReplicatesAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "max_replicates (IntelFPGA)";
+  let Content = [{
+This attribute may be attached to a variable or struct member declaration and
+instructs the backend to replicate the memory generated for the variable or
+struct member no more than the specified maximum number of times to enable
+simultaneous accesses from different load/store sites in the program.
+  }];
+}
+
+def IntelFPGASimpleDualPortAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "simple_dual_port (IntelFPGA)";
+  let Content = [{
+This attribute may be attached to a variable or struct member declaration and
+instructs the backend to implement the variable or struct member in a memory
+with simple dual port configuration (no memory port services both stores and
+loads).
+  }];
+}
+
 def SYCLIntelFPGAIVDepAttrDocs : Documentation {
   let Category = DocCatVariable;
   let Heading = "ivdep";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -130,6 +130,8 @@ def err_intel_fpga_memory_arg_invalid : Error<
   "%0 attribute requires either no argument or one of: %1">;
 def err_intel_fpga_loop_attr_duplication : Error<
   "duplicate Intel FPGA loop attribute '%0'">;
+def err_intel_fpga_merge_dir_invalid : Error<
+  "merge direction must be 'depth' or 'width'">;
 
 // C99 variable-length arrays
 def ext_vla : Extension<"variable length arrays are a C99 feature">,

clang/lib/CodeGen/CGLoopInfo.cpp

@@ -446,13 +446,12 @@ MDNode *LoopInfo::createMetadata(
 
 LoopAttributes::LoopAttributes(bool IsParallel)
     : IsParallel(IsParallel), VectorizeEnable(LoopAttributes::Unspecified),
-      SYCLIVDepEnable(false), SYCLIVDepSafelen(0), SYCLIInterval(0),
-      SYCLMaxConcurrencyNThreads(0),
       UnrollEnable(LoopAttributes::Unspecified),
       UnrollAndJamEnable(LoopAttributes::Unspecified), VectorizeWidth(0),
-      InterleaveCount(0), UnrollCount(0), UnrollAndJamCount(0),
-      DistributeEnable(LoopAttributes::Unspecified), PipelineDisabled(false),
-      PipelineInitiationInterval(0) {}
+      InterleaveCount(0), SYCLIVDepEnable(false), SYCLIVDepSafelen(0),
+      SYCLIInterval(0), SYCLMaxConcurrencyNThreads(0), UnrollCount(0),
+      UnrollAndJamCount(0), DistributeEnable(LoopAttributes::Unspecified),
+      PipelineDisabled(false), PipelineInitiationInterval(0) {}
 
 void LoopAttributes::clear() {
   IsParallel = false;

clang/lib/CodeGen/CodeGenModule.cpp

@@ -3800,6 +3800,16 @@ void CodeGenModule::generateIntelFPGAAnnotation(
     llvm::APSInt BWAInt = NBA->getValue()->EvaluateKnownConstInt(getContext());
     Out << '{' << NBA->getSpelling() << ':' << BWAInt << '}';
   }
+  if (const auto *MRA = D->getAttr<IntelFPGAMaxReplicatesAttr>()) {
+    llvm::APSInt MRAInt = MRA->getValue()->EvaluateKnownConstInt(getContext());
+    Out << '{' << MRA->getSpelling() << ':' << MRAInt << '}';
+  }
+  if (const auto *MA = D->getAttr<IntelFPGAMergeAttr>()) {
+    Out << '{' << MA->getSpelling() << ':' << MA->getName() << ':'
+        << MA->getDirection() << '}';
+  }
+  if (D->hasAttr<IntelFPGASimpleDualPortAttr>())
+    Out << "{simple_dual_port:1}";
 }
 
 void CodeGenModule::addGlobalIntelFPGAAnnotation(const VarDecl *VD,

clang/lib/Sema/SemaDeclAttr.cpp

@@ -5091,6 +5091,12 @@ static bool checkIntelFPGARegisterAttrCompatibility(Sema &S, Decl *D,
     if (!NBA->isImplicit() &&
         checkAttrMutualExclusion<IntelFPGANumBanksAttr>(S, D, Attr))
       InCompat = true;
+  if (checkAttrMutualExclusion<IntelFPGAMaxReplicatesAttr>(S, D, Attr))
+    InCompat = true;
+  if (checkAttrMutualExclusion<IntelFPGASimpleDualPortAttr>(S, D, Attr))
+    InCompat = true;
+  if (checkAttrMutualExclusion<IntelFPGAMergeAttr>(S, D, Attr))
+    InCompat = true;
 
   return InCompat;
 }
@@ -5124,6 +5130,65 @@ static void handleOneConstantPowerTwoValueAttr(Sema &S, Decl *D,
       Attr.getAttributeSpellingListIndex());
 }
 
+static void handleIntelFPGASimpleDualPortAttr(Sema &S, Decl *D,
+                                              const ParsedAttr &Attr) {
+  checkForDuplicateAttribute<IntelFPGASimpleDualPortAttr>(S, D, Attr);
+
+  if (checkAttrMutualExclusion<IntelFPGARegisterAttr>(S, D, Attr))
+    return;
+
+  if (!D->hasAttr<IntelFPGAMemoryAttr>())
+    D->addAttr(IntelFPGAMemoryAttr::CreateImplicit(
+        S.Context, IntelFPGAMemoryAttr::Default));
+
+  D->addAttr(::new (S.Context)
+                 IntelFPGASimpleDualPortAttr(Attr.getRange(), S.Context, 0));
+}
+
+static void handleIntelFPGAMaxReplicatesAttr(Sema &S, Decl *D,
+                                             const ParsedAttr &Attr) {
+  checkForDuplicateAttribute<IntelFPGAMaxReplicatesAttr>(S, D, Attr);
+
+  if (checkAttrMutualExclusion<IntelFPGARegisterAttr>(S, D, Attr))
+    return;
+
+  S.AddOneConstantValueAttr<IntelFPGAMaxReplicatesAttr>(
+      Attr.getRange(), D, Attr.getArgAsExpr(0),
+      Attr.getAttributeSpellingListIndex());
+}
+
+/// Handle the merge attribute.
+/// This requires two string arguments.  The first argument is a name, the
+/// second is a direction.  The direction must be "depth" or "width".
+/// This is incompatible with the register attribute.
+static void handleIntelFPGAMergeAttr(Sema &S, Decl *D, const ParsedAttr &Attr) {
+  checkForDuplicateAttribute<IntelFPGAMergeAttr>(S, D, Attr);
+
+  if (checkAttrMutualExclusion<IntelFPGARegisterAttr>(S, D, Attr))
+    return;
+
+  SmallVector<StringRef, 2> Results;
+  for (int I = 0; I < 2; I++) {
+    StringRef Str;
+    if (!S.checkStringLiteralArgumentAttr(Attr, I, Str))
+      return;
+
+    if (I == 1 && Str != "depth" && Str != "width") {
+      S.Diag(Attr.getLoc(), diag::err_intel_fpga_merge_dir_invalid) << Attr;
+      return;
+    }
+    Results.push_back(Str);
+  }
+
+  if (!D->hasAttr<IntelFPGAMemoryAttr>())
+    D->addAttr(IntelFPGAMemoryAttr::CreateImplicit(
+        S.Context, IntelFPGAMemoryAttr::Default));
+
+  D->addAttr(::new (S.Context) IntelFPGAMergeAttr(
+      Attr.getRange(), S.Context, Results[0], Results[1],
+      Attr.getAttributeSpellingListIndex()));
+}
+
 static void handleIntelFPGAMaxPrivateCopiesAttr(Sema &S, Decl *D,
                                               const ParsedAttr &Attr) {
   checkForDuplicateAttribute<IntelFPGAMaxPrivateCopiesAttr>(S, D, Attr);
@@ -7509,6 +7574,15 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_IntelFPGAMaxPrivateCopies:
     handleIntelFPGAMaxPrivateCopiesAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_IntelFPGAMaxReplicates:
+    handleIntelFPGAMaxReplicatesAttr(S, D, AL);
+    break;
+  case ParsedAttr::AT_IntelFPGASimpleDualPort:
+    handleIntelFPGASimpleDualPortAttr(S, D, AL);
+    break;
+  case ParsedAttr::AT_IntelFPGAMerge:
+    handleIntelFPGAMergeAttr(S, D, AL);
+    break;
 
   case ParsedAttr::AT_AnyX86NoCallerSavedRegisters:
     handleSimpleAttribute<AnyX86NoCallerSavedRegistersAttr>(S, D, AL);

clang/lib/Sema/SemaSYCL.cpp

@@ -395,12 +395,23 @@ class KernelBodyTransform : public TreeTransform<KernelBodyTransform> {
       auto NewDecl = MappingPair.second;
       return DeclRefExpr::Create(
           SemaRef.getASTContext(), DRE->getQualifierLoc(),
-          DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(),
+          DRE->getTemplateKeywordLoc(), NewDecl, false,
+          DeclarationNameInfo(DRE->getNameInfo().getName(), SourceLocation(),
+                              DRE->getNameInfo().getInfo()),
           NewDecl->getType(), DRE->getValueKind());
     }
     return DRE;
   }
 
+  StmtResult RebuildCompoundStmt(SourceLocation LBraceLoc,
+                                 MultiStmtArg Statements,
+                                 SourceLocation RBraceLoc,
+                                 bool IsStmtExpr) {
+    // Build a new compound statement but clear the source locations.
+    return getSema().ActOnCompoundStmt(SourceLocation(), SourceLocation(),
+                                       Statements, IsStmtExpr);
+  }
+
 private:
   std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair;
   Sema &SemaRef;
@@ -530,8 +541,8 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S,
         auto ME = MemberExpr::Create(
             S.Context, SpecialObjME, false, SourceLocation(),
             NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP,
-            InitMethod->getNameInfo(), nullptr, InitMethod->getType(),
-            VK_LValue, OK_Ordinary);
+            DeclarationNameInfo(InitMethod->getDeclName(), SourceLocation()),
+            nullptr, InitMethod->getType(), VK_LValue, OK_Ordinary);
 
         // Not referenced -> not emitted
         S.MarkFunctionReferenced(SourceLocation(), InitMethod, true);

clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp

@@ -1,30 +1,38 @@
-// RUN: %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s --check-prefixes CHECK
+// RUN: %clang --sycl %s -S -I %S/Inputs -emit-llvm -g -o - | FileCheck %s
 //
-// Verify the SYCL kernel routine is marked artificial.
+// Verify the SYCL kernel routine is marked artificial and has no source
+// correlation.
 //
-// Since it has no source correlation of its own, the SYCL kernel needs to be
-// marked artificial or it will inherit source correlation from the surrounding
-// code.
+// The SYCL kernel should have no source correlation of its own, so it needs
+// to be marked artificial or it will inherit source correlation from the
+// surrounding code.
 //
 
-template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+#include <sycl.hpp>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
   kernelFunc();
 }
 
 int main() {
-  int value = 0;
-  int* captured = &value;
-  kernel_single_task<class kernel_function>([=]() {
-      *captured = 1;
-      });
+  cl::sycl::sampler Sampler;
+  kernel<class use_kernel_for_test>([=]() {
+    Sampler.use();
+  });
   return 0;
 }
 
-
-// CHECK: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32 addrspace(1)*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
+// CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
+// CHECK: getelementptr inbounds %"class.{{.*}}.anon"{{.*}} !dbg [[LINE_A0:![0-9]+]]
+// CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
+// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"{{.*}} !dbg [[LINE_B0:![0-9]+]]
+// CHECK: ret void, !dbg [[LINE_A0]]
 // CHECK: [[FILE:![0-9]+]] = !DIFile(filename: "{{.*}}debug-info-srcpos-kernel.cpp"{{.*}})
-// CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "_ZTSZ4mainE15kernel_function"
+// CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "{{.*}}19use_kernel_for_test"
 // CHECK-SAME: scope: [[FILE]]
 // CHECK-SAME: file: [[FILE]]
 // CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped
+// CHECK: [[LINE_A0]] = !DILocation(line: 0
+// CHECK: [[LINE_B0]] = !DILocation(line: 0
+