Skip to content

Commit 3c0d2a0

Browse files
committed
Clarify restrictions on device global variables
Clarify the restrictions for declaring variables of type `device_global`. These restrictions are similar to those we have already for `specialization_id` variables. The restriction against arrays of `device_global` variables is not strictly necessary, but it would take extra effort to implement and we do not have a compelling use case. We could lift the restriction later if we find a need.
1 parent 7e11e48 commit 3c0d2a0

File tree

1 file changed

+88
-2
lines changed

1 file changed

+88
-2
lines changed

sycl/doc/extensions/DeviceGlobal/SYCL_INTEL_device_global.asciidoc

Lines changed: 88 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -184,8 +184,6 @@ Since _T_ is restricted to types with trivial default constructors in this versi
184184

185185
The allocation of type _T_ is zero-initialized on each device prior to the first access to that `device_global` variable.
186186

187-
`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.
188-
189187
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:
190188

191189
[source,c++]
@@ -398,6 +396,94 @@ Available only if `sycl::is_property_of_v<propertyT, sycl::ext::oneapi::device_g
398396

399397
|===
400398

399+
=== Restrictions on creating device global objects
400+
401+
There are restrictions on how the application can create objects of type
402+
`device_global`. Applications that violate these restrictions are ill-formed.
403+
404+
* The application may declare a variable of type `device_global` in the
405+
following ways:
406+
+
407+
--
408+
** As a variable at namespace scope, or
409+
** As a static member variable, but only if the member variable is publicly
410+
accessible from namespace scope.
411+
--
412+
+
413+
The application must not create an object of type `device_global` in any other
414+
way. (E.g. variables with automatic storage duration or objects created via
415+
`new` are not allowed.)
416+
417+
* The `device_global` variable must not itself be an array. The underlying
418+
type _T_ may be an array type, but the `device_global` variable itself must
419+
not be an array.
420+
421+
* The `device_global` variable must not be shadowed by another identifier _X_
422+
which has the same name and is declared in an inline namespace, such that the
423+
`device_global` variable is no longer accessible after the declaration of
424+
_X_.
425+
426+
* If the `device_global` variable is declared in a namespace, none of the
427+
enclosing namespace names _N_ may be shadowed by another identifier _X_ which
428+
has the same name as _N_ and is declared in an inline namespace, such that
429+
_N_ is no longer accessible after the declaration of _X_.
430+
431+
[NOTE]
432+
====
433+
The expectation is that some implementations may conceptually insert code at
434+
the end of a translation unit which references each `device_global` variable
435+
that is declared in that translation unit. The restrictions listed above make
436+
this possible by ensuring that these variables are accessible at the end of the
437+
translation unit.
438+
====
439+
440+
The following example illustrates some of these restrictions:
441+
442+
[source, c++]
443+
----
444+
#include <sycl/sycl.hpp>
445+
using namespace sycl::ext::oneapi;
446+
447+
device_global<int> a; // OK
448+
static device_global<int> b; // OK
449+
inline device_global<int> c; // OK
450+
451+
struct Foo {
452+
static device_global<int> d; // OK
453+
};
454+
device_global<int> Foo::d;
455+
456+
struct Bar {
457+
device_global<int> e; // ILLEGAL: non-static member variable not
458+
}; // allowed
459+
460+
struct Baz {
461+
private:
462+
static device_global<int> f; // ILLEGAL: not publicly accessible from
463+
}; // namespace scope
464+
device_global<int> Baz::f;
465+
466+
device_global<int[4]> g; // OK
467+
device_global<int> h[4]; // ILLEGAL: array of "device_global" not
468+
// allowed
469+
470+
device_global<int> same_name; // OK
471+
namespace foo {
472+
device_global<int> same_name; // OK
473+
}
474+
namespace {
475+
device_global<int> same_name; // OK
476+
}
477+
inline namespace other {
478+
device_global<int> same_name; // ILLEGAL: shadows "device_global" variable
479+
} // with same name in enclosing namespace scope
480+
inline namespace {
481+
namespace foo { // ILLEGAL: namespace name shadows "::foo"
482+
} // namespace which contains "device_global"
483+
// variable.
484+
}
485+
----
486+
401487
=== Relax language restrictions for SYCL device functions
402488

403489
SYCL 2020 restrictions must be relaxed to allow `device_global` to be used within

0 commit comments

Comments
 (0)