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

sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc

@@ -244,7 +244,14 @@ 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);
+
+  // Available if has_property<device_image_scope> is false and OtherT is
+  //convertible to T
+  template <typename OtherT, typename OtherProps>
+  constexpr device_global(const device_global<OtherT, OtherProps> &other) {}
+
   device_global(const device_global &&) = delete;
   device_global &operator=(const device_global &) = delete;
   device_global &operator=(const device_global &&) = delete;
@@ -324,6 +331,37 @@ 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++]
+----
+template <typename OtherT, typename OtherProps>
+constexpr device_global(const device_global<OtherT, OtherProps> &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 storage in `other`.
+
+`OtherT` must be convertible to `T` and `T` must be trivially destructible.
+
 // --- ROW BREAK ---
 a|
 [source,c++]

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

@@ -49,6 +49,12 @@ struct HasArrowOperator<T,
                         std::void_t<decltype(std::declval<T>().operator->())>>
     : std::true_type {};
 
+// Checks that T is a reference to either device_global or
+// device_global_base. This is used by the variadic ctor to allow copy ctors to
+// take preference.
+template <typename T>
+struct IsDeviceGlobalOrBaseRef : std::false_type {};
+
 // Base class for device_global.
 template <typename T, typename PropertyListT, typename = void>
 class device_global_base {
@@ -65,12 +71,34 @@ class device_global_base {
 
 public:
 #if __cpp_consteval
-  template <typename... Args>
+  // The SFINAE is to allow the copy constructors to take priority.
+  template <
+      typename... Args,
+      std::enable_if_t<
+          sizeof...(Args) != 1 ||
+              (!IsDeviceGlobalOrBaseRef<std::remove_cv_t<Args>>::value && ...),
+          int> = 0>
   consteval explicit device_global_base(Args &&...args) : init_val{args...} {}
 #else
   device_global_base() = default;
 #endif // __cpp_consteval
 
+#ifndef __SYCL_DEVICE_ONLY__
+  template <typename OtherT, typename OtherProps,
+            typename = std::enable_if_t<std::is_convertible_v<OtherT, T>>>
+  constexpr device_global_base(
+      const device_global_base<OtherT, OtherProps> &DGB)
+      : init_val{DGB.init_val} {}
+  constexpr device_global_base(const device_global_base &DGB)
+      : init_val{DGB.init_val} {}
+#else
+  template <typename OtherT, typename OtherProps,
+            typename = std::enable_if_t<std::is_convertible_v<OtherT, T>>>
+  constexpr device_global_base(const device_global_base<OtherT, OtherProps> &) {
+  }
+  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 {
@@ -102,12 +130,24 @@ class device_global_base<
 
 public:
 #if __cpp_consteval
-  template <typename... Args>
+  // The SFINAE is to allow the copy constructors to take priority.
+  template <
+      typename... Args,
+      std::enable_if_t<
+          sizeof...(Args) != 1 ||
+              (!IsDeviceGlobalOrBaseRef<std::remove_cv_t<Args>>::value && ...),
+          int> = 0>
   consteval explicit device_global_base(Args &&...args) : val{args...} {}
 #else
   device_global_base() = default;
 #endif // __cpp_consteval
 
+  template <typename OtherT, typename OtherProps,
+            typename = std::enable_if_t<std::is_convertible_v<OtherT, T>>>
+  constexpr device_global_base(const device_global_base<OtherT, OtherProps> &) =
+      delete;
+  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 {
@@ -124,6 +164,11 @@ class device_global_base<
                               const T>(this->get_ptr());
   }
 };
+
+template <typename T, typename PropertyListT>
+struct IsDeviceGlobalOrBaseRef<const device_global_base<T, PropertyListT> &>
+    : std::true_type {};
+
 } // namespace detail
 
 template <typename T, typename PropertyListT = empty_properties_t>
@@ -151,6 +196,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 +213,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;
@@ -244,6 +291,12 @@ class
   }
 };
 
+namespace detail {
+template <typename T, typename PropertyListT>
+struct IsDeviceGlobalOrBaseRef<device_global<T, PropertyListT> &>
+    : std::true_type {};
+} // namespace detail
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

sycl/test-e2e/DeviceGlobal/device_global_copy.cpp

@@ -0,0 +1,65 @@
+// 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> DGInit1{3};
+oneapiext::device_global<const int> DGCopy1{DGInit1};
+
+oneapiext::device_global<int> DGInit2{4};
+oneapiext::device_global<int> DGCopy2{DGInit2};
+
+oneapiext::device_global<float> DGInit3{5.0f};
+oneapiext::device_global<int> DGCopy3{DGInit3};
+
+oneapiext::device_global<const int, decltype(oneapiext::properties{oneapiext::device_image_scope})> DGInit4{6};
+oneapiext::device_global<const int> DGCopy4{DGInit4};
+
+oneapiext::device_global<const int> DGInit5{7};
+oneapiext::device_global<const int, decltype(oneapiext::properties{oneapiext::host_access_read})> DGCopy5{DGInit5};
+
+int main() {
+  sycl::queue Q;
+
+  int ReadVals[2] = {0, 0};
+  {
+    sycl::buffer<int, 10> ReadValsBuff{ReadVals, 2};
+
+    Q.submit([&](sycl::handler &CGH) {
+       sycl::accessor ReadValsAcc{ReadValsBuff, CGH, sycl::write_only};
+       CGH.single_task([=]() {
+         ReadValsAcc[0] = DGInit1.get();
+         ReadValsAcc[1] = DGCopy1.get();
+         ReadValsAcc[2] = DGInit2.get();
+         ReadValsAcc[3] = DGCopy2.get();
+         ReadValsAcc[4] = DGInit3.get();
+         ReadValsAcc[5] = DGCopy3.get();
+         ReadValsAcc[6] = DGInit4.get();
+         ReadValsAcc[7] = DGCopy4.get();
+         ReadValsAcc[8] = DGInit5.get();
+         ReadValsAcc[9] = DGCopy5.get();
+       });
+     }).wait_and_throw();
+  }
+
+  assert(ReadVals[0] == 3);
+  assert(ReadVals[1] == 3);
+  assert(ReadVals[2] == 4);
+  assert(ReadVals[3] == 4);
+  assert(ReadVals[4] == 5);
+  assert(ReadVals[5] == 5);
+  assert(ReadVals[6] == 6);
+  assert(ReadVals[7] == 6);
+  assert(ReadVals[8] == 7);
+  assert(ReadVals[9] == 7);
+
+  return 0;
+}

sycl/test/extensions/device_global/device_global_copy_negative.cpp

@@ -0,0 +1,30 @@
+// 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});
+
+// expected-error@sycl/ext/oneapi/device_global/device_global.hpp:* {{call to deleted constructor}}
+oneapiext::device_global<const int, device_image_properties> DGInit1{3};
+oneapiext::device_global<const int, device_image_properties> DGCopy1{DGInit1};
+
+// expected-error@sycl/ext/oneapi/device_global/device_global.hpp:* {{call to deleted constructor}}
+oneapiext::device_global<int, device_image_properties> DGInit2{3};
+oneapiext::device_global<int, device_image_properties> DGCopy2{DGInit2};
+
+// expected-error@+2 {{call to deleted constructor}}
+oneapiext::device_global<int, device_image_properties> DGInit3{3};
+oneapiext::device_global<float, device_image_properties> DGCopy3{DGInit3};
+
+// expected-error@+2 {{call to deleted constructor}}
+oneapiext::device_global<const int> DGInit4{3};
+oneapiext::device_global<const int, device_image_properties> DGCopy4{DGInit4};
+
+
+int main() { return 0; }