[SYCL][E2E] Enable device_global tests

sycl/test-e2e/DeviceGlobal/common.hpp

@@ -0,0 +1,14 @@
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+// Property list that contains device_image_scope if USE_DEVICE_IMAGE_SCOPE is
+// defined.
+#ifdef USE_DEVICE_IMAGE_SCOPE
+using TestProperties = decltype(properties{device_image_scope});
+#else
+using TestProperties = decltype(properties{});
+#endif

sycl/test-e2e/DeviceGlobal/device_global_arrow.cpp

@@ -1,70 +1,16 @@
-// TODO: device_global without the device_image_scope property is not currently
-//       initialized on device. Enable the following test cases when it is
-//       supported.
-// RUNx: %{build} -o %t.out
-// RUNx: %{run} %t.out
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
 //
-// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t_dev_img_scope.out
-// RUN: %{run} %t_dev_img_scope.out
+// The HIP and OpenCL GPU backends do not currently support device_global
+// backend calls.
+// UNSUPPORTED: hip || (opencl && gpu)
 //
-// CPU and accelerators are not currently guaranteed to support the required
-// extensions they are disabled until they are.
-// UNSUPPORTED: cpu, accelerator
+// Temporarily disabled for OpenCL CPU while we wait for CPU driver bump. Same
+// applies to the FPGA emulator.
+// UNSUPPORTED: opencl
 //
 // Tests operator-> on device_global.
-// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
-//       we set -fsycl-device-code-split=per_source.
 
-#include <sycl/sycl.hpp>
+#include "device_global_arrow.hpp"
 
-using namespace sycl;
-using namespace sycl::ext::oneapi::experimental;
-
-struct StructWithMember {
-  int x;
-  int getX() { return x; }
-};
-
-struct StructWithDeref {
-  StructWithMember y[1];
-  StructWithMember *operator->() { return y; }
-};
-
-#ifdef USE_DEVICE_IMAGE_SCOPE
-device_global<StructWithMember *, decltype(properties{device_image_scope})>
-    DeviceGlobalVar1;
-device_global<StructWithDeref, decltype(properties{device_image_scope})>
-    DeviceGlobalVar2;
-#else
-device_global<StructWithMember *> DeviceGlobalVar1;
-device_global<StructWithDeref> DeviceGlobalVar2;
-#endif
-
-int main() {
-  queue Q;
-
-  StructWithMember *DGMem = malloc_device<StructWithMember>(1, Q);
-
-  Q.single_task([=]() {
-     DeviceGlobalVar1 = DGMem;
-     DeviceGlobalVar1->x = 1234;
-     DeviceGlobalVar2->x = 4321;
-   }).wait();
-
-  int Out[2] = {0, 0};
-  {
-    buffer<int, 1> OutBuf{Out, 2};
-    Q.submit([&](handler &CGH) {
-      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
-      CGH.single_task([=]() {
-        OutAcc[0] = DeviceGlobalVar1->getX();
-        OutAcc[1] = DeviceGlobalVar2->getX();
-      });
-    });
-  }
-  free(DGMem, Q);
-
-  assert(Out[0] == 1234 && "First value does not match.");
-  assert(Out[1] == 4321 && "Second value does not match.");
-  return 0;
-}
+int main() { return test(); }

sycl/test-e2e/DeviceGlobal/device_global_arrow.hpp

@@ -0,0 +1,45 @@
+#pragma once
+
+#include "common.hpp"
+
+struct StructWithMember {
+  int x;
+  int getX() { return x; }
+};
+
+struct StructWithDeref {
+  StructWithMember y[1];
+  StructWithMember *operator->() { return y; }
+};
+
+device_global<StructWithMember *, TestProperties> DeviceGlobalVar1;
+device_global<StructWithDeref, TestProperties> DeviceGlobalVar2;
+
+int test() {
+  queue Q;
+
+  StructWithMember *DGMem = malloc_device<StructWithMember>(1, Q);
+
+  Q.single_task([=]() {
+     DeviceGlobalVar1 = DGMem;
+     DeviceGlobalVar1->x = 1234;
+     DeviceGlobalVar2->x = 4321;
+   }).wait();
+
+  int Out[2] = {0, 0};
+  {
+    buffer<int, 1> OutBuf{Out, 2};
+    Q.submit([&](handler &CGH) {
+      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
+      CGH.single_task([=]() {
+        OutAcc[0] = DeviceGlobalVar1->getX();
+        OutAcc[1] = DeviceGlobalVar2->getX();
+      });
+    });
+  }
+  free(DGMem, Q);
+
+  assert(Out[0] == 1234 && "First value does not match.");
+  assert(Out[1] == 4321 && "Second value does not match.");
+  return 0;
+}

sycl/test-e2e/DeviceGlobal/device_global_arrow_dis.cpp

@@ -0,0 +1,14 @@
+// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out
+// RUN: %{run} %t.out
+//
+// The HIP and OpenCL GPU backends do not currently support device_global
+// backend calls.
+// UNSUPPORTED: hip || (opencl && gpu)
+//
+// Tests operator-> on device_global with device_image_scope.
+// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
+//       we set -fsycl-device-code-split=per_source.
+
+#include "device_global_arrow.hpp"
+
+int main() { return test(); }

sycl/test-e2e/DeviceGlobal/device_global_device_only.cpp

@@ -1,46 +1,16 @@
-// TODO: device_global without the device_image_scope property is not currently
-//       initialized on device. Enable the following test cases when it is
-//       supported.
-// RUNx: %{build} -o %t.out
-// RUNx: %{run} %t.out
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
 //
-// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t_dev_img_scope.out
-// RUN: %{run} %t_dev_img_scope.out
+// The HIP and OpenCL GPU backends do not currently support device_global
+// backend calls.
+// UNSUPPORTED: hip || (opencl && gpu)
 //
-// CPU and accelerators are not currently guaranteed to support the required
-// extensions they are disabled until they are.
-// UNSUPPORTED: cpu, accelerator
+// Temporarily disabled for OpenCL CPU while we wait for CPU driver bump. Same
+// applies to the FPGA emulator.
+// UNSUPPORTED: opencl
 //
 // Tests basic device_global access through device kernels.
-// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
-//       we set -fsycl-device-code-split=per_source.
 
-#include <sycl/sycl.hpp>
+#include "device_global_device_only.hpp"
 
-using namespace sycl;
-using namespace sycl::ext::oneapi::experimental;
-
-#ifdef USE_DEVICE_IMAGE_SCOPE
-device_global<int[4], decltype(properties{device_image_scope})> DeviceGlobalVar;
-#else
-device_global<int[4]> DeviceGlobalVar;
-#endif
-
-int main() {
-  queue Q;
-
-  Q.single_task([=]() { DeviceGlobalVar.get()[0] = 42; });
-  // Make sure that the write happens before subsequent read
-  Q.wait();
-
-  int OutVal = 0;
-  {
-    buffer<int, 1> OutBuf(&OutVal, 1);
-    Q.submit([&](handler &CGH) {
-      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
-      CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; });
-    });
-  }
-  assert(OutVal == 42 && "Read value does not match.");
-  return 0;
-}
+int main() { return test(); }

sycl/test-e2e/DeviceGlobal/device_global_device_only.hpp

@@ -0,0 +1,24 @@
+#pragma once
+
+#include "common.hpp"
+
+device_global<int[4], TestProperties> DeviceGlobalVar;
+
+int test() {
+  queue Q;
+
+  Q.single_task([=]() { DeviceGlobalVar.get()[0] = 42; });
+  // Make sure that the write happens before subsequent read
+  Q.wait();
+
+  int OutVal = 0;
+  {
+    buffer<int, 1> OutBuf(&OutVal, 1);
+    Q.submit([&](handler &CGH) {
+      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
+      CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; });
+    });
+  }
+  assert(OutVal == 42 && "Read value does not match.");
+  return 0;
+}

sycl/test-e2e/DeviceGlobal/device_global_device_only_dis.cpp

@@ -0,0 +1,15 @@
+// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out
+// RUN: %{run} %t.out
+//
+// The HIP and OpenCL GPU backends do not currently support device_global
+// backend calls.
+// UNSUPPORTED: hip || (opencl && gpu)
+//
+// Tests basic device_global with device_image_scope access through device
+// kernels.
+// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
+//       we set -fsycl-device-code-split=per_source.
+
+#include "device_global_device_only.hpp"
+
+int main() { return test(); }

sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.cpp

@@ -1,50 +1,16 @@
-// TODO: device_global without the device_image_scope property is not currently
-//       initialized on device. Enable the following test cases when it is
-//       supported.
-// RUNx: %{build} -o %t.out
-// RUNx: %{run} %t.out
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
 //
-// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t_dev_img_scope.out
-// RUN: %{run} %t_dev_img_scope.out
+// The HIP and OpenCL GPU backends do not currently support device_global
+// backend calls.
+// UNSUPPORTED: hip || (opencl && gpu)
 //
-// CPU and accelerators are not currently guaranteed to support the required
-// extensions they are disabled until they are.
-// UNSUPPORTED: cpu, accelerator
+// Temporarily disabled for OpenCL CPU while we wait for CPU driver bump. Same
+// applies to the FPGA emulator.
+// UNSUPPORTED: opencl
 //
 // Tests the passthrough of operators on device_global.
-// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
-//       we set -fsycl-device-code-split=per_source.
 
-#include <sycl/sycl.hpp>
+#include "device_global_operator_passthrough.hpp"
 
-using namespace sycl;
-using namespace sycl::ext::oneapi::experimental;
-
-#ifdef USE_DEVICE_IMAGE_SCOPE
-device_global<int, decltype(properties{device_image_scope})> DeviceGlobalVar;
-#else
-device_global<int> DeviceGlobalVar;
-#endif
-
-int main() {
-  queue Q;
-
-  Q.single_task([]() {
-     DeviceGlobalVar = 2;
-     DeviceGlobalVar += 3;
-     DeviceGlobalVar = DeviceGlobalVar * DeviceGlobalVar;
-     DeviceGlobalVar = DeviceGlobalVar - 3;
-     DeviceGlobalVar = 25 - DeviceGlobalVar;
-   }).wait();
-
-  int Out = 0;
-  {
-    buffer<int, 1> OutBuf{&Out, 1};
-    Q.submit([&](handler &CGH) {
-      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
-      CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar; });
-    });
-  }
-  assert(Out == 3 && "Read value does not match.");
-  return 0;
-}
+int main() { return test(); }

sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough.hpp

@@ -0,0 +1,28 @@
+#pragma once
+
+#include "common.hpp"
+
+device_global<int, TestProperties> DeviceGlobalVar;
+
+int test() {
+  queue Q;
+
+  Q.single_task([]() {
+     DeviceGlobalVar = 2;
+     DeviceGlobalVar += 3;
+     DeviceGlobalVar = DeviceGlobalVar * DeviceGlobalVar;
+     DeviceGlobalVar = DeviceGlobalVar - 3;
+     DeviceGlobalVar = 25 - DeviceGlobalVar;
+   }).wait();
+
+  int Out = 0;
+  {
+    buffer<int, 1> OutBuf{&Out, 1};
+    Q.submit([&](handler &CGH) {
+      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
+      CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar; });
+    });
+  }
+  assert(Out == 3 && "Read value does not match.");
+  return 0;
+}

sycl/test-e2e/DeviceGlobal/device_global_operator_passthrough_dis.cpp

@@ -0,0 +1,14 @@
+// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out
+// RUN: %{run} %t.out
+//
+// The HIP and OpenCL GPU backends do not currently support device_global
+// backend calls.
+// UNSUPPORTED: hip || (opencl && gpu)
+//
+// Tests the passthrough of operators on device_global with device_image_scope.
+// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so
+//       we set -fsycl-device-code-split=per_source.
+
+#include "device_global_operator_passthrough.hpp"
+
+int main() { return test(); }