[SYCl][L0] Buffer for multi-device context is using device allocation now (#976)

Author: smaslov-intel

SYCL/Basic/buffer/buffer_create.cpp

@@ -15,20 +15,17 @@ int main() {
   buffer<::cl_int, 1> Buffer(Size);
   Queue.submit([&](handler &cgh) {
     accessor Accessor{Buffer, cgh, read_write};
-    if (NumOfDevices > 1)
-      // Currently the Level Zero plugin uses host allocations for multi-device
-      // contexts because such allocations are accessible by all devices.
-      std::cerr << "Multi GPU should use zeMemAllocHost\n";
-    else if (D.get_info<info::device::host_unified_memory>())
+    if (D.get_info<info::device::host_unified_memory>())
       std::cerr << "Integrated GPU should use zeMemAllocHost\n";
     else
       std::cerr << "Discrete GPU should use zeMemAllocDevice\n";
-    cgh.parallel_for<class CreateBuffer>(range<1>(Size), [=](id<1> ID) {});
+    cgh.parallel_for<class CreateBuffer>(range<1>(Size),
+                                         [=](id<1> ID) { Accessor[ID] = 0; });
   });
   Queue.wait();
 
   return 0;
 }
 
-// CHECK: {{Integrated|Multi|Discrete}} GPU should use [[API:zeMemAllocHost|zeMemAllocHost|zeMemAllocDevice]]
+// CHECK: {{Integrated|Discrete}} GPU should use [[API:zeMemAllocHost|zeMemAllocDevice]]
 // CHECK: ZE ---> [[API]](

SYCL/Basic/buffer/buffer_migrate.cpp

@@ -0,0 +1,47 @@
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=2 %GPU_RUN_PLACEHOLDER %t.out
+// RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 %GPU_RUN_PLACEHOLDER %t.out
+//
+// Test for buffer use in a context with multiple devices (all found
+// root-devices)
+//
+
+#include <CL/sycl.hpp>
+using namespace cl::sycl;
+
+int main() {
+
+  int Data = 0;
+  int Result = 0;
+  buffer<int, 1> Buffer(&Data, range<1>(1));
+
+  const auto &Devices =
+      platform(gpu_selector{}).get_devices(info::device_type::gpu);
+  std::cout << Devices.size() << " devices found" << std::endl;
+  context C(Devices);
+
+  int Index = 0;
+  for (auto D : Devices) {
+    std::cout << "Using on device " << Index << ": "
+              << D.get_info<info::device::name>() << std::endl;
+    Result |= (1 << Index);
+
+    queue Q(C, D);
+    Q.submit([&](handler &cgh) {
+      accessor Accessor{Buffer, cgh, read_write};
+      cgh.parallel_for<class MigrateBuffer>(
+          range<1>(1), [=](id<1> ID) { Accessor[ID] |= (1 << Index); });
+    });
+    Q.wait();
+    ++Index;
+  }
+
+  auto HostAcc = Buffer.get_host_access();
+  auto Passed = (HostAcc[0] == Result);
+  std::cout << "Checking result on host: " << (Passed ? "passed" : "FAILED")
+            << std::endl;
+  std::cout << HostAcc[0] << " ?= " << Result << std::endl;
+  return !Passed;
+}

SYCL/Plugin/interop-level-zero-buffer-ownership.cpp

@@ -9,6 +9,9 @@
 // 2. User-provided memory allocation is freed by DPCPP RT if
 // "transfer" ownership is specified.
 
+// NOTE: SYCL RT will see unbalanced count of alloc/free,
+// so this test will fail with ZE_DEBUG=4.
+
 // Keep ownership
 // CHECK: zeMemFree
 

SYCL/Plugin/interop-level-zero-buffer.cpp

@@ -90,11 +90,6 @@ int main() {
     Queue.wait();
 
     {
-      char *Ptr = (char *)HostBuffer1;
-      for (int i = 0; i < 10; i++) {
-        assert(Ptr[i] == 'a');
-      }
-
       auto HostAcc1 = HostBufferInterop1.get_host_access();
       for (int i = 0; i < 10; i++) {
         assert(HostAcc1[i] == 'a');