[PATCH] D82496: [amdgpu] Add codegen support for HIP dynamic shared memory.

[Michael Liao via Phabricator via llvm-commits]

hliao created this revision.
hliao added reviewers: arsenm, yaxunl, kpyzhov, b-sumner.
Herald added subscribers: llvm-commits, kerbowa, hiraditya, t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely, kzhuravl.
Herald added a project: LLVM.

- HIP uses an unsized extern array `extern __shared__ T s[]` to declare the dynamic shared memory, which size is not known at the compile time.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D82496

Files:
  llvm/lib/Target/AMDGPU/SIISelLowering.cpp
  llvm/test/CodeGen/AMDGPU/hip.extern.shared.array.ll


Index: llvm/test/CodeGen/AMDGPU/hip.extern.shared.array.ll
===================================================================
--- /dev/null
+++ llvm/test/CodeGen/AMDGPU/hip.extern.shared.array.ll
@@ -0,0 +1,58 @@
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck %s
+
+ at lds0 = addrspace(3) global [512 x float] undef, align 4
+ at lds1 = addrspace(3) global [256 x float] undef, align 4
+ at large = addrspace(3) global [4096 x float] undef, align 4
+
+ at dynamic_shared = external addrspace(3) global [0 x float], align 4
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_0:
+; CHECK: v_add_u32_e32 v{{[0-9]+}}, 0x800, v{{[0-9]+}}
+define amdgpu_kernel void @dynamic_shared_array_0(float addrspace(1)* %out) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %arrayidx0 = getelementptr inbounds [512 x float], [512 x float] addrspace(3)* @lds0, i32 0, i32 %tid.x
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared, i32 0, i32 %tid.x
+  store float %val0, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_1:
+; CHECK: v_mov_b32_e32 v{{[0-9]+}}, 0xc00
+define amdgpu_kernel void @dynamic_shared_array_1(float addrspace(1)* %out, i32 %cond) {
+entry:
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %idx.0 = add nsw i32 %tid.x, 64
+  %tmp = icmp eq i32 %cond, 0
+  br i1 %tmp, label %if, label %else
+
+if:                                               ; preds = %entry
+  %arrayidx0 = getelementptr inbounds [512 x float], [512 x float] addrspace(3)* @lds0, i32 0, i32 %idx.0
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  br label %endif
+
+else:                                             ; preds = %entry
+  %arrayidx1 = getelementptr inbounds [256 x float], [256 x float] addrspace(3)* @lds1, i32 0, i32 %idx.0
+  %val1 = load float, float addrspace(3)* %arrayidx1, align 4
+  br label %endif
+
+endif:                                            ; preds = %else, %if
+  %val = phi float [ %val0, %if ], [ %val1, %else ]
+  %arrayidx = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared, i32 0, i32 %tid.x
+  store float %val, float addrspace(3)* %arrayidx, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}dynamic_shared_array_2:
+; CHECK: v_mov_b32_e32 v{{[0-9]+}}, 0x4000
+define amdgpu_kernel void @dynamic_shared_array_2(i32 %idx) {
+  %tid.x = tail call i32 @llvm.amdgcn.workitem.id.x()
+  %vidx = add i32 %tid.x, %idx
+  %arrayidx0 = getelementptr inbounds [4096 x float], [4096 x float] addrspace(3)* @large, i32 0, i32 %vidx
+  %val0 = load float, float addrspace(3)* %arrayidx0, align 4
+  %arrayidx1 = getelementptr inbounds [0 x float], [0 x float] addrspace(3)* @dynamic_shared, i32 0, i32 %tid.x
+  store float %val0, float addrspace(3)* %arrayidx1, align 4
+  ret void
+}
+
+declare i32 @llvm.amdgcn.workitem.id.x()
Index: llvm/lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -5510,15 +5510,25 @@
                                              SDValue Op,
                                              SelectionDAG &DAG) const {
   GlobalAddressSDNode *GSD = cast<GlobalAddressSDNode>(Op);
+  SDLoc DL(GSD);
+  EVT PtrVT = Op.getValueType();
+
   const GlobalValue *GV = GSD->getGlobal();
   if ((GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS &&
        shouldUseLDSConstAddress(GV)) ||
       GSD->getAddressSpace() == AMDGPUAS::REGION_ADDRESS ||
-      GSD->getAddressSpace() == AMDGPUAS::PRIVATE_ADDRESS)
+      GSD->getAddressSpace() == AMDGPUAS::PRIVATE_ADDRESS) {
+    if (GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS &&
+        GV->hasExternalLinkage()) {
+      ArrayType *ATy = dyn_cast<ArrayType>(GV->getValueType());
+      if (ATy && ATy->getNumElements() == 0) {
+        assert(PtrVT == MVT::i32 && "32-bit pointer is expected.");
+        return SDValue(
+            DAG.getMachineNode(AMDGPU::GET_GROUPSTATICSIZE, DL, PtrVT), 0);
+      }
+    }
     return AMDGPUTargetLowering::LowerGlobalAddress(MFI, Op, DAG);
-
-  SDLoc DL(GSD);
-  EVT PtrVT = Op.getValueType();
+  }
 
   if (GSD->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
     SDValue GA = DAG.getTargetGlobalAddress(GV, DL, MVT::i32, GSD->getOffset(),


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D82496.273132.patch
Type: text/x-patch
Size: 4455 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20200624/8278b426/attachment.bin>