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

[Artem Belevich via Phabricator via cfe-commits]

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