Skip to content

Commit 4a9e9a0

Browse files
steffenlarsensteffenlarsen
authored
[SYCL][Docs] Update sycl_ext_intel_usm_address_spaces and fix ctors (#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 <[email protected]>
1 parent 166bbc3 commit 4a9e9a0

File tree

10 files changed

+300
-85
lines changed

10 files changed

+300
-85
lines changed

‎sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc

+138-74
Original file line numberDiff line numberDiff line change
@@ -1,139 +1,203 @@
11
= sycl_ext_intel_usm_address_spaces
22

3-
== Introduction
4-
This extension introduces two new address spaces and their corresponding multi_ptr specializations.
5-
These address spaces are subsets of the global address space and are added to enable users to provide more optimization information to their compiler.
3+
:source-highlighter: coderay
4+
:coderay-linenums-mode: table
65

7-
IMPORTANT: This specification is a draft.
6+
// This section needs to be after the document title.
7+
:doctype: book
8+
:toc2:
9+
:toc: left
10+
:encoding: utf-8
11+
:lang: en
12+
:dpcpp: pass:[DPC++]
813

9-
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.
14+
// Set the default source code type in this document to C++,
15+
// for syntax highlighting purposes. This is needed because
16+
// docbook uses c++ and html5 uses cpp.
17+
:language: {basebackend@docbook:c++:cpp}
1018

11-
NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons.
12-
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.
1319

1420
== Notice
15-
Copyright (c) 2020 Intel Corporation. All rights reserved.
1621

17-
== Status
22+
[%hardbreaks]
23+
Copyright (C) 2022 Intel Corporation. All rights reserved.
24+
25+
Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
26+
of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by
27+
permission by Khronos.
1828

19-
Draft
2029

21-
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.
30+
== Contact
2231

23-
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.
32+
To report problems with this extension, please open a new issue at:
2433

25-
== Version
34+
https://github.com/intel/llvm/issues
2635

27-
Built On: {docdate} +
28-
Revision: 2
2936

3037
== Dependencies
3138

32-
This extension is written against the SYCL 2020 specification, Revision 3.
39+
This extension is written against the SYCL 2020 revision 6 specification. All
40+
references below to the "core SYCL specification" or to section numbers in the
41+
SYCL specification refer to that revision.
42+
43+
44+
== Status
45+
46+
This extension is implemented and fully supported by {dpcpp}.
47+
48+
49+
== Overview
50+
51+
This extension adds two new address spaces: device and host that are subsets of
52+
the global address space.
53+
New interfaces for `multi_ptr` are added for each of these address spaces.
54+
55+
The goal of this division of the global address space is to enable users to
56+
explicitly tell the compiler which address space a pointer resides in for the
57+
purposes of enabling optimization.
58+
While automatic address space inference is often possible for accessors, it is
59+
harder for USM pointers as it requires inter-procedural optimization with the
60+
host code.
61+
This additional information can be particularly beneficial on FPGA targets where
62+
knowing that a pointer only ever accesses host or device memory can allow
63+
compilers to produce more area efficient memory-accessing hardware.
64+
3365

34-
If SPIR-V is used by the implementation, this extension also requires support for the SPV_INTEL_usm_storage_classes SPIR-V extension.
66+
== Specification
3567

36-
== Feature Test Macro
68+
=== Feature test macro
3769

3870
This extension provides a feature-test macro as described in the core SYCL
39-
specification section 6.3.3 "Feature test macros". Therefore, an
40-
implementation supporting this extension must predefine the macro
41-
`SYCL_EXT_INTEL_USM_ADDRESS_SPACES` to one of the values defined in the table below.
42-
Applications can test for the existence of this macro to determine if the
43-
implementation supports this feature, or applications can test the macro's
44-
value to determine which of the extension's APIs the implementation supports.
71+
specification. An implementation supporting this extension must predefine the
72+
macro `SYCL_EXT_INTEL_USM_ADDRESS_SPACES` to one of the values defined in the table
73+
below. Applications can test for the existence of this macro to determine if
74+
the implementation supports this feature, or applications can test the macro's
75+
value to determine which of the extension's features the implementation
76+
supports.
4577

4678
[%header,cols="1,5"]
4779
|===
48-
|Value |Description
49-
|1 |Initial extension version. Base features are supported.
50-
|===
80+
|Value
81+
|Description
5182

52-
== Overview
83+
|1
84+
|Initial version of this extension.
5385

54-
This extension adds two new address spaces: device and host that are subsets of the global address space.
55-
New specializations of multi_ptr are added for each of these address spaces.
86+
|2
87+
|Adds `sycl::ext::intel::host_ptr`, `sycl::ext::intel::raw_host_ptr`,
88+
`sycl::ext::intel::decorated_host_ptr`, `sycl::ext::intel::device_ptr`,
89+
`sycl::ext::intel::raw_device_ptr` and `sycl::ext::intel::decorated_device_ptr`.
90+
`sycl::host_ptr` and `sycl::device_ptr` are deprecated.
91+
|===
5692

57-
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.
58-
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.
59-
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.
93+
== Modifications to SYCL 2020
6094

61-
== Modifications to the SYCL Specification, Version 2020 revision 3
95+
The following sections contain the related changes and additions to the SYCL
96+
2020 specification relating to this extension.
6297

63-
=== Section 3.8.2 SYCL Device Memory Model
98+
=== SYCL Device Memory Model
6499

65100
Add to the end of the definition of global memory:
66-
Global memory is a virtual address space which overlaps the device and host address spaces.
101+
Global memory is a virtual address space which overlaps the device and host
102+
address spaces.
67103

68104
Add two new memory regions as follows:
69105

70-
*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.
71-
72-
*Host memory* is a sub-region of global memory. USM pointers allocated with the host alloc type reside in this address space.
106+
*Device memory* is a sub-region of global memory that is not directly accessible
107+
by the host. Buffer accessors and USM allocations whose kind is
108+
`usm::alloc::device` reside in this address space.
73109

74-
=== Section 3.8.2.1 Access to memory
110+
*Host memory* is a sub-region of global memory. USM allocations whose kind is
111+
`usm::alloc::host` reside in this address space.
75112

76-
In the second last paragraph, add sycl::device_ptr and sycl::host_ptr to the list of explicit pointer classes.
77113

78-
=== Section 4.7.7.1 Multi-pointer Class
114+
=== Multi-pointer Class
79115

80-
In the overview of the multi_ptr class replace the address_space enum with the following:
116+
Add the following enumerations to the `access::address_space` enum:
81117
```c++
82-
enum class address_space : int {
83-
global_space,
84-
local_space,
85-
constant_space, // Deprecated in SYCL 2020
86-
private_space,
87-
generic_space,
118+
enum class address_space : /* unspecified */ {
119+
...
88120
ext_intel_global_device_space,
89121
ext_intel_global_host_space
90122
};
91123
```
92124

93-
Add the following new conversion operator:
125+
Add the following new conversion operator to the `multi_ptr` class:
94126
```c++
95127
// Explicit conversion to global_space
96128
// Only available if Space == address_space::ext_intel_global_device_space || Space == address_space::ext_intel_global_host_space
97-
explicit operator multi_ptr<ElementType, access::address_space::global_space>() const;
129+
explicit operator multi_ptr<ElementType, access::address_space::global_space, DecorateAddress>() const;
98130
```
99131

100-
Add a new row to Table 91: Constructors of the SYCL multi_ptr class template, as follows:
132+
Change the `multi_ptr` constructor taking an accessor with `target::device` to
133+
also allow `access::address_space::ext_intel_global_device_space` as follows:
101134

102135
--
103136
[options="header"]
104137
|===
105138
| Constructor | Description
106139
a|
107140
```c++
108-
template<typename ElementType, access::
109-
address_space Space = access::address_space::
110-
ext_intel_global_device_space>
111-
template <int dimensions, access::mode Mode>
112-
multi_ptr(
113-
accessor<ElementType, dimensions, Mode, access::
114-
target::global_buffer>)
115-
``` | Constructs a multi_ptr<ElementType, access::address_space::ext_intel_global_device_space> from an accessor of access::target::global_buffer.
141+
template <int Dimensions, access_mode Mode, access::placeholder IsPlaceholder>
142+
multi_ptr(
143+
accessor<ElementType, Dimensions, Mode, target::device, IsPlaceholder>);
144+
```
145+
| Available only when:
146+
`Space == access::address_space::global_space \|\| Space == access::address_space::ext_intel_global_device_space \|\| Space == access::address_space::generic_space`.
147+
148+
Constructs a `multi_ptr` from an accessor of `target::device`.
149+
150+
This constructor may only be called from within a command.
116151
|===
117152
--
118153

119-
=== Section 4.7.7.2 Explicit Pointer Aliases
120154

121-
Add device_ptr and host_ptr aliases to the list of multi_ptr aliases as follows:
155+
=== Explicit Pointer Aliases
156+
157+
Add `device_ptr` and `host_ptr` aliases to the list of `multi_ptr` aliases as
158+
follows:
122159
```c++
160+
namespace sycl {
161+
162+
// Deprecated.
163+
template<typename ElementType,
164+
access::decorated IsDecorated = access::decorated::legacy>
165+
using device_ptr =
166+
multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
167+
IsDecorated>
168+
169+
// Deprecated.
170+
template<typename ElementType,
171+
access::decorated IsDecorated = access::decorated::legacy>
172+
using host_ptr =
173+
multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
174+
IsDecorated>
175+
176+
namespace ext {
177+
namespace intel {
178+
123179
template<typename ElementType>
124-
using device_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_device_space>
180+
using raw_device_ptr =
181+
multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
182+
access::decorated::no>
125183

126184
template<typename ElementType>
127-
using host_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_host_space>
128-
```
185+
using raw_host_ptr =
186+
multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
187+
access::decorated::no>
129188

130-
== Revision History
189+
template<typename ElementType>
190+
using decorated_device_ptr =
191+
multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
192+
access::decorated::yes>
193+
194+
template<typename ElementType>
195+
using decorated_host_ptr =
196+
multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
197+
access::decorated::yes>
198+
199+
} // namespace intel
200+
} // namespace ext
201+
} // namespace sycl
202+
```
131203

132-
[cols="5,15,15,70"]
133-
[grid="rows"]
134-
[options="header"]
135-
|========================================
136-
|Rev|Date|Author|Changes
137-
|1|2020-06-18|Joe Garvey|Initial public draft
138-
|2|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions
139-
|========================================

‎sycl/include/sycl/access/access.hpp

+11
Original file line numberDiff line numberDiff line change
@@ -337,6 +337,17 @@ template <typename ToT, typename FromT> inline ToT cast_AS(FromT from) {
337337
return reinterpret_cast<ToT>(from);
338338
#endif // defined(__NVPTX__) || defined(__AMDGCN__)
339339
} else
340+
#ifdef __ENABLE_USM_ADDR_SPACE__
341+
if constexpr (FromAS == access::address_space::global_space &&
342+
(ToAS ==
343+
access::address_space::ext_intel_global_device_space ||
344+
ToAS ==
345+
access::address_space::ext_intel_global_host_space)) {
346+
// Casting from global address space to the global device and host address
347+
// spaces is allowed.
348+
return (ToT)from;
349+
} else
350+
#endif // __ENABLE_USM_ADDR_SPACE__
340351
#endif // __SYCL_DEVICE_ONLY__
341352
{
342353
return reinterpret_cast<ToT>(from);

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

+64
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,64 @@
1+
//==-------- usm_pointers.hpp - Extended SYCL pointers classes -------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
#include <sycl/access/access.hpp>
11+
12+
namespace sycl {
13+
__SYCL_INLINE_VER_NAMESPACE(_V1) {
14+
15+
template <typename ElementType, access::address_space Space,
16+
access::decorated DecorateAddress>
17+
class multi_ptr;
18+
19+
namespace ext {
20+
namespace intel {
21+
22+
template <typename ElementType,
23+
access::decorated IsDecorated = access::decorated::legacy>
24+
using device_ptr =
25+
multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
26+
IsDecorated>;
27+
28+
template <typename ElementType,
29+
access::decorated IsDecorated = access::decorated::legacy>
30+
using host_ptr =
31+
multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
32+
IsDecorated>;
33+
34+
// Template specialization aliases for different pointer address spaces.
35+
// The interface exposes non-decorated pointer while keeping the
36+
// address space information internally.
37+
38+
template <typename ElementType>
39+
using raw_device_ptr =
40+
multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
41+
access::decorated::no>;
42+
43+
template <typename ElementType>
44+
using raw_host_ptr =
45+
multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
46+
access::decorated::no>;
47+
48+
// Template specialization aliases for different pointer address spaces.
49+
// The interface exposes decorated pointer.
50+
51+
template <typename ElementType>
52+
using decorated_device_ptr =
53+
multi_ptr<ElementType, access::address_space::ext_intel_global_device_space,
54+
access::decorated::yes>;
55+
56+
template <typename ElementType>
57+
using decorated_host_ptr =
58+
multi_ptr<ElementType, access::address_space::ext_intel_global_host_space,
59+
access::decorated::yes>;
60+
61+
} // namespace intel
62+
} // namespace ext
63+
} // __SYCL_INLINE_VER_NAMESPACE(_V1)
64+
} // namespace sycl

‎sycl/include/sycl/feature_test.hpp.in

+1-1
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
6262
#define SYCL_EXT_INTEL_FPGA_REG 1
6363
#define SYCL_EXT_INTEL_KERNEL_ARGS_RESTRICT 1
6464
#define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1
65-
#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1
65+
#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 2
6666
#define SYCL_EXT_INTEL_RUNTIME_BUFFER_LOCATION 1
6767
#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 3
6868
#define SYCL_EXT_ONEAPI_USM_DEVICE_READ_ONLY 1

0 commit comments

Comments
 (0)