Skip to content

[SYCL][Doc] Update address_cast specification #12689

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Feb 28, 2024
Merged

[SYCL][Doc] Update address_cast specification #12689

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
dea2d4e
[SYCL][Doc] Default address_cast to no decoration
Pennycook Feb 12, 2024
9410347
[SYCL][Doc] Add overloads for generic multi_ptr
Pennycook Feb 12, 2024
6e45f98
[SYCL][Doc] Add address_space shorthands
Pennycook Feb 12, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
73 changes: 60 additions & 13 deletions sycl/doc/extensions/proposed/sycl_ext_oneapi_address_cast.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -93,31 +93,64 @@ implementation supports.
----
namespace sycl::ext::oneapi::experimental {

template <access::address_space Space, access::decorated DecorateAddress,
// Shorthands for address space names
constexpr inline address_space global_space = sycl::access::address_space::global_space;
constexpr inline address_space local_space = sycl::access::address_space::local_space;
constexpr inline address_space private_space = sycl::access::address_space::private_space;
constexpr inline address_space generic_space = sycl::access::address_space::generic_space;

template <access::address_space Space,
typename ElementType>
multi_ptr<ElementType, Space, DecorateAddress>
multi_ptr<ElementType, Space, access::decorated::no>
static_address_cast(ElementType* ptr);
Comment on lines +102 to 105
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

out of curiosity, why removing the possibility to select the yes/no decoration ?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The logic was that there shouldn't be two ways to do the same thing, because that's harder to teach.

Before this change, *_address_cast functions could be used to cast between address spaces and to introduce decoration. Now, they can only be used to cast between address spaces, which is aligned with their name.

multi_ptr already provides interfaces to move between decorated and undecorated pointers (via implicit casts), so we don't need to handle it here as well.


template <access::address_space Space, access::decorated DecorateAddress,
typename ElementType>
template <access::address_space Space,
typename ElementType,
access::decorated DecorateAddress>
multi_ptr<ElementType, Space, DecorateAddress>
static_address_cast(multi_ptr<ElementType, generic_space, DecorateAddress> ptr);

template <access::address_space Space,
typename ElementType>
multi_ptr<ElementType, Space, access::decorated::no>
dynamic_address_cast(ElementType* ptr);

template <access::address_space Space,
typename ElementType,
access::decorated DecorateAddress>
multi_ptr<ElementType, Space, DecorateAddress>
dynamic_address_cast(multi_ptr<ElementType, generic_space, DecorateAddress> ptr);

} // namespace sycl::ext::oneapi::experimental
----


[source,c++]
----
template <access::address_space Space, access::decorated DecorateAddress,
template <access::address_space Space,
typename ElementType>
multi_ptr<ElementType, Space, DecorateAddress>
multi_ptr<ElementType, Space, access::decorated::no>
static_address_cast(ElementType* ptr);
----
_Preconditions_: `ptr` points to an object allocated in the address space
designated by `Space`.

_Returns_: A `multi_ptr` with the specified address space and decoration that
points to the same object as `ptr`.
_Returns_: A `multi_ptr` with the specified address space that points to the
same object as `ptr`.

[source,c++]
----
template <access::address_space Space,
typename ElementType,
access::decorated DecorateAddress>
multi_ptr<ElementType, Space, DecorateAddress>
static_address_cast(multi_ptr<ElementType, generic_space, DecorateAddress> ptr);
----
_Preconditions_: `ptr` points to an object allocated in the address space
designated by `Space`.

_Returns_: A `multi_ptr` with the specified address space that points to the
same object as `ptr` and has the same decoration.

[NOTE]
====
Expand All @@ -129,16 +162,30 @@ does not point to an object allocated in the address space designated by

[source,c++]
----
template <access::address_space Space, access::decorated DecorateAddress,
template <access::address_space Space,
typename ElementType>
multi_ptr<ElementType, Space, DecorateAddress>
multi_ptr<ElementType, Space, access::decorated::no>
dynamic_address_cast(ElementType* ptr);
----
_Preconditions_: The memory at `ptr` is accessible to the calling work-item.

_Returns_: A `multi_ptr` with the specified address space and decoration that
points to the same object as `ptr` if `ptr` points to an object allocated in
the address space designated by `Space`, and `nullptr` otherwise.
_Returns_: A `multi_ptr` with the specified address space that points to the
same object as `ptr` if `ptr` points to an object allocated in the address
space designated by `Space`, and `nullptr` otherwise.

[source,c++]
----
template <access::address_space Space,
typename ElementType,
access::decorated DecorateAddress>
multi_ptr<ElementType, Space, DecorateAddress>
dynamic_address_cast(multi_ptr<ElementType, generic_space, DecorateAddress> ptr);
----
_Preconditions_: The memory at `ptr` is accessible to the calling work-item.

_Returns_: A `multi_ptr` with the specified address space that points to the
same object as `ptr` and has the same decoration, if `ptr` points to an object
allocated in the address space designated by `Space`, and `nullptr` otherwise.

[NOTE]
====
Expand Down