[sycl][cuda][hip] Expose const addrsp via device_global<T, decltype(properties{device_constant})> #16001
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -40,8 +40,6 @@ namespace sycl { | |
| inline namespace _V1 { | ||
| namespace ext::oneapi::experimental { | ||
|
|
||
| namespace detail { | ||
| // Type-trait for checking if a type defines `operator->`. | ||
| template <typename T, typename = void> | ||
|
|
@@ -61,8 +59,8 @@ 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, typename Cond> | ||
| struct IsDeviceGlobalOrBaseRef<device_global<T, PropertyListT, Cond> &> | ||
| : std::true_type {}; | ||
|
|
||
| // Base class for device_global. | ||
|
|
@@ -141,11 +139,10 @@ class device_global_base { | |
|
|
||
| // Specialization of device_global base class for when device_image_scope is in | ||
| // the property list. | ||
| template <typename T, typename Props> | ||
| class device_global_base< | ||
| T, Props, | ||
| std::enable_if_t<Props::template has_property<device_image_scope_key>()>> { | ||
| protected: | ||
| T val{}; | ||
| T *get_ptr() noexcept { return &val; } | ||
|
|
@@ -192,7 +189,8 @@ class device_global_base< | |
|
|
||
| } // namespace detail | ||
|
|
||
| template <typename T, typename PropertyListT = empty_properties_t, | ||
| typename Cond = void> | ||
| class | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| // FIXME: Temporary work-around. Remove when fixed. | ||
|
|
@@ -204,112 +202,139 @@ class | |
| "Property list is invalid."); | ||
| }; | ||
|
|
||
| // Common code for device_global with and without | ||
| // __sycl_detail__::device_constant attribute | ||
| // Inherit the base class' constructors | ||
| #define DEVICE_GLOBAL_COMMON() \ | ||
| using property_list_t = detail::properties_t<Props...>; \ | ||
| using base_t = detail::device_global_base<T, property_list_t>; \ | ||
| using element_type = std::remove_extent_t<T>; \ | ||
| static_assert(std::is_trivially_destructible_v<T>, \ | ||
| "Type T must be trivially destructible."); \ | ||
| static_assert(is_property_list<property_list_t>::value, \ | ||
| "Property list is invalid."); \ | ||
| 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; \ | ||
| const T &get() const noexcept { \ | ||
| __SYCL_HOST_NOT_SUPPORTED("get()") \ | ||
| return *this->get_ptr(); \ | ||
| } \ | ||
| \ | ||
| operator const T &() const noexcept { \ | ||
| __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of device_global to T") \ | ||
| return get(); \ | ||
| } \ | ||
| \ | ||
| template <class RelayT = T> \ | ||
| std::remove_reference_t< \ | ||
| decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> & \ | ||
| operator[](std::ptrdiff_t idx) noexcept { \ | ||
| __SYCL_HOST_NOT_SUPPORTED("Subscript operator") \ | ||
| return (*this->get_ptr())[idx]; \ | ||
| } \ | ||
| \ | ||
| template <class RelayT = T> \ | ||
| const std::remove_reference_t< \ | ||
| decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> & \ | ||
| operator[](std::ptrdiff_t idx) const noexcept { \ | ||
| __SYCL_HOST_NOT_SUPPORTED("Subscript operator") \ | ||
| return (*this->get_ptr())[idx]; \ | ||
| } \ | ||
| \ | ||
| template <class RelayT = T> \ | ||
| std::enable_if_t<detail::HasArrowOperator<RelayT>::value || \ | ||
| std::is_pointer_v<RelayT>, \ | ||
| RelayT> & \ | ||
| operator->() noexcept { \ | ||
| __SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global") \ | ||
| return *this->get_ptr(); \ | ||
| } \ | ||
| \ | ||
| template <class RelayT = T> \ | ||
| std::enable_if_t<detail::HasArrowOperator<RelayT>::value || \ | ||
| std::is_pointer_v<RelayT>, \ | ||
| const RelayT> & \ | ||
| operator->() const noexcept { \ | ||
| __SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global") \ | ||
| return *this->get_ptr(); \ | ||
| } \ | ||
| \ | ||
| template <typename propertyT> static constexpr bool has_property() { \ | ||
| return property_list_t::template has_property<propertyT>(); \ | ||
| } \ | ||
| \ | ||
| template <typename propertyT> static constexpr auto get_property() { \ | ||
| return property_list_t::template get_property<propertyT>(); \ | ||
| } | ||
|
|
||
| template <typename T, typename... Props> | ||
| class | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global, | ||
|
There was a problem hiding this comment. I was wondering if There was a problem hiding this comment. We could even do that with the current property. It should be as simple as doing a As for the member availability, this could be done through either conditionally picking base classes or SFINAE. There was a problem hiding this comment.
Doing address space declarations directly in source code is currently not allowed by SEMA: I get e.g. (same if field is pointer type): I looked into changing this behaviour, but I didn't think there was a simple solution. such that when the property device_constant is used we add a clang attribute to the class :
That the compiler then uses to manually set the address space to .const only for cuda/hip backends. This I think in theory should be compatible with the partial specializations of the @steffenlarsen @Naghasan Maybe there is a better solution though? Once this is done there should also probably be a few more tests added to check all valid combined functionality of the device_constant property with the device_image_scope property. There should also be a test checking that the compiler does not allow writing to a device_global with the device_constant property (apart from via sycl::hander::/queue::memcpy etc). There was a problem hiding this comment. If we need the attribute, could we maybe make it apply to the field instead of the device-global class then? I.e. that way we can use the conditional solution we previously discussed, but with the new attribute. As a thought experiment, can we think of a case where a const global variable (including const fields of non-const global variables) would not want the variables to be in the There was a problem hiding this comment.
Yeah this sounds like it might be a better solution. I'll look into doing this. Thanks for the input
I am not sure if it can be an issue for NVPTX and AMDGCN backends, but in theory you can run out of .const memory space (It is normally limited to 4kb for NVPTX), and hence a user may want to be careful about which variables to put in .const space. If it is not a real issue in these backends, it could be an issue in other as yet unsupported backends. This is why we do not wan to have
|
||
| __sycl_detail__::device_constant, | ||
| __sycl_detail__::add_ir_attributes_global_variable( | ||
| "sycl-device-global-size", | ||
| __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::name..., sizeof(T), | ||
| __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::value...)]] | ||
| #endif | ||
| device_global<T, detail::properties_t<Props...>, | ||
| std::enable_if_t<detail::properties_t< | ||
| Props...>::template has_property<device_constant_key>()>> | ||
| : public detail::device_global_base<T, detail::properties_t<Props...>> { | ||
|
|
||
| public: | ||
| #if !__cpp_consteval | ||
| static_assert(std::is_trivially_default_constructible_v<T>, | ||
| "Type T must be trivially default constructable (until C++20 " | ||
| "consteval is supported and enabled.)"); | ||
| #endif // !__cpp_consteval | ||
|
|
||
| DEVICE_GLOBAL_COMMON() | ||
| }; | ||
|
|
||
| typedef void abvk; | ||
|
|
||
| template <typename T, typename... Props> | ||
| class | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global, | ||
| __sycl_detail__::add_ir_attributes_global_variable( | ||
| "sycl-device-global-size", | ||
| __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::name..., sizeof(T), | ||
| __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::value...)]] | ||
| #endif | ||
| device_global<T, detail::properties_t<Props...>, | ||
| std::enable_if_t<!( | ||
| detail::properties_t<Props...>::template has_property< | ||
| device_constant_key>()), abvk>> | ||
| : public detail::device_global_base<T, detail::properties_t<Props...>> { | ||
| public: | ||
| #if !__cpp_consteval | ||
| static_assert(std::is_trivially_default_constructible_v<T>, | ||
| "Type T must be trivially default constructable (until C++20 " | ||
| "consteval is supported and enabled.)"); | ||
| #endif // !__cpp_consteval | ||
|
|
||
| DEVICE_GLOBAL_COMMON() | ||
|
|
||
| T &get() noexcept { | ||
| __SYCL_HOST_NOT_SUPPORTED("get()") | ||
| return *this->get_ptr(); | ||
| } | ||
| operator T &() noexcept { | ||
| __SYCL_HOST_NOT_SUPPORTED("Implicit conversion of device_global to T") | ||
| return get(); | ||
| } | ||
| device_global &operator=(const T &newValue) noexcept { | ||
| __SYCL_HOST_NOT_SUPPORTED("Assignment operator") | ||
| *this->get_ptr() = newValue; | ||
| return *this; | ||
| } | ||
| }; | ||
|
|
||
| } // namespace ext::oneapi::experimental | ||
|
|
||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This doesn't seem like it corresponds to what is implemented. It could be done in tandem by making all conditional behavior dependent on the new property dependent on the disjunction of that and
std::is_const_v<T>. That said, if that is the case we also need to document how that affects the members of thedevice_globalclass.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I haven't finalised the implementation yet (I'm just testing the draft requested changes atm), so I haven't updated the documentation which you are right is completely out of date. I described this here: #16001 (comment)
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Can you mention me with a comment like "spec ready for review" when it is ready?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Sure, no problem.