Skip to content

Commit

Permalink
[SYCL][DOC] Update device global design shadowing (#5525)
Browse files Browse the repository at this point in the history
Update the design for device global variables for variables that are
"shadowed" in an unnamed namespace.  This change was suggested by
@Fznamznon because it allows us to reuse existing machinery we have
for handling shadowed `specialization_id` variables.
  • Loading branch information
gmlueck authored Feb 11, 2022
1 parent f376826 commit e211d73
Showing 1 changed file with 16 additions and 14 deletions.
30 changes: 16 additions & 14 deletions sycl/doc/DeviceGlobal.md
Original file line number Diff line number Diff line change
Expand Up @@ -323,37 +323,39 @@ no way to reference the variable in the unnamed namespace using fully qualified
lookup.

Such programs are still legal, though. The integration footer can support
cases like this by defining a temporary variable that holds the address of the
cases like this by defining a shim function that returns a reference to the
shadowed device global:

```
namespace {
const void *__sycl_UNIQUE_STRING = &FuBar; // References 'FuBar' in the
// unnamed namespace
namespace __sycl_detail {
static constexpr decltype(FuBar) &__shim_1() {
return FuBar; // References 'FuBar' in the unnamed namespace
}
} // namespace __sycl_detail
} // namespace (unnamed)
namespace sycl::detail {
namespace {
__sycl_device_global_registration::__sycl_device_global_registration() noexcept {
device_global_map::add(&::FuBar,
/* same string returned from __builtin_sycl_unique_stable_id(::FuBar) */);
device_global_map::add(::__sycl_UNIQUE_STRING,
device_global_map::add(&::__sycl_detail::__shim_1(),
/* same string returned from __builtin_sycl_unique_stable_id(::(unnamed)::FuBar) */);
}
} // namespace (unnamed)
} // namespace sycl::detail
```

The `__sycl_UNIQUE_STRING` variable is defined in the same namespace as the
second `FuBar` device global, so it can reference the variable through
unqualified name lookup. Furthermore, the name of the temporary variable
(`__sycl_UNIQUE_STRING`) is globally unique, so it is guaranteed not to be
shadowed by any other name in the translation unit. This problem with variable
shadowing is also a problem for the integration footer we use for
specialization constants. See the [specialization constant design document][5]
for more details on this topic.
The `__shim_1()` function is defined in the same namespace as the second
`FuBar` device global, so it can reference the variable through unqualified
name lookup. Furthermore, the name of the shim function is globally unique, so
it is guaranteed not to be shadowed by any other name in the translation unit.
This problem with variable shadowing is also a problem for the integration
footer we use for specialization constants. See the [specialization constant
design document][5] for more details on this topic.

[5]: <SpecializationConstants.md>

Expand Down

0 comments on commit e211d73

Please sign in to comment.