tra added a comment.
In D88345#2298688 <https://reviews.llvm.org/D88345#2298688>, @jlebar wrote:
> 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.
Yes. `__shared__` is an odd duck. It is implicitly static, so whether we
explicitly specify `static` makes no difference.
We're not changing anything about how it's implemented.
> - 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?
No. `__constant__` is not allowed on non-static local variables as it can't be
allocated on stack.
> - And `__device__` does not behave the same whether or not it's static?
Correct.
> In function scope `__device__ int x;` is a variable local to the thread.
Correct. `__device__` in a device function is effectively a no-op and can be
placed on stack as a regular local variable.
> Whereas in global scope `__device__ int x;` is a global variable that lives
> in device memory (?).
Correct.
> In function scope `static __device__ int x;` is equivalent to `static int x;`
> which is equivalent to `int x;` in namespace scope?
Yes, assuming you mean a `__device__` function and `__device__ int x;` in the
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;`?
No. Accoring to PTX spec: `Variables in .const and .global state spaces are
initialized to zero by default.`
Those are the address spaces `__constant__` and `__device__` variables map to.
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#state-spaces
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
