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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
