From 2df61942af38fcd4498c69eccf83efb26c8c8d6a Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Wed, 14 Aug 2024 08:02:15 -0700
Subject: [PATCH 01/16] [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 <steffen.larsen@intel.com>
---
 .../sycl_ext_oneapi_device_global.asciidoc    | 19 +++++++++-
 .../oneapi/device_global/device_global.hpp    | 17 +++++++--
 .../DeviceGlobal/device_global_copy.cpp       | 37 +++++++++++++++++++
 .../device_global_copy_negative.cpp           | 18 +++++++++
 4 files changed, 87 insertions(+), 4 deletions(-)
 create mode 100644 sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
 create mode 100644 sycl/test/extensions/device_global/device_global_copy_negative.cpp

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 804b30a1a979e..140c5588eb06e 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/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++]
diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index eb26ec5410709..31097e3888a94 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/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;
diff --git a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
new file mode 100644
index 0000000000000..78117e544a6cf
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test/extensions/device_global/device_global_copy_negative.cpp b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
new file mode 100644
index 0000000000000..8af69578a2817
--- /dev/null
+++ b/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; }

From bbf50c4ee87d765307bf6678021269590953de3f Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 15 Aug 2024 02:46:57 -0700
Subject: [PATCH 02/16] Allow cases for copying device_global with different
 template arguments

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../sycl_ext_oneapi_device_global.asciidoc    | 21 +++++++++
 .../oneapi/device_global/device_global.hpp    | 46 ++++++++++++++++++-
 .../DeviceGlobal/device_global_copy.cpp       | 38 +++++++++++++--
 .../device_global_copy_negative.cpp           | 16 ++++++-
 4 files changed, 112 insertions(+), 9 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 140c5588eb06e..7ec2ca4c732f5 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -247,6 +247,11 @@ public:
   // 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;
@@ -341,6 +346,22 @@ The storage on each device for `T` is initialized with a copy of the storage in
 
 `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++]
diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index 31097e3888a94..8cdf9185e8487 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/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,16 +71,31 @@ 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__
 
@@ -109,12 +130,22 @@ 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>
@@ -133,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>
@@ -255,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
diff --git a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
index 78117e544a6cf..bda92b189e1d4 100644
--- a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
+++ b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
@@ -11,27 +11,55 @@
 
 namespace oneapiext = sycl::ext::oneapi::experimental;
 
-oneapiext::device_global<const int> DGInit{3};
-oneapiext::device_global<const int> DGCopy{DGInit};
+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, 1> ReadValsBuff{ReadVals, 2};
+    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] = DGInit.get();
-         ReadValsAcc[1] = DGCopy.get();
+         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;
 }
diff --git a/sycl/test/extensions/device_global/device_global_copy_negative.cpp b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
index 8af69578a2817..c60a60e8f3948 100644
--- a/sycl/test/extensions/device_global/device_global_copy_negative.cpp
+++ b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
@@ -10,9 +10,21 @@ 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}}
+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; }

From c83c716550225cb6432e01ac43a4d089fd1a0ce6 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 15 Aug 2024 03:54:12 -0700
Subject: [PATCH 03/16] Fix buffer size

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/test-e2e/DeviceGlobal/device_global_copy.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
index bda92b189e1d4..9f26bc0438797 100644
--- a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
+++ b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
@@ -29,9 +29,9 @@ oneapiext::device_global<const int, decltype(oneapiext::properties{oneapiext::ho
 int main() {
   sycl::queue Q;
 
-  int ReadVals[2] = {0, 0};
+  int ReadVals[10] = {0, 0};
   {
-    sycl::buffer<int, 10> ReadValsBuff{ReadVals, 2};
+    sycl::buffer<int, 1> ReadValsBuff{ReadVals, 10};
 
     Q.submit([&](sycl::handler &CGH) {
        sycl::accessor ReadValsAcc{ReadValsBuff, CGH, sycl::write_only};

From 27b212952b54bf1f22dfab48d59a4f59c740c0f1 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 15 Aug 2024 03:59:00 -0700
Subject: [PATCH 04/16] Fix formatting

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/test-e2e/DeviceGlobal/device_global_copy.cpp | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

diff --git a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
index 9f26bc0438797..a9b0fc95638cc 100644
--- a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
+++ b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
@@ -20,11 +20,15 @@ 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, 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};
+oneapiext::device_global<const int, decltype(oneapiext::properties{
+                                        oneapiext::host_access_read})>
+    DGCopy5{DGInit5};
 
 int main() {
   sycl::queue Q;

From fdaf3f63200da80dc604f0fe40470e5d248bffe5 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 15 Aug 2024 04:07:54 -0700
Subject: [PATCH 05/16] More formatting fixes

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/include/sycl/ext/oneapi/device_global/device_global.hpp   | 3 +--
 .../extensions/device_global/device_global_copy_negative.cpp   | 1 -
 2 files changed, 1 insertion(+), 3 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index 8cdf9185e8487..84e9c0ee174cd 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
@@ -52,8 +52,7 @@ struct HasArrowOperator<T,
 // 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 {};
+template <typename T> struct IsDeviceGlobalOrBaseRef : std::false_type {};
 
 // Base class for device_global.
 template <typename T, typename PropertyListT, typename = void>
diff --git a/sycl/test/extensions/device_global/device_global_copy_negative.cpp b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
index c60a60e8f3948..d097162544dbc 100644
--- a/sycl/test/extensions/device_global/device_global_copy_negative.cpp
+++ b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
@@ -26,5 +26,4 @@ oneapiext::device_global<float, device_image_properties> DGCopy3{DGInit3};
 oneapiext::device_global<const int> DGInit4{3};
 oneapiext::device_global<const int, device_image_properties> DGCopy4{DGInit4};
 
-
 int main() { return 0; }

From d35e68e6ceb5baeea3df01ed4387524a24583905 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 15 Aug 2024 05:38:57 -0700
Subject: [PATCH 06/16] Fix extraction of values from device_global with
 different template args

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../ext/oneapi/device_global/device_global.hpp  | 17 ++++++++++++++++-
 1 file changed, 16 insertions(+), 1 deletion(-)

diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index 84e9c0ee174cd..b484af8acf6ba 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
@@ -68,6 +68,19 @@ class device_global_base {
   pointer_t get_ptr() noexcept { return usmptr; }
   pointer_t get_ptr() const noexcept { return usmptr; }
 
+  template <typename, typename, typename> friend class device_global_base;
+
+#ifndef __SYCL_DEVICE_ONLY__
+  template <typename OtherT, typename OtherProps>
+  static constexpr const T &
+  ExtractInitialVal(const device_global_base<OtherT, OtherProps> &Other) {
+    if constexpr (OtherProps::template has_property<device_image_scope_key>())
+      return Other.val;
+    else
+      return Other.init_val;
+  }
+#endif // __SYCL_DEVICE_ONLY__
+
 public:
 #if __cpp_consteval
   // The SFINAE is to allow the copy constructors to take priority.
@@ -87,7 +100,7 @@ class device_global_base {
             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} {}
+      : init_val{ExtractInitialVal(DGB)} {}
   constexpr device_global_base(const device_global_base &DGB)
       : init_val{DGB.init_val} {}
 #else
@@ -127,6 +140,8 @@ class device_global_base<
   T *get_ptr() noexcept { return &val; }
   const T *get_ptr() const noexcept { return &val; }
 
+  template <typename, typename, typename> friend class device_global_base;
+
 public:
 #if __cpp_consteval
   // The SFINAE is to allow the copy constructors to take priority.

From bf0a40d7818b071a494a81b622248bf280eaa11b Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Mon, 28 Oct 2024 08:47:35 -0700
Subject: [PATCH 07/16] Fix wording and requirements

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../sycl_ext_oneapi_device_global.asciidoc    | 21 ++++++++++---------
 1 file changed, 11 insertions(+), 10 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 7ec2ca4c732f5..8b1fd6896eb20 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -232,6 +232,8 @@ public:
 
   // device_global initializes underlying T with the args argument
 #if __cpp_consteval
+  // Available only if sizeof...(Args) > 1 or the one argument in args is not a
+  // device_global.
   template <typename... Args>
   consteval explicit device_global(Args&&... args);
 #else
@@ -244,11 +246,11 @@ public:
   device_global() = default;
 #endif // __cpp_consteval
 
-  // Available if has_property<device_image_scope> is false
+  // Available if PropertyListT::has_property<device_image_scope_key>() is false.
   constexpr device_global(const device_global &other);
 
-  // Available if has_property<device_image_scope> is false and OtherT is
-  //convertible to T
+  // Available if PropertyListT::has_property<device_image_scope_key>() is false
+  // and OtherT is convertible to T.
   template <typename OtherT, typename OtherProps>
   constexpr device_global(const device_global<OtherT, OtherProps> &other) {}
 
@@ -325,6 +327,8 @@ template <typename... Args>
 consteval explicit device_global(Args&&... args);
 ----
 |
+Available only if sizeof...(Args) != 1 or the one argument in args is not a device_global.
+
 Constructs a `device_global` object, and implicit storage for `T` in the global address space on each device that may access it.
 
 The object of type `T` is initialized from the `args` parameter pack using list initialization as defined in the {cpp} specification.
@@ -338,14 +342,12 @@ a|
 constexpr device_global(const device_global &other);
 ----
 |
-Available if `has_property<device_image_scope> == false`.
+Available if `PropertyListT::has_property<device_image_scope_key>() == 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++]
@@ -354,13 +356,12 @@ template <typename OtherT, typename OtherProps>
 constexpr device_global(const device_global<OtherT, OtherProps> &other) {}
 ----
 |
-Available if `has_property<device_image_scope> == false`.
+Available if `PropertyListT::has_property<device_image_scope_key>() == false` and
+`std::is_convertible_v<OtherT, T> == true`;
 
 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.
+The storage on each device for `T` is initialized with a copy of the storage in `other`.
 
 // --- ROW BREAK ---
 a|

From d1c4ac2ee5f8532a9c4150336f1474d652352238 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Mon, 28 Oct 2024 09:01:04 -0700
Subject: [PATCH 08/16] Remove redundant ctors

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../sycl_ext_oneapi_device_global.asciidoc      | 17 ++---------------
 .../ext/oneapi/device_global/device_global.hpp  |  6 ------
 2 files changed, 2 insertions(+), 21 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 8b1fd6896eb20..1e10cb6d81c7d 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -252,7 +252,7 @@ public:
   // Available if PropertyListT::has_property<device_image_scope_key>() is false
   // and OtherT is convertible to T.
   template <typename OtherT, typename OtherProps>
-  constexpr device_global(const device_global<OtherT, OtherProps> &other) {}
+  constexpr device_global(const device_global<OtherT, OtherProps> &other);
 
   device_global(const device_global &&) = delete;
   device_global &operator=(const device_global &) = delete;
@@ -335,25 +335,12 @@ 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 `PropertyListT::has_property<device_image_scope_key>() == 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`.
-
 // --- ROW BREAK ---
 a|
 [source,c++]
 ----
 template <typename OtherT, typename OtherProps>
-constexpr device_global(const device_global<OtherT, OtherProps> &other) {}
+constexpr device_global(const device_global<OtherT, OtherProps> &other);
 ----
 |
 Available if `PropertyListT::has_property<device_image_scope_key>() == false` and
diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index b484af8acf6ba..2d0d535131749 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
@@ -101,14 +101,11 @@ class device_global_base {
   constexpr device_global_base(
       const device_global_base<OtherT, OtherProps> &DGB)
       : init_val{ExtractInitialVal(DGB)} {}
-  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>
@@ -229,9 +226,6 @@ class
   // Inherit the base class' constructors
   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 &operator=(const device_global &) = delete;
   device_global &operator=(const device_global &&) = delete;

From 09c7d4b6e3301cc8b1de293c2b4cda2c8ffd6e13 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Mon, 28 Oct 2024 10:47:10 -0700
Subject: [PATCH 09/16] Fix expected errors

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../extensions/device_global/device_global_copy_negative.cpp  | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/test/extensions/device_global/device_global_copy_negative.cpp b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
index d097162544dbc..c03357c7ec483 100644
--- a/sycl/test/extensions/device_global/device_global_copy_negative.cpp
+++ b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
@@ -10,11 +10,11 @@ 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}}
+// expected-error@+2 {{call to implicitly-deleted copy constructor of}}
 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}}
+// expected-error@+2 {{call to implicitly-deleted copy constructor of}}
 oneapiext::device_global<int, device_image_properties> DGInit2{3};
 oneapiext::device_global<int, device_image_properties> DGCopy2{DGInit2};
 

From 376185b9ce827af69e32bd0ca5c01b1fe76b5f7d Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 29 Oct 2024 02:53:15 -0700
Subject: [PATCH 10/16] Add unsupported tracker to new test

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/test-e2e/DeviceGlobal/device_global_copy.cpp | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
index a9b0fc95638cc..2d9ea347179ce 100644
--- a/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
+++ b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp
@@ -1,9 +1,8 @@
 // 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
+// UNSUPPORTED-TRACKER: GSD-4287
 //
 // Tests the copy ctor on device_global without device_image_scope.
 

From eca9cad058faacf5f8cfd4ab23518e3b82070298 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 29 Oct 2024 03:43:12 -0700
Subject: [PATCH 11/16] Make post-modification ctor undefined behavior

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../experimental/sycl_ext_oneapi_device_global.asciidoc         | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 1e10cb6d81c7d..61e075aa9572f 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -348,7 +348,7 @@ Available if `PropertyListT::has_property<device_image_scope_key>() == false` an
 
 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`.
+The storage on each device for `T` is initialized with a copy of the initial value of `other`. The behavior is undefined if `other` has been written to prior to a call this constructor.
 
 // --- ROW BREAK ---
 a|

From a9b3acd9e6e740749482bae3fc7333d6fc2df2ac Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 29 Oct 2024 04:42:31 -0700
Subject: [PATCH 12/16] Readd copy ctor

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../sycl_ext_oneapi_device_global.asciidoc          | 13 +++++++++++++
 .../sycl/ext/oneapi/device_global/device_global.hpp |  3 +++
 .../device_global/device_global_copy_negative.cpp   |  4 ++--
 3 files changed, 18 insertions(+), 2 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 61e075aa9572f..fb9f5b4c17b7e 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -335,6 +335,19 @@ 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 `PropertyListT::has_property<device_image_scope_key>() == 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`.
+
 // --- ROW BREAK ---
 a|
 [source,c++]
diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index 2d0d535131749..db08c33002738 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
@@ -226,6 +226,9 @@ class
   // Inherit the base class' constructors
   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 &operator=(const device_global &) = delete;
   device_global &operator=(const device_global &&) = delete;
diff --git a/sycl/test/extensions/device_global/device_global_copy_negative.cpp b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
index c03357c7ec483..d097162544dbc 100644
--- a/sycl/test/extensions/device_global/device_global_copy_negative.cpp
+++ b/sycl/test/extensions/device_global/device_global_copy_negative.cpp
@@ -10,11 +10,11 @@ namespace oneapiext = sycl::ext::oneapi::experimental;
 using device_image_properties =
     decltype(oneapiext::properties{oneapiext::device_image_scope});
 
-// expected-error@+2 {{call to implicitly-deleted copy constructor of}}
+// 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@+2 {{call to implicitly-deleted copy constructor of}}
+// 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};
 

From 10d207e2526d6afa1ae4bbe816172c388df0793b Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 29 Oct 2024 07:29:44 -0700
Subject: [PATCH 13/16] Add UB to readded ctor

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../experimental/sycl_ext_oneapi_device_global.asciidoc         | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index fb9f5b4c17b7e..093b0edb2dd1d 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -346,7 +346,7 @@ Available if `PropertyListT::has_property<device_image_scope_key>() == 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`.
+The storage on each device for `T` is initialized with a copy of the initial value of `other`. The behavior is undefined if `other` has been written to prior to a call this constructor.
 
 // --- ROW BREAK ---
 a|

From a531a180bb5aa616cbb7755024085791734d9fe3 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Wed, 30 Oct 2024 00:32:52 -0700
Subject: [PATCH 14/16] Add back base ctors to avoid llvm-spirv issues

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/include/sycl/ext/oneapi/device_global/device_global.hpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index db08c33002738..b484af8acf6ba 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
@@ -101,11 +101,14 @@ class device_global_base {
   constexpr device_global_base(
       const device_global_base<OtherT, OtherProps> &DGB)
       : init_val{ExtractInitialVal(DGB)} {}
+  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>

From c66d01dae9c3ed200c3d8a06ada5a144501607d7 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Wed, 30 Oct 2024 00:58:11 -0700
Subject: [PATCH 15/16] Forward declare and specialize trait in one place

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 .../oneapi/device_global/device_global.hpp    | 21 ++++++++++---------
 1 file changed, 11 insertions(+), 10 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
index b484af8acf6ba..0e74036f4012c 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
@@ -40,6 +40,8 @@ namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi::experimental {
 
+template <typename T, typename PropertyListT> class device_global;
+
 namespace detail {
 // Type-trait for checking if a type defines `operator->`.
 template <typename T, typename = void>
@@ -49,10 +51,19 @@ struct HasArrowOperator<T,
                         std::void_t<decltype(std::declval<T>().operator->())>>
     : std::true_type {};
 
+template <typename T, typename PropertyListT, typename>
+class device_global_base;
+
 // 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 {};
+template <typename T, typename PropertyListT>
+struct IsDeviceGlobalOrBaseRef<device_global_base<T, PropertyListT, void> &>
+    : std::true_type {};
+template <typename T, typename PropertyListT>
+struct IsDeviceGlobalOrBaseRef<device_global<T, PropertyListT> &>
+    : std::true_type {};
 
 // Base class for device_global.
 template <typename T, typename PropertyListT, typename = void>
@@ -179,10 +190,6 @@ class device_global_base<
   }
 };
 
-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>
@@ -305,12 +312,6 @@ 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

From 67df6e8ad0390aa95f99ab70bd1d75a4819c757e Mon Sep 17 00:00:00 2001
From: Steffen Larsen <steffen.larsen@intel.com>
Date: Wed, 30 Oct 2024 10:54:31 +0100
Subject: [PATCH 16/16] Apply suggestions from code review

Co-authored-by: John Pennycook <john.pennycook@intel.com>
---
 .../experimental/sycl_ext_oneapi_device_global.asciidoc       | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
index 093b0edb2dd1d..fd7fdb0fe65fb 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc
@@ -346,7 +346,7 @@ Available if `PropertyListT::has_property<device_image_scope_key>() == 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 initial value of `other`. The behavior is undefined if `other` has been written to prior to a call this constructor.
+The storage on each device for `T` is initialized with a copy of the initial value of `other`. The behavior is undefined if `other` has been written to prior to a call to this constructor.
 
 // --- ROW BREAK ---
 a|
@@ -361,7 +361,7 @@ Available if `PropertyListT::has_property<device_image_scope_key>() == false` an
 
 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 initial value of `other`. The behavior is undefined if `other` has been written to prior to a call this constructor.
+The storage on each device for `T` is initialized with a copy of the initial value of `other`. The behavior is undefined if `other` has been written to prior to a call to this constructor.
 
 // --- ROW BREAK ---
 a|