Add a section about shadowed variables and integration footer

Author: AlexeySachkov

sycl/doc/design/MappingHostAddressesToDeviceEntities.md

@@ -201,6 +201,19 @@ __sycl_device_global_registration __sycl_device_global_registrar;
 } // namespace sycl::detail
 ```
 
+Examples below are written for the following code snippet:
+
+```
+#include <sycl/sycl.hpp>
+
+static sycl::device_global<int> Foo;
+namespace inner {
+  sycl::device_global<double[2]> Bar;
+} // namespace inner
+
+// ...
+```
+
 The integration footer generated by the compiler contains the definition of the
 constructor, which calls a function in the DPC++ runtime, which registers
 needed mappings:
@@ -213,7 +226,7 @@ __sycl_device_global_registration::__sycl_device_global_registration() noexcept
   __register_uniquely_identifiable_object(
     &::Foo,
     /* same string returned from __builtin_sycl_unique_stable_id(::Foo) */,
-    "specialization_id");
+    "device_global");
   __register_uniquely_identifiable_object(
     &::inner::Bar,
     /* same string returned from __builtin_sycl_unique_stable_id(::inner::Bar) */,
@@ -224,6 +237,67 @@ __sycl_device_global_registration::__sycl_device_global_registration() noexcept
 } // namespace sycl::detail
 ```
 
+### Handling shadowed variables
+
+The example above shows a simple case where the user's device global variables
+can all be uniquely referenced via fully qualified lookup (e.g.
+`::inner::Bar`).  However, it is possible for users to construct applications
+where this is not the case, for example:
+
+```
+sycl::device_global<int> FuBar;
+namespace {
+  sycl::device_global<int> FuBar;
+}
+```
+
+In this example, the `FuBar` variable in the global namespace shadows a
+variable with the same name in the unnamed namespace.  The integration footer
+can reference the variable in the global namespace as `::FuBar`, but there is
+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 shim function that returns a reference to the
+shadowed device global:
+
+```
+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 {
+
+__sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+  __register_uniquely_identifiable_object(
+    &::FuBar,
+    /* same string returned from __builtin_sycl_unique_stable_id(::FuBar) */,
+    "device_global");
+  __register_uniquely_identifiable_object(
+    &::__sycl_detail::__shim_1(),
+    /* same string returned from __builtin_sycl_unique_stable_id(::(unnamed)::FuBar) */,
+    "device_global");
+}
+
+} // namespace sycl::detail
+```
+
+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>
+
 ## Custom host compiler approach
 
 With this approach, we simply schedule a one more pass in the optimization