Merge from 'sycl' to 'sycl-web' (intel#11)

Author: iclsrc

buildbot/configure.py

@@ -26,13 +26,14 @@ def do_configure(args):
                  "-DLLVM_EXTERNAL_PROJECTS=sycl;llvm-spirv",
                  "-DLLVM_EXTERNAL_SYCL_SOURCE_DIR={}".format(sycl_dir),
                  "-DLLVM_EXTERNAL_LLVM_SPIRV_SOURCE_DIR={}".format(spirv_dir),
-                 "-DLLVM_ENABLE_PROJECTS=clang;llvm-spirv;sycl",
+                 "-DLLVM_ENABLE_PROJECTS=clang;sycl;llvm-spirv",
                  "-DOpenCL_INCLUDE_DIR={}".format(ocl_header_dir),
                  "-DOpenCL_LIBRARY={}".format(icd_loader_lib),
                  "-DLLVM_BUILD_TOOLS=OFF",
                  "-DSYCL_ENABLE_WERROR=ON",
                  "-DLLVM_ENABLE_ASSERTIONS=ON",
                  "-DCMAKE_INSTALL_PREFIX={}".format(install_dir),
+                 "-DSYCL_INCLUDE_TESTS=ON", # Explicitly include all kinds of SYCL tests.
                  llvm_dir]
 
     print(cmake_cmd)

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -540,6 +540,11 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
   if (!FD->hasAttr<OpenCLKernelAttr>())
     return;
 
+  // TODO Module identifier is not reliable for this purpose since two modules
+  // can have the same ID, needs improvement
+  if (getLangOpts().SYCLIsDevice)
+    Fn->addFnAttr("sycl-module-id", Fn->getParent()->getModuleIdentifier());
+
   llvm::LLVMContext &Context = getLLVMContext();
 
   CGM.GenOpenCLArgMetadata(Fn, FD, this);

clang/lib/Driver/ToolChains/MSVC.cpp

@@ -324,7 +324,7 @@ void visualstudio::Linker::constructMSVCLibCommand(Compilation &C,
   CmdArgs.push_back(
       C.getArgs().MakeArgString(Twine("-OUT:") + Output.getFilename()));
 
-  SmallString<128> ExecPath(getToolChain().GetProgramPath("lib"));
+  SmallString<128> ExecPath(getToolChain().GetProgramPath("lib.exe"));
   const char *Exec = C.getArgs().MakeArgString(ExecPath);
   C.addCommand(std::make_unique<Command>(JA, *this, Exec, CmdArgs, None));
 }

clang/test/CodeGenSYCL/module-id.cpp

@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  kernel_single_task<class kernel>([]() {});
+  return 0;
+}
+// CHECK: define spir_kernel void @{{.*}}kernel{{.*}}() #[[KERN_ATTR:[0-9]+]]
+
+// CHECK: #[[KERN_ATTR]] = { {{.*}}"sycl-module-id"="{{.*}}module-id.cpp"{{.*}} }

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

@@ -30,6 +30,17 @@
 // CHK-FPGA-IMAGE: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocx]]" "[[OUTPUT4]]" "-sycl"
 // CHK-FPGA-LINK: {{lib|llvm-ar}}{{.*}}
 
+/// -fintelfpga -fsycl-link clang-cl specific
+// RUN:  touch %t.obj
+// RUN:  %clang_cl -### -clang:-target -clang:x86_64-pc-windows-msvc -fsycl -fintelfpga -fsycl-link %t.obj 2>&1 \
+// RUN:  | FileCheck -check-prefixes=CHK-FPGA-LINK-WIN %s
+// CHK-FPGA-LINK-WIN: clang-offload-bundler{{.*}} "-type=o" "-targets=host-x86_64-pc-windows-msvc,sycl-spir64_fpga-unknown-{{.*}}-sycldevice{{.*}}" "-inputs=[[INPUT:.+\.obj]]" "-outputs=[[OUTPUT1:.+\.obj]],[[OUTPUT2:.+\.obj]]" "-unbundle"
+// CHK-FPGA-LINK-WIN: llvm-link{{.*}} "[[OUTPUT2]]" "-o" "[[OUTPUT3:.+\.bc]]"
+// CHK-FPGA-LINK-WIN: llvm-spirv{{.*}} "-spirv-max-version=1.1" "-spirv-ext=+all" "-o" "[[OUTPUT4:.+\.spv]]" "[[OUTPUT3]]"
+// CHK-FPGA-LINK-WIN: aoc{{.*}} "-o" "[[OUTPUT5:.+\.aocr]]" "[[OUTPUT4]]" "-sycl" "-rtl"
+// CHK-FPGA-LINK-WIN: lib.exe{{.*}}
+
+
 /// Check -fintelfpga -fsycl-link with an FPGA archive
 // Create the dummy archive
 // RUN:  echo "Dummy AOCR image" > %t.aocr

sycl/CMakeLists.txt

@@ -159,9 +159,23 @@ add_custom_target( sycl-toolchain
   COMMENT "Building SYCL compiler toolchain..."
 )
 
-add_subdirectory( test )
-add_subdirectory( unittests )
-add_subdirectory( tools )
+if (NOT DEFINED LLVM_INCLUDE_TESTS)
+  set(LLVM_INCLUDE_TESTS ON)
+endif()
+
+option(SYCL_INCLUDE_TESTS
+  "Generate build targets for the SYCL unit tests."
+  ${LLVM_INCLUDE_TESTS})
+
+add_subdirectory(tools)
+
+if(SYCL_INCLUDE_TESTS)
+  if(EXISTS ${LLVM_MAIN_SRC_DIR}/utils/unittest/googletest/include/gtest/gtest.h)
+    add_subdirectory(unittests)
+    list(APPEND SYCL_TEST_DEPS SYCLUnitTests)
+  endif()
+  add_subdirectory(test)
+endif()
 
 # Package deploy support
 # Listed here are component names contributing the package

sycl/doc/extensions/USM/USM.adoc

@@ -528,9 +528,9 @@ float* a = static_cast<float*>(malloc_shared(10*sizeof(float), dev, ctxt));
 float* b = static_cast<float*>(malloc_shared(10*sizeof(float), dev, ctxt));
 float* c = static_cast<float*>(malloc_shared(10*sizeof(float), dev, ctxt));
 
-queue Q;
-auto e = Q.submit([&](handler& cgh) {
-  cgh.parallel_for<class vec_add>(range<1> {10}, [=](id<1> i) {
+auto e = q.submit([&](handler& cgh) {
+  cgh.parallel_for<class vec_add>(range<1> {10}, [=](id<1> ID) {
+    size_t i = ID[0];
     c[i] = a[i] + b[i];
   });
 });

sycl/include/CL/sycl/detail/pi.h

@@ -56,6 +56,7 @@ typedef enum {
   PI_INVALID_CONTEXT = CL_INVALID_CONTEXT,
   PI_INVALID_PLATFORM = CL_INVALID_PLATFORM,
   PI_INVALID_DEVICE = CL_INVALID_DEVICE,
+  PI_INVALID_BINARY = CL_INVALID_BINARY,
   PI_MISALIGNED_SUB_BUFFER_OFFSET = CL_MISALIGNED_SUB_BUFFER_OFFSET,
   PI_OUT_OF_HOST_MEMORY = CL_OUT_OF_HOST_MEMORY
 } _pi_result;
@@ -245,9 +246,20 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4;
 /// Target identification strings for
 /// pi_device_binary_struct.DeviceTargetSpec
 ///
+/// A device type represented by a particular target
+/// triple requires specific binary images. We need
+/// to map the image type onto the device target triple
+///
 #define PI_DEVICE_BINARY_TARGET_UNKNOWN "<unknown>"
+/// SPIR-V 32-bit image <-> 32-bit OpenCL device
 #define PI_DEVICE_BINARY_TARGET_SPIRV32 "spir"
-#define PI_DEVICE_BINARY_TARGET_SPIRV64 "spir64";
+/// SPIR-V 64-bit image <-> 64-bit OpenCL device
+#define PI_DEVICE_BINARY_TARGET_SPIRV64 "spir64"
+/// Device-specific binary images produced from SPIR-V 64-bit <->
+/// various triples for specific 64-bit OpenCL devices
+#define PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64 "spir64_x86_64"
+#define PI_DEVICE_BINARY_TARGET_SPIRV64_GEN "spir64_gen"
+#define PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA "spir64_fpga"
 
 /// This struct is a record of the device binary information. If the Kind field
 /// denotes a portable binary type (SPIRV or LLVMIR), the DeviceTargetSpec field
@@ -264,6 +276,13 @@ struct pi_device_binary_struct {
   /// format of the binary data - SPIRV, LLVMIR bitcode,...
   uint8_t Format;
   /// null-terminated string representation of the device's target architecture
+  /// which holds one of:
+  /// PI_DEVICE_BINARY_TARGET_UNKNOWN - unknown
+  /// PI_DEVICE_BINARY_TARGET_SPIRV32 - general value for 32-bit OpenCL devices
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64 - general value for 64-bit OpenCL devices
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64 - 64-bit OpenCL CPU device
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64_GEN - GEN GPU device (64-bit OpenCL)
+  /// PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA - 64-bit OpenCL FPGA device
   const char *DeviceTargetSpec;
   /// a null-terminated string; target- and compiler-specific options
   /// which are suggested to use to "build" program at runtime
@@ -405,7 +424,7 @@ pi_result piDevicePartition(
 /// and the IR characteristics.
 ///
 pi_result piextDeviceSelectBinary(
-  pi_device           device, // TODO: does this need to be context?
+  pi_device           device,
   pi_device_binary *  binaries,
   pi_uint32           num_binaries,
   pi_device_binary *  selected_binary);

sycl/include/CL/sycl/detail/scheduler/commands.hpp

@@ -66,7 +66,8 @@ class Command {
     RELEASE,
     MAP_MEM_OBJ,
     UNMAP_MEM_OBJ,
-    UPDATE_REQUIREMENT
+    UPDATE_REQUIREMENT,
+    EMPTY_TASK
   };
 
   Command(CommandType Type, QueueImplPtr Queue, bool UseExclusiveQueue = false);
@@ -123,6 +124,23 @@ class Command {
   std::atomic<bool> MEnqueued;
 };
 
+// The command does nothing during enqueue. The task can be used to implement
+// lock in the graph, or to merge several nodes into one.
+class EmptyCommand : public Command {
+public:
+  EmptyCommand(QueueImplPtr Queue, Requirement *Req)
+      : Command(CommandType::EMPTY_TASK, std::move(Queue)),
+        MStoredRequirement(*Req) {}
+
+  Requirement *getStoredRequirement() { return &MStoredRequirement; }
+
+private:
+  cl_int enqueueImp() override { return CL_SUCCESS; }
+  void printDot(std::ostream &Stream) const override;
+
+  Requirement MStoredRequirement;
+};
+
 // The command enqueues release instance of memory allocated on Host or
 // underlying framework.
 class ReleaseCommand : public Command {

sycl/include/CL/sycl/detail/usm_impl.hpp

@@ -19,8 +19,8 @@ namespace usm {
 void *alignedAlloc(size_t Alignment, size_t Bytes, const context &Ctxt,
                    const device &Dev, cl::sycl::usm::alloc Kind);
 
-void *alignedAlloc(size_t Alignment, size_t Bytes, const context &Ctxt,
-                   cl::sycl::usm::alloc Kind);
+void *alignedAllocHost(size_t Alignment, size_t Bytes, const context &Ctxt,
+                       cl::sycl::usm::alloc Kind);
 
 void free(void *Ptr, const context &Ctxt);
 

sycl/include/CL/sycl/usm/usm_allocator.hpp

@@ -109,7 +109,7 @@ class usm_allocator {
       usm::alloc AllocT = AllocKind,
       typename std::enable_if<AllocT == usm::alloc::host, int>::type = 0>
   pointer allocate(size_t Size) {
-    auto Result = reinterpret_cast<pointer>(detail::usm::alignedAlloc(
+    auto Result = reinterpret_cast<pointer>(detail::usm::alignedAllocHost(
         getAlignment(), Size * sizeof(value_type), mContext, AllocKind));
     if (!Result) {
       throw memory_allocation_error();

sycl/source/detail/pi_opencl.cpp

@@ -59,13 +59,16 @@ pi_result OCL(piDevicesGet)(pi_platform      platform,
   return cast<pi_result>(result);
 }
 
-pi_result OCL(piextDeviceSelectBinary)(
-  pi_device           device, // TODO: does this need to be context?
-  pi_device_binary *  images,
-  pi_uint32           num_images,
-  pi_device_binary *  selected_image) {
-
-  // TODO dummy implementation.
+pi_result OCL(piextDeviceSelectBinary)(pi_device device,
+                                       pi_device_binary *images,
+                                       pi_uint32 num_images,
+                                       pi_device_binary *selected_image) {
+
+  // TODO: this is a bare-bones implementation for choosing a device image
+  // that would be compatible with the targeted device. An AOT-compiled
+  // image is preferred over SPIRV for known devices (i.e. Intel devices)
+  // The implementation makes no effort to differentiate between multiple images
+  // for the given device, and simply picks the first one compatible
   // Real implementaion will use the same mechanism OpenCL ICD dispatcher
   // uses. Somthing like:
   //   PI_VALIDATE_HANDLE_RETURN_HANDLE(ctx, PI_INVALID_CONTEXT);
@@ -74,8 +77,56 @@ pi_result OCL(piextDeviceSelectBinary)(
   // where context->dispatch is set to the dispatch table provided by PI
   // plugin for platform/device the ctx was created for.
 
-  *selected_image = num_images > 0 ? images[0] : nullptr;
-  return PI_SUCCESS;
+  // Choose the binary target for the provided device
+  const char *image_target = nullptr;
+  // Get the type of the device
+  cl_device_type device_type;
+  cl_int ret_err = clGetDeviceInfo(cast<cl_device_id>(device), CL_DEVICE_TYPE,
+                                   sizeof(cl_device_type), &device_type, nullptr);
+  if (ret_err != CL_SUCCESS) {
+    *selected_image = nullptr;
+    return cast<pi_result>(ret_err);
+  }
+
+  switch (device_type) {
+  // TODO: Factor out vendor specifics into a separate source
+  // E.g. sycl/source/detail/vendor/intel/detail/pi_opencl.cpp?
+
+  // We'll attempt to find an image that was AOT-compiled
+  // from a SPIR-V image into an image specific for:
+
+  case CL_DEVICE_TYPE_CPU: // OpenCL 64-bit CPU
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64;
+    break;
+  case CL_DEVICE_TYPE_GPU: // OpenCL 64-bit GEN GPU
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64_GEN;
+    break;
+  case CL_DEVICE_TYPE_ACCELERATOR: // OpenCL 64-bit FPGA
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA;
+    break;
+  default:
+    // Otherwise, we'll attempt to find and JIT-compile
+    // a device-independent SPIR-V image
+    image_target = PI_DEVICE_BINARY_TARGET_SPIRV64;
+    break;
+  }
+
+  // Find the appropriate device image, fallback to spirv if not found
+  pi_device_binary fallback = nullptr;
+  for (size_t i = 0; i < num_images; ++i) {
+    if (strcmp(images[i]->DeviceTargetSpec, image_target) == 0) {
+      *selected_image = images[i];
+      return PI_SUCCESS;
+    }
+    if (strcmp(images[i]->DeviceTargetSpec, PI_DEVICE_BINARY_TARGET_SPIRV64) ==
+        0)
+      fallback = images[i];
+  }
+  // Points to a spirv image, if such indeed was found
+  if ((*selected_image = fallback))
+    return PI_SUCCESS;
+  // No image can be loaded for the given device
+  return PI_INVALID_BINARY;
 }
 
 pi_result OCL(piQueueCreate)(pi_context context, pi_device device,
@@ -290,7 +341,7 @@ _PI_CL(piDeviceRetain,       clRetainDevice)
 _PI_CL(piDeviceRelease,      clReleaseDevice)
 _PI_CL(piextDeviceSelectBinary, OCL(piextDeviceSelectBinary))
 _PI_CL(piextGetDeviceFunctionPointer, OCL(piextGetDeviceFunctionPointer))
-  // Context
+// Context
 _PI_CL(piContextCreate,     clCreateContext)
 _PI_CL(piContextGetInfo,    clGetContextInfo)
 _PI_CL(piContextRetain,     clRetainContext)

sycl/source/detail/program_manager/program_manager.cpp

@@ -263,6 +263,7 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
               << getRawSyclObjImpl(Context) << ")\n";
   }
 
+  const RT::PiContext &Ctx = getRawSyclObjImpl(Context)->getHandleRef();
   DeviceImage *Img = nullptr;
   bool UseKernelSpv = false;
   const std::string UseSpvEnv("SYCL_USE_KERNEL_SPV");
@@ -309,6 +310,10 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
       std::cerr << "loaded device image from " << Fname << "\n";
     }
   } else {
+    // TODO: There may be cases with cl::sycl::program class usage in source code
+    // that will result in a multi-device context. This case needs to be handled
+    // here or at the program_impl class level
+
     // Take all device images in module M and ask the native runtime under the
     // given context to choose one it prefers.
     auto ImgIt = m_DeviceImages.find(M);
@@ -318,8 +323,8 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
     }
     std::vector<DeviceImage *> *Imgs = (ImgIt->second).get();
 
-    PI_CALL(RT::piextDeviceSelectBinary(
-      0, Imgs->data(), (cl_uint)Imgs->size(), &Img));
+    PI_CALL(RT::piextDeviceSelectBinary(getFirstDevice(Ctx), Imgs->data(),
+                                        (cl_uint)Imgs->size(), &Img));
 
     if (DbgProgMgr > 0) {
       std::cerr << "available device images:\n";
@@ -400,7 +405,6 @@ RT::PiProgram ProgramManager::loadProgram(OSModuleHandle M,
   // Load the selected image
   if (!is_device_binary_type_supported(Context, Format))
     throw feature_not_supported("Online compilation is not supported in this context");
-  const RT::PiContext &Ctx = getRawSyclObjImpl(Context)->getHandleRef();
   RT::PiProgram Res = nullptr;
   Res = Format == PI_DEVICE_BINARY_TYPE_SPIRV
             ? createSpirvProgram(Ctx, Img->BinaryStart, ImgSize)

sycl/source/detail/scheduler/commands.cpp

@@ -424,9 +424,7 @@ void UpdateHostRequirementCommand::printDot(std::ostream &Stream) const {
   for (const auto &Dep : MDeps) {
     Stream << "  \"" << this << "\" -> \"" << Dep.MDepCommand << "\""
            << " [ label = \"Access mode: "
-           << accessModeToString(
-                  Dep.MAllocaCmd->getAllocationReq()->MAccessMode)
-           << "\\n"
+           << accessModeToString(Dep.MReq->MAccessMode) << "\\n"
            << "MemObj: " << Dep.MAllocaCmd->getSYCLMemObj() << " \" ]"
            << std::endl;
   }
@@ -467,6 +465,23 @@ cl_int MemCpyCommandHost::enqueueImp() {
   return CL_SUCCESS;
 }
 
+void EmptyCommand::printDot(std::ostream &Stream) const {
+  Stream << "\"" << this << "\" [style=filled, fillcolor=\"#8d8f29\", label=\"";
+
+  Stream << "ID = " << this << "\n";
+  Stream << "EMPTY NODE"
+         << "\\n";
+
+  Stream << "\"];" << std::endl;
+
+  for (const auto &Dep : MDeps) {
+    Stream << "  \"" << this << "\" -> \"" << Dep.MDepCommand << "\""
+           << " [ label = \"Access mode: "
+           << accessModeToString(Dep.MReq->MAccessMode) << "\\n"
+           << "MemObj: " << Dep.MReq->MSYCLMemObj << " \" ]" << std::endl;
+  }
+}
+
 void MemCpyCommandHost::printDot(std::ostream &Stream) const {
   Stream << "\"" << this << "\" [style=filled, fillcolor=\"#B6A2EB\", label=\"";
 

sycl/source/detail/scheduler/graph_builder.cpp

@@ -167,10 +167,8 @@ UpdateHostRequirementCommand *Scheduler::GraphBuilder::insertUpdateHostReqCmd(
     UpdateCommand->addDep(DepDesc{Dep, StoredReq, AllocaCmd});
     Dep->addUser(UpdateCommand);
   }
-  // access::mode::read_write is always used here regardless of requieremnt
-  // access mode because this node shouldn't be skipped.
-  UpdateLeafs(Deps, Record, access::mode::read_write);
-  AddNodeToLeafs(Record, UpdateCommand, access::mode::read_write);
+  UpdateLeafs(Deps, Record, Req->MAccessMode);
+  AddNodeToLeafs(Record, UpdateCommand, Req->MAccessMode);
   return UpdateCommand;
 }