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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl marked 4 inline comments as done.
yaxunl added inline comments.


================
Comment at: clang/lib/AST/ExprConstant.cpp:2227
+            !Var->hasAttr<CUDAConstantAttr>() &&
+            !Var->hasAttr<CUDASharedAttr>() &&
+            !Var->getType()->isCUDADeviceBuiltinSurfaceType() &&
----------------
tra wrote:
> 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.
Will do. This is similar to a TLS var in C++, which is treated as non-constant by clang.


================
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
+
----------------
tra wrote:
> 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?
> 
will do.


================
Comment at: clang/test/SemaCUDA/const-var.cu:17
+const int *const B::p = &a;
+__device__ const int *const B::p2 = &a;
+
----------------
tra wrote:
> 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.
will do.

B::p == B::p2 evaluates to 1 on both device and host side.


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

https://reviews.llvm.org/D118153