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