Merge from 'sycl' to 'sycl-web' (triSYCL#7)

sycl/include/CL/sycl/detail/accessor_impl.hpp

@@ -160,6 +160,8 @@ class LocalAccessorImplHost {
   }
 };
 
+using LocalAccessorImplPtr = std::shared_ptr<LocalAccessorImplHost>;
+
 class LocalAccessorBaseHost {
 public:
   LocalAccessorBaseHost(sycl::range<3> Size, int Dims, int ElemSize) {

sycl/include/CL/sycl/detail/platform_impl.hpp

@@ -10,7 +10,6 @@
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/force_device.hpp>
 #include <CL/sycl/detail/pi.hpp>
-#include <CL/sycl/detail/platform_info.hpp>
 #include <CL/sycl/info/info_desc.hpp>
 #include <CL/sycl/stl.hpp>
 
@@ -41,15 +40,7 @@ class platform_impl {
   ///
   /// @param ExtensionName is a string containing extension name.
   /// @return true if platform supports specified extension.
-  bool has_extension(const string_class &ExtensionName) const {
-    if (is_host())
-      return false;
-
-    string_class AllExtensionNames =
-        get_platform_info<string_class, info::platform::extensions>::get(
-            MPlatform);
-    return (AllExtensionNames.find(ExtensionName) != std::string::npos);
-  }
+  bool has_extension(const string_class &ExtensionName) const;
 
   /// Returns all SYCL devices associated with this platform.
   ///
@@ -68,14 +59,7 @@ class platform_impl {
   /// The return type depends on information being queried.
   template <info::platform param>
   typename info::param_traits<info::platform, param>::return_type
-  get_info() const {
-    if (is_host())
-      return get_platform_info_host<param>();
-
-    return get_platform_info<
-        typename info::param_traits<info::platform, param>::return_type,
-        param>::get(this->getHandleRef());
-  }
+  get_info() const;
 
   /// @return true if this SYCL platform is a host platform.
   bool is_host() const { return MHostPlatform; };

sycl/include/CL/sycl/handler.hpp

@@ -143,6 +143,7 @@ class handler {
   // we exit the method they are passed in.
   std::vector<std::vector<char>> MArgsStorage;
   std::vector<detail::AccessorImplPtr> MAccStorage;
+  std::vector<detail::LocalAccessorImplPtr> MLocalAccStorage;
   std::vector<std::shared_ptr<detail::stream_impl>> MStreamStorage;
   std::vector<std::shared_ptr<const void>> MSharedPtrStorage;
   // The list of arguments for the kernel.
@@ -221,6 +222,10 @@ class handler {
           detail::AccessorBaseHost *AccBase =
               static_cast<detail::AccessorBaseHost *>(Ptr);
           Ptr = detail::getSyclObjImpl(*AccBase).get();
+        } else if (AccTarget == access::target::local) {
+          detail::LocalAccessorBaseHost *LocalAccBase =
+              static_cast<detail::LocalAccessorBaseHost *>(Ptr);
+          Ptr = detail::getSyclObjImpl(*LocalAccBase).get();
         }
       }
       processArg(Ptr, Kind, Size, I, IndexShift, IsKernelCreatedFromSource);
@@ -292,20 +297,17 @@ class handler {
         break;
       }
       case access::target::local: {
-        detail::LocalAccessorBaseHost *LAcc =
-            static_cast<detail::LocalAccessorBaseHost *>(Ptr);
+        detail::LocalAccessorImplHost *LAcc =
+            static_cast<detail::LocalAccessorImplHost *>(Ptr);
         // Stream implementation creates local accessor with size per work item
         // in work group. Number of work items is not available during stream
         // construction, that is why size of the accessor is updated here using
         // information about number of work items in the work group.
-        if (detail::getSyclObjImpl(*LAcc)->PerWI) {
-          auto LocalAccImpl = detail::getSyclObjImpl(*LAcc);
-          LocalAccImpl->resize(MNDRDesc.LocalSize.size(),
-                               MNDRDesc.GlobalSize.size());
-        }
-        range<3> &Size = LAcc->getSize();
-        const int Dims = LAcc->getNumOfDims();
-        int SizeInBytes = LAcc->getElementSize();
+        if (LAcc->PerWI)
+          LAcc->resize(MNDRDesc.LocalSize.size(), MNDRDesc.GlobalSize.size());
+        range<3> &Size = LAcc->MSize;
+        const int Dims = LAcc->MDims;
+        int SizeInBytes = LAcc->MElemSize;
         for (int I = 0; I < Dims; ++I)
           SizeInBytes *= Size[I];
         MArgs.emplace_back(kind_std_layout, nullptr, SizeInBytes,
@@ -480,7 +482,11 @@ class handler {
                              IsPlaceholder> &&Arg) {
     detail::LocalAccessorBaseHost *LocalAccBase =
         (detail::LocalAccessorBaseHost *)&Arg;
-    MArgs.emplace_back(detail::kernel_param_kind_t::kind_accessor, LocalAccBase,
+    detail::LocalAccessorImplPtr LocalAccImpl =
+        detail::getSyclObjImpl(*LocalAccBase);
+    detail::LocalAccessorImplHost *Req = LocalAccImpl.get();
+    MLocalAccStorage.push_back(std::move(LocalAccImpl));
+    MArgs.emplace_back(detail::kernel_param_kind_t::kind_accessor, Req,
                        static_cast<int>(access::target::local), ArgIndex);
   }
 

sycl/include/CL/sycl/intel/builtins.hpp

@@ -18,7 +18,7 @@ extern int __spirv_ocl_printf(const CONSTANT_AS char *__format, ...);
 #define CONSTANT_AS
 #endif
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 namespace intel {
 namespace experimental {

sycl/source/detail/platform_impl.cpp

@@ -8,6 +8,7 @@
 
 #include <CL/sycl/detail/device_impl.hpp>
 #include <CL/sycl/detail/platform_impl.hpp>
+#include <CL/sycl/detail/platform_info.hpp>
 #include <CL/sycl/device.hpp>
 #include <detail/config.hpp>
 
@@ -227,6 +228,34 @@ platform_impl::get_devices(info::device_type DeviceType) const {
 
   return Res;
 }
+
+bool platform_impl::has_extension(const string_class &ExtensionName) const {
+  if (is_host())
+    return false;
+
+  string_class AllExtensionNames =
+      get_platform_info<string_class, info::platform::extensions>::get(
+          MPlatform);
+  return (AllExtensionNames.find(ExtensionName) != std::string::npos);
+}
+
+template <info::platform param>
+typename info::param_traits<info::platform, param>::return_type
+platform_impl::get_info() const {
+  if (is_host())
+    return get_platform_info_host<param>();
+
+  return get_platform_info<
+      typename info::param_traits<info::platform, param>::return_type,
+      param>::get(this->getHandleRef());
+}
+
+#define PARAM_TRAITS_SPEC(param_type, param, ret_type)                         \
+  template ret_type platform_impl::get_info<info::param_type::param>() const;
+
+#include <CL/sycl/info/platform_traits.def>
+#undef PARAM_TRAITS_SPEC
+
 } // namespace detail
 } // namespace sycl
 } // namespace cl

sycl/source/detail/program_manager/program_manager.cpp

@@ -197,23 +197,30 @@ static bool isDeviceBinaryTypeSupported(const context &C,
   if (Format != PI_DEVICE_BINARY_TYPE_SPIRV)
     return true;
 
+  vector_class<device> Devices = C.get_devices();
+
+  // Program type is SPIR-V, so we need a device compiler to do JIT.
+  for (const device &D : Devices) {
+    if (!D.get_info<info::device::is_compiler_available>())
+      return false;
+  }
+
   // OpenCL 2.1 and greater require clCreateProgramWithIL
   if (pi::useBackend(pi::SYCL_BE_PI_OPENCL) &&
       C.get_platform().get_info<info::platform::version>() >= "2.1")
     return true;
 
-  // Otherwise we need cl_khr_il_program extension to be present
-  // and we can call clCreateProgramWithILKHR using the extension
-  for (const device &D : C.get_devices()) {
+  for (const device &D : Devices) {
+    // We need cl_khr_il_program extension to be present
+    // and we can call clCreateProgramWithILKHR using the extension
     vector_class<string_class> Extensions =
         D.get_info<info::device::extensions>();
-    if (std::find(Extensions.begin(), Extensions.end(),
-                  string_class("cl_khr_il_program")) != Extensions.end())
-      return true;
+    if (Extensions.end() ==
+        std::find(Extensions.begin(), Extensions.end(), "cl_khr_il_program"))
+      return false;
   }
 
-  // This device binary type is not supported.
-  return false;
+  return true;
 }
 
 static const char *getFormatStr(RT::PiDeviceBinaryType Format) {
@@ -618,14 +625,6 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context,
   }
   const char *Opts = std::getenv("SYCL_PROGRAM_BUILD_OPTIONS");
 
-  for (const auto &DeviceId : Devices) {
-    if (!createSyclObjFromImpl<device>(std::make_shared<device_impl>(DeviceId))
-             .get_info<info::device::is_compiler_available>()) {
-      throw feature_not_supported(
-          "Online compilation is not supported by this device");
-    }
-  }
-
   if (!Opts)
     Opts = Options.c_str();
 

sycl/source/program.cpp

@@ -11,7 +11,7 @@
 
 #include <vector>
 
-namespace cl {
+__SYCL_INLINE namespace cl {
 namespace sycl {
 
 program::program(const context &context)

sycl/test/aot/accelerator.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: aoc
+// REQUIRES: aoc, accelerator
 
 // RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga-unknown-unknown-sycldevice %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out

sycl/test/aot/cpu.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: opencl-aot
+// REQUIRES: opencl-aot, cpu
 
 // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out

sycl/test/aot/gpu.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: ocloc
+// REQUIRES: ocloc, gpu
 
 // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen-unknown-unknown-sycldevice -Xsycl-target-backend=spir64_gen-unknown-unknown-sycldevice "-device skl" %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out

sycl/test/aot/multiple-devices.cpp

@@ -6,7 +6,7 @@
 //
 //===------------------------------------------------------------------------===//
 
-// REQUIRES: opencl-aot, ocloc, aoc
+// REQUIRES: opencl-aot, ocloc, aoc, cpu, gpu, accelerator
 
 // 1-command compilation case
 // Targeting CPU, GPU, FPGA

sycl/test/basic_tests/set_arg_interop.cpp

@@ -16,12 +16,14 @@ int main() {
 
     cl_context ClContext = Context.get();
 
-    const size_t CountSources = 2;
+    const size_t CountSources = 3;
     const char *Sources[CountSources] = {
         "kernel void foo1(global float* Array, global int* Value) { *Array = "
         "42; *Value = 1; }\n",
         "kernel void foo2(global float* Array) { int id = get_global_id(0); "
         "Array[id] = id; }\n",
+        "kernel void foo3(global float* Array, local float* LocalArray) { "
+        "(void)LocalArray; (void)Array; }\n",
     };
 
     cl_int Err;
@@ -38,11 +40,15 @@ int main() {
     cl_kernel SecondCLKernel = clCreateKernel(ClProgram, "foo2", &Err);
     assert(Err == CL_SUCCESS);
 
+    cl_kernel ThirdCLKernel = clCreateKernel(ClProgram, "foo3", &Err);
+    assert(Err == CL_SUCCESS);
+
     const size_t Count = 100;
     float Array[Count];
 
     kernel FirstKernel(FirstCLKernel, Context);
     kernel SecondKernel(SecondCLKernel, Context);
+    kernel ThirdKernel(ThirdCLKernel, Context);
     int Value;
     {
       buffer<float, 1> FirstBuffer(Array, range<1>(1));
@@ -92,9 +98,24 @@ int main() {
       }
     }
 
+    {
+      buffer<float, 1> FirstBuffer(Array, range<1>(Count));
+      Queue.submit([&](handler &CGH) {
+        auto Acc = FirstBuffer.get_access<access::mode::read_write>(CGH);
+        CGH.set_arg(0, FirstBuffer.get_access<access::mode::read_write>(CGH));
+        CGH.set_arg(
+            1, cl::sycl::accessor<float, 1, cl::sycl::access::mode::read_write,
+                                  cl::sycl::access::target::local>(
+                   cl::sycl::range<1>(Count), CGH));
+        CGH.parallel_for(range<1>{Count}, ThirdKernel);
+      });
+    }
+    Queue.wait_and_throw();
+
     clReleaseContext(ClContext);
     clReleaseKernel(FirstCLKernel);
     clReleaseKernel(SecondCLKernel);
+    clReleaseKernel(ThirdCLKernel);
     clReleaseProgram(ClProgram);
   }
   return 0;

sycl/test/device-code-split/aot-accelerator.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: aoc
+// REQUIRES: aoc, accelerator
 
 // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source -fsycl-targets=spir64_fpga-unknown-unknown-sycldevice -I %S/Inputs -o %t.out %S/split-per-source-main.cpp %S/Inputs/split-per-source-second-file.cpp
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

sycl/test/device-code-split/aot-cpu.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: ioc64
+// REQUIRES: opencl-aot, cpu
 
 // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice -I %S/Inputs -o %t.out %S/split-per-source-main.cpp %S/Inputs/split-per-source-second-file.cpp
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

sycl/test/device-code-split/aot-gpu.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: ocloc
+// REQUIRES: ocloc, gpu
 
 // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source -fsycl-targets=spir64_gen-unknown-unknown-sycldevice -Xsycl-target-backend=spir64_gen-unknown-unknown-sycldevice "-device skl" -I %S/Inputs -o %t.out %S/split-per-source-main.cpp %S/Inputs/split-per-source-second-file.cpp
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

sycl/test/lit.cfg.py

@@ -150,6 +150,7 @@ def getDeviceCount(device_type):
     print("Found available accelerator device")
     acc_run_substitute = " env SYCL_DEVICE_TYPE=ACC "
     acc_check_substitute = "| FileCheck %s"
+    config.available_features.add('accelerator')
 config.substitutions.append( ('%ACC_RUN_PLACEHOLDER',  acc_run_substitute) )
 config.substitutions.append( ('%ACC_CHECK_PLACEHOLDER',  acc_check_substitute) )