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