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