[PATCH] D73979: [HIP] Allow non-incomplete array type for extern shared var

[Jon Chesterfield via Phabricator via cfe-commits]

JonChesterfield added a comment.

In D73979#1857736 <https://reviews.llvm.org/D73979#1857736>, @yaxunl wrote:

> BTW this is requested by HIP users, who have similar code for CUDA and HIP. They found it surprised that nvcc allows it but hip-clang does not.


I think I'm one of the HIP users here, but the above change is not what I was hoping for.

I'd like:

  __shared__ int x;
  __shared__ int y;
  __device__ void foo()
  {
    assert(&x != &y);
    x = 2 * y;
  }

to compile and behave as it does on cuda, i.e. the 'x' variable gets allocated in __shared__ memory for each kernel which accesses it, and so does 'y'.

The 'extern __shared__' feature where nvcc builds a union out of all the things it sees and the user indexes into it at runtime is totally unappealing. That cuda uses the 'extern' keyword to opt into this magic union also seems undesirable.


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

https://reviews.llvm.org/D73979