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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In D73979#1907965 <https://reviews.llvm.org/D73979#1907965>, @JonChesterfield wrote:

> 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.


Clang emits correct IR for this code. If you use x and y in a kernel directly, amdgcn backend can generate correct ISA where &x and &y are different. What the backend does is to accumulate all shared memory and get the total shared memory usage and assign address to different shared variables. Therefore x and y get different addresses.

Currently amdgcn backend emits a diagnostic message if shared variable is used in non-kernel function:

https://github.com/llvm/llvm-project/blob/6085593c128e91fd7db998c5441ebe120c7e4f04/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp#L1232

https://github.com/llvm/llvm-project/blob/3fda1fde8f7bdf3b90d8700f5a386f63409b4313/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp#L1952

This is unreasonable if the backend is able to calculate total shared memory usage, so this is a bug. With this bug fixed, you should be able to use shared variables in device functions.


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

https://reviews.llvm.org/D73979