From e211d733989e4e6680585a68f0fc4f6c236aaa99 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Fri, 11 Feb 2022 01:42:37 -0500 Subject: [PATCH] [SYCL][DOC] Update device global design shadowing (#5525) 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. --- sycl/doc/DeviceGlobal.md | 30 ++++++++++++++++-------------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/sycl/doc/DeviceGlobal.md b/sycl/doc/DeviceGlobal.md index dcf74f6489b9e..975d5dfeb1720 100644 --- a/sycl/doc/DeviceGlobal.md +++ b/sycl/doc/DeviceGlobal.md @@ -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]: