Merge remote-tracking branch 'upstream/sycl' into glyons-name

glyons-intel · glyons-intel · commit df201b132433 · 2020-07-16T06:56:56.000-07:00
diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp
@@ -42,15 +42,29 @@ template <typename Group> bool GroupAny(bool pred) {
 }
 
 // Broadcast with scalar local index
+// Work-group supports any integral type
+// Sub-group currently supports only uint32_t
 template <typename Group, typename T, typename IdT>
-detail::enable_if_t<std::is_integral<IdT>::value, T>
+detail::enable_if_t<is_group<Group>::value && std::is_integral<IdT>::value, T>
 GroupBroadcast(T x, IdT local_id) {
   using OCLT = detail::ConvertToOpenCLType_t<T>;
   using OCLIdT = detail::ConvertToOpenCLType_t<IdT>;
   OCLT ocl_x = detail::convertDataToType<T, OCLT>(x);
   OCLIdT ocl_id = detail::convertDataToType<IdT, OCLIdT>(local_id);
   return __spirv_GroupBroadcast(group_scope<Group>::value, ocl_x, ocl_id);
 }
+template <typename Group, typename T, typename IdT>
+detail::enable_if_t<is_sub_group<Group>::value && std::is_integral<IdT>::value,
+                    T>
+GroupBroadcast(T x, IdT local_id) {
+  using SGIdT = uint32_t;
+  SGIdT sg_local_id = static_cast<SGIdT>(local_id);
+  using OCLT = detail::ConvertToOpenCLType_t<T>;
+  using OCLIdT = detail::ConvertToOpenCLType_t<SGIdT>;
+  OCLT ocl_x = detail::convertDataToType<T, OCLT>(x);
+  OCLIdT ocl_id = detail::convertDataToType<SGIdT, OCLIdT>(sg_local_id);
+  return __spirv_GroupBroadcast(group_scope<Group>::value, ocl_x, ocl_id);
+}
 
 // Broadcast with vector local index
 template <typename Group, typename T, int Dimensions>
diff --git a/sycl/include/CL/sycl/detail/type_traits.hpp b/sycl/include/CL/sycl/detail/type_traits.hpp
@@ -17,6 +17,10 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+template <int Dimensions> class group;
+namespace intel {
+struct sub_group;
+} // namespace intel
 namespace detail {
 namespace half_impl {
 class half;
@@ -302,6 +306,20 @@ template <access::address_space AS, class DataT>
 using const_if_const_AS = DataT;
 #endif
 
+template <typename T> struct is_group : std::false_type {};
+
+template <int Dimensions>
+struct is_group<group<Dimensions>> : std::true_type {};
+
+template <typename T> struct is_sub_group : std::false_type {};
+
+template <> struct is_sub_group<intel::sub_group> : std::true_type {};
+
+template <typename T>
+struct is_generic_group
+    : std::integral_constant<bool,
+                             is_group<T>::value || is_sub_group<T>::value> {};
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -77,20 +77,6 @@ template <> inline id<3> linear_id_to_id(range<3> r, size_t linear_id) {
   return result;
 }
 
-template <typename T> struct is_group : std::false_type {};
-
-template <int Dimensions>
-struct is_group<group<Dimensions>> : std::true_type {};
-
-template <typename T> struct is_sub_group : std::false_type {};
-
-template <> struct is_sub_group<intel::sub_group> : std::true_type {};
-
-template <typename T>
-struct is_generic_group
-    : std::integral_constant<bool,
-                             is_group<T>::value || is_sub_group<T>::value> {};
-
 template <typename T, class BinaryOperation> struct identity {};
 
 template <typename T, typename V> struct identity<T, intel::plus<V>> {
diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp
@@ -678,7 +678,7 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
 
   std::vector<ze_device_memory_properties_t> ZeDeviceMemoryProperties;
   try {
-    ZeDeviceMemoryProperties.reserve(ZeAvailMemCount);
+    ZeDeviceMemoryProperties.resize(ZeAvailMemCount);
   } catch (const std::bad_alloc &) {
     return PI_OUT_OF_HOST_MEMORY;
   } catch (...) {
diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp
@@ -25,15 +25,16 @@ namespace detail {
 context_impl::context_impl(const device &Device, async_handler AsyncHandler,
                            bool UseCUDAPrimaryContext)
     : MAsyncHandler(AsyncHandler), MDevices(1, Device), MContext(nullptr),
-      MPlatform(), MPluginInterop(false), MHostContext(true),
+      MPlatform(), MHostContext(Device.is_host()),
       MUseCUDAPrimaryContext(UseCUDAPrimaryContext) {
   MKernelProgramCache.setContextPtr(this);
 }
 
 context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
-                           async_handler AsyncHandler, bool UseCUDAPrimaryContext)
+                           async_handler AsyncHandler,
+                           bool UseCUDAPrimaryContext)
     : MAsyncHandler(AsyncHandler), MDevices(Devices), MContext(nullptr),
-      MPlatform(), MPluginInterop(true), MHostContext(false),
+      MPlatform(), MHostContext(false),
       MUseCUDAPrimaryContext(UseCUDAPrimaryContext) {
   MPlatform = detail::getSyclObjImpl(MDevices[0].get_platform());
   vector_class<RT::PiDevice> DeviceIds;
@@ -48,14 +49,15 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
         static_cast<pi_context_properties>(PI_CONTEXT_PROPERTIES_CUDA_PRIMARY),
         static_cast<pi_context_properties>(UseCUDAPrimaryContext), 0};
 
-    getPlugin().call<PiApiKind::piContextCreate>(props, DeviceIds.size(), 
-	  	  DeviceIds.data(), nullptr, nullptr, &MContext);
+    getPlugin().call<PiApiKind::piContextCreate>(
+        props, DeviceIds.size(), DeviceIds.data(), nullptr, nullptr, &MContext);
 #else
     cl::sycl::detail::pi::die("CUDA support was not enabled at compilation time");
 #endif
   } else {
-    getPlugin().call<PiApiKind::piContextCreate>(nullptr, DeviceIds.size(), 
-	  	  DeviceIds.data(), nullptr, nullptr, &MContext);
+    getPlugin().call<PiApiKind::piContextCreate>(nullptr, DeviceIds.size(),
+                                                 DeviceIds.data(), nullptr,
+                                                 nullptr, &MContext);
   }
 
   MKernelProgramCache.setContextPtr(this);
@@ -64,7 +66,7 @@ context_impl::context_impl(const vector_class<cl::sycl::device> Devices,
 context_impl::context_impl(RT::PiContext PiContext, async_handler AsyncHandler,
                            const plugin &Plugin)
     : MAsyncHandler(AsyncHandler), MDevices(), MContext(PiContext), MPlatform(),
-      MPluginInterop(true), MHostContext(false) {
+      MHostContext(false) {
 
   vector_class<RT::PiDevice> DeviceIds;
   size_t DevicesNum = 0;
@@ -92,7 +94,7 @@ context_impl::context_impl(RT::PiContext PiContext, async_handler AsyncHandler,
 }
 
 cl_context context_impl::get() const {
-  if (MPluginInterop) {
+  if (!MHostContext) {
     // TODO catch an exception and put it to list of asynchronous exceptions
     getPlugin().call<PiApiKind::piContextRetain>(MContext);
     return pi::cast<cl_context>(MContext);
@@ -102,17 +104,17 @@ cl_context context_impl::get() const {
       PI_INVALID_CONTEXT);
 }
 
-bool context_impl::is_host() const { return MHostContext || !MPluginInterop; }
+bool context_impl::is_host() const { return MHostContext; }
 
 context_impl::~context_impl() {
-  if (MPluginInterop) {
-    // TODO catch an exception and put it to list of asynchronous exceptions
-    getPlugin().call<PiApiKind::piContextRelease>(MContext);
-  }
   for (auto LibProg : MCachedLibPrograms) {
     assert(LibProg.second && "Null program must not be kept in the cache");
     getPlugin().call<PiApiKind::piProgramRelease>(LibProg.second);
   }
+  if (!MHostContext) {
+    // TODO catch an exception and put it to list of asynchronous exceptions
+    getPlugin().call<PiApiKind::piContextRelease>(MContext);
+  }
 }
 
 const async_handler &context_impl::get_async_handler() const {
diff --git a/sycl/source/detail/context_impl.hpp b/sycl/source/detail/context_impl.hpp
@@ -153,7 +153,6 @@ class context_impl {
   vector_class<device> MDevices;
   RT::PiContext MContext;
   PlatformImplPtr MPlatform;
-  bool MPluginInterop;
   bool MHostContext;
   bool MUseCUDAPrimaryContext;
   std::map<DeviceLibExt, RT::PiProgram> MCachedLibPrograms;
diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp
@@ -353,10 +353,12 @@ void copyD2D(SYCLMemObjI *SYCLMemObj, RT::PiMem SrcMem, QueueImplPtr SrcQueue,
           DepEvents.size(), &DepEvents[0], &OutEvent);
     } else {
       size_t SrcRowPitch = (1 == DimSrc) ? 0 : SrcSize[0];
-      size_t SrcSlicePitch = (3 == DimSrc) ? SrcSize[0] * SrcSize[1] : 0;
+      size_t SrcSlicePitch =
+          (DimSrc > 1) ? SrcSize[0] * SrcSize[1] : SrcSize[0];
 
       size_t DstRowPitch = (1 == DimDst) ? 0 : DstSize[0];
-      size_t DstSlicePitch = (3 == DimDst) ? DstSize[0] * DstSize[1] : 0;
+      size_t DstSlicePitch =
+          (DimDst > 1) ? DstSize[0] * DstSize[1] : DstSize[0];
 
       Plugin.call<PiApiKind::piEnqueueMemBufferCopyRect>(
           Queue, SrcMem, DstMem, &SrcOffset[0], &DstOffset[0],
diff --git a/sycl/test/basic_tests/buffer/buffer_full_copy.cpp b/sycl/test/basic_tests/buffer/buffer_full_copy.cpp
@@ -6,6 +6,8 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t2.out
 // RUN: %ACC_RUN_PLACEHOLDER %t2.out
 
+// XFAIL: level0
+
 //==------------- buffer_full_copy.cpp - SYCL buffer basic test ------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -18,7 +20,7 @@
 #include <cassert>
 
 void check_copy_device_to_host(cl::sycl::queue &Queue) {
-  const int size = 6, offset = 2;
+  constexpr int size = 6, offset = 2;
 
   // Create 2d buffer
   cl::sycl::buffer<int, 2> simple_buffer(cl::sycl::range<2>(size, size));
@@ -56,11 +58,11 @@ void check_copy_device_to_host(cl::sycl::queue &Queue) {
 }
 
 void check_fill(cl::sycl::queue &Queue) {
-  const int size = 6, offset = 2;
+  constexpr int size = 6, offset = 2;
   cl::sycl::buffer<float, 1> buf_1(size);
   cl::sycl::buffer<float, 1> buf_2(size / 2);
-  std::vector<float> expected_res_1(size);
-  std::vector<float> expected_res_2(size / 2);
+  std::array<float, size> expected_res_1;
+  std::array<float, size / 2> expected_res_2;
 
   // fill with offset
   {
@@ -86,11 +88,11 @@ void check_fill(cl::sycl::queue &Queue) {
 }
 
 void check_copy_host_to_device(cl::sycl::queue &Queue) {
-  const int size = 6, offset = 2;
+  constexpr int size = 6, offset = 2;
   cl::sycl::buffer<float, 1> buf_1(size);
   cl::sycl::buffer<float, 1> buf_2(size / 2);
-  std::vector<float> expected_res_1(size);
-  std::vector<float> expected_res_2(size / 2);
+  std::array<float, size> expected_res_1;
+  std::array<float, size / 2> expected_res_2;
 
   // copy acc 2 acc with offset
   {
@@ -126,24 +128,23 @@ void check_copy_host_to_device(cl::sycl::queue &Queue) {
 
   cl::sycl::buffer<float, 2> buf_3({size, size});
   cl::sycl::buffer<float, 2> buf_4({size / 2, size / 2});
-  std::vector<float> expected_res_3(size * size);
-  std::vector<float> expected_res_4(size * size / 4);
+  std::array<std::array<float, size>, size> expected_res_3;
+  std::array<std::array<float, size / 2>, size / 2> expected_res_4;
 
   // copy acc 2 acc with offset for 2D buffers
   {
     auto acc = buf_3.get_access<cl::sycl::access::mode::write>();
     for (int i = 0; i < size; ++i) {
       for (int j = 0; j < size; ++j) {
         acc[i][j] = i * size + j + 1;
-        expected_res_3[i * size + j] = i * size + j + 1;
+        expected_res_3[i][j] = acc[i][j];
       }
     }
   }
 
   for (int i = 0; i < size / 2; ++i)
     for (int j = 0; j < size / 2; ++j)
-      expected_res_4[i * size / 2 + j] =
-          expected_res_3[(i + offset) * size + j + offset];
+      expected_res_4[i][j] = expected_res_3[i + offset][j + offset];
 
   e = Queue.submit([&](cl::sycl::handler &cgh) {
     auto a = buf_3.get_access<cl::sycl::access::mode::read>(
@@ -161,12 +162,61 @@ void check_copy_host_to_device(cl::sycl::queue &Queue) {
     // check that there was no data corruption/loss
     for (int i = 0; i < size; ++i) {
       for (int j = 0; j < size; ++j)
-        assert(expected_res_3[i * size + j] == acc_1[i][j]);
+        assert(expected_res_3[i][j] == acc_1[i][j]);
     }
 
     for (int i = 0; i < size / 2; ++i)
       for (int j = 0; j < size / 2; ++j)
-        assert(expected_res_4[i * size / 2 + j] == acc_2[i][j]);
+        assert(expected_res_4[i][j] == acc_2[i][j]);
+  }
+
+  cl::sycl::buffer<float, 3> buf_5({size, size, size});
+  cl::sycl::buffer<float, 3> buf_6({size / 2, size / 2, size / 2});
+  std::array<std::array<std::array<float, size>, size>, size> expected_res_5;
+  std::array<std::array<std::array<float, size / 2>, size / 2>, size / 2> expected_res_6;
+
+  // copy acc 2 acc with offset for 3D buffers
+  {
+    auto acc = buf_5.get_access<cl::sycl::access::mode::write>();
+    for (int i = 0; i < size; ++i) {
+      for (int j = 0; j < size; ++j) {
+        for (int k = 0; k < size; ++k) {
+          acc[i][j][k] = (i * size * size) + (j * size) + k + 1;
+          expected_res_5[i][j][k] = (i * size * size) + (j * size) + k + 1;
+        }
+      }
+    }
+  }
+
+  for (int i = 0; i < size / 2; ++i)
+    for (int j = 0; j < size / 2; ++j)
+      for (int k = 0; k < size / 2; ++k)
+        expected_res_6[i][j][k] = expected_res_5[i + offset][j + offset][k + offset];
+
+  e = Queue.submit([&](cl::sycl::handler &cgh) {
+    auto a = buf_5.get_access<cl::sycl::access::mode::read>(
+        cgh, {size / 2, size / 2, size / 2}, {offset, offset, offset});
+    auto b = buf_6.get_access<cl::sycl::access::mode::write>(
+        cgh, {size / 2, size / 2, size / 2});
+    cgh.copy(a, b);
+  });
+  e.wait();
+
+  {
+    auto acc_1 = buf_5.get_access<cl::sycl::access::mode::read>();
+    auto acc_2 = buf_6.get_access<cl::sycl::access::mode::read>();
+
+    // check that there was no data corruption/loss
+    for (int i = 0; i < size; ++i)
+      for (int j = 0; j < size; ++j)
+        for (int k = 0; k < size; ++k)
+          assert(expected_res_5[i][j][k] == acc_1[i][j][k]);
+
+    for (int i = 0; i < size / 2; ++i) {
+      for (int j = 0; j < size / 2; ++j)
+        for (int k = 0; k < size / 2; ++k)
+          assert(expected_res_6[i][j][k] == acc_2[i][j][k]);
+    }
   }
 }
 
@@ -178,6 +228,7 @@ int main() {
     check_fill(Queue);
   } catch (cl::sycl::exception &ex) {
     std::cerr << ex.what() << std::endl;
+    return 1;
   }
 
   return 0;
diff --git a/sycl/test/regression/pi_release.cpp b/sycl/test/regression/pi_release.cpp
@@ -0,0 +1,13 @@
+// REQUIRES: cpu
+// RUN: %clangxx -fsycl  -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_PI_TRACE=-1 %CPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s
+
+#include <CL/sycl.hpp>
+
+int main() {
+  sycl::queue q;
+  return 0;
+}
+
+// CHECK: piQueueRelease
+// CHECK: piContextRelease
