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

[Justin Lebar via Phabricator via cfe-commits]

jlebar accepted this revision.
jlebar added inline comments.
This revision is now accepted and ready to land.


================
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<
----------------
`__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?


================
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);
----------------
So just to check, in our new world, `__constant__` variables don't have to be const.  That matches nvcc, fine.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D88345