Skip to content

[SYCL][DOC] Update device global design shadowing #5525

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Feb 11, 2022
Merged
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
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.
Comment on lines +357 to +358
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is a bit confusing that SpecializationConstants.md document doesn't actually describe this topic (only #3331 does). Tagging @AlexeySachkov .


[5]: <SpecializationConstants.md>

Expand Down