[AMDGPU] added writelane intrinsic

[tpr]

Summary

For use by LLPC SPV_AMD_shader_ballot extension.

The v_writelane instruction was already implemented for use by SGPR
spilling, but I had to add an extra dummy operand tied to the
destination, to represent that all lanes except the selected one keep
the old value of the destination register.

.ll test changes were due to schedule changes caused by that new
operand.

Change-Id: I26fcb56acb60cf3a9cce63cc673f704fad24cde2

[dstuttard]

include/llvm/IR/IntrinsicsAMDGPU.td

776

It is by no means consistent, but the usual way to document the intrinsic arguments is different.
See e.g. int_amdgcn_exp for an example of this

lib/Target/AMDGPU/SIRegisterInfo.cpp

742

I notice that you've left this as the pseudo whereas the original lowered to an actual MCOpcode - I guess it's more consistent to leave that final lowering until later.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

779

It would make senses for this to be a mangled type so you could use f32/f16/i16/v2i16/v2f16, however this was already a mistake with readlane so maybe it should be consistent.

lib/Target/AMDGPU/SIInstrInfo.cpp

3155–3178

All of this code should be unnecessary. The instruction is already broken if it gets here. VGPRs are not allowed in src0/src1, and required for src2. The selection pattern should already be respecting these constraints. Since you're using a straightforward tablegen pattern, there shouldn't be an issue. The main purpose of the legalizeOperands* functions is to handle legalizing operand constraints that require looking at the entire instruction, and not just a single operand (i.e. the constant bus restriction)

lib/Target/AMDGPU/SIRegisterInfo.cpp

735–739

This is broken. You shouldn't need to do this. The kill flags need to be correct. The VGPRs used for SGPR spilling are hackily reserved. Depending on where this is called from it's possible that the way that is done is broken. I think bad things could still happen from this undef if the verifier is unhappy.

742

There is a reason for this, but I don't remember what it was exactly. I think there was some related change between SI and VI. I think you should leave this as a separate change. I know I've tried to fix this before and hit some issue.

lib/Target/AMDGPU/VOP2Instructions.td

329

This isn't a src2, so should be named otherwise to indicate it's tied (like vdst_in or something)

[tpr]

lib/Target/AMDGPU/SIRegisterInfo.cpp

735–739

So what do you think I should do? What bad things could happen? Post-RA scheduling going wrong?

742

Yes; gfx6 and gfx7 have vop2 and gfx8 has vop3. But I saw that my new use of writelane does not lower the pseudo until later, and my testing with this change (with llpc only, not other users of the backend) passed. But good idea to make it a separate change.

[tpr]

lib/Target/AMDGPU/SIInstrInfo.cpp

3155–3178

Hi Matt. I tried removing this code, and the verifier falls over due to src1 being a vgpr_32 in the test I am trying. Also I see that readlane has analogous special case code here in legalizeOperandsVOP2.

[tpr]

lib/Target/AMDGPU/SIRegisterInfo.cpp

735–739

Also it's not any more broken than it was before my change, when the destination register is not mentioned as an input at all. Fixing any brokenness of sgpr spilling should not be part of this change. I'll change the comment to reflect suspected brokenness.

[arsenm]

lib/Target/AMDGPU/SIInstrInfo.cpp

3155–3178

The src2 verification shouldn't be necessary.

The readlane handling is surprising to me, I would expect we would need to handle this case by waterfalling

lib/Target/AMDGPU/SIRegisterInfo.cpp

735–739

My main concern would be the register scavenger finding this as an available register to use

[tpr]

V2: Addressed some review comments.

[tpr]

V3: Addressed review comment about not needing to check src2 in
legalizeOperandsVop2.

[tpr]

Ping. Can I land this on the basis that vgpr allocation for sgpr spilling is no more broken than it was before?

[nhaehnle]

I think there should be two more tests: one where the vdst_in comes from a uniform value in an SGPR (e.g. a shader where it comes from an inreg) and one where it's a constant.

lib/Target/AMDGPU/SIInstrInfo.cpp

3155–3178

For context, the GLSL readlane intrinsic is defined to require a uniform lane selection. So, like other intrinsics which have the same requirement, not doing automatic waterfalling is the right thing to do here.

lib/Target/AMDGPU/SIRegisterInfo.cpp

742

What is actually breaking without the RegState::Undef?

When we allocate a VGPR for spilling, we add it as LiveIn to all basic blocks, and the V_READLANE_B32 generated for restoring a spilled SGPR should never set the VGPR to killed. So I don't see where the machine verifier would complain. An example would be good.

[tpr]

lib/Target/AMDGPU/SIRegisterInfo.cpp

742

Just re-checked, and I think the problem (shown by test/CodeGen/AMDGPU/byval-frame-setup.ll) is when it happens to pick a callee saves register for spilling sgprs into. After inserting the writelanes for sgpr spills, it runs the prologue/epilogue insertion pass, which adds a store of the callee saves register to stack frame, with a kill.

Any ideas on a good way to fix that? Anything less hacky than scanning the first basic block to find the first writelane and mark just that one as undef?

[tpr]

V4: Mark the input undef only on the first sgpr spill in the first basic

block for that vgpr. Added the suggested two tests.

[nhaehnle]

Thank you for taking another look. I understand now, and I agree that there doesn't seem to be a super clean solution. Not to mention that the function prolog/epilog code here is clearly incorrect, because it is saving/restoring the VGPR with a potentially insufficient EXEC mask.

Anyway, those issues are unrelated, and I think your approach to solving it is reasonable now. I'd just ask you to inline that one helper function as mentioned in the comment, apart from that the change looks good to me.

lib/Target/AMDGPU/SIMachineFunctionInfo.h

656–664 ↗

(On Diff #135922)

I don't like this function. The name is misleading: it suggests a pure function ("is XXX"), but in reality there is a side effect; and the list of parameters is pretty surprising given the name as well.

Please just inline the function (including its comments) in the one single location where it is used, I think that will be much cleaner.

[tpr]

lib/Target/AMDGPU/SIMachineFunctionInfo.h

656–664 ↗

(On Diff #135922)

Will inline before landing.

include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_readfirstlane :
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 
 // The lane argument must be uniform across the currently active threads of the
 // current wave. Otherwise, the result is undefined.
 def int_amdgcn_readlane :
   GCCBuiltin<"__builtin_amdgcn_readlane">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 
+// The value to write and lane select arguments must be uniform across the
+// currently active threads of the current wave. Otherwise, the result is
+// undefined.
+def int_amdgcn_writelane :
+  GCCBuiltin<"__builtin_amdgcn_writelane">,
+  Intrinsic<[llvm_i32_ty], [
+    llvm_i32_ty,    // uniform value to write: returned by the selected lane
+    llvm_i32_ty,    // uniform lane select
+    llvm_i32_ty     // returned by all lanes other than the selected one
+  ],
+  [IntrNoMem, IntrConvergent]
+>;
+
 def int_amdgcn_alignbit : Intrinsic<[llvm_i32_ty],
   [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]
 >;
 
 def int_amdgcn_alignbyte : Intrinsic<[llvm_i32_ty],
   [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
   [IntrNoMem, IntrSpeculatable]

lib/Target/AMDGPU/SIInstrInfo.cpp

if (DstUnused && DstUnused->isImm() &&
       } else if (TargetRegisterInfo::isPhysicalRegister(TiedMO.getReg()) &&
                  Dst.getReg() != TiedMO.getReg()) {
         ErrInfo = "Dst register should use same physical register as preserved";
         return false;
       }
     }
   }
 
-  // Verify VOP*
-  if (isVOP1(MI) || isVOP2(MI) || isVOP3(MI) || isVOPC(MI) || isSDWA(MI)) {
+  // Verify VOP*. Ignore multiple sgpr operands on writelane.
+  if (Desc.getOpcode() != AMDGPU::V_WRITELANE_B32
+      && (isVOP1(MI) || isVOP2(MI) || isVOP3(MI) || isVOPC(MI) || isSDWA(MI))) {
     // Only look at the true operands. Only a real operand can use the constant
     // bus, and we don't want to check pseudo-operands like the source modifier
     // flags.
     const int OpIndices[] = { Src0Idx, Src1Idx, Src2Idx };
 
     unsigned ConstantBusCount = 0;
 
     if (AMDGPU::getNamedOperandIdx(Opcode, AMDGPU::OpName::imm) != -1)
void SIInstrInfo::legalizeOperandsVOP2(MachineRegisterInfo &MRI,
   if (HasImplicitSGPR) {
     int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
     MachineOperand &Src0 = MI.getOperand(Src0Idx);
 
     if (Src0.isReg() && RI.isSGPRReg(MRI, Src0.getReg()))
       legalizeOpWithMove(MI, Src0Idx);
   }
 
+  // Special case: V_WRITELANE_B32 accepts only immediate or SGPR operands for
+  // both the value to write (src0) and lane select (src1), and the value to
+  // use for other lanes (src2) must be VGPR. Fix up non-SGPR src0/src1 with
+  // V_READFIRSTLANE.
+  if (Opc == AMDGPU::V_WRITELANE_B32) {
+    int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
+    MachineOperand &Src0 = MI.getOperand(Src0Idx);
+    const DebugLoc &DL = MI.getDebugLoc();
+    if (Src0.isReg() && RI.isVGPR(MRI, Src0.getReg())) {
+      unsigned Reg = MRI.createVirtualRegister(&AMDGPU::SReg_32_XM0RegClass);
+      BuildMI(*MI.getParent(), MI, DL, get(AMDGPU::V_READFIRSTLANE_B32), Reg)
+          .add(Src0);
+      Src0.ChangeToRegister(Reg, false);
+    }
+    if (Src1.isReg() && RI.isVGPR(MRI, Src1.getReg())) {
+      unsigned Reg = MRI.createVirtualRegister(&AMDGPU::SReg_32_XM0RegClass);
+      const DebugLoc &DL = MI.getDebugLoc();
+      BuildMI(*MI.getParent(), MI, DL, get(AMDGPU::V_READFIRSTLANE_B32), Reg)
+          .add(Src1);
+      Src1.ChangeToRegister(Reg, false);
+    }
+    int VdstInIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::vdst_in);
+    MachineOperand &VdstIn = MI.getOperand(VdstInIdx);
+    if (VdstIn.isReg() && RI.isSGPRReg(MRI, VdstIn.getReg()))
+      legalizeOpWithMove(MI, VdstInIdx);
+    return;
+  }
+
   // VOP2 src0 instructions support all operand types, so we don't need to check
   // their legality. If src1 is already legal, we don't need to do anything.
   if (isLegalRegOperand(MRI, InstrDesc.OpInfo[Src1Idx], Src1))
     return;
 
   // Special case: V_READLANE_B32 accepts only immediate or SGPR operands for
   // lane select. Fix up using V_READFIRSTLANE, since we assume that the lane
   // select is uniform.

lib/Target/AMDGPU/SIRegisterInfo.cpp

if (SpillToSMEM) {
         .addMemOperand(MMO);
 
       continue;
     }
 
     if (SpillToVGPR) {
       SIMachineFunctionInfo::SpilledReg Spill = VGPRSpills[i];
 
+      // FIXME: Using Undef on the "old value of VGPR" operand here is a hack
+      // to avoid MachineVerifier objecting when that same physical VGPR was
+      // killed by an earlier instruction. The way that VGPRs are allocated for
+      // SGPR spilling is hacky, and the undef flags are incorrect, and could
+      // possibly cause subsequent problems. Although it is no more incorrect
+      // than it was before adding this extra operand to V_WRITELANE_B32.
       BuildMI(*MBB, MI, DL,
               TII->getMCOpcodeFromPseudo(AMDGPU::V_WRITELANE_B32),
               Spill.VGPR)
         .addReg(SubReg, getKillRegState(IsKill))
-        .addImm(Spill.Lane);
+        .addImm(Spill.Lane)
+        .addReg(Spill.VGPR, RegState::Undef);
 
       // FIXME: Since this spills to another register instead of an actual
       // frame index, we should delete the frame index when all references to
       // it are fixed.
     } else {
       // XXX - Can to VGPR spill fail for some subregisters but not others?
       if (OnlyToVGPR)
         return false;

lib/Target/AMDGPU/VOP2Instructions.td

def VOP_READLANE : VOPProfile<[i32, i32, i32]> {
   let Ins32 = (ins VGPR_32:$src0, SCSrc_b32:$src1);
   let Ins64 = Ins32;
   let Asm32 = " $vdst, $src0, $src1";
   let Asm64 = Asm32;
   let HasExt = 0;
   let HasSDWA9 = 0;
 }
 
-def VOP_WRITELANE : VOPProfile<[i32, i32, i32]> {
+def VOP_WRITELANE : VOPProfile<[i32, i32, i32, i32]> {
   let Outs32 = (outs VGPR_32:$vdst);
   let Outs64 = Outs32;
-  let Ins32 = (ins SCSrc_b32:$src0, SCSrc_b32:$src1);
+  let Ins32 = (ins SCSrc_b32:$src0, SCSrc_b32:$src1, VGPR_32:$vdst_in);
   let Ins64 = Ins32;
   let Asm32 = " $vdst, $src0, $src1";
   let Asm64 = Asm32;
   let HasExt = 0;
   let HasSDWA9 = 0;
+  let HasSrc2 = 0;
+  let HasSrc2Mods = 0;
 }
 
 //===----------------------------------------------------------------------===//
 // VOP2 Instructions
 //===----------------------------------------------------------------------===//
 
 let SubtargetPredicate = isGCN in {
 

 
 } // End isCommutable = 1
 
 // These are special and do not read the exec mask.
 let isConvergent = 1, Uses = []<Register> in {
 def V_READLANE_B32 : VOP2_Pseudo<"v_readlane_b32", VOP_READLANE,
   [(set i32:$vdst, (int_amdgcn_readlane i32:$src0, i32:$src1))], "">;
 
-def V_WRITELANE_B32 : VOP2_Pseudo<"v_writelane_b32", VOP_WRITELANE, [], "">;
+let Constraints = "$vdst = $vdst_in", DisableEncoding="$vdst_in" in {
+def V_WRITELANE_B32 : VOP2_Pseudo<"v_writelane_b32", VOP_WRITELANE,
+  [(set i32:$vdst, (int_amdgcn_writelane i32:$src0, i32:$src1, i32:$vdst_in))], "">;
+} // End $vdst = $vdst_in, DisableEncoding $vdst_in
 } // End isConvergent = 1
 
 defm V_BFM_B32 : VOP2Inst <"v_bfm_b32", VOP_NO_EXT<VOP_I32_I32_I32>>;
 defm V_BCNT_U32_B32 : VOP2Inst <"v_bcnt_u32_b32", VOP_NO_EXT<VOP_I32_I32_I32>>;
 defm V_MBCNT_LO_U32_B32 : VOP2Inst <"v_mbcnt_lo_u32_b32", VOP_NO_EXT<VOP_I32_I32_I32>, int_amdgcn_mbcnt_lo>;
 defm V_MBCNT_HI_U32_B32 : VOP2Inst <"v_mbcnt_hi_u32_b32", VOP_NO_EXT<VOP_I32_I32_I32>, int_amdgcn_mbcnt_hi>;
 defm V_LDEXP_F32 : VOP2Inst <"v_ldexp_f32", VOP_NO_EXT<VOP_F32_F32_I32>, AMDGPUldexp>;
 defm V_CVT_PKACCUM_U8_F32 : VOP2Inst <"v_cvt_pkaccum_u8_f32", VOP_NO_EXT<VOP_I32_F32_I32>>; // TODO: set "Uses = dst"

 defm V_SUB_I32            : VOP2be_Real_e32e64_si <0x26>;
 defm V_SUBREV_I32         : VOP2be_Real_e32e64_si <0x27>;
 defm V_ADDC_U32           : VOP2be_Real_e32e64_si <0x28>;
 defm V_SUBB_U32           : VOP2be_Real_e32e64_si <0x29>;
 defm V_SUBBREV_U32        : VOP2be_Real_e32e64_si <0x2a>;
 
 defm V_READLANE_B32       : VOP2_Real_si <0x01>;
 
-let InOperandList = (ins SSrc_b32:$src0, SCSrc_b32:$src1) in {
+let InOperandList = (ins SSrc_b32:$src0, SCSrc_b32:$src1, VSrc_b32:$vdst_in) in {
 defm V_WRITELANE_B32      : VOP2_Real_si <0x02>;
 }
 
 defm V_MAC_LEGACY_F32     : VOP2_Real_e32e64_si <0x6>;
 defm V_MIN_LEGACY_F32     : VOP2_Real_e32e64_si <0xd>;
 defm V_MAX_LEGACY_F32     : VOP2_Real_e32e64_si <0xe>;
 defm V_LSHR_B32           : VOP2_Real_e32e64_si <0x15>;
 defm V_ASHR_I32           : VOP2_Real_e32e64_si <0x17>;

test/CodeGen/AMDGPU/byval-frame-setup.ll

entry:
   ret void
 }
 
 ; GCN-LABEL: {{^}}void_func_byval_struct_non_leaf:
 ; GCN: s_mov_b32 s5, s32
 ; GCN-DAG: buffer_store_dword v32
 ; GCN-DAG: buffer_store_dword v33
 ; GCN-NOT: v_writelane_b32 v{{[0-9]+}}, s32
-; GCN: v_writelane_b32
-
+; GCN-DAG: v_writelane_b32
 ; GCN-DAG: s_add_u32 s32, s32, 0xb00{{$}}
-
 ; GCN-DAG: buffer_load_dword [[LOAD0:v[0-9]+]], off, s[0:3], s5 offset:4{{$}}
-; GCN: v_add_{{[iu]}}32_e32 [[ADD0:v[0-9]+]], vcc, 1, [[LOAD0]]
-; GCN: buffer_store_dword [[ADD0]], off, s[0:3], s5 offset:4{{$}}
+; GCN-DAG: v_add_{{[iu]}}32_e32 [[ADD0:v[0-9]+]], vcc, 1, [[LOAD0]]
+; GCN-DAG: buffer_store_dword [[ADD0]], off, s[0:3], s5 offset:4{{$}}
 
-; GCN: buffer_load_dword [[LOAD1:v[0-9]+]], off, s[0:3], s5 offset:20{{$}}
-; GCN: v_add_{{[iu]}}32_e32 [[ADD1:v[0-9]+]], vcc, 2, [[LOAD1]]
+; GCN-DAG: buffer_load_dword [[LOAD1:v[0-9]+]], off, s[0:3], s5 offset:20{{$}}
+; GCN-DAG: v_add_{{[iu]}}32_e32 [[ADD1:v[0-9]+]], vcc, 2, [[LOAD1]]
 
 ; GCN: s_swappc_b64
 
 ; GCN: buffer_store_dword [[ADD1]], off, s[0:3], s5 offset:20{{$}}
 
 ; GCN: v_readlane_b32
 ; GCN-NOT: v_readlane_b32 s32
 ; GCN: buffer_load_dword v32,

 ; GCN-DAG: v_writelane_b32
 
 ; GCN-DAG: v_mov_b32_e32 [[NINE:v[0-9]+]], 9
 ; GCN-DAG: v_mov_b32_e32 [[THIRTEEN:v[0-9]+]], 13
 
 ; GCN-DAG: buffer_store_dword [[NINE]], off, s[0:3], s5 offset:8
 ; GCN-DAG: buffer_store_dword [[THIRTEEN]], off, s[0:3], s5 offset:24
 
-; GCN: buffer_load_dword [[LOAD0:v[0-9]+]], off, s[0:3], s5 offset:8
-; GCN: buffer_load_dword [[LOAD1:v[0-9]+]], off, s[0:3], s5 offset:12
-; GCN: buffer_load_dword [[LOAD2:v[0-9]+]], off, s[0:3], s5 offset:16
-; GCN: buffer_load_dword [[LOAD3:v[0-9]+]], off, s[0:3], s5 offset:20
+; GCN-DAG: buffer_load_dword [[LOAD0:v[0-9]+]], off, s[0:3], s5 offset:8
+; GCN-DAG: buffer_load_dword [[LOAD1:v[0-9]+]], off, s[0:3], s5 offset:12
+; GCN-DAG: buffer_load_dword [[LOAD2:v[0-9]+]], off, s[0:3], s5 offset:16
+; GCN-DAG: buffer_load_dword [[LOAD3:v[0-9]+]], off, s[0:3], s5 offset:20
 
 ; GCN-NOT: s_add_u32 s32, s32, 0x800
 
 ; GCN-DAG: buffer_store_dword [[LOAD0]], off, s[0:3], s32 offset:4{{$}}
 ; GCN-DAG: buffer_store_dword [[LOAD1]], off, s[0:3], s32 offset:8
 ; GCN-DAG: buffer_store_dword [[LOAD2]], off, s[0:3], s32 offset:12
 ; GCN-DAG: buffer_store_dword [[LOAD3]], off, s[0:3], s32 offset:16
 

test/CodeGen/AMDGPU/callee-frame-setup.ll

 ; GCN: s_mov_b32 s5, s32
 ; GCN: buffer_store_dword v32, off, s[0:3], s5 offset:8
 
 ; GCN-DAG: v_writelane_b32 v32, s33,
 ; GCN-DAG: v_writelane_b32 v32, s34,
 ; GCN-DAG: v_writelane_b32 v32, s35,
 ; GCN-DAG: s_add_u32 s32, s32, 0x300{{$}}
 ; GCN-DAG: v_mov_b32_e32 v0, 0{{$}}
-; GCN: buffer_store_dword v0, off, s[0:3], s5 offset:4{{$}}
+; GCN-DAG: buffer_store_dword v0, off, s[0:3], s5 offset:4{{$}}
 ; GCN-DAG: s_mov_b32 s33, s5
 
 
 ; GCN: s_swappc_b64
 ; GCN: s_mov_b32 s5, s33
 ; GCN-DAG: v_readlane_b32 s35,
 ; GCN-DAG: v_readlane_b32 s34,
 ; GCN-DAG: v_readlane_b32 s33,

test/CodeGen/AMDGPU/inserted-wait-states.mir

 body: |
   bb.0:
     $vgpr0,$sgpr0_sgpr1 = V_ADD_I32_e64 $vgpr1, $vgpr2, implicit $vcc, implicit $exec
     $sgpr4 = V_READLANE_B32 $vgpr4, $sgpr0
     S_BRANCH %bb.1
 
   bb.1:
     $vgpr0,$sgpr0_sgpr1 = V_ADD_I32_e64 $vgpr1, $vgpr2, implicit $vcc, implicit $exec
-    $vgpr4 = V_WRITELANE_B32 $sgpr0, $sgpr0
+    $vgpr4 = V_WRITELANE_B32 $sgpr0, $sgpr0, $vgpr4
     S_BRANCH %bb.2
 
   bb.2:
     $vgpr0,implicit $vcc = V_ADD_I32_e32 $vgpr1, $vgpr2, implicit $vcc, implicit $exec
     $sgpr4 = V_READLANE_B32 $vgpr4, $vcc_lo
     S_BRANCH %bb.3
 
   bb.3:
     $vgpr0,implicit $vcc = V_ADD_I32_e32 $vgpr1, $vgpr2, implicit $vcc, implicit $exec
-    $vgpr4 = V_WRITELANE_B32 $sgpr4, $vcc_lo
+    $vgpr4 = V_WRITELANE_B32 $sgpr4, $vcc_lo, $vgpr4
     S_ENDPGM
 
 ...
 
 ...
 ---
 
 # GCN-LABEL: name: rfe

test/CodeGen/AMDGPU/llvm.amdgcn.writelane.ll

+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx802 -verify-machineinstrs < %s | FileCheck %s
+
+declare i32 @llvm.amdgcn.writelane(i32, i32, i32) #0
+
+; CHECK-LABEL: {{^}}test_writelane_sreg:
+; CHECK: v_writelane_b32 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}
+define amdgpu_kernel void @test_writelane_sreg(i32 addrspace(1)* %out, i32 %src0, i32 %src1) #1 {
+  %oldval = load i32, i32 addrspace(1)* %out
+  %writelane = call i32 @llvm.amdgcn.writelane(i32 %src0, i32 %src1, i32 %oldval)
+  store i32 %writelane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_imm_sreg:
+; CHECK: v_writelane_b32 v{{[0-9]+}}, 32, s{{[0-9]+}}
+define amdgpu_kernel void @test_writelane_imm_sreg(i32 addrspace(1)* %out, i32 %src1) #1 {
+  %oldval = load i32, i32 addrspace(1)* %out
+  %writelane = call i32 @llvm.amdgcn.writelane(i32 32, i32 %src1, i32 %oldval)
+  store i32 %writelane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_vreg_lane:
+; CHECK: v_readfirstlane_b32 [[LANE:s[0-9]+]], v{{[0-9]+}}
+; CHECK: v_writelane_b32 v{{[0-9]+}}, 12, [[LANE]]
+define amdgpu_kernel void @test_writelane_vreg_lane(i32 addrspace(1)* %out, <2 x i32> addrspace(1)* %in) #1 {
+  %tid = call i32 @llvm.amdgcn.workitem.id.x()
+  %gep.in = getelementptr <2 x i32>, <2 x i32> addrspace(1)* %in, i32 %tid
+  %args = load <2 x i32>, <2 x i32> addrspace(1)* %gep.in
+  %oldval = load i32, i32 addrspace(1)* %out
+  %lane = extractelement <2 x i32> %args, i32 1
+  %writelane = call i32 @llvm.amdgcn.writelane(i32 12, i32 %lane, i32 %oldval)
+  store i32 %writelane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; TODO: m0 should be folded.
+; CHECK-LABEL: {{^}}test_writelane_m0_sreg:
+; CHECK: s_mov_b32 m0, -1
+; CHECK: s_mov_b32 [[COPY_M0:s[0-9]+]], m0
+; CHECK: v_writelane_b32 v{{[0-9]+}}, [[COPY_M0]], s{{[0-9]+}}
+define amdgpu_kernel void @test_writelane_m0_sreg(i32 addrspace(1)* %out, i32 %src1) #1 {
+  %oldval = load i32, i32 addrspace(1)* %out
+  %m0 = call i32 asm "s_mov_b32 m0, -1", "={M0}"()
+  %writelane = call i32 @llvm.amdgcn.writelane(i32 %m0, i32 %src1, i32 %oldval)
+  store i32 %writelane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+; CHECK-LABEL: {{^}}test_writelane_imm:
+; CHECK: v_writelane_b32 v{{[0-9]+}}, s{{[0-9]+}}, 32
+define amdgpu_kernel void @test_writelane_imm(i32 addrspace(1)* %out, i32 %src0) #1 {
+  %oldval = load i32, i32 addrspace(1)* %out
+  %writelane = call i32 @llvm.amdgcn.writelane(i32 %src0, i32 32, i32 %oldval) #0
+  store i32 %writelane, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+declare i32 @llvm.amdgcn.workitem.id.x() #2
+
+attributes #0 = { nounwind readnone convergent }
+attributes #1 = { nounwind }
+attributes #2 = { nounwind readnone }

test/CodeGen/AMDGPU/sibling-call.ll

 ; GCN: buffer_store_dword v34, off, s[0:3], s5 offset:12
 ; GCN: buffer_store_dword v32, off, s[0:3], s5 offset:8 ; 4-byte Folded Spill
 ; GCN: buffer_store_dword v33, off, s[0:3], s5 offset:4 ; 4-byte Folded Spill
 ; GCN-DAG: v_writelane_b32 v34, s33, 0
 ; GCN-DAG: v_writelane_b32 v34, s34, 1
 ; GCN-DAG: v_writelane_b32 v34, s35, 2
 ; GCN-DAG: s_add_u32 s32, s32, 0x400
 
-; GCN: s_getpc_b64
+; GCN-DAG: s_getpc_b64
 ; GCN: s_swappc_b64
 
 ; GCN: s_getpc_b64 s[6:7]
 ; GCN: s_add_u32 s6, s6, sibling_call_i32_fastcc_i32_i32@rel32@lo+4
 ; GCN: s_addc_u32 s7, s7, sibling_call_i32_fastcc_i32_i32@rel32@hi+4
 
 ; GCN-DAG: v_readlane_b32 s33, v34, 0
 ; GCN-DAG: v_readlane_b32 s34, v34, 1