[SYCL][DOC] Add SPV_INTEL_usm_storage_classes and SYCL_INTEL_usm_addr…

…ess_spaces extensions (#1840)

SPV_INTEL_usm_storage_classes extension introduces two new storage classes that
are sub classes of the CrossWorkgroup storage class. Using these more specific
storage classes provides additional information that can enable optimization.

SYCL_INTEL_usm_address_spaces 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.

sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc

@@ -0,0 +1,160 @@
+= SPV_INTEL_usm_storage_classes
+
+== Name Strings
+
+SPV_INTEL_usm_storage_classes
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/KhronosGroup/SPIRV-Headers
+
+== Contributors
+
+- Joe Garvey, Intel
+
+== Notice
+
+Copyright (c) 2020 Intel Corporation.  All rights reserved.
+
+== Status
+
+Draft
+
+== Version
+
+[width="40%",cols="25,25"]
+|========================================
+| Last Modified Date | {docdate}
+| Revision           | A
+|========================================
+
+== Dependencies
+
+This extension is written against the SPIR-V Specification,
+Version 1.5 Revision 3.
+
+This extension requires SPIR-V 1.0.
+
+== Overview
+
+This extension introduces two new storage classes that are sub classes of the CrossWorkgroup storage class.  
+Using these more specific storage classes provides additional information that can enable optimization. 
+The extension also introduces two new conversion instructions to enable converting pointers from and to these storage classes.  
+
+== Extension Name
+To use this extension within a SPIR-V module, the following *OpExtension* must be present in the module:
+
+----
+OpExtension "SPV_INTEL_usm_storage_classes"
+----
+
+== New capabilities
+This extension introduces a new capability:
+
+----
+USMStorageClassesINTEL
+----
+
+== Token Number Assignments
+
+[width="40%"]
+[cols="70%,30%"]
+[grid="rows"]
+|====
+| OpPtrCastToCrossWorkgroupINTEL | 5934
+| USMStorageClassesINTEL  | 5935
+| DeviceOnlyINTEL | 5936
+| HostOnlyINTEL | 5937
+| OpCrossWorkgroupCastToPtrINTEL | 5938
+|====
+
+== Modifications to the SPIR-V Specification, Version 1.5 Revision 3
+
+=== Storage Class
+
+Modify Section 3.7, Storage Class, adding these rows to the table:
+
+--
+[options="header"]
+|===
+2+^| Storage Class | Enabling Capabilities
+| 5936 | *DeviceOnlyINTEL* +
+A subset of the *CrossWorkgroup* *Storage Class*.  Stored on the device. SYCL device allocations. | *USMStorageClassesINTEL*
+| 5937 | *HostOnlyINTEL* +
+A subset of the *CrossWorkgroup* *Storage Class*.  Stored on the host. SYCL host allocations. | *USMStorageClassesINTEL*
+|===
+--
+
+=== Capability
+
+Modify Section 3.31, Capability, adding a row to the Capability table:
+--
+[options="header"]
+|====
+2+^| Capability ^| Implicitly Declares
+| 5935 | *USMStorageClassesINTEL* | *Kernel*
+|====
+--
+
+=== Instructions
+
+Modify Section 3.36.11, Conversion Instructions, adding two new instructions as follows:
+
+|===
+3+^| *OpPtrCastToCrossWorkgroupINTEL*
+
+Converts a pointer's Storage Class from a more specific class to *CrossWorkgroup*.
+
+_Result Type_ must be an OpTypePointer.  Its Storage Class must be *CrossWorkgroup*.  
+
+_Pointer_ must point to the *DeviceOnlyINTEL* or *HostOnlyINTEL* Storage Class.  
+
+_Result Type_ and _Pointer_ must point to the same type. 2+^| Capability: +
+*USMStorageClassesINTEL*
+| 4 | 5934 | <id> +
+Result Type | Result <id> | <id> +
+Pointer
+|===
+
+|===
+3+^| *OpCrossWorkgroupCastToPtrINTEL*
+
+Convert a pointer's Storage Class from *CrossWorkgroup* to a more specific class.
+
+_Result Type_ must be an OpTypePointer.  _Result Type_'s Storage Class must be *DeviceOnlyINTEL* or *HostOnlyINTEL*.
+
+_Pointer_ must point to the *CrossWorkgroup* Storage Class.
+
+_Result Type_ and _Pointer_ must point to the same type.  2+^| Capability: +
+*USMStorageClassesINTEL*
+| 4 | 5938 | <id> +
+Result Type | Result <id> | <id> +
+Pointer
+|===
+
+=== Validation Rules
+
+None.
+
+== Issues
+
+None.
+
+//. Issue.
+//+
+//--
+//*RESOLVED*: Resolution.
+//--
+
+== Revision History
+
+[cols="5,15,15,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|Rev|Date|Author|Changes
+|A|2020-06-08|Joe Garvey|*First public draft*
+|B|2020-06-18|Joe Garvey|*Added new conversion instructions*
+|======================================== 

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
+|========================================