Added SYCL_INTEL_usm_address_spaces, an extension that adds two new a…

…ddress spaces to SYCL to enable additional optimization.

sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc

@@ -0,0 +1,124 @@
+= SYCL_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.  
+
+IMPORTANT: This specification is a draft.
+
+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.
+
+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.  
+
+== Name Strings
++SYCL_INTEL_usm_address_spaces+
+
+== Notice
+Copyright (c) 2020 Intel Corporation.  All rights reserved.
+
+== Status
+
+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.
+
+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.
+
+== Version
+
+Built On: {docdate} +
+Revision: 1
+
+== Dependencies
+
+This extension is written against the SYCL 1.2.1 specification, Revision 7.  It requires the Unified Shared Memory SYCL proposal.
+
+If SPIR-V is used by the implementation, this extension also requires support for the SPV_INTEL_usm_storage_classes SPIR-V extension.
+
+== Overview
+
+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.  
+
+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 the SYCL Specification, Version 1.2.1 revision 7
+
+=== Section 3.5.2 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.  
+
+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.  
+
+=== Section 3.5.2.1 Access to memory
+
+In the second last paragraph, add cl::sycl::device_ptr and cl::sycl::host_ptr to the list of explicit pointer classes.  
+
+=== Section 4.7.7.1 Multi-pointer Class
+
+In the overview of the multi_ptr class replace the address_space enum with the following:
+```c++
+enum class address_space : int {
+    global_space,
+    local_space,
+    constant_space,
+    private_space,
+    device_space,
+    host_space
+};
+```
+
+Add the following new conversion operator:
+```c++
+// Explicit conversion to global_space
+// Only available if Space == address_space::device_space || Space == address_space::host_space
+explicit operator multi_ptr<ElementType, access::address_space::global_space>() const; 
+```
+
+Add a new row to Table 4.54: Constructors of the SYCL multi_ptr class template, as follows:
+
+--
+[options="header"]
+|===
+| Constructor | Description
+a|
+```c++
+template<typename ElementType, access::
+address_space Space = access::address_space::
+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::device_space> from an accessor of access::target::global_buffer.
+|===
+--
+
+=== Section 4.7.7.2 Explicit Pointer Aliases
+
+Add device_ptr and host_ptr aliases to the list of multi_ptr aliases as follows:
+```c++
+template<typename ElementType>
+using device_ptr = multi_ptr<ElementType, access::address_space::device_space>
+
+template<typename ElementType>
+using host_ptr = multi_ptr<ElementType, access::address_space::host_space>
+```
+
+== Revision History
+
+[cols="5,15,15,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|Rev|Date|Author|Changes
+|A|2020-06-18|Joe Garvey|Initial public draft
+|========================================