[SYCL][Docs] Allow copy-construction of device_global

This commit makes it possible to copy-construct device_global variables
if they do not have the device_image_scope property. The restriction on
device_image_scope is due to static construction not being allowed in
device code, which they would require, while other device_globals have
USM storage which will be initialized by the host code, so the
constructor on the device is a simple zero-initialization.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc

@@ -244,7 +244,9 @@ public:
   device_global() = default;
 #endif // __cpp_consteval
 
-  device_global(const device_global &) = delete;
+  // Available if has_property<device_image_scope> is false
+  constexpr device_global(const device_global &other);
+
   device_global(const device_global &&) = delete;
   device_global &operator=(const device_global &) = delete;
   device_global &operator=(const device_global &&) = delete;
@@ -324,6 +326,21 @@ The object of type `T` is initialized from the `args` parameter pack using list
 
 `T` must be trivially destructible.
 
+// --- ROW BREAK ---
+a|
+[source,c++]
+----
+constexpr device_global(const device_global &other);
+----
+|
+Available if `has_property<device_image_scope> == false`.
+
+Constructs a `device_global` object, and implicit storage for `T` in the global address space on each device that may access it.
+
+The storage on each device for `T` is initialized with a copy of the storage in `other`.
+
+`T` must be copy constructible and trivially destructible.
+
 // --- ROW BREAK ---
 a|
 [source,c++]

sycl/include/sycl/ext/oneapi/device_global/device_global.hpp

@@ -71,6 +71,13 @@ class device_global_base {
   device_global_base() = default;
 #endif // __cpp_consteval
 
+#ifndef __SYCL_DEVICE_ONLY__
+  constexpr device_global_base(const device_global_base &DGB)
+      : init_val{DGB.init_val} {}
+#else
+  constexpr device_global_base(const device_global_base &) {}
+#endif // __SYCL_DEVICE_ONLY__
+
   template <access::decorated IsDecorated>
   multi_ptr<T, access::address_space::global_space, IsDecorated>
   get_multi_ptr() noexcept {
@@ -108,6 +115,8 @@ class device_global_base<
   device_global_base() = default;
 #endif // __cpp_consteval
 
+  constexpr device_global_base(const device_global_base &) = delete;
+
   template <access::decorated IsDecorated>
   multi_ptr<T, access::address_space::global_space, IsDecorated>
   get_multi_ptr() noexcept {
@@ -151,6 +160,7 @@ class
     : public detail::device_global_base<T, detail::properties_t<Props...>> {
 
   using property_list_t = detail::properties_t<Props...>;
+  using base_t = detail::device_global_base<T, property_list_t>;
 
 public:
   using element_type = std::remove_extent_t<T>;
@@ -167,10 +177,11 @@ class
                 "Property list is invalid.");
 
   // Inherit the base class' constructors
-  using detail::device_global_base<
-      T, detail::properties_t<Props...>>::device_global_base;
+  using detail::device_global_base<T, property_list_t>::device_global_base;
+
+  constexpr device_global(const device_global &DG)
+      : base_t(static_cast<const base_t &>(DG)) {}
 
-  device_global(const device_global &) = delete;
   device_global(const device_global &&) = delete;
   device_global &operator=(const device_global &) = delete;
   device_global &operator=(const device_global &&) = delete;

sycl/test-e2e/DeviceGlobal/device_global_copy.cpp

@@ -0,0 +1,37 @@
+// RUN: %{build} -std=c++23 -o %t.out
+// RUN: %{run} %t.out
+//
+// The OpenCL GPU backends do not currently support device_global backend
+// calls.
+// UNSUPPORTED: opencl && gpu
+//
+// Tests the copy ctor on device_global without device_image_scope.
+
+#include <sycl/detail/core.hpp>
+
+namespace oneapiext = sycl::ext::oneapi::experimental;
+
+oneapiext::device_global<const int> DGInit{3};
+oneapiext::device_global<const int> DGCopy{DGInit};
+
+int main() {
+  sycl::queue Q;
+
+  int ReadVals[2] = {0, 0};
+  {
+    sycl::buffer<int, 1> ReadValsBuff{ReadVals, 2};
+
+    Q.submit([&](sycl::handler &CGH) {
+       sycl::accessor ReadValsAcc{ReadValsBuff, CGH, sycl::write_only};
+       CGH.single_task([=]() {
+         ReadValsAcc[0] = DGInit.get();
+         ReadValsAcc[1] = DGCopy.get();
+       });
+     }).wait_and_throw();
+  }
+
+  assert(ReadVals[0] == 3);
+  assert(ReadVals[1] == 3);
+
+  return 0;
+}

sycl/test/extensions/device_global/device_global_copy_negative.cpp

@@ -0,0 +1,18 @@
+// RUN: %clangxx -std=c++23 -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+//
+// Tests that the copy ctor on device_global with device_image_scope is
+// unavailable.
+
+#include <sycl/sycl.hpp>
+
+namespace oneapiext = sycl::ext::oneapi::experimental;
+
+using device_image_properties =
+    decltype(oneapiext::properties{oneapiext::device_image_scope});
+
+oneapiext::device_global<const int, device_image_properties> DGInit{3};
+oneapiext::device_global<const int, device_image_properties> DGCopy{DGInit};
+
+// expected-error@sycl/ext/oneapi/device_global/device_global.hpp:* {{call to deleted constructor}}
+
+int main() { return 0; }