[SYCL][Docs] Update sycl_ext_intel_usm_address_spaces and fix ctors (i…

…ntel#7680)

This commit updates the sycl_ext_intel_usm_address_spaces extension to
adhere to SYCL 2020 `multi_ptr` and updates the extension specification
to use the new extension template.

Additionally this commit fixes the `multi_ptr` constructors for the
extension address spaces.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc

@@ -1,139 +1,203 @@
 = sycl_ext_intel_usm_address_spaces
 
-== Introduction
-This extension introduces two new address spaces and their corresponding multi_ptr specializations.  
-These address spaces are subsets of the global address space and are added to enable users to provide more optimization information to their compiler.  
+:source-highlighter: coderay
+:coderay-linenums-mode: table
 
-IMPORTANT: This specification is a draft.
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
 
-NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos.
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
 
-NOTE: This document is better viewed when rendered as html with asciidoctor.  GitHub does not render image icons.
-This document describes an extension to the SYCL USM extension that adds new explicit address spaces for the possible locations that USM pointers can be allocated.  Users can create pointers that point into these address spaces explicitly in order to pass additional information to their compiler so as to enable optimizations.  
 
 == Notice
-Copyright (c) 2020 Intel Corporation.  All rights reserved.
 
-== Status
+[%hardbreaks]
+Copyright (C) 2022 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
 
-Draft
 
-This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension.
+== Contact
 
-Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products.
+To report problems with this extension, please open a new issue at:
 
-== Version
+https://github.com/intel/llvm/issues
 
-Built On: {docdate} +
-Revision: 2
 
 == Dependencies
 
-This extension is written against the SYCL 2020 specification, Revision 3.
+This extension is written against the SYCL 2020 revision 6 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+
+== Status
+
+This extension is implemented and fully supported by {dpcpp}.
+
+
+== Overview
+
+This extension adds two new address spaces: device and host that are subsets of
+the global address space.
+New interfaces for `multi_ptr` are added for each of these address spaces.
+
+The goal of this division of the global address space is to enable users to
+explicitly tell the compiler which address space a pointer resides in for the
+purposes of enabling optimization.
+While automatic address space inference is often possible for accessors, it is
+harder for USM pointers as it requires inter-procedural optimization with the
+host code.
+This additional information can be particularly beneficial on FPGA targets where
+knowing that a pointer only ever accesses host or device memory can allow
+compilers to produce more area efficient memory-accessing hardware.
+
 
-If SPIR-V is used by the implementation, this extension also requires support for the SPV_INTEL_usm_storage_classes SPIR-V extension.
+== Specification
 
-== Feature Test Macro
+=== Feature test macro
 
 This extension provides a feature-test macro as described in the core SYCL
-specification section 6.3.3 "Feature test macros".  Therefore, an
-implementation supporting this extension must predefine the macro
-`SYCL_EXT_INTEL_USM_ADDRESS_SPACES` to one of the values defined in the table below.
-Applications can test for the existence of this macro to determine if the
-implementation supports this feature, or applications can test the macro's
-value to determine which of the extension's APIs the implementation supports.
+specification.  An implementation supporting this extension must predefine the
+macro `SYCL_EXT_INTEL_USM_ADDRESS_SPACES` to one of the values defined in the table
+below.  Applications can test for the existence of this macro to determine if
+the implementation supports this feature, or applications can test the macro's
+value to determine which of the extension's features the implementation
+supports.
 
 [%header,cols="1,5"]
 |===
-|Value |Description
-|1     |Initial extension version.  Base features are supported.
-|===
+|Value
+|Description
 
-== Overview
+|1
+|Initial version of this extension.
 
-This extension adds two new address spaces: device and host that are subsets of the global address space.  
-New specializations of multi_ptr are added for each of these address spaces.  
+|2
+|Adds `sycl::ext::intel::host_ptr`, `sycl::ext::intel::raw_host_ptr`,
+`sycl::ext::intel::decorated_host_ptr`, `sycl::ext::intel::device_ptr`,
+`sycl::ext::intel::raw_device_ptr` and `sycl::ext::intel::decorated_device_ptr`.
+`sycl::host_ptr` and `sycl::device_ptr` are deprecated.
+|===
 
-The goal of this division of the global address space is to enable users to explicitly tell the compiler which address space a pointer resides in for the purposes of enabling optimization.  
-While automatic address space inference is often possible for accessors, it is harder for USM pointers as it requires inter-procedural optimization with the host code.
-This additional information can be particularly beneficial on FPGA targets where knowing that a pointer only ever accesses host or device memory can allow compilers to produce more area efficient memory-accessing hardware.  
+== Modifications to SYCL 2020
 
-== Modifications to the SYCL Specification, Version 2020 revision 3
+The following sections contain the related changes and additions to the SYCL
+2020 specification relating to this extension.
 
-=== Section 3.8.2 SYCL Device Memory Model
+=== SYCL Device Memory Model
 
 Add to the end of the definition of global memory:
-Global memory is a virtual address space which overlaps the device and host address spaces.  
+Global memory is a virtual address space which overlaps the device and host
+address spaces.
 
 Add two new memory regions as follows:
 
-*Device memory* is a sub-region of global memory that is not directly accessible by the host.  Global accessors and USM allocations of the device alloc type reside in this address space.
-
-*Host memory* is a sub-region of global memory.  USM pointers allocated with the host alloc type reside in this address space.  
+*Device memory* is a sub-region of global memory that is not directly accessible
+by the host.  Buffer accessors and USM allocations whose kind is
+`usm::alloc::device` reside in this address space.
 
-=== Section 3.8.2.1 Access to memory
+*Host memory* is a sub-region of global memory.  USM allocations whose kind is
+`usm::alloc::host` reside in this address space.
 
-In the second last paragraph, add sycl::device_ptr and sycl::host_ptr to the list of explicit pointer classes.  
 
-=== Section 4.7.7.1 Multi-pointer Class
+=== Multi-pointer Class
 
-In the overview of the multi_ptr class replace the address_space enum with the following:
+Add the following enumerations to the `access::address_space` enum:
 ```c++
-enum class address_space : int {
-  global_space,
-  local_space,
-  constant_space, // Deprecated in SYCL 2020
-  private_space,
-  generic_space,
+enum class address_space : /* unspecified */ {
+  ...
   ext_intel_global_device_space,
   ext_intel_global_host_space
 };
 ```
 
-Add the following new conversion operator:
+Add the following new conversion operator to the `multi_ptr` class:
 ```c++
 // Explicit conversion to global_space
 // Only available if Space == address_space::ext_intel_global_device_space || Space == address_space::ext_intel_global_host_space
-explicit operator multi_ptr<ElementType, access::address_space::global_space>() const; 
+explicit operator multi_ptr<ElementType, access::address_space::global_space, DecorateAddress>() const;
 ```
 
-Add a new row to Table 91: Constructors of the SYCL multi_ptr class template, as follows:
+Change the `multi_ptr` constructor taking an accessor with `target::device` to
+also allow `access::address_space::ext_intel_global_device_space` as follows:
 
 --
 [options="header"]
 |===
 | Constructor | Description
 a|
 ```c++
-template<typename ElementType, access::
-address_space Space = access::address_space::
-ext_intel_global_device_space>
-    template <int dimensions, access::mode Mode>
-    multi_ptr(
-    accessor<ElementType, dimensions, Mode, access::
-    target::global_buffer>)
-``` | Constructs a multi_ptr<ElementType, access::address_space::ext_intel_global_device_space> from an accessor of access::target::global_buffer.
+template <int Dimensions, access_mode Mode, access::placeholder IsPlaceholder>
+multi_ptr(
+    accessor<ElementType, Dimensions, Mode, target::device, IsPlaceholder>);
+```
+| Available only when:
+`Space == access::address_space::global_space \|\| Space == access::address_space::ext_intel_global_device_space \|\| Space == access::address_space::generic_space`.
+
+Constructs a `multi_ptr` from an accessor of `target::device`.
+
+This constructor may only be called from within a command.
 |===
 --
 
-=== Section 4.7.7.2 Explicit Pointer Aliases
 
-Add device_ptr and host_ptr aliases to the list of multi_ptr aliases as follows:
+=== Explicit Pointer Aliases
+
+Add `device_ptr` and `host_ptr` aliases to the list of `multi_ptr` aliases as
+follows:
 ```c++
+namespace sycl {
+
+// Deprecated.
+template<typename ElementType,
+         access::decorated IsDecorated = access::decorated::legacy>
+using device_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
+              IsDecorated>
+
+// Deprecated.
+template<typename ElementType,
+         access::decorated IsDecorated = access::decorated::legacy>
+using host_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
+              IsDecorated>
+
+namespace ext {
+namespace intel {
+
 template<typename ElementType>
-using device_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_device_space>
+using raw_device_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
+              access::decorated::no>
 
 template<typename ElementType>
-using host_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_host_space>
-```
+using raw_host_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
+              access::decorated::no>
 
-== Revision History
+template<typename ElementType>
+using decorated_device_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
+              access::decorated::yes>
+
+template<typename ElementType>
+using decorated_host_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
+              access::decorated::yes>
+
+} // namespace intel
+} // namespace ext
+} // namespace sycl
+```
 
-[cols="5,15,15,70"]
-[grid="rows"]
-[options="header"]
-|========================================
-|Rev|Date|Author|Changes
-|1|2020-06-18|Joe Garvey|Initial public draft
-|2|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions
-|======================================== 

sycl/include/sycl/access/access.hpp

@@ -337,6 +337,17 @@ template <typename ToT, typename FromT> inline ToT cast_AS(FromT from) {
       return reinterpret_cast<ToT>(from);
 #endif // defined(__NVPTX__) || defined(__AMDGCN__)
   } else
+#ifdef __ENABLE_USM_ADDR_SPACE__
+      if constexpr (FromAS == access::address_space::global_space &&
+                    (ToAS ==
+                         access::address_space::ext_intel_global_device_space ||
+                     ToAS ==
+                         access::address_space::ext_intel_global_host_space)) {
+    // Casting from global address space to the global device and host address
+    // spaces is allowed.
+    return (ToT)from;
+  } else
+#endif // __ENABLE_USM_ADDR_SPACE__
 #endif // __SYCL_DEVICE_ONLY__
   {
     return reinterpret_cast<ToT>(from);

sycl/include/sycl/ext/intel/usm_pointers.hpp

@@ -0,0 +1,64 @@
+//==-------- usm_pointers.hpp - Extended SYCL pointers classes -------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+#include <sycl/access/access.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+
+template <typename ElementType, access::address_space Space,
+          access::decorated DecorateAddress>
+class multi_ptr;
+
+namespace ext {
+namespace intel {
+
+template <typename ElementType,
+          access::decorated IsDecorated = access::decorated::legacy>
+using device_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
+              IsDecorated>;
+
+template <typename ElementType,
+          access::decorated IsDecorated = access::decorated::legacy>
+using host_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
+              IsDecorated>;
+
+// Template specialization aliases for different pointer address spaces.
+// The interface exposes non-decorated pointer while keeping the
+// address space information internally.
+
+template <typename ElementType>
+using raw_device_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
+              access::decorated::no>;
+
+template <typename ElementType>
+using raw_host_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
+              access::decorated::no>;
+
+// Template specialization aliases for different pointer address spaces.
+// The interface exposes decorated pointer.
+
+template <typename ElementType>
+using decorated_device_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
+              access::decorated::yes>;
+
+template <typename ElementType>
+using decorated_host_ptr =
+    multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
+              access::decorated::yes>;
+
+} // namespace intel
+} // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/feature_test.hpp.in

@@ -62,7 +62,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 #define SYCL_EXT_INTEL_FPGA_REG 1
 #define SYCL_EXT_INTEL_KERNEL_ARGS_RESTRICT 1
 #define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1
-#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1
+#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 2
 #define SYCL_EXT_INTEL_RUNTIME_BUFFER_LOCATION 1
 #define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 3
 #define SYCL_EXT_ONEAPI_USM_DEVICE_READ_ONLY 1