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