[PATCH] D103225: [AMDGPU] Replace non-kernel function uses of LDS globals by pointers.

[Mahesha S via Phabricator via llvm-commits]

hsmhsm marked an inline comment as done.
hsmhsm added a comment.

In D103225#2805419 <https://reviews.llvm.org/D103225#2805419>, @b-sumner wrote:

> Two approaches for limiting the stores to lane 0 of each wave:
>
> 1. Write 1 to exec mask, store, and write -1 to exec mask.  This works since the exec mask at the start of the wave when this happens is -1
> 2. Check for lane == 0 and branch.  The lane can be computed by a) wave64: __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u)) b) wave32: __builtin_amdgcn_mbcnt_lo(~0u, 0u)

OK, let's see, which one is more feasible from the implementation point of view.



================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp:152
+
+    // FIXME: Any other scenarios which disqualify LDS from replacement?
+  }
----------------
rampitec wrote:
> JonChesterfield wrote:
> > Don't replace it when the variable is used by all (possibly most) kernels that can access any LDS. For example, when the variable is part of a language runtime that has been linked into all kernels.
> > 
> > Such a variable will still be allocated in all the kernels so putting it behind a pointer wastes two bytes of LDS everywhere and introduces code to do the indirection with zero benefit.
> > Don't replace it when the variable is used by all (possibly most) kernels that can access any LDS. For example, when the variable is part of a language runtime that has been linked into all kernels.
> > 
> > Such a variable will still be allocated in all the kernels so putting it behind a pointer wastes two bytes of LDS everywhere and introduces code to do the indirection with zero benefit.
> 
> Sound like a good TODO at the very least.
It is in my TODO list. It requires some logical thinking, before carefully handling this case. That is the reason, I had not yet replied to Jon's comment.


================
Comment at: llvm/lib/Target/AMDGPU/AMDGPUReplaceLDSUseWithPointer.cpp:150
+    LDSToNonKernels[GV] = AMDGPU::collectNonKernelAccessorsOfLDS(GV);
+    return LDSToNonKernels[GV].empty();
+
----------------
rampitec wrote:
> If it is empty it probably will not be returned by findVariablesToLower()?
That is true.  But, there is a catch. Assume the below case. Here "@lds.1" used within non-kernel function scope, but also it is used as initializer to global @gptr.1. So, we cannot pointer replace it.  In this case, collectNonKernelAccessorsOfLDS() returns empty set.

```
@lds.1 = internal unnamed_addr addrspace(3) global [1 x i8] undef, align 1

@gptr.1 = addrspace(1) global i64* addrspacecast ([1 x i8] addrspace(3)* @lds.1 to i64*), align 8

define void @f0() {
  %bc = bitcast [1 x i8] addrspace(3)* @lds.1 to i8 addrspace(3)*
  store i8 1, i8 addrspace(3)* %bc, align 1
  ret void
}

define amdgpu_kernel void @k0() {  
  call void @f0()
  ret void
}
```


================
Comment at: llvm/test/CodeGen/AMDGPU/replace-lds-by-ptr-call-selected_functions.ll:41
+; CHECK:   %0 = load i16, i16 addrspace(3)* @lds_used_within_function_2.ptr, align 2
+; CHECK:   %1 = getelementptr i8, i8 addrspace(3)* null, i16 %0
+; CHECK:   %2 = bitcast i8 addrspace(3)* %1 to [2 x i32] addrspace(3)*
----------------
rampitec wrote:
> We had problems with null before. Instruction selection will happily allocate it overlapping with first non-null variable. That's how AMDGPULowerModuleLDSPass turned to use struct instead of null. Did you check we are not introducing this again in the final ISA?
The null represent the base address of LDS memory.  My understanding is that, earlier, instruction selection had gone for toss because of the bug where we were not actually allocating memory at all for llvm.amdgcn.module.lds, but we were accessing it at address 0. And, instruction selection was allocating address 0 for some other lds (since this address was not occupied), so there was an issue. 

But, nevertheless, let me look into ISA and see if we are fine or any issue because of null.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D103225