================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
