tra added a comment.
In D88345#2296118 <https://reviews.llvm.org/D88345#2296118>, @jlebar wrote:
> wha... As you know, `const` doesn't mean anything, that can be const-casted
> away. And then you'll be able to observe that this nominally-static variable
> is just a normal variable.
Yes, I'm aware of that. It is irrelevant in this case.
The point is that `__constant__` and `__device__` are not allowed for the local
`static` variables. I believe originally only `__shared__` was allowed for
local static vars is because `__shared__` is already implicitly static, while
`__constant__` and `__device__` are not.
On the other hand, nothing stops us **giving** `static __constant__` of
`__device__` variables static storage class and that's what NVIDIA has
apparently done, though without updating the docs.
The checks on the constructor still apply -- trivialness is enforced, `const`
must have an initializer, etc.
Looks like the const-ness check should not be there, either. NVCC allows
non-const statics. We may get rid of this check altogether and allow everything.
> Since this doesn't make sense and contradicts their documentation, I'm
> tempted to say this should only apply to the nvidia headers. Is that
> technically possible? And then we file a bug against nvidia and/or ask Bryce?
IMO the relaxation is sensible. A `static` var is just a global var in an odd
namespace and we do support that.
E.g. these produce the same PTX for `c` (modulo mangling of the names):
namespace a {
__constant__ int c = 4;
}
__device__ void foo() {
static __constant__ int c = 4;
}
================
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;
----------------
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.
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
