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..fd7fdb0fe65fb 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 consteval explicit device_global(Args&&... args); #else @@ -244,7 +246,14 @@ public: device_global() = default; #endif // __cpp_consteval - device_global(const device_global &) = delete; + // Available if PropertyListT::has_property() is false. + constexpr device_global(const device_global &other); + + // Available if PropertyListT::has_property() is false + // and OtherT is convertible to T. + template + 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; @@ -318,12 +327,42 @@ template 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. `T` must be trivially destructible. +// --- ROW BREAK --- +a| +[source,c++] +---- +constexpr device_global(const device_global &other); +---- +| +Available if `PropertyListT::has_property() == 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 to this constructor. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +constexpr device_global(const device_global &other); +---- +| +Available if `PropertyListT::has_property() == false` and +`std::is_convertible_v == 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 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| [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..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 class device_global; + namespace detail { // Type-trait for checking if a type defines `operator->`. template @@ -49,6 +51,20 @@ struct HasArrowOperator().operator->())>> : std::true_type {}; +template +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 struct IsDeviceGlobalOrBaseRef : std::false_type {}; +template +struct IsDeviceGlobalOrBaseRef &> + : std::true_type {}; +template +struct IsDeviceGlobalOrBaseRef &> + : std::true_type {}; + // Base class for device_global. template class device_global_base { @@ -63,14 +79,49 @@ class device_global_base { pointer_t get_ptr() noexcept { return usmptr; } pointer_t get_ptr() const noexcept { return usmptr; } + template friend class device_global_base; + +#ifndef __SYCL_DEVICE_ONLY__ + template + static constexpr const T & + ExtractInitialVal(const device_global_base &Other) { + if constexpr (OtherProps::template has_property()) + return Other.val; + else + return Other.init_val; + } +#endif // __SYCL_DEVICE_ONLY__ + public: #if __cpp_consteval - template + // The SFINAE is to allow the copy constructors to take priority. + template < + typename... Args, + std::enable_if_t< + sizeof...(Args) != 1 || + (!IsDeviceGlobalOrBaseRef>::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 >> + constexpr device_global_base( + const device_global_base &DGB) + : init_val{ExtractInitialVal(DGB)} {} + constexpr device_global_base(const device_global_base &DGB) + : init_val{DGB.init_val} {} +#else + template >> + constexpr device_global_base(const device_global_base &) { + } + constexpr device_global_base(const device_global_base &) {} +#endif // __SYCL_DEVICE_ONLY__ + template multi_ptr get_multi_ptr() noexcept { @@ -100,14 +151,28 @@ class device_global_base< T *get_ptr() noexcept { return &val; } const T *get_ptr() const noexcept { return &val; } + template friend class device_global_base; + public: #if __cpp_consteval - template + // The SFINAE is to allow the copy constructors to take priority. + template < + typename... Args, + std::enable_if_t< + sizeof...(Args) != 1 || + (!IsDeviceGlobalOrBaseRef>::value && ...), + int> = 0> consteval explicit device_global_base(Args &&...args) : val{args...} {} #else device_global_base() = default; #endif // __cpp_consteval + template >> + constexpr device_global_base(const device_global_base &) = + delete; + constexpr device_global_base(const device_global_base &) = delete; + template multi_ptr get_multi_ptr() noexcept { @@ -124,6 +189,7 @@ class device_global_base< const T>(this->get_ptr()); } }; + } // namespace detail template @@ -151,6 +217,7 @@ class : public detail::device_global_base> { using property_list_t = detail::properties_t; + using base_t = detail::device_global_base; public: using element_type = std::remove_extent_t; @@ -167,10 +234,11 @@ class "Property list is invalid."); // Inherit the base class' constructors - using detail::device_global_base< - T, detail::properties_t>::device_global_base; + using detail::device_global_base::device_global_base; + + constexpr device_global(const device_global &DG) + : base_t(static_cast(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..2d9ea347179ce --- /dev/null +++ b/sycl/test-e2e/DeviceGlobal/device_global_copy.cpp @@ -0,0 +1,68 @@ +// RUN: %{build} -std=c++23 -o %t.out +// RUN: %{run} %t.out +// +// UNSUPPORTED: opencl && gpu +// UNSUPPORTED-TRACKER: GSD-4287 +// +// Tests the copy ctor on device_global without device_image_scope. + +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +oneapiext::device_global DGInit1{3}; +oneapiext::device_global DGCopy1{DGInit1}; + +oneapiext::device_global DGInit2{4}; +oneapiext::device_global DGCopy2{DGInit2}; + +oneapiext::device_global DGInit3{5.0f}; +oneapiext::device_global DGCopy3{DGInit3}; + +oneapiext::device_global + DGInit4{6}; +oneapiext::device_global DGCopy4{DGInit4}; + +oneapiext::device_global DGInit5{7}; +oneapiext::device_global + DGCopy5{DGInit5}; + +int main() { + sycl::queue Q; + + int ReadVals[10] = {0, 0}; + { + sycl::buffer ReadValsBuff{ReadVals, 10}; + + 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; +} 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..d097162544dbc --- /dev/null +++ b/sycl/test/extensions/device_global/device_global_copy_negative.cpp @@ -0,0 +1,29 @@ +// 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 + +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 DGInit1{3}; +oneapiext::device_global DGCopy1{DGInit1}; + +// expected-error@sycl/ext/oneapi/device_global/device_global.hpp:* {{call to deleted constructor}} +oneapiext::device_global DGInit2{3}; +oneapiext::device_global DGCopy2{DGInit2}; + +// expected-error@+2 {{call to deleted constructor}} +oneapiext::device_global DGInit3{3}; +oneapiext::device_global DGCopy3{DGInit3}; + +// expected-error@+2 {{call to deleted constructor}} +oneapiext::device_global DGInit4{3}; +oneapiext::device_global DGCopy4{DGInit4}; + +int main() { return 0; }