Add sycl_khr_addrspace_cast extension

This extension splits sycl::address_space_cast into two functions:

- khr::static_addrspace_cast, which casts with no run-time checks.
- khr::dynamic_addrspace_cast, which casts with run-time checks.

adoc/extensions/index.adoc

@@ -11,3 +11,5 @@ specification, but their design is subject to change.
 
 // leveloffset=2 allows extensions to be written as standalone documents
 // include::sycl_khr_extension_name.adoc[leveloffset=2]
+
+include::sycl_khr_addrspace_cast.adoc[leveloffset=2]

adoc/extensions/sycl_khr_addrspace_cast.adoc

@@ -0,0 +1,155 @@
+[[sec:khr-addrspace-cast]]
+= SYCL_KHR_ADDRSPACE_CAST
+
+[code]#sycl::address_space_cast# does two things: first, it checks whether a
+given raw pointer can be cast to a specific address space; and second, it
+performs the requested cast.
+In cases where a developer is asserting that a raw pointer points to an
+object in a specific address space, run-time checks are not required and may
+have undesirable performance impact.
+
+This extension splits [code]#sycl::address_space_cast# into two functions:
+
+- [code]#static_addrspace_cast#, which casts with no run-time checks.
+- [code]#dynamic_addrspace_cast#, which casts with run-time checks.
+
+[[sec:khr-addrspace-cast-dependencies]]
+== Dependencies
+
+This extension has no dependencies on other extensions.
+
+[[sec:khr-addrspace-cast-feature-test]]
+== Feature test macro
+
+An implementation supporting this extension must predefine the macro
+[code]#SYCL_KHR_ADDRSPACE_CAST# to one of the values defined in the table
+below.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+[[sec:khr-addrspace-cast-functions]]
+== Address space cast functions
+
+[source,role=synopsis]
+----
+
+namespace sycl {
+namespace khr {
+
+template <addrspace Space, typename ElementType>
+multi_ptr<ElementType, Space, access::decorated::no>
+static_addrspace_cast(ElementType* ptr);
+
+template <addrspace Space, typename ElementType, access::decorated DecorateAddress>
+multi_ptr<ElementType, Space, DecorateAddress>
+static_addrspace_cast(multi_ptr<ElementType, addrspace_generic, DecorateAddress> ptr);
+
+template <addrspace Space, typename ElementType>
+multi_ptr<ElementType, Space, access::decorated::no>
+dynamic_addrspace_cast(ElementType* ptr);
+
+template <addrspace Space, typename ElementType, access::decorated DecorateAddress>
+multi_ptr<ElementType, Space, DecorateAddress>
+dynamic_addrspace_cast(multi_ptr<ElementType, addrspace_generic, DecorateAddress> ptr);
+
+} // namespace khr
+} // namespace sycl
+----
+
+.[apidef]#static_addrspace_cast#
+[source,role=synopsis,id=api:khr-addrspace-cast-static_addrspace_cast]
+----
+template <addrspace Space, typename ElementType>
+multi_ptr<ElementType, Space, access::decorated::no>
+static_addrspace_cast(ElementType* ptr);
+
+template <addrspace Space, typename ElementType, access::decorated DecorateAddress>
+multi_ptr<ElementType, Space, DecorateAddress>
+static_addrspace_cast(multi_ptr<ElementType, addrspace_generic, DecorateAddress> ptr);
+----
+
+_Preconditions_: [code]#ptr# points to an object allocated in the address space
+designated by [code]#Space#.
+
+_Returns_: A [code]#multi_ptr# with the specified address space that points to
+the same object as [code]#ptr#.
+If [code]#ptr# is a [code]#multi_ptr# then the return value has the same
+decoration.
+
+{note}Implementations may choose to issue a diagnostic if they can prove that
+[code]#ptr# does not point to an object allocated in the address space
+designated by [code]#Space#.{endnote}
+
+'''
+
+.[apidef]#dynamic_addrspace_cast#
+[source,role=synopsis,id=api:khr-addrspace-cast-dynamic_addrspace_cast]
+----
+template <addrspace Space, typename ElementType>
+multi_ptr<ElementType, Space, access::decorated::no>
+dynamic_addrspace_cast(ElementType* ptr);
+
+template <addrspace Space, typename ElementType, access::decorated DecorateAddress>
+multi_ptr<ElementType, Space, DecorateAddress>
+dynamic_addrspace_cast(multi_ptr<ElementType, addrspace_generic, DecorateAddress> ptr);
+----
+
+_Preconditions_: The memory at [code]#ptr# can be accessed by the calling
+work-item.
+
+_Returns_: A [code]#multi_ptr# with the specified address space that points to
+the same object as [code]#ptr# if [code]#ptr# points to an object allocated in
+the address space designated by [code]#Space#, and [code]#nullptr# otherwise.
+If [code]#ptr# is a [code]#multi_ptr# then the return value has the same
+decoration.
+
+{note}The precondition prevents reasoning about the address space of pointers
+originating from another work-item (in the case of [code]#private# pointers) or
+another work-group (in the case of [code]#local# pointers).
+Such pointers could not be dereferenced by the calling work-item, and it is
+thus unclear that being able to reason about the address space would be useful.
+Limiting usage to accessible pointers is expected to result in simpler and
+faster implementations.{endnote}
+
+[[sec:khr-addrspace-cast-example]]
+== Example
+
+The example below demonstrates the usage of this extension.
+
+[source,,linenums]
+----
+#include <sycl/sycl.hpp>
+using namespace sycl; // (optional) avoids need for "sycl::" before SYCL name
+
+// This function accepts raw pointers, but assumes a global address space.
+template <typename T>
+void update_global(T* ptr, int x)
+{
+  // Assert that implementation can treat ptr as pointing to global.
+  auto mptr = khr::static_addrspace_cast<addrspace_global>(ptr);
+
+  *mptr += x;
+}
+
+// This function accepts raw pointers, but dispatches internally.
+template <typename T>
+void update(T* ptr, int x)
+{
+  // If cast to global returns non-null, call the global version.
+  if (khr::dynamic_addrspace_cast<addrspace_global>(ptr)) {
+    update_global(ptr, x);
+  }
+
+  // If cast to local returns non-null, call the local version.
+  else if (khr::dynamic_addrspace_cast<addrspace_local>(ptr)) {
+    update_local(ptr, x);
+  }
+}
+----