[clang] [clang][HIP] Do not treat address of managed variable as a constant expression (PR #217047)

[Yaxun Liu via cfe-commits]

================
@@ -161,6 +161,26 @@ __device__ __host__ int load4() {
   return ex;
 }
 
+namespace gh198079 {
+__managed__ int x = 0;
+struct S {
+    int *p;
+};
+__attribute__((device)) void f() {
+  S s{&x};
+}
+// DEV-LABEL: define {{.*}}@{{.*}}gh198079{{.*}}f{{.*}}()
+// DEV: %p = getelementptr inbounds nuw %"struct.gh198079::S", ptr %s.ascast, i32 0, i32 0
+// DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @_ZN8gh1980791xE, align 4
+// DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr
+// DEV: store ptr %0, ptr %p, align 8
+int *hostglob = &x;
+
----------------
yxsamliu wrote:

Could we diagnose this in Sema instead of adding a positive CodeGen test? The CUDA programming guide does not allow the address or value of a managed variable in static or thread-local initialization. With the current HIP initialization order, this initializer may also run before the managed variable is registered, leaving `hostglob` null.

https://github.com/llvm/llvm-project/pull/217047