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

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Fri Oct 2 11:30:22 PDT 2020


tra added inline comments.


================
Comment at: clang/test/SemaCUDA/device-var-init.cu:404
 __host__ __device__ void hd_sema() {
   static int x = 42;
 }
----------------
yaxunl wrote:
> tra wrote:
> > yaxunl wrote:
> > > how does this work in device compilation? Is this equivalent to `static __device__ int x = 42`?
> > Correct. 
> so static variable without `__device__/__constant__` attribute in host device function implies `__device__` attribute in device compilation.
> 
> Is this also true in device function? We need Sema and CodeGen tests for these cases.
> 
> Also, can we document these changes? It is easily forgotten.
I think of it as a static variable in a `__device__` function. There should be no host-side shadow for it, which would normally be created for a `__device__` variable.

The tests at the beginning of df_sema() in SemaCUDA/device-var-init.cu already check that static w/o attributes is allowed in `__device__`.
I'll update CodeGen tests to verify that we generate correct code.



Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D88345



More information about the cfe-commits mailing list