Merge from 'sycl' to 'sycl-web' (#13)

Author: iclsrc

clang/include/clang/Basic/Attr.td

@@ -1735,6 +1735,13 @@ def IntelFPGASimpleDualPort : Attr {
   let Documentation = [IntelFPGASimpleDualPortAttrDocs];
 }
 
+def SYCLFPGAPipe : TypeAttr {
+  let Spellings = [GNU<"pipe">];
+  let Args = [StringArgument<"Mode">];
+  let LangOpts = [SYCLIsDevice];
+  let Documentation = [SYCLFPGAPipeDocs];
+}
+
 def Naked : InheritableAttr {
   let Spellings = [GCC<"naked">, Declspec<"naked">];
   let Subjects = SubjectList<[Function]>;

clang/include/clang/Basic/AttrDocs.td

@@ -1885,6 +1885,16 @@ function, and no effect otherwise.
   }];
 }
 
+def SYCLFPGAPipeDocs : Documentation {
+  let Category = DocCatStmt;
+  let Heading = "pipe (read_only, write_only)";
+  let Content = [{
+This attribute applies to a type to indicate that it is a pipe type.
+It requires a string argument that specifies if the pipe is read only or
+write only. Expected to be used only in SYCL headers.
+  }];
+}
+
 def SYCLIntelFPGAIVDepAttrDocs : Documentation {
   let Category = DocCatVariable;
   let Heading = "ivdep";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -137,6 +137,8 @@ def err_sycl_loop_attr_duplication : Error<
   "duplicate %select{unroll|Intel FPGA}0 loop attribute '%1'">;
 def err_loop_unroll_compatibility : Error<
   "incompatible loop unroll instructions: '%0' and '%1'">;
+def err_pipe_attribute_arg_not_allowed : Error<
+  "'%0' mode for pipe attribute is not supported. Allowed modes: 'read_only', 'write_only'">;
 
 // C99 variable-length arrays
 def ext_vla : Extension<"variable length arrays are a C99 feature">,

clang/include/clang/Basic/TokenKinds.def

@@ -578,7 +578,6 @@ KEYWORD(vec_step                    , KEYOPENCLC | KEYALTIVEC | KEYZVECTOR)
 KEYWORD(__builtin_omp_required_simd_align, KEYALL)
 
 KEYWORD(pipe                        , KEYOPENCLC | KEYOPENCLCXX)
-KEYWORD(__pipe                      , KEYSYCL)
 
 // Borland Extensions.
 KEYWORD(__pascal                    , KEYALL)

clang/lib/AST/TypePrinter.cpp

@@ -1509,6 +1509,10 @@ void TypePrinter::printAttributedAfter(const AttributedType *T,
     // AttributedType nodes for them.
     break;
 
+  case attr::SYCLFPGAPipe:
+    OS << "pipe";
+    break;
+
   case attr::LifetimeBound:
   case attr::TypeNonNull:
   case attr::TypeNullable:

clang/lib/Driver/Driver.cpp

@@ -3299,6 +3299,11 @@ class OffloadingActionBuilder final {
       if (auto *IA = dyn_cast<InputAction>(HostAction)) {
         SYCLDeviceActions.clear();
 
+        // Options that are considered LinkerInput are not valid input actions
+        // to the device tool chain.
+        if (IA->getInputArg().getOption().hasFlag(options::LinkerInput))
+          return ABRT_Inactive;
+
         std::string InputName = IA->getInputArg().getAsString(Args);
         // Objects should already be consumed with -foffload-static-lib
         if (Args.hasArg(options::OPT_foffload_static_lib_EQ) &&
@@ -3319,6 +3324,11 @@ class OffloadingActionBuilder final {
       if (auto *UA = dyn_cast<OffloadUnbundlingJobAction>(HostAction)) {
         SYCLDeviceActions.clear();
         if (auto *IA = dyn_cast<InputAction>(UA->getInputs().back())) {
+          // Options that are considered LinkerInput are not valid input actions
+          // to the device tool chain.
+          if (IA->getInputArg().getOption().hasFlag(options::LinkerInput))
+            return ABRT_Inactive;
+
           std::string FileName = IA->getInputArg().getAsString(Args);
           // Check if the type of the file is the same as the action. Do not
           // unbundle it if it is not. Do not unbundle .so files, for example,
@@ -3772,6 +3782,7 @@ class OffloadingActionBuilder final {
     // the input is not a bundle.
     if (CanUseBundler && isa<InputAction>(HostAction) &&
         InputArg->getOption().getKind() == llvm::opt::Option::InputClass &&
+        !InputArg->getOption().hasFlag(options::LinkerInput) &&
         !types::isSrcFile(HostAction->getType())) {
       std::string InputName = InputArg->getAsString(Args);
       // Do not create an unbundling action for an object when we know a fat

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5649,12 +5649,17 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     // Host-side SYCL compilation receives the integration header file as
     // Inputs[1].  Include the header with -include
     if (!IsSYCLOffloadDevice && SYCLDeviceInput) {
+      SmallString<128> RealPath;
+      // Fixup the header path name in case there are discrepancies in the
+      // string used for the temporary directory environment variable and
+      // actual path expectations.
+      llvm::sys::fs::real_path(SYCLDeviceInput->getFilename(), RealPath);
       CmdArgs.push_back("-include");
-      CmdArgs.push_back(SYCLDeviceInput->getFilename());
+      CmdArgs.push_back(Args.MakeArgString(RealPath));
       // When creating dependency information, filter out the generated
       // header file.
       CmdArgs.push_back("-dependency-filter");
-      CmdArgs.push_back(SYCLDeviceInput->getFilename());
+      CmdArgs.push_back(Args.MakeArgString(RealPath));
       // Let the FE know we are doing a SYCL offload compilation, but we are
       // doing the host pass.
       CmdArgs.push_back("-fsycl-is-host");

clang/lib/Parse/ParseDecl.cpp

@@ -3830,11 +3830,6 @@ void Parser::ParseDeclarationSpecifiers(DeclSpec &DS,
       }
       isInvalid = DS.SetTypePipe(true, Loc, PrevSpec, DiagID, Policy);
       break;
-    case tok::kw___pipe:
-      if (getLangOpts().SYCLIsDevice)
-        // __pipe keyword is defined only for SYCL kernel language
-        isInvalid = DS.SetTypePipe(true, Loc, PrevSpec, DiagID, Policy);
-      break;
 #define GENERIC_IMAGE_TYPE(ImgType, Id) \
   case tok::kw_##ImgType##_t: \
     isInvalid = DS.SetTypeSpecType(DeclSpec::TST_##ImgType##_t, Loc, PrevSpec, \
@@ -4967,9 +4962,8 @@ bool Parser::isDeclarationSpecifier(bool DisambiguatingWithExpression) {
   default: return false;
 
   case tok::kw_pipe:
-  case tok::kw___pipe:
     return (getLangOpts().OpenCL && getLangOpts().OpenCLVersion >= 200) ||
-            getLangOpts().OpenCLCPlusPlus || getLangOpts().SYCLIsDevice;
+           getLangOpts().OpenCLCPlusPlus;
 
   case tok::identifier:   // foo::bar
     // Unfortunate hack to support "Class.factoryMethod" notation.
@@ -5463,9 +5457,8 @@ static bool isPtrOperatorToken(tok::TokenKind Kind, const LangOptions &Lang,
   if (Kind == tok::star || Kind == tok::caret)
     return true;
 
-  if ((Kind == tok::kw_pipe || Kind == tok::kw___pipe) &&
-      ((Lang.OpenCL && Lang.OpenCLVersion >= 200) || Lang.OpenCLCPlusPlus ||
-       Lang.SYCLIsDevice))
+  if ((Kind == tok::kw_pipe) &&
+      ((Lang.OpenCL && Lang.OpenCLVersion >= 200) || Lang.OpenCLCPlusPlus))
     return true;
 
   if (!Lang.CPlusPlus)

clang/lib/Parse/ParseTentative.cpp

@@ -1467,8 +1467,6 @@ Parser::isCXXDeclarationSpecifier(Parser::TPResult BracedCastResult,
   case tok::kw___read_write:
     // OpenCL pipe
   case tok::kw_pipe:
-    // SYCL pipe
-  case tok::kw___pipe:
 
     // GNU
   case tok::kw_restrict:

clang/lib/Sema/SemaType.cpp

@@ -6103,6 +6103,66 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
   }
 }
 
+static void HandleSYCLFPGAPipeAttribute(QualType &Type, const ParsedAttr &Attr,
+                                        TypeProcessingState &State) {
+  Sema &S = State.getSema();
+  ASTContext &Ctx = S.Context;
+
+  // Check the attribute arguments.
+  if (Attr.getNumArgs() != 1) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments)
+        << Attr << 1;
+    Attr.setInvalid();
+    return;
+  }
+
+  if (!Attr.isArgExpr(0)) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_type)
+        << Attr << AANT_ArgumentString;
+    Attr.setInvalid();
+    return;
+  }
+
+  StringRef Str;
+  if (auto *SL = dyn_cast<StringLiteral>(Attr.getArgAsExpr(0))) {
+    Str = SL->getString();
+  } else {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_type)
+        << Attr << AANT_ArgumentString;
+    Attr.setInvalid();
+    return;
+  }
+
+  bool isReadOnlyPipe;
+  if (Str == "write_only")
+    isReadOnlyPipe = false;
+  else if (Str == "read_only")
+    isReadOnlyPipe = true;
+  else {
+    S.Diag(Attr.getLoc(), diag::err_pipe_attribute_arg_not_allowed) << Str;
+    Attr.setInvalid();
+    return;
+  }
+
+  auto *PipeAttr = ::new (Ctx) SYCLFPGAPipeAttr(Ctx, Attr, Str);
+
+  // Apply pipe qualifiers just to the equivalent type, as the expression is not
+  // value dependent (not templated).
+  QualType EquivType = isReadOnlyPipe
+                           ? S.BuildReadPipeType(Type, Attr.getLoc())
+                           : S.BuildWritePipeType(Type, Attr.getLoc());
+  if (EquivType.isNull()) {
+    Attr.setInvalid();
+    return;
+  }
+
+  QualType T = State.getAttributedType(PipeAttr, Type, EquivType);
+  if (!T.isNull())
+    Type = T;
+  else
+    Attr.setInvalid();
+}
+
 /// handleObjCOwnershipTypeAttr - Process an objc_ownership
 /// attribute on the specified type.
 ///
@@ -7670,6 +7730,10 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
       HandleOpenCLAccessAttr(type, attr, state.getSema());
       attr.setUsedAsTypeAttr();
       break;
+    case ParsedAttr::AT_SYCLFPGAPipe:
+      HandleSYCLFPGAPipeAttribute(type, attr, state);
+      attr.setUsedAsTypeAttr();
+      break;
     case ParsedAttr::AT_LifetimeBound:
       if (TAL == TAL_DeclChunk)
         HandleLifetimeBoundAttr(state, type, attr);

clang/test/CodeGenSYCL/fpga_pipes.cpp

@@ -0,0 +1,26 @@
+// RUN: %clang %s -S -emit-llvm --sycl -o - | FileCheck %s
+#include "CL/sycl.hpp"
+// CHECK: %opencl.pipe_wo_t
+// CHECK: %opencl.pipe_ro_t
+
+class SomePipe;
+void foo() {
+  using Pipe = cl::sycl::pipe<SomePipe, int>;
+  // CHECK: %WPipe = alloca %opencl.pipe_wo_t
+  Pipe::write(42);
+  // CHECK: %RPipe = alloca %opencl.pipe_ro_t
+  int a = Pipe::read();
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  kernel_single_task<class kernel_function>([]() {
+    foo();
+  });
+  return 0;
+}
+

clang/test/CodeGenSYCL/intel-restrict.cpp

@@ -53,7 +53,7 @@ int foo(int X) {
       cgh.single_task<class kernel_restrict_struct>(f);
     });
 
-    // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_other_params{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}, i32 %_arg_9)
+    // CHECK: define {{.*}} spir_kernel {{.*}}kernel_restrict_other_params{{.*}}(i32 addrspace(1)* noalias %{{.*}} i32 addrspace(1)* noalias %{{.*}}, i32 %{{[^,]*}})
     int num = 42;
     Q.submit([&](cl::sycl::handler& cgh) {
       auto AccA = BufA.get_access<sycl_read_write, sycl_global_buffer>(cgh);

clang/test/Driver/sycl-offload-header-check.cpp

@@ -0,0 +1,14 @@
+// REQUIRES: system-windows
+// Test the integration header file name generation.  Name should match the
+// actual path name and not the environment variable setting
+// RUN: mkdir -p %t_DiRnAmE
+// invoke the compiler overriding output temp location
+// RUN: env TMP=%t_dirname  \
+// RUN: %clang_cl -### -fsycl %s 2>&1 | \
+// RUN: FileCheck --check-prefix=CHECK-HEADER %s
+// RUN: env TMP=%t_dirname  \
+// RUN: %clang -### -fsycl %s 2>&1 | \
+// RUN: FileCheck --check-prefix=CHECK-HEADER %s
+// CHECK-HEADER: clang{{.*}} "-fsycl-int-header=[[HEADER:.+\.h]]"
+// CHECK-HEADER-NOT: clang{{.*}} "-include" "[[HEADER]]"
+// CHECK-HEADER: clang{{.*}} "-include" "{{.*}}_DiRnAmE{{.+}}.h"

clang/test/Driver/sycl-offload.c

@@ -182,6 +182,13 @@
 // CHK-PHASES-LIB: 15: clang-offload-wrapper, {14}, object, (device-sycl)
 // CHK-PHASES-LIB: 16: offload, "host-sycl (x86_64-unknown-linux-gnu)" {10}, "device-sycl (spir64-unknown-unknown-sycldevice)" {15}, image
 
+/// Compilation check with -lstdc++ (treated differently than regular lib)
+// RUN:   %clang -### -target x86_64-unknown-linux-gnu -lstdc++ -fsycl %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-LIB-STDCXX %s
+// CHK-LIB-STDCXX: ld{{.*}} "-lstdc++"
+// CHK-LIB-STDCXX-NOT: clang-offload-bundler{{.*}}
+// CHK-LIB-STDCXX-NOT: llvm-link{{.*}} "-lstdc++"
+
 /// ###########################################################################
 
 /// Check the phases when using and multiple source files

clang/test/SemaSYCL/fpga_pipes.cpp

@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -x c++ -fsycl-is-device -std=c++11 -fsyntax-only -verify -pedantic %s
+
+// no error expected
+using type1 = __attribute__((pipe("read_only"))) const int;
+
+// no error expected
+using type2 = __attribute__((pipe("write_only"))) const int;
+
+// expected-error@+1 {{'42' mode for pipe attribute is not supported. Allowed modes: 'read_only', 'write_only'}}
+using type3 = __attribute__((pipe("42"))) const int;
+
+// expected-error@+1{{'pipe' attribute requires a string}}
+using type4 = __attribute__((pipe(0))) const int;
+
+// expected-error@+1{{'pipe' attribute takes one argument}}
+using type5 = __attribute__((pipe)) const int;

clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp

@@ -11,7 +11,6 @@
 /// as input and creates wrapper bitcode file containing target binaries
 /// packaged as data. Wrapper bitcode also includes initialization code which
 /// registers target binaries in offloading runtime at program startup.
-/// TODO Add Windows support.
 ///
 //===----------------------------------------------------------------------===//
 
@@ -874,6 +873,9 @@ template <typename... Tys> class ListArgsSequencer {
   int Cur = -1;
 
   /// Class IDs of all options from all lists. Filled in the constructor.
+  /// Can also be seen as a map from command line position to the option class
+  /// ID. If there is no option participating in one of the sequenced lists at
+  /// given position, then it is mapped to -1 marker value.
   std::unique_ptr<std::vector<int>> OptListIDs;
 
   using tuple_of_iters_t = std::tuple<typename Tys::iterator...>;
@@ -894,8 +896,11 @@ template <typename... Tys> class ListArgsSequencer {
   /// Args - the cl::list objects to sequence elements of
   ListArgsSequencer(size_t Sz, Tys &... Args)
       : Prevs(Args.end()...), Iters(Args.begin()...) {
-    assert(Sz >= sizeof...(Tys));
+    // make OptListIDs big enough to hold IDs of all options coming from the
+    // command line and initialize all IDs to default class -1
     OptListIDs.reset(new std::vector<int>(Sz, -1));
+    // map command line positions where sequenced options occur to appropriate
+    // class IDs
     addLists<sizeof...(Tys) - 1, 0>(Args...);
   }
 
@@ -942,9 +947,12 @@ template <typename... Tys> class ListArgsSequencer {
 
   /// Does the actual sequencing of options found in given list.
   template <int ID, typename T> void addListImpl(T &L) {
+    // iterate via all occurences of an option of given list class
     for (auto It = L.begin(); It != L.end(); It++) {
+      // calculate its sequential position in the command line
       unsigned Pos = L.getPosition(It - L.begin());
       assert((*OptListIDs)[Pos] == -1);
+      // ... and fill the corresponding spot in the list with the class ID
       (*OptListIDs)[Pos] = ID;
     }
   }

sycl/CMakeLists.txt

@@ -178,6 +178,7 @@ add_subdirectory( source )
 # SYCL toolchain builds all components: compiler, libraries, headers, etc.
 add_custom_target( sycl-toolchain
   DEPENDS ${SYCL_RT_LIBS}
+          pi_opencl
           clang
           clang-offload-wrapper
           clang-offload-bundler
@@ -236,6 +237,7 @@ set( SYCL_TOOLCHAIN_DEPLOY_COMPONENTS
      opencl-headers
      sycl-headers
      sycl
+     pi_opencl
 )
 
 # Use it as fake dependency in order to force another command(s) to execute.

sycl/doc/GetStartedWithSYCLCompiler.md

@@ -87,9 +87,7 @@ cmake -G "Ninja" -DCMAKE_BUILD_TYPE=Release -DLLVM_TARGETS_TO_BUILD="X86" ^
 -DLLVM_ENABLE_PROJECTS="clang;llvm-spirv;sycl" ^
 -DLLVM_EXTERNAL_SYCL_SOURCE_DIR="%SYCL_HOME%\llvm\sycl" ^
 -DLLVM_EXTERNAL_LLVM_SPIRV_SOURCE_DIR="%SYCL_HOME%\llvm\llvm-spirv" ^
--DCMAKE_C_COMPILER=cl -DCMAKE_CXX_COMPILER=cl -DCMAKE_C_FLAGS="/GS" ^
--DCMAKE_CXX_FLAGS="/GS" -DCMAKE_EXE_LINKER_FLAGS="/NXCompat /DynamicBase" ^
--DCMAKE_SHARED_LINKER_FLAGS="/NXCompat /DynamicBase" ^
+-DCMAKE_C_COMPILER=cl -DCMAKE_CXX_COMPILER=cl ^
 "%SYCL_HOME%\llvm\llvm"
 
 ninja sycl-toolchain