[PATCH] D118153: [CUDA][HIP] Do not treat host var address as constant in device compilation

[Artem Belevich via Phabricator via cfe-commits]

tra added a subscriber: rsmith.
tra added a comment.

@rsmith -- is there anything else we need to worry about when it comes to treating pointers as constant values (or not)?



================
Comment at: clang/lib/AST/ExprConstant.cpp:2227
+            !Var->hasAttr<CUDAConstantAttr>() &&
+            !Var->hasAttr<CUDASharedAttr>() &&
+            !Var->getType()->isCUDADeviceBuiltinSurfaceType() &&
----------------
Does it mean that we currently treat address of a `__shared__` variable as a constant? Looks like we do:  https://godbolt.org/z/eG4vG1rbf

```
__shared__ int s;
__device__ const int *const p2 = &s;

__device__ bool f() { return p2 == &s; }
```

`f()` currently always returns true, even though there will be multiple instances of `s` and we probably should not have allowed to init `p2` to start with.

I think we should always return false for `__shared__` vars.


================
Comment at: clang/test/SemaCUDA/const-var.cu:1-2
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -fsyntax-only -verify
+
----------------
Would it make sense to add a host-side test to illustrate a case where some pointer expressions will be const on the host side, but not on the device side?



================
Comment at: clang/test/SemaCUDA/const-var.cu:17
+const int *const B::p = &a;
+__device__ const int *const B::p2 = &a;
+
----------------
I'd add a comment that global const variables are treated as `__constant__`.  Otherwise no error on taking address of `a` on device side looks wrong.

Makes me wonder what would be result of `B::p == B::p2`? On device side, it should be possible to const evaluate it to `true`.
However, on the host, `B::p` would have a real value, but `B::p2` will be represented by a proxy var.  Off the top of my head I'm not sure whether it will be initialized or left undefined. 

I'd add a codegen test to make sure `B::p == B::p2` does get evaluated consistently on host and device.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D118153/new/

https://reviews.llvm.org/D118153