================
@@ -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:
Thanks for adding this check. I agree that following calls is not practical,
especially when a function is defined in another translation unit.
Could we use `ConstEvaluatedExprVisitor` instead of
`DynamicRecursiveASTVisitor`? It visits only potentially evaluated parts of the
initializer. This avoids rejecting valid cases such as `sizeof(x)` or a
reference to `x` inside an uncalled lambda body.
The check can diagnose direct uses, for example:
```cpp
int *p = &x;
int v = x;
```
It will not diagnose every indirect use:
```cpp
extern int *getX();
int *p = getX();
```
Clang cannot know whether `getX()` accesses a managed variable. I think that
limitation is acceptable, but it is important to document it in
`clang/docs/HIPSupport.md`. The documentation should say that Clang diagnoses
some direct uses but does not follow function calls, constructors, destructors,
function pointers, or definitions in other translation units. Accessing a
managed variable during static or thread-local initialization or destruction is
still undefined behavior, even if Clang does not emit a diagnostic.
https://github.com/llvm/llvm-project/pull/217047
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
