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

Justin Lebar via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Mon Sep 28 10:46:20 PDT 2020


jlebar added a comment.

OK, backing up, what are the semantics of `static` on `__constant__`, `__device__`, and `__shared__`?

- My understanding is that `__shared__` behaves the same whether or not it's static.  It's not equivalent to `namespace a { __shared__ int c = 4; }`, because that's illegal.
- Does `__constant__` behave the same whether or not it's static?  A static `__constant__` is equivalent to `namespace a { __constant__ int c = 4; }`, and a non-static `__constant__` is *also* equivalent to that?
- And `__device__` does not behave the same whether or not it's static?  In function scope `__device__ int x;` is a variable local to the thread.  Whereas in global scope `__device__ int x;` is a global variable that lives in device memory (?).  In function scope `static __device__ int x;` is equivalent to `static int x;` which is equivalent to `int x;` in namespace scope?

Should we mandate that you initialize `static __constant__` variables in function scope?  That is, if you write `static __constant__ int x;` in a function, then x is always uninitialized (right)?  You should do `static __constant__ int x = 42;`?


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