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



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
  • [PATCH] D73979: [HIP] All... Jon Chesterfield via Phabricator via cfe-commits

Reply via email to