Allow cases for copying device_global with different template arguments

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

Author: steffenlarsen

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++]

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

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;
 }

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; }