Skip to content

Commit fac9b20

Browse files
gmlueckgmlueck
authored
Merge pull request KhronosGroup#650 from Pennycook/khr_addrspace_cast
Add sycl_khr_addrspace_cast extension
2 parents a009088 + c2972dd commit fac9b20

File tree

3 files changed

+178
-0
lines changed

3 files changed

+178
-0
lines changed

‎adoc/extensions/index.adoc‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,3 +16,5 @@ include::sycl_khr_group_interface.adoc[leveloffset=2]
1616
include::sycl_khr_max_work_group_queries.adoc[leveloffset=2]
1717
include::sycl_khr_queue_flush.adoc[leveloffset=2]
1818
include::sycl_khr_work_item_queries.adoc[leveloffset=2]
19+
include::sycl_khr_static_addrspace_cast.adoc[leveloffset=2]
20+
include::sycl_khr_dynamic_addrspace_cast.adoc[leveloffset=2]

‎adoc/extensions/sycl_khr_dynamic_addrspace_cast.adoc‎

Lines changed: 91 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,91 @@
1+
[[sec:khr-dynamic-addrspace-cast]]
2+
= sycl_khr_dynamic_addrspace_cast
3+
4+
This extension introduces a [code]#dynamic_addrspace_cast# function with the
5+
same semantics as [code]#sycl::address_space_cast# to align with the
6+
[code]#static_addrspace_cast# function defined by the
7+
SYCL_KHR_STATIC_ADDRSPACE_CAST extension, and clarifies the expected behavior of
8+
a dynamic address space cast.
9+
10+
[[sec:khr-dynamic-addrspace-cast-dependencies]]
11+
== Dependencies
12+
13+
This extension has no dependencies on other extensions.
14+
15+
[[sec:khr-dynamic-addrspace-cast-feature-test]]
16+
== Feature test macro
17+
18+
An implementation supporting this extension must predefine the macro
19+
[code]#SYCL_KHR_DYNAMIC_ADDRSPACE_CAST# to one of the values defined in the
20+
table below.
21+
22+
[%header,cols="1,5"]
23+
|===
24+
|Value
25+
|Description
26+
27+
|1
28+
|Initial version of this extension.
29+
|===
30+
31+
[[sec:khr-dynamic-addrspace-cast-functions]]
32+
== Dynamic address space cast functions
33+
34+
.[apidef]#dynamic_addrspace_cast#
35+
[source,role=synopsis,id=api:khr-dynamic-addrspace-cast-dynamic_addrspace_cast]
36+
----
37+
namespace sycl::khr {
38+
39+
template <access::address_space Space, typename ElementType>
40+
multi_ptr<ElementType, Space, access::decorated::no>
41+
dynamic_addrspace_cast(ElementType* ptr);
42+
43+
template <access::address_space Space, typename ElementType, access::decorated DecorateAddress>
44+
multi_ptr<ElementType, Space, DecorateAddress>
45+
dynamic_addrspace_cast(multi_ptr<ElementType, addrspace_generic, DecorateAddress> ptr);
46+
47+
} // namespace khr::sycl
48+
----
49+
50+
_Preconditions_: The memory at [code]#ptr# can be accessed by the calling
51+
work-item.
52+
53+
_Returns_: A [code]#multi_ptr# with the specified address space that points to
54+
the same object as [code]#ptr# if [code]#ptr# points to an object allocated in
55+
the address space designated by [code]#Space#, and [code]#nullptr# otherwise.
56+
If [code]#ptr# is a [code]#multi_ptr# then the return value has the same
57+
decoration.
58+
59+
{note}The precondition prevents reasoning about the address space of pointers
60+
originating from another work-item (in the case of [code]#private# pointers) or
61+
another work-group (in the case of [code]#local# pointers).
62+
Such pointers could not be dereferenced by the calling work-item, and it is thus
63+
unclear that being able to reason about the address space would be useful.
64+
Limiting usage to accessible pointers is expected to result in simpler and
65+
faster implementations.{endnote}
66+
67+
[[sec:khr-dynamic-addrspace-cast-example]]
68+
== Example
69+
70+
The example below demonstrates the usage of this extension.
71+
72+
[source,,linenums]
73+
----
74+
#include <sycl/sycl.hpp>
75+
using namespace sycl; // (optional) avoids need for "sycl::" before SYCL name
76+
77+
// This function accepts raw pointers, but dispatches internally.
78+
template <typename T>
79+
void update(T* ptr, int x)
80+
{
81+
// If cast to global returns non-null, call the global version.
82+
if (khr::dynamic_addrspace_cast<access::address_space::global_space>(ptr).get() != nullptr) {
83+
update_global(ptr, x);
84+
}
85+
86+
// If cast to local returns non-null, call the local version.
87+
else if (khr::dynamic_addrspace_cast<access::address_space::local_space>(ptr).get() != nullptr) {
88+
update_local(ptr, x);
89+
}
90+
}
91+
----

‎adoc/extensions/sycl_khr_static_addrspace_cast.adoc‎

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
[[sec:khr-static-addrspace-cast]]
2+
= sycl_khr_static_addrspace_cast
3+
4+
[code]#sycl::address_space_cast# does two things: first, it checks whether a
5+
given raw pointer can be cast to a specific address space; and second, it
6+
performs the requested cast.
7+
In cases where a developer is asserting that a raw pointer points to an object
8+
in a specific address space, run-time checks are not required and may have
9+
undesirable performance impact.
10+
This extension defines [code]#static_addrspace_cast# to provide developers a
11+
mechanism which casts with no run-time checks, enabling address space casts
12+
without any performance overhead.
13+
14+
[[sec:khr-static-addrspace-cast-dependencies]]
15+
== Dependencies
16+
17+
This extension has no dependencies on other extensions.
18+
19+
[[sec:khr-static-addrspace-cast-feature-test]]
20+
== Feature test macro
21+
22+
An implementation supporting this extension must predefine the macro
23+
[code]#SYCL_KHR_STATIC_ADDRSPACE_CAST# to one of the values defined in the table
24+
below.
25+
26+
[%header,cols="1,5"]
27+
|===
28+
|Value
29+
|Description
30+
31+
|1
32+
|Initial version of this extension.
33+
|===
34+
35+
[[sec:khr-static-addrspace-cast-functions]]
36+
== Static address space cast functions
37+
38+
.[apidef]#static_addrspace_cast#
39+
[source,role=synopsis,id=api:khr-static-addrspace-cast-static_addrspace_cast]
40+
----
41+
namespace sycl::khr {
42+
43+
template <access::address_space Space, typename ElementType>
44+
multi_ptr<ElementType, Space, access::decorated::no>
45+
static_addrspace_cast(ElementType* ptr);
46+
47+
template <access::address_space Space, typename ElementType, access::decorated DecorateAddress>
48+
multi_ptr<ElementType, Space, DecorateAddress>
49+
static_addrspace_cast(multi_ptr<ElementType, addrspace_generic, DecorateAddress> ptr);
50+
51+
} // namespace khr::sycl
52+
----
53+
54+
_Preconditions_: [code]#ptr# points to an object allocated in the address space
55+
designated by [code]#Space#.
56+
57+
_Returns_: A [code]#multi_ptr# with the specified address space that points to
58+
the same object as [code]#ptr#.
59+
If [code]#ptr# is a [code]#multi_ptr# then the return value has the same
60+
decoration.
61+
62+
{note}Implementations may choose to issue a diagnostic if they can prove that
63+
[code]#ptr# does not point to an object allocated in the address space
64+
designated by [code]#Space#.{endnote}
65+
66+
[[sec:khr-static-addrspace-cast-example]]
67+
== Example
68+
69+
The example below demonstrates the usage of this extension.
70+
71+
[source,,linenums]
72+
----
73+
#include <sycl/sycl.hpp>
74+
using namespace sycl; // (optional) avoids need for "sycl::" before SYCL name
75+
76+
// This function accepts raw pointers, but assumes a global address space.
77+
template <typename T>
78+
void update_global(T* ptr, int x)
79+
{
80+
// Assert that implementation can treat ptr as pointing to global.
81+
auto mptr = khr::static_addrspace_cast<access::address_space::global_space>(ptr);
82+
83+
*mptr += x;
84+
}
85+
----

0 commit comments

Comments
 (0)