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

  CONFLICT (content): Merge conflict in sycl/plugins/opencl/pi_opencl.cpp

Author: bader

clang/lib/Driver/Driver.cpp

@@ -5527,9 +5527,11 @@ const char *Driver::GetNamedOutputPath(Compilation &C, const JobAction &JA,
   if ((!AtTopLevel && !isSaveTempsEnabled() &&
        (!C.getArgs().hasArg(options::OPT__SLASH_Fo) ||
         // FIXME - The use of /Fo is limited when offloading is enabled.  When
-        // compiling to exe use of /Fo does not produce the named obj
+        // compiling to exe use of /Fo does not produce the named obj.  We also
+        // should not use the named output when performing unbundling.
         (C.getArgs().hasArg(options::OPT__SLASH_Fo) &&
          (!JA.isOffloading(Action::OFK_None) ||
+          isa<OffloadUnbundlingJobAction>(JA) ||
           JA.getOffloadingHostActiveKinds() > Action::OFK_Host)))) ||
       CCGenDiagnostics) {
     StringRef Name = llvm::sys::path::filename(BaseInput);
@@ -5567,7 +5569,8 @@ const char *Driver::GetNamedOutputPath(Compilation &C, const JobAction &JA,
   // Determine what the derived output name should be.
   const char *NamedOutput;
 
-  if ((JA.getType() == types::TY_Object || JA.getType() == types::TY_LTO_BC) &&
+  if ((JA.getType() == types::TY_Object || JA.getType() == types::TY_LTO_BC ||
+       JA.getType() == types::TY_Archive) &&
       C.getArgs().hasArg(options::OPT__SLASH_Fo, options::OPT__SLASH_o)) {
     // The /Fo or /o flag decides the object filename.
     StringRef Val =

clang/lib/Sema/SemaSYCL.cpp

@@ -1285,6 +1285,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
   FunctionDecl *OpenCLKernel =
       CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs);
 
+  ContextRAII FuncContext(*this, OpenCLKernel);
+
   // Let's copy source location of a functor/lambda to emit nicer diagnostics
   OpenCLKernel->setLocation(LE->getLocation());
 

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

@@ -16,30 +16,44 @@
 
 /// -fintelfpga -fsycl-link tests
 // RUN:  touch %t.o
-// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link %t.o 2>&1 \
+// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link %t.o -o libfoo.a 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK,CHK-FPGA-EARLY %s
-// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link=early %t.o 2>&1 \
+// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link=early %t.o -o libfoo.a 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK,CHK-FPGA-EARLY %s
-// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link=image %t.o 2>&1 \
+// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link=image %t.o -o libfoo.a 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK,CHK-FPGA-IMAGE %s
 // CHK-FPGA-LINK-NOT: clang-offload-bundler{{.*}} "-check-section"
 // CHK-FPGA-LINK: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64_fpga-unknown-unknown-sycldevice" "-inputs=[[INPUT:.+\.o]]" "-outputs=[[OUTPUT1:.+\.o]]" "-unbundle"
+// CHK-FPGA-LINK-NOT: clang-offload-bundler{{.*}}
 // CHK-FPGA-LINK: llvm-link{{.*}} "[[OUTPUT1]]" "-o" "[[OUTPUT2:.+\.bc]]"
 // CHK-FPGA-LINK: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.spv]]" "-spirv-max-version=1.1" "-spirv-ext=+all" "[[OUTPUT2]]"
 // CHK-FPGA-EARLY: aoc{{.*}} "-o" "[[OUTPUT4:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl"
 // CHK-FPGA-IMAGE: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocx]]" "[[OUTPUT3]]" "-sycl"
-// CHK-FPGA-LINK: {{lib|llvm-ar}}{{.*}} "[[INPUT]]"
+// CHK-FPGA-LINK: llvm-ar{{.*}} "cr" "libfoo.a" "[[INPUT]]"
+
+// Output designation should not be used for unbundling step
+// RUN:  touch %t.o
+// RUN:  touch %t.obj
+// RUN:  %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -fsycl-link %t.o -o libfoo.a 2>&1 \
+// RUN:  | FileCheck -check-prefix=CHK-FPGA-LINK-OUT %s
+// RUN:  %clang_cl -### -fsycl -fintelfpga -fsycl-link %t.obj -Folibfoo.a 2>&1 \
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK-OUT %s
+// RUN:  %clang_cl -### -fsycl -fintelfpga -fsycl-link %t.obj -o libfoo.a 2>&1 \
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK-OUT %s
+// CHK-FPGA-LINK-OUT-NOT: clang-offload-bundler{{.*}} "-outputs=libfoo.a" "-unbundle"
 
 /// -fintelfpga -fsycl-link clang-cl specific
 // RUN:  touch %t.obj
-// RUN:  %clang_cl -### -fsycl -fintelfpga -fsycl-link %t.obj 2>&1 \
+// RUN:  %clang_cl -### -fsycl -fintelfpga -fsycl-link %t.obj -Folibfoo.lib 2>&1 \
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK-WIN %s
+// RUN:  %clang_cl -### -fsycl -fintelfpga -fsycl-link %t.obj -o libfoo.lib 2>&1 \
 // RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK-WIN %s
 // CHK-FPGA-LINK-WIN: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64_fpga-unknown-unknown-sycldevice{{.*}}" "-inputs=[[INPUT:.+\.obj]]" "-outputs=[[OUTPUT1:.+\.obj]]" "-unbundle"
+// CHK-FPGA-LINK-WIN-NOT: clang-offload-bundler{{.*}}
 // CHK-FPGA-LINK-WIN: llvm-link{{.*}} "[[OUTPUT1]]" "-o" "[[OUTPUT2:.+\.bc]]"
 // CHK-FPGA-LINK-WIN: llvm-spirv{{.*}} "-o" "[[OUTPUT3:.+\.spv]]" "-spirv-max-version=1.1" "-spirv-ext=+all" "[[OUTPUT2]]"
 // CHK-FPGA-LINK-WIN: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocr]]" "[[OUTPUT3]]" "-sycl" "-rtl"
-// CHK-FPGA-LINK-WIN: lib.exe{{.*}} "[[INPUT]]"
-
+// CHK-FPGA-LINK-WIN: lib.exe{{.*}} "[[INPUT]]" {{.*}} "-OUT:libfoo.lib"
 
 /// Check -fintelfpga -fsycl-link with an FPGA archive
 // Create the dummy archive

sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

@@ -165,11 +165,14 @@ template <typename AllocatorT> class SYCLMemObjT : public SYCLMemObjI {
 
   template <typename Destination>
   EnableIfOutputPointerT<Destination> set_final_data(Destination FinalData) {
-    MUploadDataFunctor = [this, FinalData]() {
-      EventImplPtr Event = updateHostMemory(FinalData);
-      if (Event)
-        Event->wait(Event);
-    };
+    if (!FinalData)
+      MUploadDataFunctor = nullptr;
+    else
+      MUploadDataFunctor = [this, FinalData]() {
+        EventImplPtr Event = updateHostMemory(FinalData);
+        if (Event)
+          Event->wait(Event);
+      };
   }
 
   template <typename Destination>

sycl/plugins/opencl/pi_opencl.cpp

@@ -30,17 +30,28 @@ template <class To, class From> To cast(From value) {
   return (To)(value);
 }
 
+// Older versions of GCC don't like "const" here
+#if defined(__GNUC__) && (__GNUC__ < 7 || (__GNU__C == 7 && __GNUC_MINOR__ < 2))
+#define CONSTFIX constexpr
+#else
+#define CONSTFIX const
+#endif
+
 // Names of USM functions that are queried from OpenCL
-constexpr char clHostMemAllocName[] = "clHostMemAllocINTEL";
-constexpr char clDeviceMemAllocName[] = "clDeviceMemAllocINTEL";
-constexpr char clSharedMemAllocName[] = "clSharedMemAllocINTEL";
-constexpr char clMemFreeName[] = "clMemFreeINTEL";
-constexpr char clSetKernelArgMemPointerName[] = "clSetKernelArgMemPointerINTEL";
-constexpr char clEnqueueMemsetName[] = "clEnqueueMemsetINTEL";
-constexpr char clEnqueueMemcpyName[] = "clEnqueueMemcpyINTEL";
-constexpr char clEnqueueMigrateMemName[] = "clEnqueueMigrateMemINTEL";
-constexpr char clEnqueueMemAdviseName[] = "clEnqueueMemAdviseINTEL";
-constexpr char clGetMemAllocInfoName[] = "clGetMemAllocInfoINTEL";
+CONSTFIX char clHostMemAllocName[] = "clHostMemAllocINTEL";
+CONSTFIX char clDeviceMemAllocName[] = "clDeviceMemAllocINTEL";
+CONSTFIX char clSharedMemAllocName[] = "clSharedMemAllocINTEL";
+CONSTFIX char clMemFreeName[] = "clMemFreeINTEL";
+CONSTFIX char clSetKernelArgMemPointerName[] = "clSetKernelArgMemPointerINTEL";
+CONSTFIX char clEnqueueMemsetName[] = "clEnqueueMemsetINTEL";
+CONSTFIX char clEnqueueMemcpyName[] = "clEnqueueMemcpyINTEL";
+CONSTFIX char clEnqueueMigrateMemName[] = "clEnqueueMigrateMemINTEL";
+CONSTFIX char clEnqueueMemAdviseName[] = "clEnqueueMemAdviseINTEL";
+CONSTFIX char clGetMemAllocInfoName[] = "clGetMemAllocInfoINTEL";
+
+#undef CONSTFIX
+
+
 
 // USM helper function to get an extension function pointer
 template <const char *FuncName, typename T>

sycl/source/detail/scheduler/graph_builder.cpp

@@ -232,16 +232,31 @@ Command *Scheduler::GraphBuilder::insertMemoryMove(MemObjRecord *Record,
 
   AllocaCommandBase *AllocaCmdSrc =
       findAllocaForReq(Record, Req, Record->MCurContext);
+  if (!AllocaCmdSrc && IsSuitableSubReq(Req)) {
+    // Since no alloca command for the sub buffer requirement was found in the
+    // current context, need to find a parent alloca command for it (it must be
+    // there)
+    auto IsSuitableAlloca = [Record, Req](AllocaCommandBase *AllocaCmd) {
+      bool Res = sameCtx(AllocaCmd->getQueue()->get_context_impl(),
+                         Record->MCurContext) &&
+                 // Looking for a parent buffer alloca command
+                 AllocaCmd->getType() == Command::CommandType::ALLOCA;
+      return Res;
+    };
+    const auto It =
+        std::find_if(Record->MAllocaCommands.begin(),
+                     Record->MAllocaCommands.end(), IsSuitableAlloca);
+    AllocaCmdSrc = (Record->MAllocaCommands.end() != It) ? *It : nullptr;
+  }
   if (!AllocaCmdSrc)
     throw runtime_error("Cannot find buffer allocation");
   // Get parent allocation of sub buffer to perform full copy of whole buffer
   if (IsSuitableSubReq(Req)) {
     if (AllocaCmdSrc->getType() == Command::CommandType::ALLOCA_SUB_BUF)
       AllocaCmdSrc =
           static_cast<AllocaSubBufCommand *>(AllocaCmdSrc)->getParentAlloca();
-    else
-      assert(
-          !"Inappropriate alloca command. AllocaSubBufCommand was expected.");
+    else if (AllocaCmdSrc->getSYCLMemObj() != Req->MSYCLMemObj)
+      assert(!"Inappropriate alloca command.");
   }
 
   Command *NewCmd = nullptr;

sycl/test/basic_tests/buffer/subbuffer.cpp

@@ -247,11 +247,37 @@ void copyBlock() {
   }
 }
 
+void checkMultipleContexts() {
+  constexpr int N = 64;
+  int a[N] = {0};
+  {
+    sycl::queue queue1;
+    sycl::queue queue2;
+    sycl::buffer<int, 1> buf(a, sycl::range<1>(N));
+    sycl::buffer<int, 1> subbuf1(buf, sycl::id<1>(0), sycl::range<1>(N / 2));
+    sycl::buffer<int, 1> subbuf2(buf, sycl::id<1>(N / 2),
+                                 sycl::range<1>(N / 2));
+    queue1.submit([&](sycl::handler &cgh) {
+      auto bufacc = subbuf1.get_access<sycl::access::mode::read_write>(cgh);
+      cgh.parallel_for<class sub_buffer_1>(
+          sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = 1; });
+    });
+
+    queue2.submit([&](sycl::handler &cgh) {
+      auto bufacc = subbuf2.get_access<sycl::access::mode::read_write>(cgh);
+      cgh.parallel_for<class sub_buffer_2>(
+          sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = 2; });
+    });
+  }
+  assert(a[N / 2 - 1] == 1 && a[N / 2] == 2 && "Sub buffer data loss");
+}
+
 int main() {
   cl::sycl::queue q;
   check1DSubBuffer(q);
   checkHostAccessor(q);
   checkExceptions();
   copyBlock();
+  checkMultipleContexts();
   return 0;
 }

sycl/test/basic_tests/image.cpp

@@ -73,6 +73,31 @@ int main() {
     });
   }
 
+  // image with write accessor to it in kernel
+  {
+    int NX = 32;
+    int NY = 32;
+
+    sycl::image<2> Img(sycl::image_channel_order::rgba,
+                       sycl::image_channel_type::fp32,
+                       sycl::range<2>(NX, NY));
+
+    sycl::queue Q;
+    Q.submit([&](sycl::handler &CGH) {
+        auto ImgAcc = Img.get_access<sycl::float4, sycl::access::mode::write>(
+              CGH);
+
+        sycl::nd_range<2> Rng(sycl::range<2>(NX, NY), sycl::range<2>(16, 16));
+
+        CGH.parallel_for<class sample>(Rng, [=](sycl::nd_item<2> Item) {
+            sycl::id<2> Idx = Item.get_global_id();
+            sycl::float4 C(0.5f, 0.5f, 0.2f, 1.0f);
+            ImgAcc.write(sycl::int2(Idx[0], Idx[1]), C);
+        });
+
+    }).wait_and_throw();
+  }
+
   std::cout << "Success" << std::endl;
   return 0;
 }

sycl/test/basic_tests/stream/auto_flush.cpp

@@ -3,7 +3,6 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
 // RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
-// TODO: SYCL specific fail - analyze and enable
 //==-------------- copy.cpp - SYCL stream obect auto flushing test ---------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

sycl/test/basic_tests/stream/stream.cpp

@@ -3,7 +3,6 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
 // RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
-// TODO: SYCL specific fail - analyze and enable
 //==------------------ stream.cpp - SYCL stream basic test -----------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

sycl/test/hier_par/hier_par_basic.cpp

@@ -8,7 +8,6 @@
 
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out
-// TODO: SYCL specific fail - analyze and enable on Windows
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

sycl/test/hier_par/hier_par_wgscope.cpp

@@ -8,7 +8,6 @@
 
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out
-// TODO: SYCL specific fail - analyze and enable on Windows
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out