[AMDGPU] Legalize initialized LDS variables

[cdevadas]

Summary

We don't allow an initializer for LDS variables
and there is an early abort during instruction
selection. This patch legalizes them by ignoring
the init values. During assembly emission, proper
error reporting already exists for such instances.

[arsenm]

llvm/test/CodeGen/AMDGPU/legalize-lds-initializer.ll

2

Do we not already have a test for this? I would prefer to run end-to-end to ensure the diagnostic is in fact emitted for the initializer in the end

[cdevadas]

Addressed review comments.

[cdevadas]

Ping

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

1381–1383

Should explicitly note this will error later

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

2423–2428

Ditto

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

if (!MFI->isModuleEntryFunction() &&
       return DAG.getUNDEF(Op.getValueType());
     }
 
     // XXX: What does the value of G->getOffset() mean?
     assert(G->getOffset() == 0 &&
          "Do not know what to do with an non-zero offset");
 
     // TODO: We could emit code to handle the initialization somewhere.
-    if (!hasDefinedInitializer(GV)) {
-      unsigned Offset = MFI->allocateLDSGlobal(DL, *cast<GlobalVariable>(GV));
-      return DAG.getConstant(Offset, SDLoc(Op), Op.getValueType());
-    }
-  }
-
-  const Function &Fn = DAG.getMachineFunction().getFunction();
-  DiagnosticInfoUnsupported BadInit(
-      Fn, "unsupported initializer for address space", SDLoc(Op).getDebugLoc());
-  DAG.getContext()->diagnose(BadInit);
+    // We ignore the initializer for now and legalize it to allow selection.
   return SDValue();
 }
 
 SDValue AMDGPUTargetLowering::LowerCONCAT_VECTORS(SDValue Op,
                                                   SelectionDAG &DAG) const {
   SmallVector<SDValue, 8> Args;
 
   EVT VT = Op.getValueType();

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

if (!MFI->isModuleEntryFunction() &&
       // and a trap, since there should be no callable path here.
       B.buildIntrinsic(Intrinsic::trap, ArrayRef<Register>(), true);
       B.buildUndef(DstReg);
       MI.eraseFromParent();
       return true;
     }
 
     // TODO: We could emit code to handle the initialization somewhere.
-    if (!AMDGPUTargetLowering::hasDefinedInitializer(GV)) {
-      const SITargetLowering *TLI = ST.getTargetLowering();
-      if (!TLI->shouldUseLDSConstAddress(GV)) {
-        MI.getOperand(1).setTargetFlags(SIInstrInfo::MO_ABS32_LO);
-        return true; // Leave in place;
-      }
+    // We ignore the initializer for now and legalize it to allow selection.
 
-      if (AS == AMDGPUAS::LOCAL_ADDRESS && GV->hasExternalLinkage()) {
-        Type *Ty = GV->getValueType();
-        // HIP uses an unsized array `extern __shared__ T s[]` or similar
-        // zero-sized type in other languages to declare the dynamic shared
-        // memory which size is not known at the compile time. They will be
-        // allocated by the runtime and placed directly after the static
-        // allocated ones. They all share the same offset.
-        if (B.getDataLayout().getTypeAllocSize(Ty).isZero()) {
-          // Adjust alignment for that dynamic shared memory array.
-          MFI->setDynLDSAlign(B.getDataLayout(), *cast<GlobalVariable>(GV));
-          LLT S32 = LLT::scalar(32);
-          auto Sz =
-              B.buildIntrinsic(Intrinsic::amdgcn_groupstaticsize, {S32}, false);
-          B.buildIntToPtr(DstReg, Sz);
-          MI.eraseFromParent();
-          return true;
-        }
-      }
 
-      B.buildConstant(
-          DstReg,
-          MFI->allocateLDSGlobal(B.getDataLayout(), *cast<GlobalVariable>(GV)));
-      MI.eraseFromParent();
-      return true;
-    }
+    B.buildConstant(DstReg, MFI->allocateLDSGlobal(B.getDataLayout(),
+                                                   *cast<GlobalVariable>(GV)));
 
-    const Function &Fn = MF.getFunction();
-    DiagnosticInfoUnsupported BadInit(
-      Fn, "unsupported initializer for address space", MI.getDebugLoc());
-    Fn.getContext().diagnose(BadInit);
-    return true;
-  }
-
   const SITargetLowering *TLI = ST.getTargetLowering();
 
   if (TLI->shouldEmitFixup(GV)) {
     buildPCRelGlobalAddress(DstReg, Ty, B, GV, 0);
     MI.eraseFromParent();
     return true;
   }
 

llvm/test/CodeGen/AMDGPU/GlobalISel/legalize-lds-initializer.ll

+; NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -global-isel -stop-after=instruction-select -verify-machineinstrs -o - %s | FileCheck %s
+
+; Even if we don't emit an initializer for LDS variables,
+; allow compiler to legalize them instead of an abort.
+
+@lds = addrspace(3) global i32 10
+
+define amdgpu_kernel void @test_lds_initializer() {
+  ; CHECK-LABEL: name: test_lds_initializer
+  ; CHECK: bb.1.entry:
+  ; CHECK:   [[S_MOV_B32_:%[0-9]+]]:sreg_32 = S_MOV_B32 0
+  ; CHECK:   [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_MOV_B32 20
+  ; CHECK:   [[COPY:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_1]]
+  ; CHECK:   [[COPY1:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_]]
+  ; CHECK:   DS_WRITE_B32_gfx9 [[COPY1]], [[COPY]], 0, 0, implicit $exec :: (volatile store (s32) into @lds, addrspace 3)
+  ; CHECK:   S_ENDPGM 0
+entry:
+  store volatile i32 20, i32 addrspace(3)* @lds, align 4
+  ret void
+}

llvm/test/CodeGen/AMDGPU/legalize-lds-initializer.ll

+; NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -stop-after=amdgpu-isel -verify-machineinstrs -o - %s | FileCheck %s
+
+; Even if we don't emit an initializer for LDS variables,
+; allow compiler to legalize them instead of an abort.
+
+@lds = addrspace(3) global i32 10
+
+define amdgpu_kernel void @test_lds_initializer() {
+  ; CHECK-LABEL: name: test_lds_initializer
+  ; CHECK: bb.0.entry:
+  ; CHECK:   [[V_MOV_B32_e32_:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 20, implicit $exec
+  ; CHECK:   [[V_MOV_B32_e32_1:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 0, implicit $exec
+  ; CHECK:   DS_WRITE_B32_gfx9 killed [[V_MOV_B32_e32_1]], killed [[V_MOV_B32_e32_]], 0, 0, implicit $exec :: (volatile store (s32) into @lds, addrspace 3)
+  ; CHECK:   S_ENDPGM 0
+entry:
+  store volatile i32 20, i32 addrspace(3)* @lds, align 4
+  ret void
+}