From 1475d09281fad0a42250f9950efe2d2eff5a52aa Mon Sep 17 00:00:00 2001 From: GarveyJoe <49960392+GarveyJoe@users.noreply.github.com> Date: Mon, 8 Jun 2020 18:46:08 -0400 Subject: [PATCH 1/3] Added a SPIR-V spec that introduces two new storage classes. --- .../SPV_INTEL_usm_storage_classes.asciidoc | 147 ++++++++++++++++++ 1 file changed, 147 insertions(+) create mode 100644 sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc diff --git a/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc b/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc new file mode 100644 index 0000000000000..7d1d5fbaf9c83 --- /dev/null +++ b/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc @@ -0,0 +1,147 @@ += 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. + +== 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"] +|==== +| USMStorageClassesINTEL | 5935 +| DeviceOnlyINTEL | 5936 +| HostOnlyINTEL | 5937 +|==== + +== 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, replacing the descriptions of *OpPtrCastToGeneric* and *OpGenericCastToPtr* with the following. Their required capabilities and operands remain unchanged. + +|=== +*OpPtrCastToGeneric* + +Converts a pointer's Storage Class to a more generic class. + +_Result Type_ must be an OpTypePointer. Its Storage Class must be *Generic* or *CrossWorkgroup*. + +If _Result Type_'s Storage Class is *Generic*, _Pointer_ must point to the *Workgroup*, *CrossWorkgroup*, *Function*, *DeviceOnlyINTEL*, or *HostOnlyINTEL* Storage Class. If _Result Type_'s Storage Class is *CrossWorkgroup*, _Pointer_ must point to the *DeviceOnlyINTEL* or *HostOnlyINTEL* Storage Class. + +_Result Type_ and _Pointer_ must point to the same type. +|=== + +|=== +*OpGenericCastToPtr* + +Convert a pointer's Storage Class to a more specific class. + +_Result Type_ must be an OpTypePointer. If _Pointer_ points to the *Generic* Storage Class, _Result Type_'s Storage Class must be *Workgroup*, *CrossWorkgroup*, *Function*, *DeviceOnlyINTEL*, or *HostOnlyINTEL*. If _Pointer_ points to the *CrossWorkgroup* storage class, _Result Type_'s Storage Class must be *DeviceOnlyINTEL* or *HostOnlyINTEL*. + +_Pointer_ must point to the *Generic* or *CrossWorkgroup* Storage Class. + +_Result Type_ and _Pointer_ must point to the same type. +|=== + +=== 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* +|======================================== From decaa0eb15e579edae793ba1060cd6e57a7ca0b2 Mon Sep 17 00:00:00 2001 From: GarveyJoe <49960392+GarveyJoe@users.noreply.github.com> Date: Thu, 18 Jun 2020 14:52:12 -0400 Subject: [PATCH 2/3] Added new conversion instructions instead of changing the behaviour of the existing ones. --- .../SPV_INTEL_usm_storage_classes.asciidoc | 37 +++++++++++++------ 1 file changed, 25 insertions(+), 12 deletions(-) diff --git a/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc b/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc index 7d1d5fbaf9c83..e13b44332d30a 100644 --- a/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc +++ b/sycl/doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc @@ -39,7 +39,9 @@ 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. +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: @@ -61,9 +63,11 @@ USMStorageClassesINTEL [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 @@ -96,30 +100,38 @@ Modify Section 3.31, Capability, adding a row to the Capability table: === Instructions -Modify Section 3.36.11, Conversion Instructions, replacing the descriptions of *OpPtrCastToGeneric* and *OpGenericCastToPtr* with the following. Their required capabilities and operands remain unchanged. +Modify Section 3.36.11, Conversion Instructions, adding two new instructions as follows: |=== -*OpPtrCastToGeneric* +3+^| *OpPtrCastToCrossWorkgroupINTEL* -Converts a pointer's Storage Class to a more generic class. +Converts a pointer's Storage Class from a more specific class to *CrossWorkgroup*. -_Result Type_ must be an OpTypePointer. Its Storage Class must be *Generic* or *CrossWorkgroup*. +_Result Type_ must be an OpTypePointer. Its Storage Class must be *CrossWorkgroup*. -If _Result Type_'s Storage Class is *Generic*, _Pointer_ must point to the *Workgroup*, *CrossWorkgroup*, *Function*, *DeviceOnlyINTEL*, or *HostOnlyINTEL* Storage Class. If _Result Type_'s Storage Class is *CrossWorkgroup*, _Pointer_ must point to the *DeviceOnlyINTEL* or *HostOnlyINTEL* Storage Class. +_Pointer_ must point to the *DeviceOnlyINTEL* or *HostOnlyINTEL* Storage Class. -_Result Type_ and _Pointer_ must point to the same type. +_Result Type_ and _Pointer_ must point to the same type. 2+^| Capability: + +*USMStorageClassesINTEL* +| 4 | 5934 | + +Result Type | Result | + +Pointer |=== |=== -*OpGenericCastToPtr* +3+^| *OpCrossWorkgroupCastToPtrINTEL* -Convert a pointer's Storage Class to a more specific class. +Convert a pointer's Storage Class from *CrossWorkgroup* to a more specific class. -_Result Type_ must be an OpTypePointer. If _Pointer_ points to the *Generic* Storage Class, _Result Type_'s Storage Class must be *Workgroup*, *CrossWorkgroup*, *Function*, *DeviceOnlyINTEL*, or *HostOnlyINTEL*. If _Pointer_ points to the *CrossWorkgroup* storage class, _Result Type_'s Storage Class must be *DeviceOnlyINTEL* or *HostOnlyINTEL*. +_Result Type_ must be an OpTypePointer. _Result Type_'s Storage Class must be *DeviceOnlyINTEL* or *HostOnlyINTEL*. -_Pointer_ must point to the *Generic* or *CrossWorkgroup* Storage Class. +_Pointer_ must point to the *CrossWorkgroup* Storage Class. -_Result Type_ and _Pointer_ must point to the same type. +_Result Type_ and _Pointer_ must point to the same type. 2+^| Capability: + +*USMStorageClassesINTEL* +| 4 | 5938 | + +Result Type | Result | + +Pointer |=== === Validation Rules @@ -144,4 +156,5 @@ None. |======================================== |Rev|Date|Author|Changes |A|2020-06-08|Joe Garvey|*First public draft* +|B|2020-06-18|Joe Garvey|*Added new conversion instructions* |======================================== From 66bdf4a6c753448376aec1d7411dc626c2c5cf1c Mon Sep 17 00:00:00 2001 From: GarveyJoe <49960392+GarveyJoe@users.noreply.github.com> Date: Thu, 18 Jun 2020 16:07:06 -0400 Subject: [PATCH 3/3] Added SYCL_INTEL_usm_address_spaces, an extension that adds two new address spaces to SYCL to enable additional optimization. --- .../usm_address_spaces.asciidoc | 124 ++++++++++++++++++ 1 file changed, 124 insertions(+) create mode 100644 sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc diff --git a/sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc b/sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc new file mode 100644 index 0000000000000..73e9de4758904 --- /dev/null +++ b/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() 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 + template + multi_ptr( + accessor) +``` | Constructs a multi_ptr 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 +using device_ptr = multi_ptr + +template +using host_ptr = multi_ptr +``` + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|A|2020-06-18|Joe Garvey|Initial public draft +|========================================