Skip to content

Commit f4d0feb

Browse files
author
Michael Kinsner
committed
Incorporate suggestions from @artemrad
1 parent fc9fbae commit f4d0feb

File tree

1 file changed

+1
-22
lines changed

1 file changed

+1
-22
lines changed

sycl/doc/extensions/DeviceGlobal/SYCL_INTEL_device_global.asciidoc

Lines changed: 1 addition & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -186,27 +186,6 @@ The allocation of type _T_ is zero-initialized on each device prior to the first
186186

187187
`device_global` may only be declared with static storage duration at namespace scope or class scope. If a `device_global` is declared with any other storage duration or scope, the program is ill-formed.
188188

189-
The example below creates two global namespace scope `device_global` objects named `dm1` and `dm2`. `dm1` contains one object of type `MyClass` on each device, and the `device_global` object has external linkage. `dm2` contains an array of four integers on each device, and the `device_global` object has internal linkage. In both cases, the `MyClass` and `int[4]` allocations on each device are zero-initialized before any non-initialization accesses occur.
190-
191-
[source,c++]
192-
----
193-
using namespace sycl::ext::oneapi;
194-
195-
device_global<MyClass> dm1;
196-
static device_global<int[4]> dm2;
197-
198-
int main () {
199-
sycl::queue Q;
200-
Q.submit([&](sycl::handler& h) {
201-
h.single_task([=]() {
202-
int x = 5;
203-
if (dm1.get().flag)
204-
x = dm2[0];
205-
});
206-
});
207-
}
208-
----
209-
210189
Properties may be specified for a `device_global` to provide semantic modification or optimization hint information to the compiler. Specific properties are defined in other extensions, but example uses of a property (with a "no copy" attribute described by another extension) are:
211190

212191
[source,c++]
@@ -763,7 +742,7 @@ int main () {
763742

764743
`device_global` prioritizes usability over simplicity of implementation, and therefore adds requirements such as (1) that contents and addresses of the allocation on each device remain stable across changes to specialization constant values, and (2) that the allocation be accessible across `device_image` on the same device. These requirements mean that the semantics of `device_global` do not match the semantics of SPIR-V module scope variables, and therefore may not be implementable exclusively using the SPIR-V feature in existing SPIR-V consuming implementations.
765744

766-
Also note that there are no restrictions on passing (and subsequent dereferencing) of pointers obtained on a device from a `device_global`, between kernels on a device, including through storage to memory.
745+
Also note that there are no restrictions on passing (and subsequent dereferencing) of pointers obtained on a device from a `device_global`, between kernels on the same device, including through storage to memory.
767746

768747
== Issues
769748

0 commit comments

Comments
 (0)