[AMDGPU][GISel] Add inverse ballot intrinsic

[OutOfCache]

Summary

The inverse ballot intrinsic takes in a boolean mask for all lanes and
returns the boolean for the current lane. See SPIR-V's
subgroupInverseBallot() in the GL_KHR_shader_subgroup extension.
This allows decision making via branch and select instructions with a manually
manipulated mask.

Implemented in GlobalISel and SelectionDAG, since currently both are supported.
The SelectionDAG required pseudo instructions to use the custom inserter.

The boolean mask needs to be uniform for all lanes.
Therefore we expect SGPR input. In case the source is in a
VGPR, we insert one or more v_readfirstlane instructions.

[OutOfCache]

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.inverse.ballot.i32.ll

21 ↗

(On Diff #506042)

This s_sendmsg instruction is the only difference between gfx10 and gfx11 in every test. Should I remove the gfx10 test?

[tsymalla]

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.inverse.ballot.i32.ll

2 ↗

(On Diff #506042)

Why do you require -risky-select?

21 ↗

(On Diff #506042)

I'd add just a single test run if the differences are minor.

[arsenm]

Drop the "[GISel]" from the description, this isn't adding existing functionality to it. Obsoletes D57748?

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

2540

Should let it fall through to fail to select instead

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

1363–1364

This isn't enforced by the verifier

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4480

Why do you have to go through selecting the pseudo? Can't you do this split during the select?

[foad]

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

1380

Emitting the readfirstlanes directly during instruction selection is pretty unusual. Usually it is done by legalizeOperands during SIFixSGPRCopies. Is there a reason that inverse ballot has to be done differently?

[tsymalla]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4489

Use a better name, like Is32BitMask.

4504

Add a comment on why 0 is used in 32-bit mode.

4513

Could also use a comment.

llvm/lib/Target/AMDGPU/SIInstructions.td

194

Can unknown be changed to a more constrained register class instead of allowing arbitrary inputs?

194

Rather call $src $mask.

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.inverse.ballot.i32.ll

153 ↗

(On Diff #506042)

Nit: Missing whitespace after %entry

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.inverse.ballot.i64.ll

52 ↗

(On Diff #506042)

Can you add tests for more complex bitmask inputs?

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.inverse.ballot.i64.ll

16

Can be optimized, as s0-s2 and v2-v3 are zero. Probably add another test to show the MIR your code reduces the intrinsic call to

[rovka]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4497

Nit: Most of this looks the same for GlobalISel, could you try to extract a helper function to use for both?

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.inverse.ballot.i32.ll

3 ↗

(On Diff #506042)

I have 2 nits here:

1. Why go all the way to assembly? You could -stop-after=finalize-isel (or even before it). You'll probably have to use update_mir_test_checks.py though.
2. There's precedent for having a single test with -global-isel=0 and -global-isel=1 RUN lines, so you don't really need separate files in this directory.

[OutOfCache]

• Simplify GISel Implementation by using legalizeOperands
• Remove unnecessary checks.
• Remove redundant tests and move GISel ones to SDAG tests
• Increase readability.

[OutOfCache]

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

1380

Turns out there isn't! Thank you for the suggestion. It is much cleaner now.

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.inverse.ballot.i32.ll

3 ↗

(On Diff #506042)

Thank you for the suggestion. Combining the GISel and SelDAG Tests into one file is much better.

I run into an issue when trying the -stop-after=finalize-isel however. The CHECK lines seem to be generated fine with update_mir_test_checks.py, but running the tests with llvm-lit fails the tests. Changing back to asm tests works again, with no other changes to the code. I will need to investigate further.

[OutOfCache]

• Fix typos
• Rename variable

[OutOfCache]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4480

I have attempted simply selecting the node to a copy, similar to the current GISel implementation. However, I have run into issues in case the mask is inside a VGPR. The following passes changed the VGPR->SGPR copy into a VGPR->VGPR copy instead of producing the v_readfirstlanes. Do you have any other suggestion of things I could try?

[nhaehnle]

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

2540

Probably better to error out explicitly, no? A fall-through is easily missed by future changes.

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4502–4523

This entire sequence could probably be replaced by just calling SIInstrInfo::readlaneVGPRToSGPR.

llvm/lib/Target/AMDGPU/SIInstructions.td

193

Remove the Defs = [VCC]. Defs is used to designate implicit defs of physical registers, but these instructions don't do that, they only use virtual registers.

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.inverse.ballot.i32.ll

119–130

This code is wrong. That's a problem that's orthogonal to your change and has to do with GlobalISel not correctly taking uniformity analysis into account (yet). But please add a comment to make a note of it (same in the corresponding wave64 test).

The SDAG code is correct.

[OutOfCache]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.inverse.ballot.i32.ll

119–130

What is incorrect here? I only see the additional v_mov in the SelDAG code, but every other instruction seems to be the same.

[OutOfCache]

• Simplify SelDAG with call to SIInstrInfo::readlaneVGPRToSGPR
• Add comment to clarify that current GISel test is incorrect
• remove [VCC] form pseudo instruction def

[OutOfCache]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4502–4523

Good call. Would it make sense to also use this method for other pseudos that generate readfirstlanes here? Like S_ADD_CO_PSEUDO. In another change, of course.

[nhaehnle]

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

4502–4523

I think so, yes. And yes, it should be a separate change.

[nhaehnle]

What about the comment(s) in AMDGPUISelDAGToDAG.cpp?

[OutOfCache]

• add internal error if the mask size is not 32 or 64

[nhaehnle]

LGTM

[OutOfCache]

• Change llvm_unreachable_internal to llvm_unreachable

I decided to keep the error output since it is more
descriptive than the Invalid Opcode error we
would otherwise output.

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_fcmp :
   Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
             [IntrNoMem, IntrConvergent,
              ImmArg<ArgIndex<2>>, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 def int_amdgcn_ballot :
   Intrinsic<[llvm_anyint_ty], [llvm_i1_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
+def int_amdgcn_inverse_ballot :
+  Intrinsic<[llvm_i1_ty], [llvm_anyint_ty],
+            [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
+
 def int_amdgcn_readfirstlane :
   ClangBuiltin<"__builtin_amdgcn_readfirstlane">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 // The lane argument must be uniform across the currently active threads of the
 // current wave. Otherwise, the result is undefined.
 def int_amdgcn_readlane :

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

case Intrinsic::amdgcn_strict_wwm:
     Opcode = AMDGPU::STRICT_WWM;
     break;
   case Intrinsic::amdgcn_strict_wqm:
     Opcode = AMDGPU::STRICT_WQM;
     break;
   case Intrinsic::amdgcn_interp_p1_f16:
     SelectInterpP1F16(N);
     return;
+  case Intrinsic::amdgcn_inverse_ballot:
+    switch (N->getOperand(1).getValueSizeInBits()) {
+    case 32:
+      Opcode = AMDGPU::S_INVERSE_BALLOT_U32;
+      break;
+    case 64:
+      Opcode = AMDGPU::S_INVERSE_BALLOT_U64;
+      break;
+    }
+    break;
   default:
     SelectCode(N);
     return;
   }
 
   SDValue Src = N->getOperand(1);
   CurDAG->SelectNodeTo(N, Opcode, N->getVTList(), {Src});
 }

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.h

private:
   bool selectG_INSERT(MachineInstr &I) const;
   bool selectG_SBFX_UBFX(MachineInstr &I) const;
 
   bool selectInterpP1F16(MachineInstr &MI) const;
   bool selectWritelane(MachineInstr &MI) const;
   bool selectDivScale(MachineInstr &MI) const;
   bool selectIntrinsicCmp(MachineInstr &MI) const;
   bool selectBallot(MachineInstr &I) const;
+  bool selectInverseBallot(MachineInstr &I) const;
   bool selectRelocConstant(MachineInstr &I) const;
   bool selectGroupStaticSize(MachineInstr &I) const;
   bool selectReturnAddress(MachineInstr &I) const;
   bool selectG_INTRINSIC(MachineInstr &I) const;
 
   bool selectEndCfIntrinsic(MachineInstr &MI) const;
   bool selectDSOrderedIntrinsic(MachineInstr &MI, Intrinsic::ID IID) const;
   bool selectDSGWSIntrinsic(MachineInstr &MI, Intrinsic::ID IID) const;

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

case Intrinsic::amdgcn_div_scale:
     return selectDivScale(I);
   case Intrinsic::amdgcn_icmp:
   case Intrinsic::amdgcn_fcmp:
     if (selectImpl(I, *CoverageInfo))
       return true;
     return selectIntrinsicCmp(I);
   case Intrinsic::amdgcn_ballot:
     return selectBallot(I);
+  case Intrinsic::amdgcn_inverse_ballot:
+    return selectInverseBallot(I);
   case Intrinsic::amdgcn_reloc_constant:
     return selectRelocConstant(I);
   case Intrinsic::amdgcn_groupstaticsize:
     return selectGroupStaticSize(I);
   case Intrinsic::returnaddress:
     return selectReturnAddress(I);
   case Intrinsic::amdgcn_smfmac_f32_16x16x32_f16:
   case Intrinsic::amdgcn_smfmac_f32_32x32x16_f16:
if (Arg) {
     Register SrcReg = I.getOperand(2).getReg();
     BuildMI(*BB, &I, DL, TII.get(AMDGPU::COPY), DstReg).addReg(SrcReg);
   }
 
   I.eraseFromParent();
   return true;
 }
 
+bool AMDGPUInstructionSelector::selectInverseBallot(MachineInstr &I) const {
+  MachineBasicBlock *BB = I.getParent();
+  const DebugLoc &DL = I.getDebugLoc();
+  const Register DstReg = I.getOperand(0).getReg();
+  const Register MaskReg = I.getOperand(2).getReg();
+
+  BuildMI(*BB, &I, DL, TII.get(AMDGPU::COPY), DstReg).addReg(MaskReg);
+  I.eraseFromParent();
+  return true;
+}
+
 bool AMDGPUInstructionSelector::selectRelocConstant(MachineInstr &I) const {
   Register DstReg = I.getOperand(0).getReg();
   const RegisterBank *DstBank = RBI.getRegBank(DstReg, *MRI, TRI);
   const TargetRegisterClass *DstRC = TRI.getRegClassForSizeOnBank(32, *DstBank);
   if (!DstRC || !RBI.constrainGenericRegister(DstReg, *DstRC, *MRI))
     return false;
 
   const bool IsVALU = DstBank->getID() == AMDGPU::VGPRRegBankID;
 
   Module *M = MF->getFunction().getParent();
   const MDNode *Metadata = I.getOperand(2).getMetadata();
   auto SymbolName = cast<MDString>(Metadata->getOperand(0))->getString();
   auto RelocSymbol = cast<GlobalVariable>(
     M->getOrInsertGlobal(SymbolName, Type::getInt32Ty(M->getContext())));
 
   MachineBasicBlock *BB = I.getParent();
   BuildMI(*BB, &I, I.getDebugLoc(),
           TII.get(IsVALU ? AMDGPU::V_MOV_B32_e32 : AMDGPU::S_MOV_B32), DstReg)
     .addGlobalAddress(RelocSymbol, 0, SIInstrInfo::MO_ABS32_LO);
 
   I.eraseFromParent();
   return true;

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

case Intrinsic::amdgcn_permlanex16: {
       return;
     }
     case Intrinsic::amdgcn_sbfe:
       applyMappingBFE(OpdMapper, true);
       return;
     case Intrinsic::amdgcn_ubfe:
       applyMappingBFE(OpdMapper, false);
       return;
+    case Intrinsic::amdgcn_inverse_ballot:
+      applyDefaultMapping(OpdMapper);
+      constrainOpWithReadfirstlane(MI, MRI, 2); // Mask
+      return;
     case Intrinsic::amdgcn_ballot:
       // Use default handling and insert copy to vcc source.
       break;
     }
     break;
   }
   case AMDGPU::G_AMDGPU_INTRIN_IMAGE_LOAD:
   case AMDGPU::G_AMDGPU_INTRIN_IMAGE_LOAD_D16:
case AMDGPU::G_INTRINSIC: {
     }
     case Intrinsic::amdgcn_ballot: {
       unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
       unsigned SrcSize = MRI.getType(MI.getOperand(2).getReg()).getSizeInBits();
       OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::SGPRRegBankID, DstSize);
       OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VCCRegBankID, SrcSize);
       break;
     }
+    case Intrinsic::amdgcn_inverse_ballot: {
+      // This must be an SGPR, but accept a VGPR.
+      Register MaskReg = MI.getOperand(2).getReg();
+      unsigned MaskSize = MRI.getType(MaskReg).getSizeInBits();
+      unsigned MaskBank = getRegBankID(MaskReg, MRI, AMDGPU::SGPRRegBankID);
+      OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VCCRegBankID, 1);
+      OpdsMapping[2] = AMDGPU::getValueMapping(MaskBank, MaskSize);
+      break;
+    }
     }
     break;
   }
   case AMDGPU::G_AMDGPU_INTRIN_IMAGE_LOAD:
   case AMDGPU::G_AMDGPU_INTRIN_IMAGE_LOAD_D16:
   case AMDGPU::G_AMDGPU_INTRIN_IMAGE_STORE:
   case AMDGPU::G_AMDGPU_INTRIN_IMAGE_STORE_D16: {
     auto IntrID = MI.getIntrinsicID();

llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

 
 // The dummy boolean output is divergent from the IR's perspective,
 // but the mask results are uniform. These produce a divergent and
 // uniform result, so the returned struct is collectively divergent.
 // isAlwaysUniform can override the extract of the uniform component.
 def : SourceOfDivergence<int_amdgcn_if>;
 def : SourceOfDivergence<int_amdgcn_else>;
 def : SourceOfDivergence<int_amdgcn_loop>;
+def : SourceOfDivergence<int_amdgcn_inverse_ballot>;
 
 foreach intr = AMDGPUImageDimAtomicIntrinsics in
 def : SourceOfDivergence<intr>;
 
 class AlwaysUniform<Intrinsic intr> {
   Intrinsic Intr = intr;
 }
 

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

case AMDGPU::S_SETREG_B32: {
 
     // If only FP bits are touched, used the no side effects pseudo.
     if ((SetMask & (AMDGPU::Hwreg::FP_ROUND_MASK |
                     AMDGPU::Hwreg::FP_DENORM_MASK)) == SetMask)
       MI.setDesc(TII->get(AMDGPU::S_SETREG_B32_mode));
 
     return BB;
   }
+  case AMDGPU::S_INVERSE_BALLOT_U32:
+  case AMDGPU::S_INVERSE_BALLOT_U64: {
+    MachineRegisterInfo &MRI = BB->getParent()->getRegInfo();
+    const GCNSubtarget &ST = MF->getSubtarget<GCNSubtarget>();
+    const SIRegisterInfo *TRI = ST.getRegisterInfo();
+    const DebugLoc &DL = MI.getDebugLoc();
+    const Register DstReg = MI.getOperand(0).getReg();
+    Register MaskReg = MI.getOperand(1).getReg();
+
+    const bool IsVALU = TRI->isVectorRegister(MRI, MaskReg);
+
+    if (IsVALU) {
+      MaskReg = TII->readlaneVGPRToSGPR(MaskReg, MI, MRI);
+    }
+
+    BuildMI(*BB, &MI, DL, TII->get(AMDGPU::COPY), DstReg).addReg(MaskReg);
+    MI.eraseFromParent();
+    return BB;
+  }
   default:
     return AMDGPUTargetLowering::EmitInstrWithCustomInserter(MI, BB);
   }
 }
 
 bool SITargetLowering::hasAtomicFaddRtnForTy(SDValue &Op) const {
   switch (Op.getValue(0).getSimpleValueType().SimpleTy) {
   case MVT::f32:
     return Subtarget->hasAtomicFaddRtnInsts();
   case MVT::v2f16:
   case MVT::f64:
     return Subtarget->hasGFX90AInsts();
   default:
     return false;
   }
 }
 
 bool SITargetLowering::enableAggressiveFMAFusion(EVT VT) const {
   // This currently forces unfolding various combinations of fsub into fma with
   // free fneg'd operands. As long as we have fast FMA (controlled by
   // isFMAFasterThanFMulAndFAdd), we should perform these.
 
   // When fma is quarter rate, for f64 where add / sub are at best half rate,
   // most of these combines appear to be cycle neutral but save on instruction
   // count / code size.
   return true;
 }
 
 bool SITargetLowering::enableAggressiveFMAFusion(LLT Ty) const { return true; }
 
 EVT SITargetLowering::getSetCCResultType(const DataLayout &DL, LLVMContext &Ctx,
                                          EVT VT) const {
   if (!VT.isVector()) {

llvm/lib/Target/AMDGPU/SIInstructions.td

 }
 
 def EXIT_STRICT_WQM : SPseudoInstSI <(outs SReg_1:$sdst), (ins SReg_1:$src0)> {
   let hasSideEffects = 0;
   let mayLoad = 0;
   let mayStore = 0;
 }
 
+let usesCustomInserter = 1 in {
+def S_INVERSE_BALLOT_U32 : SPseudoInstSI <(outs SReg_32:$sdst), (ins SSrc_b32:$mask)>;
+
+def S_INVERSE_BALLOT_U64 : SPseudoInstSI <(outs SReg_64:$sdst), (ins SSrc_b64:$mask)>;
+} // End usesCustomInserter = 1
+
 // PSEUDO_WM is treated like STRICT_WWM/STRICT_WQM without exec changes.
 def ENTER_PSEUDO_WM : SPseudoInstSI <(outs), (ins)> {
   let Uses = [EXEC];
   let Defs = [EXEC];
   let hasSideEffects = 0;
   let mayLoad = 0;
   let mayStore = 0;
 }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.inverse.ballot.i32.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 2
+; RUN: llc -march=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -mattr=+wavefrontsize32,-wavefrontsize64 -global-isel=1 -verify-machineinstrs < %s | FileCheck  -check-prefixes=GFX11,GISEL %s
+; RUN: llc -march=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -mattr=+wavefrontsize32,-wavefrontsize64 -global-isel=0 -verify-machineinstrs < %s | FileCheck  -check-prefixes=GFX11,SDAG %s
+
+declare i1 @llvm.amdgcn.inverse.ballot(i32)
+
+; Test ballot(0)
+define amdgpu_cs void @constant_false_inverse_ballot(ptr addrspace(1) %out) {
+; GFX11-LABEL: constant_false_inverse_ballot:
+; GFX11:       ; %bb.0: ; %entry
+; GFX11-NEXT:    s_mov_b32 s0, 0
+; GFX11-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s0
+; GFX11-NEXT:    global_store_b32 v[0:1], v2, off
+; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GFX11-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 0)
+  %sel    = select i1 %ballot, i32 1, i32 0
+  store i32 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test ballot(1)
+
+define amdgpu_cs void @constant_true_inverse_ballot(ptr addrspace(1) %out) {
+; GFX11-LABEL: constant_true_inverse_ballot:
+; GFX11:       ; %bb.0: ; %entry
+; GFX11-NEXT:    s_mov_b32 s0, -1
+; GFX11-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s0
+; GFX11-NEXT:    global_store_b32 v[0:1], v2, off
+; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GFX11-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 u0xFFFFFFFF)
+  %sel    = select i1 %ballot, i32 1, i32 0
+  store i32 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+define amdgpu_cs void @constant_mask_inverse_ballot(ptr addrspace(1) %out) {
+; GFX11-LABEL: constant_mask_inverse_ballot:
+; GFX11:       ; %bb.0: ; %entry
+; GFX11-NEXT:    s_movk_i32 s0, 0x1000
+; GFX11-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s0
+; GFX11-NEXT:    global_store_b32 v[0:1], v2, off
+; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GFX11-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 u0x00001000)
+  %sel    = select i1 %ballot, i32 1, i32 0
+  store i32 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test inverse ballot using a vgpr as input
+
+define amdgpu_cs void @vgpr_inverse_ballot(i32 %input, ptr addrspace(1) %out) {
+; GFX11-LABEL: vgpr_inverse_ballot:
+; GFX11:       ; %bb.0: ; %entry
+; GFX11-NEXT:    v_readfirstlane_b32 s0, v0
+; GFX11-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s0
+; GFX11-NEXT:    global_store_b32 v[1:2], v0, off
+; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GFX11-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 %input)
+  %sel = select i1 %ballot, i32 1, i32 0
+  store i32 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+define amdgpu_cs void @sgpr_inverse_ballot(i32 inreg %input, ptr addrspace(1) %out) {
+; GFX11-LABEL: sgpr_inverse_ballot:
+; GFX11:       ; %bb.0: ; %entry
+; GFX11-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s0
+; GFX11-NEXT:    global_store_b32 v[0:1], v2, off
+; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GFX11-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 %input)
+  %sel = select i1 %ballot, i32 1, i32 0
+  store i32 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test ballot after phi
+define amdgpu_cs void @phi_uniform(i32 inreg %s0_1, i32 inreg %s2, ptr addrspace(1) %out) {
+; GFX11-LABEL: phi_uniform:
+; GFX11:       ; %bb.0: ; %entry
+; GFX11-NEXT:    s_cmp_lg_u32 s1, 0
+; GFX11-NEXT:    s_cbranch_scc1 .LBB5_2
+; GFX11-NEXT:  ; %bb.1: ; %if
+; GFX11-NEXT:    s_add_i32 s0, s0, 1
+; GFX11-NEXT:  .LBB5_2: ; %endif
+; GFX11-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s0
+; GFX11-NEXT:    global_store_b32 v[0:1], v2, off
+; GFX11-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GFX11-NEXT:    s_endpgm
+entry:
+  %cc = icmp ne i32 %s2, 0
+  br i1 %cc, label %endif, label %if
+
+if:
+  %tmp = add i32 %s0_1, 1
+  br label %endif
+
+endif:
+  %input = phi i32 [ %s0_1, %entry ], [ %tmp, %if ]
+
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 %input)
+  %sel = select i1 %ballot, i32 1, i32 0
+  store i32 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test for branching
+; GISel implementation is currently incorrect.
+; The change in the branch affects all lanes, not just the branching ones.
+; This test will be fixed once GISel correctly takes uniformity analysis into account.
+define amdgpu_cs void @inverse_ballot_branch(i32 inreg %s0_1, i32 inreg %s2, ptr addrspace(1) %out) {
+; GISEL-LABEL: inverse_ballot_branch:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    s_xor_b32 s2, s1, -1
+; GISEL-NEXT:    s_and_saveexec_b32 s1, s2
+; GISEL-NEXT:  ; %bb.1: ; %if
+; GISEL-NEXT:    s_add_i32 s0, s0, 1
+; GISEL-NEXT:  ; %bb.2: ; %endif
+; GISEL-NEXT:    s_or_b32 exec_lo, exec_lo, s1
+; GISEL-NEXT:    v_mov_b32_e32 v2, s0
+; GISEL-NEXT:    global_store_b32 v[0:1], v2, off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: inverse_ballot_branch:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    v_mov_b32_e32 v2, s0
+; SDAG-NEXT:    s_xor_b32 s2, s1, -1
+; SDAG-NEXT:    s_and_saveexec_b32 s1, s2
+; SDAG-NEXT:  ; %bb.1: ; %if
+; SDAG-NEXT:    s_add_i32 s0, s0, 1
+; SDAG-NEXT:    v_mov_b32_e32 v2, s0
+; SDAG-NEXT:  ; %bb.2: ; %endif
+; SDAG-NEXT:    s_or_b32 exec_lo, exec_lo, s1
+; SDAG-NEXT:    global_store_b32 v[0:1], v2, off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot(i32 %s2)
+  br i1 %ballot, label %endif, label %if
+
+if:
+  %tmp = add i32 %s0_1, 1
+  br label %endif
+
+endif:
+  %input = phi i32 [ %s0_1, %entry ], [ %tmp, %if ]
+  store i32 %input, ptr addrspace(1) %out
+  ret void
+}

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.inverse.ballot.i64.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 2
+; RUN: llc -march=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -mattr=-wavefrontsize32,+wavefrontsize64 -global-isel=1 -verify-machineinstrs < %s | FileCheck -check-prefixes=GISEL %s
+; RUN: llc -march=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -mattr=-wavefrontsize32,+wavefrontsize64 -global-isel=0 -verify-machineinstrs < %s | FileCheck -check-prefixes=SDAG %s
+
+declare i1 @llvm.amdgcn.inverse.ballot.i64(i64)
+
+; Test ballot(0)
+define amdgpu_cs void @constant_false_inverse_ballot(ptr addrspace(1) %out) {
+; GISEL-LABEL: constant_false_inverse_ballot:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    s_mov_b64 s[0:1], 0
+; GISEL-NEXT:    v_mov_b32_e32 v3, 0
+; GISEL-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; GISEL-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: constant_false_inverse_ballot:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    s_mov_b32 s2, 0
+; SDAG-NEXT:    s_mov_b64 s[0:1], 0
+; SDAG-NEXT:    v_mov_b32_e32 v3, s2
+; SDAG-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; SDAG-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 0)
+  %sel    = select i1 %ballot, i64 1, i64 0
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test ballot(1)
+
+define amdgpu_cs void @constant_true_inverse_ballot(ptr addrspace(1) %out) {
+; GISEL-LABEL: constant_true_inverse_ballot:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    s_mov_b64 s[0:1], -1
+; GISEL-NEXT:    v_mov_b32_e32 v3, 0
+; GISEL-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; GISEL-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: constant_true_inverse_ballot:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    s_mov_b32 s2, 0
+; SDAG-NEXT:    s_mov_b64 s[0:1], -1
+; SDAG-NEXT:    v_mov_b32_e32 v3, s2
+; SDAG-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; SDAG-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 u0xFFFFFFFFFFFFFFFF)
+  %sel    = select i1 %ballot, i64 1, i64 0
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test ballot(u0x0040F8010000)
+
+define amdgpu_cs void @constant_mask_inverse_ballot(ptr addrspace(1) %out) {
+; GISEL-LABEL: constant_mask_inverse_ballot:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    s_mov_b32 s0, 0xf8010000
+; GISEL-NEXT:    s_mov_b32 s1, 64
+; GISEL-NEXT:    v_mov_b32_e32 v3, 0
+; GISEL-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; GISEL-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: constant_mask_inverse_ballot:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    s_mov_b32 s0, 0xf8010000
+; SDAG-NEXT:    s_mov_b32 s2, 0
+; SDAG-NEXT:    s_mov_b32 s1, 64
+; SDAG-NEXT:    v_mov_b32_e32 v3, s2
+; SDAG-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; SDAG-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 u0x0040F8010000)
+  %sel    = select i1 %ballot, i64 1, i64 0
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test inverse ballot using a vgpr as input
+
+define amdgpu_cs void @vgpr_inverse_ballot(i64 %input, ptr addrspace(1) %out) {
+; GISEL-LABEL: vgpr_inverse_ballot:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    v_readfirstlane_b32 s0, v0
+; GISEL-NEXT:    v_readfirstlane_b32 s1, v1
+; GISEL-NEXT:    v_mov_b32_e32 v1, 0
+; GISEL-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GISEL-NEXT:    global_store_b64 v[2:3], v[0:1], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: vgpr_inverse_ballot:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    v_readfirstlane_b32 s0, v0
+; SDAG-NEXT:    v_readfirstlane_b32 s1, v1
+; SDAG-NEXT:    s_mov_b32 s2, 0
+; SDAG-NEXT:    v_mov_b32_e32 v1, s2
+; SDAG-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; SDAG-NEXT:    global_store_b64 v[2:3], v[0:1], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 %input)
+  %sel    = select i1 %ballot, i64 1, i64 0
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+define amdgpu_cs void @sgpr_inverse_ballot(i64 inreg %input, ptr addrspace(1) %out) {
+; GISEL-LABEL: sgpr_inverse_ballot:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; GISEL-NEXT:    v_mov_b32_e32 v3, 0
+; GISEL-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: sgpr_inverse_ballot:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; SDAG-NEXT:    s_mov_b32 s0, 0
+; SDAG-NEXT:    s_waitcnt_depctr 0xfffe
+; SDAG-NEXT:    v_mov_b32_e32 v3, s0
+; SDAG-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 %input)
+  %sel = select i1 %ballot, i64 1, i64 0
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test ballot after phi
+define amdgpu_cs void @phi_uniform(i64 inreg %s0_1, i64 inreg %s2, ptr addrspace(1) %out) {
+; GISEL-LABEL: phi_uniform:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    s_cmp_lg_u64 s[2:3], 0
+; GISEL-NEXT:    s_cbranch_scc1 .LBB5_2
+; GISEL-NEXT:  ; %bb.1: ; %if
+; GISEL-NEXT:    s_add_u32 s0, s0, 1
+; GISEL-NEXT:    s_addc_u32 s1, s1, 0
+; GISEL-NEXT:  .LBB5_2: ; %endif
+; GISEL-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; GISEL-NEXT:    v_mov_b32_e32 v3, 0
+; GISEL-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: phi_uniform:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    s_cmp_lg_u64 s[2:3], 0
+; SDAG-NEXT:    s_cbranch_scc1 .LBB5_2
+; SDAG-NEXT:  ; %bb.1: ; %if
+; SDAG-NEXT:    s_add_u32 s0, s0, 1
+; SDAG-NEXT:    s_addc_u32 s1, s1, 0
+; SDAG-NEXT:  .LBB5_2: ; %endif
+; SDAG-NEXT:    s_mov_b32 s2, 0
+; SDAG-NEXT:    v_cndmask_b32_e64 v2, 0, 1, s[0:1]
+; SDAG-NEXT:    v_mov_b32_e32 v3, s2
+; SDAG-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %cc = icmp ne i64 %s2, 0
+  br i1 %cc, label %endif, label %if
+
+if:
+  %tmp = add  i64 %s0_1, 1
+  br label %endif
+
+endif:
+  %input = phi i64 [ %s0_1, %entry ], [ %tmp, %if ]
+
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 %input)
+  %sel = select i1 %ballot, i64 1, i64 0
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}
+
+; Test for branching
+; GISel implementation is currently incorrect.
+; The change in the branch affects all lanes, not just the branching ones.
+; This test will be fixed once GISel correctly takes uniformity analysis into account.
+define amdgpu_cs void @inverse_ballot_branch(i64 inreg %s0_1, i64 inreg %s2, ptr addrspace(1) %out) {
+; GISEL-LABEL: inverse_ballot_branch:
+; GISEL:       ; %bb.0: ; %entry
+; GISEL-NEXT:    s_xor_b64 s[4:5], s[2:3], -1
+; GISEL-NEXT:    s_and_saveexec_b64 s[2:3], s[4:5]
+; GISEL-NEXT:  ; %bb.1: ; %if
+; GISEL-NEXT:    s_add_u32 s0, s0, 1
+; GISEL-NEXT:    s_addc_u32 s1, s1, 0
+; GISEL-NEXT:  ; %bb.2: ; %endif
+; GISEL-NEXT:    s_or_b64 exec, exec, s[2:3]
+; GISEL-NEXT:    v_mov_b32_e32 v3, s1
+; GISEL-NEXT:    v_mov_b32_e32 v2, s0
+; GISEL-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; GISEL-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GISEL-NEXT:    s_endpgm
+;
+; SDAG-LABEL: inverse_ballot_branch:
+; SDAG:       ; %bb.0: ; %entry
+; SDAG-NEXT:    v_mov_b32_e32 v3, s1
+; SDAG-NEXT:    v_mov_b32_e32 v2, s0
+; SDAG-NEXT:    s_xor_b64 s[4:5], s[2:3], -1
+; SDAG-NEXT:    s_and_saveexec_b64 s[2:3], s[4:5]
+; SDAG-NEXT:  ; %bb.1: ; %if
+; SDAG-NEXT:    s_add_u32 s0, s0, 1
+; SDAG-NEXT:    s_addc_u32 s1, s1, 0
+; SDAG-NEXT:    v_mov_b32_e32 v3, s1
+; SDAG-NEXT:    v_mov_b32_e32 v2, s0
+; SDAG-NEXT:  ; %bb.2: ; %endif
+; SDAG-NEXT:    s_or_b64 exec, exec, s[2:3]
+; SDAG-NEXT:    global_store_b64 v[0:1], v[2:3], off
+; SDAG-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; SDAG-NEXT:    s_endpgm
+entry:
+  %ballot = call i1 @llvm.amdgcn.inverse.ballot.i64(i64 %s2)
+  br i1 %ballot, label %endif, label %if
+
+if:
+  %tmp = add  i64 %s0_1, 1
+  br label %endif
+
+endif:
+  %sel = phi i64 [ %s0_1, %entry ], [ %tmp, %if ]
+  store i64 %sel, ptr addrspace(1) %out
+  ret void
+}