[PATCH] D88345: [CUDA] Allow local `static const {__constant__, __device__}` variables.

[Artem Belevich via Phabricator via cfe-commits]

tra added inline comments.


================
Comment at: clang/include/clang/Basic/DiagnosticSemaKinds.td:8163
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
-def err_cuda_nonglobal_constant : Error<"__constant__ variables must be global">;
+def err_cuda_nonstatic_constdev: Error<"__constant__ and __device__ are not allowed on non-static local variables">;
 def err_cuda_ovl_target : Error<
----------------
jlebar wrote:
> `__device__` is not allowed on non-static function-scope variables?
> 
> This appears to be more restrictive than we were before.  I want to check, are we OK with the possibility that this will break user code?  https://gcc.godbolt.org/z/Y85GKe work with clang, though not with nvcc.
> 
> I notice that we even allow `__device__ int x;` in `__host__ __device__` functions, which is...questionable.  :)  https://gcc.godbolt.org/z/GjjMGx
> 
> I'm OK matching the nvcc behavior here and accepting user breakage so long as we're intentional about it.  Possibly should be called out in relnotes?
It appears to have been an oversight. AFAICT, we just ignored the `__device__` attribute of the local vars that sneaked past the `isStaticLocal()` check.



================
Comment at: clang/lib/Sema/SemaDeclAttr.cpp:4353
   const auto *VD = cast<VarDecl>(D);
-  if (!VD->hasGlobalStorage()) {
-    S.Diag(AL.getLoc(), diag::err_cuda_nonglobal_constant);
+  if (VD->hasLocalStorage()) {
+    S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
----------------
jlebar wrote:
> So just to check, in our new world, `__constant__` variables don't have to be const.  That matches nvcc, fine.
No. IMO, it's similar to explicitly putting a non-const variable into a `.rodata` section -- inadvisable, probably not very useful, and possibly implementation-defined, but not illegal.




================
Comment at: clang/test/CodeGenCUDA/static-device-var-no-rdc.cu:84
+  const static __constant__ int local_static_constant = 42;
+  const static __device__ int local_static_device = 43;
   a[0] = x;
----------------
tra wrote:
> yaxunl wrote:
> > what happens to a const static device or constant var with non-trivial constructor? can we have a test for that?
> I believe constructor trivialness check is orthogonal and will still be applied.
> I'll add a test.
This is checked in `SemaCUDA/device-var-init.cu`.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D88345