Skip to content

Commit

Permalink
[SYCL][Doc] Update address_cast specification (#12689)
Browse files Browse the repository at this point in the history
This PR contains three improvements to the specification:
- Default address_cast to no decoration.
- Add overloads for generic multi_ptr.
- Add address_space shorthands.

This allows developers using generic raw pointers to cast to `multi_ptr`
without worrying about decoration:
```c++
float* ptr = ...;
auto mptr = syclex::static_address_cast<syclex::local_space>(ptr);
```

...or to opt-in to decoration by using a generic multi-pointer, which
propagates decoration information:
```c++
multi_ptr<float, syclex::generic_space, access::decorated::yes> decorated_in = ...;
auto decorated_out = syclex::static_address_cast<syclex::local_space>(decorated_in);
```

---------

Signed-off-by: John Pennycook <john.pennycook@intel.com>
  • Loading branch information
Pennycook authored Feb 28, 2024
1 parent c232e8e commit 84a92e0
Showing 1 changed file with 60 additions and 13 deletions.
73 changes: 60 additions & 13 deletions sycl/doc/extensions/proposed/sycl_ext_oneapi_address_cast.asciidoc
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);
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

0 comments on commit 84a92e0

Please sign in to comment.