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

Author: iclsrc

clang/lib/Driver/Driver.cpp

@@ -3414,6 +3414,9 @@ class OffloadingActionBuilder final {
           ++TC;
           continue;
         }
+        if (LI.empty())
+          // Current list is empty, nothing to process.
+          continue;
 
         // Perform a check for device kernels.  This is done for FPGA when an
         // aocx or aocr based file is found.

clang/lib/Driver/ToolChains/Clang.cpp

@@ -3853,8 +3853,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     CmdArgs.push_back(Args.MakeArgString(NormalizedTriple));
 
     bool IsMSVC = AuxT.isWindowsMSVCEnvironment();
-    if (types::isCXX(Input.getType()))
-      CmdArgs.push_back(IsMSVC ? "-std=c++14" : "-std=c++11");
     if (IsMSVC) {
       CmdArgs.push_back("-fms-extensions");
       CmdArgs.push_back("-fms-compatibility");

clang/lib/Sema/SemaDeclAttr.cpp

@@ -4613,10 +4613,11 @@ bool Sema::CheckCallingConvAttr(const ParsedAttr &Attrs, CallingConv &CC,
     return true;
   }
 
+  const TargetInfo &TI = Context.getTargetInfo();
   // TODO: diagnose uses of these conventions on the wrong target.
   switch (Attrs.getKind()) {
   case ParsedAttr::AT_CDecl:
-    CC = CC_C;
+    CC = TI.getDefaultCallingConv();
     break;
   case ParsedAttr::AT_FastCall:
     CC = CC_X86FastCall;
@@ -4681,7 +4682,6 @@ bool Sema::CheckCallingConvAttr(const ParsedAttr &Attrs, CallingConv &CC,
   }
 
   TargetInfo::CallingConvCheckResult A = TargetInfo::CCCR_OK;
-  const TargetInfo &TI = Context.getTargetInfo();
   // CUDA functions may have host and/or device attributes which indicate
   // their targeted execution environment, therefore the calling convention
   // of functions in CUDA should be checked against the target deduced based

clang/lib/Sema/SemaSYCL.cpp

@@ -989,7 +989,8 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) {
 // Fields of kernel object must be initialized with SYCL kernel arguments so
 // in the following function we extract types of kernel object fields and add it
 // to the array with kernel parameters descriptors.
-static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
+// Returns true if all arguments are successfully built.
+static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
                         SmallVectorImpl<ParamDesc> &ParamDescs) {
   const LambdaCapture *Cpt = KernelObj->captures_begin();
   auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) {
@@ -1040,6 +1041,7 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
             }
           };
 
+  bool AllArgsAreValid = true;
   // Run through kernel object fields and create corresponding kernel
   // parameters descriptors. There are a several possible cases:
   //   - Kernel object field is a SYCL special object (SYCL accessor or SYCL
@@ -1054,17 +1056,22 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
     QualType ArgTy = Fld->getType();
     if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) {
       createSpecialSYCLObjParamDesc(Fld, ArgTy);
-    } else if (ArgTy->isStructureOrClassType()) {
+    } else if (!ArgTy->isStandardLayoutType()) {
       // SYCL v1.2.1 s4.8.10 p5:
       // C++ non-standard layout values must not be passed as arguments to a
       // kernel that is compiled for a device.
-      if (!ArgTy->isStandardLayoutType()) {
-        const DeclaratorDecl *V =
-            Cpt ? cast<DeclaratorDecl>(Cpt->getCapturedVar())
-                : cast<DeclaratorDecl>(Fld);
-        KernelObj->getASTContext().getDiagnostics().Report(
-            V->getLocation(), diag::err_sycl_non_std_layout_type);
-      }
+      const auto &DiagLocation =
+          Cpt ? Cpt->getLocation() : cast<DeclaratorDecl>(Fld)->getLocation();
+
+      Context.getDiagnostics().Report(DiagLocation,
+                                      diag::err_sycl_non_std_layout_type);
+
+      // Set the flag and continue processing so we can emit error for each
+      // invalid argument.
+      AllArgsAreValid = false;
+    } else if (ArgTy->isStructureOrClassType()) {
+      assert(ArgTy->isStandardLayoutType());
+
       CreateAndAddPrmDsc(Fld, ArgTy);
 
       // Create descriptors for each accessor field in the class or struct
@@ -1077,22 +1084,27 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
       PointeeTy = Context.getQualifiedType(PointeeTy.getUnqualifiedType(),
                                            Quals);
       QualType ModTy = Context.getPointerType(PointeeTy);
-      
+
       CreateAndAddPrmDsc(Fld, ModTy);
     } else if (ArgTy->isScalarType()) {
       CreateAndAddPrmDsc(Fld, ArgTy);
     } else {
       llvm_unreachable("Unsupported kernel parameter type");
     }
+
+    // Update capture iterator as we process arguments
+    if (Cpt && Cpt != KernelObj->captures_end())
+      ++Cpt;
   }
+
+  return AllArgsAreValid;
 }
 
 /// Adds necessary data describing given kernel to the integration header.
 /// \param H           the integration header object
 /// \param Name        kernel name
 /// \param NameType    type representing kernel name (first template argument
-/// of
-///                      single_task, parallel_for, etc)
+/// of single_task, parallel_for, etc)
 /// \param KernelObjTy kernel object type
 static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
                               QualType NameType, CXXRecordDecl *KernelObjTy) {
@@ -1238,15 +1250,24 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
 
   // Build list of kernel arguments
   llvm::SmallVector<ParamDesc, 16> ParamDescs;
-  buildArgTys(getASTContext(), LE, ParamDescs);
+  if (!buildArgTys(getASTContext(), LE, ParamDescs))
+    return;
 
   // Extract name from kernel caller parameters and mangle it.
   const TemplateArgumentList *TemplateArgs =
       KernelCallerFunc->getTemplateSpecializationArgs();
   assert(TemplateArgs && "No template argument info");
   QualType KernelNameType = TypeName::getFullyQualifiedType(
       TemplateArgs->get(0).getAsType(), getASTContext(), true);
-  std::string Name = constructKernelName(KernelNameType, MC);
+
+  std::string Name;
+  // TODO SYCLIntegrationHeader also computes a unique stable name. It should
+  // probably lose this responsibility and only use the name provided here.
+  if (getLangOpts().SYCLUnnamedLambda)
+    Name = PredefinedExpr::ComputeName(
+        getASTContext(), PredefinedExpr::UniqueStableNameExpr, KernelNameType);
+  else
+    Name = constructKernelName(KernelNameType, MC);
 
   // TODO Maybe don't emit integration header inside the Sema?
   populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE);

clang/test/CodeGenSYCL/fpga_pipes.cpp

@@ -6,9 +6,9 @@
 class SomePipe;
 void foo() {
   using Pipe = cl::sycl::pipe<SomePipe, int>;
-  // CHECK: %WPipe = alloca %opencl.pipe_wo_t
+  // CHECK: %{{.*}} = alloca %opencl.pipe_wo_t
   Pipe::write(42);
-  // CHECK: %RPipe = alloca %opencl.pipe_ro_t
+  // CHECK: %{{.*}} = alloca %opencl.pipe_ro_t
   int a = Pipe::read();
 }
 

clang/test/Driver/sycl-offload-intelfpga.cpp

@@ -88,23 +88,22 @@
 // RUN:  llc -filetype=obj -o %t-aocx.o %t-aocx.bc
 // RUN:  llvm-ar crv %t_aocx.a %t.o %t-aocx.o
 // RUN:  %clangxx -target x86_64-unknown-linux-gnu -fsycl -fintelfpga %t_aocx.a -ccc-print-phases 2>&1 \
-// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX-PHASES,CHK-FPGA-AOCX-PHASES-DEFAULT %s
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX-PHASES %s
 // RUN:  %clang_cl -fsycl -fintelfpga %t_aocx.a -ccc-print-phases 2>&1 \
-// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX-PHASES,CHK-FPGA-AOCX-PHASES-CL %s
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX-PHASES %s
 // CHK-FPGA-AOCX-PHASES: 0: input, "{{.*}}", object, (host-sycl)
 // CHK-FPGA-AOCX-PHASES: 1: linker, {0}, image, (host-sycl)
-// CHK-FPGA-AOCX-PHASES: 2: linker, {}, spirv, (device-sycl)
-// CHK-FPGA-AOCX-PHASES: 3: backend-compiler, {2}, fpga-aocx, (device-sycl)
-// CHK-FPGA-AOCX-PHASES: 4: clang-offload-wrapper, {3}, object, (device-sycl)
-// CHK-FPGA-AOCX-PHASES-DEFAULT: 5: offload, "host-sycl (x86_64-unknown-linux-gnu)" {1}, "device-sycl (spir64_fpga-unknown-unknown-sycldevice)" {4}, image
-// CHK-FPGA-AOCX-PHASES-CL: 5: offload, "host-sycl (x86_64-pc-windows-msvc)" {1}, "device-sycl (spir64_fpga-unknown-unknown-sycldevice-coff)" {4}, image
 
 // RUN:  %clangxx -target x86_64-unknown-linux-gnu -fsycl -fintelfpga %t_aocx.a -### 2>&1 \
-// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX %s
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX,CHK-FPGA-AOCX-LIN %s
+// RUN:  %clang_cl -fsycl -fintelfpga %t_aocx.a -### 2>&1 \
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-AOCX,CHK-FPGA-AOCX-WIN %s
 // CHK-FPGA-AOCX: clang-offload-bundler{{.*}} "-type=ao" "-targets=sycl-fpga_aocx-intel-unknown-sycldevice" "-inputs=[[LIBINPUT:.+\.a]]" "-outputs=[[BUNDLEOUT:.+\.aocx]]" "-unbundle"
-// CHK-FPGA-AOCX: clang-offload-wrapper{{.*}} "-o=[[WRAPOUT:.+\.bc]]" "-host=x86_64-unknown-linux-gnu" "-target=spir64_fpga" "-kind=sycl" "[[BUNDLEOUT]]"
-// CHK-FPGA-AOCX: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.o]]" "[[WRAPOUT]]"
-// CHK-FPGA-AOCX: ld{{.*}} "[[LIBINPUT]]" "[[LLCOUT]]"
+// CHK-FPGA-AOCX: clang-offload-wrapper{{.*}} "-o=[[WRAPOUT:.+\.bc]]" {{.*}} "-target=spir64_fpga" "-kind=sycl" "[[BUNDLEOUT]]"
+// CHK-FPGA-AOCX-LIN: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT:.+\.o]]" "[[WRAPOUT]]"
+// CHK-FPGA-AOCX-WIN: llc{{.*}} "-filetype=obj" "-o" "[[LLCOUT2:.+\.obj]]" "[[WRAPOUT]]"
+// CHK-FPGA-AOCX-LIN: ld{{.*}} "[[LIBINPUT]]" "[[LLCOUT]]"
+// CHK-FPGA-AOCX-WIN: link{{.*}} "[[LIBINPUT]]" "[[LLCOUT2]]"
 
 /// -fintelfpga -fsycl-link from source
 // RUN: touch %t.cpp

clang/test/Driver/sycl-offload.c

@@ -793,5 +793,12 @@
 // RUN:  | FileCheck -check-prefix=DUP-OPT %s
 // DUP-OPT-NOT: aoc{{.*}} "-DBLAH" {{.*}} "-DBLAH"
 
+/// passing of only a library should not create a device link
+// RUN: %clang -ccc-print-phases -fsycl -lsomelib 2>&1 \
+// RUN:  | FileCheck -check-prefix=LIB-NODEVICE %s
+// LIB-NODEVICE: 0: input, "somelib", object, (host-sycl)
+// LIB-NODEVICE: 1: linker, {0}, image, (host-sycl)
+// LIB-NODEVICE-NOT: linker, {{.*}}, spirv, (device-sycl)
+
 // TODO: SYCL specific fail - analyze and enable
 // XFAIL: windows-msvc

clang/test/Driver/sycl.c

@@ -4,10 +4,14 @@
 // RUN: %clang -### --sycl -fno-sycl-use-bitcode -c %s 2>&1 | FileCheck %s --check-prefix=NO-BITCODE
 // RUN: %clang -### -target spir64-unknown-linux-sycldevice -c -emit-llvm %s 2>&1 | FileCheck %s --check-prefix=TARGET
 // RUN: %clang -### --sycl -c -emit-llvm %s 2>&1 | FileCheck %s --check-prefix=COMBINED
+// RUN: %clangxx -### -fsycl-device-only %s 2>&1 | FileCheck %s --check-prefix=DEFAULT
+// RUN: %clang_cl -### -fsycl-device-only %s 2>&1 | FileCheck %s --check-prefix=DEFAULT
 
-// DEFAULT: "-triple" "spir64-unknown-{{.*}}-sycldevice"{{.*}} "-fsycl-is-device"{{.*}} "-emit-llvm-bc"
+// DEFAULT: "-triple" "spir64-unknown-{{.*}}-sycldevice{{.*}}" "-fsycl-is-device"{{.*}} "-emit-llvm-bc"
 // DEFAULT: "-internal-isystem" "{{.*lib.*clang.*include}}"
 // DEFAULT-NOT: "{{.*}}llvm-spirv"{{.*}} "-spirv-max-version=1.1"{{.*}} "-spirv-ext=+all"
+// DEFAULT-NOT: "-std=c++11"
+// DEFAULT-NOT: "-std=c++14"
 // NO-BITCODE: "-triple" "spir64-unknown-{{.*}}-sycldevice"{{.*}} "-fsycl-is-device"{{.*}} "-emit-llvm-bc"
 // NO-BITCODE: "{{.*}}llvm-spirv"{{.*}} "-spirv-max-version=1.1"{{.*}} "-spirv-ext=+all"
 // TARGET: "-triple" "spir64-unknown-linux-sycldevice"{{.*}} "-fsycl-is-device"{{.*}} "-emit-llvm-bc"

clang/test/SemaSYCL/mangle-unnamed-kernel.cpp

@@ -8,5 +8,5 @@ int main() {
   return 0;
 }
 
-// CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_
-// CHECK: _ZTSZZ4mainENK3$_1clERN2cl4sycl7handlerEEUlvE_
+// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE6->12clES2_EUlvE6->54{{.*}}
+// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE7->12clES2_EUlvE7->54{{.*}}

clang/test/SemaSYCL/non-std-layout-param.cpp

@@ -20,9 +20,61 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
 
 
 void test() {
-  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
   C C0;
   C0.Y=0;
-  kernel_single_task<class MyKernel>([=] { (void)C0.Y; });
+  kernel_single_task<class MyKernel>([=] {
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+    (void)C0.Y;
+  });
+}
+
+void test_capture_explicit_ref() {
+  int p = 0;
+  double q = 0;
+  float s = 0;
+  kernel_single_task<class kernel_capture_single_ref>([
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+      &p,
+      q,
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+      &s] {
+    (void) q;
+    (void) p;
+    (void) s;
+  });
 }
 
+void test_capture_implicit_refs() {
+  int p = 0;
+  double q = 0;
+  kernel_single_task<class kernel_capture_refs>([&] {
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+    (void) p;
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+    (void) q;
+  });
+}
+
+struct Kernel {
+  void operator()() {
+    (void) c1;
+    (void) c2;
+    (void) p;
+    (void) q;
+  }
+
+  int p;
+  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+  C c1;
+
+  int q;
+
+  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+  C c2;
+};
+
+void test_struct_field() {
+  Kernel k{};
+
+  kernel_single_task<class kernel_object>(k);
+}

clang/test/SemaSYCL/sycl-cconv.cpp

@@ -1,14 +1,23 @@
 // RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -fsyntax-only -verify %s
 
-//expected-warning@+1 {{'__cdecl' calling convention is not supported for this target}}
+// expected-no-warning@+1
 __inline __cdecl int printf(char const* const _Format, ...) { return 0; }
+// expected-no-warning@+1
+__inline __cdecl __attribute__((sycl_device)) int foo() { return 0; }
+// expected-no-warning@+1
+__inline __cdecl int moo() { return 0; }
 
 void bar() {
   printf("hello\n"); // expected-no-error
 }
 
 template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+// expected-no-warning@+1
+__cdecl __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  // expected-error@+1{{SYCL kernel cannot call a variadic function}}
+  printf("cannot call from here\n");
+  // expected-no-error@+1
+  moo();
   kernelFunc();
 }
 

sycl/doc/SYCLPluginInterface.md

@@ -120,17 +120,17 @@ These APIs match semantics of the corresponding OpenCL APIs exactly.
 For example:
 
 ```
-pi_program piclProgramCreateWithSource(
+pi_result piclProgramCreateWithSource(
   pi_context        context,
   pi_uint32         count,
   const char **     strings,
   const size_t *    lengths,
-  pi_result *       errcode);
+  pi_program *      ret_program);
 ```
 
 ### PI Extension mechanism
 
 TBD This section describes a mechanism for SYCL or other runtimes to detect
 availability of and obtain interfaces beyond those defined by the PI dispatch.
 
-TBD Add API to query PI version supported by plugin at runtime.
+TBD Add API to query PI version supported by plugin at runtime.