Add device code test/Update ext spec with note

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>

sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc

@@ -197,6 +197,11 @@ class device_global {
 
 _T_ is restricted to types that have a trivial destructor. _PropertyListT_ enables properties to be associated with a `device_global`.
 
+[NOTE]
+====
+If _T_ is prepended `const` then in the cuda and hip backends the compiler will use the constant address space for the `device_global` instead of the global address space. This allows for optimizations such as using the Nvidia constant cache, and is equivalent to usage of `__constant__` in the CUDA and HIP native programming languages.
+====
+
 When compiling with {cpp} versions before {cpp}20, _T_ must also have a trivial default constructor. In this case, the allocation of type _T_ for a given `device_global` is zero-initialized on a given device prior to the first access to that `device_global` on that device. For the purposes of this definition an access can be a direct access of the `device_global` in kernel code or a copy to or from that `device_global` enqueued to the given device.
 
 When compiling with {cpp}20 or later, _T_ must have a constructor that can be `constexpr` evaluated, and the parameters to the `device_global` constructor are forwarded to the _T_ constructor. In this case, the allocation of type _T_ for a given `device_global` is initialized on a given device prior to the first access to that `device_global` on that device.

sycl/test/check_device_code/device_global_const.cpp

@@ -0,0 +1,19 @@
+// RUN: %clangxx -fsycl -fsycl-device-only %if cuda %{ -fsycl-targets=nvptx64-nvidia-cuda %} %if hip-amd %{ -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx90a %} -S -emit-llvm %s -o - | FileCheck %s  %if cuda || hip-amd %{ --check-prefixes=CHECK-CONST %}
+
+// Tests that const T device_global uses const address space for cuda/hip
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+device_global<const int> DeviceGlobalVar;
+
+int main() {
+  queue Q;
+  Q.single_task([]() {
+    // CHECK-CONST: (ptr addrspace(4) @DeviceGlobalVar
+    volatile int ReadVal = DeviceGlobalVar;
+  });
+  return 0;
+}