[SYCL] Make internal device_global pointer mutable

The implementation of device_global allows it to be const. However, for
cases where the device_global is shared between programs, the value of
the device_global is defined through a USM pointer inside the
device_global. Since this pointer is seen as const when the
device_global is const-qualified, the compiler may chose to optimize it
out, potentially causing invalid accesses on device. To avoid this, this
commit marks the underlying pointer member as mutable which prevents
this potential optimization.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/ext/oneapi/device_global/device_global.hpp

@@ -44,7 +44,7 @@ template <typename T, typename PropertyListT, typename = void>
 class device_global_base {
 protected:
   using pointer_t = typename decorated_global_ptr<T>::pointer;
-  pointer_t usmptr{};
+  mutable pointer_t usmptr{};
   pointer_t get_ptr() noexcept { return usmptr; }
   const pointer_t get_ptr() const noexcept { return usmptr; }
 

sycl/test/check_device_code/device_global_ptr_use.cpp

@@ -0,0 +1,20 @@
+// RUN: %clangxx -fsycl -c -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
+
+// Tests that the underlying pointer in a const-qualified shared device_global
+// is not optimized out during access.
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+const device_global<int> DeviceGlobalVar;
+
+int main() {
+  queue Q;
+  Q.single_task([]() {
+    // CHECK: load i32 {{.*}}({{.*}}* @_ZL15DeviceGlobalVar, i64 0, i32 0)
+    volatile int ReadVal = DeviceGlobalVar;
+  });
+  return 0;
+}