================
@@ -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;
+
----------------
Fznamznon wrote:
We can try. I do have a couple of notes though. The first is that NVCC itself
(despite the programming guide saying it is not allowed), actually accepts that
case https://godbolt.org/z/xPahqYeba (unlike the constexpr case I'm disallowing
in this PR). But if that never works at runtime, we probably won't run into any
compatibility issues, right?
The second concern I have is that I don't think we are able to diagnose
nonsense like this in Sema since any callgraph analysis will negatively impact
compile time performance:
```
__managed__ int b;
int *foo() {
return &b;
}
int *p = foo();
```
note foo could call anything else that ends up returning address of b.
https://github.com/llvm/llvm-project/pull/217047
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
