[AMDGPU] Add GFX11 ds_bvh_stack_rtn_b32 instruction

[arsenm]

I don't see why this needs manual selection

[foad]

In D133928#3792158, @arsenm wrote:

I don't see why this needs manual selection

The comments say "tablegen doesn't support matching instructions with multiple outputs". Is that not true? (I'm not the original author of this patch.)

[arsenm]

In D133928#3792248, @foad wrote:

In D133928#3792158, @arsenm wrote:

I don't see why this needs manual selection

The comments say "tablegen doesn't support matching instructions with multiple outputs". Is that not true? (I'm not the original author of this patch.)

I missed the second return value somehow

llvm/lib/Target/AMDGPU/DSInstructions.td

289

Does this really have side effects?

[foad]

llvm/lib/Target/AMDGPU/DSInstructions.td

289

It reads and writes LDS in a complicated way. Previously we've used hasSideEffects=1 to model complicated GDS access, e.g. on the DS_ADD/SUB_GS_REG_RTN instructions. Is there a better way? Should we be using MMOs with addrspace 2/3 but no pointer info instead? If so I'd prefer to do that as a follow up.

[arsenm]

llvm/lib/Target/AMDGPU/DSInstructions.td

289

Yes, MMOs would be better

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_ds_add_gs_reg_rtn :
   Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
             [ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;
 
 def int_amdgcn_ds_sub_gs_reg_rtn :
   ClangBuiltin<"__builtin_amdgcn_ds_sub_gs_reg_rtn">,
   Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
             [ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;
 
+def int_amdgcn_ds_bvh_stack_rtn :
+  Intrinsic<
+    [llvm_i32_ty, llvm_i32_ty], // %vdst, %addr
+    [
+      llvm_i32_ty,   // %addr
+      llvm_i32_ty,   // %data0
+      llvm_v4i32_ty, // %data1
+      llvm_i32_ty,   // %offset
+    ],
+    [ImmArg<ArgIndex<3>>, IntrWillReturn]
+  >;
+
 // WMMA (Wave Matrix Multiply-Accumulate) intrinsics
 //
 // These operations perform a matrix multiplication and accumulation of
 // the form: D = A * B + C .
 
 class AMDGPUWmmaIntrinsic<LLVMType AB, LLVMType CD> :
   Intrinsic<
     [CD],               // %D

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.h

private:
   SDNode *getBFE32(bool IsSigned, const SDLoc &DL, SDValue Val, uint32_t Offset,
                    uint32_t Width);
   void SelectS_BFEFromShifts(SDNode *N);
   void SelectS_BFE(SDNode *N);
   bool isCBranchSCC(const SDNode *N) const;
   void SelectBRCOND(SDNode *N);
   void SelectFMAD_FMA(SDNode *N);
   void SelectDSAppendConsume(SDNode *N, unsigned IntrID);
+  void SelectDSBvhStackIntrinsic(SDNode *N);
   void SelectDS_GWS(SDNode *N, unsigned IntrID);
   void SelectInterpP1F16(SDNode *N);
   void SelectINTRINSIC_W_CHAIN(SDNode *N);
   void SelectINTRINSIC_WO_CHAIN(SDNode *N);
   void SelectINTRINSIC_VOID(SDNode *N);
 
 protected:
   // Include the pieces autogenerated from the target description.
 #include "AMDGPUGenDAGISel.inc"
 };
 
 #endif // LLVM_LIB_TARGET_AMDGPU_AMDGPUISELDAGTODAG_H

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

 
 using namespace llvm;
 
 //===----------------------------------------------------------------------===//
 // Instruction Selector Implementation
 //===----------------------------------------------------------------------===//
 
 namespace {
-
 static SDValue stripBitcast(SDValue Val) {
   return Val.getOpcode() == ISD::BITCAST ? Val.getOperand(0) : Val;
 }
 
 // Figure out if this is really an extract of the high 16-bits of a dword.
 static bool isExtractHiElt(SDValue In, SDValue &Out) {
   In = stripBitcast(In);
 
if (In.getOpcode() == ISD::TRUNCATE) {
     SDValue Src = In.getOperand(0);
     if (Src.getValueType().getSizeInBits() == 32)
       return stripBitcast(Src);
   }
 
   return In;
 }
 
-}  // end anonymous namespace
+} // end anonymous namespace
 
 INITIALIZE_PASS_BEGIN(AMDGPUDAGToDAGISel, "amdgpu-isel",
                       "AMDGPU DAG->DAG Pattern Instruction Selection", false, false)
 INITIALIZE_PASS_DEPENDENCY(AMDGPUArgumentUsageInfo)
 INITIALIZE_PASS_DEPENDENCY(AMDGPUPerfHintAnalysis)
 INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis)
 #ifdef EXPENSIVE_CHECKS
 INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass)
SDValue Ops[] = {
     Chain,
     N->getOperand(N->getNumOperands() - 1) // New glue
   };
 
   SDNode *Selected = CurDAG->SelectNodeTo(N, Opc, N->getVTList(), Ops);
   CurDAG->setNodeMemRefs(cast<MachineSDNode>(Selected), {MMO});
 }
 
+// We need to handle this here because tablegen doesn't support matching
+// instructions with multiple outputs.
+void AMDGPUDAGToDAGISel::SelectDSBvhStackIntrinsic(SDNode *N) {
+  unsigned Opc = AMDGPU::DS_BVH_STACK_RTN_B32;
+  SDValue Ops[] = {N->getOperand(2), N->getOperand(3), N->getOperand(4),
+                   N->getOperand(5), N->getOperand(0)};
+
+  MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
+  MachineMemOperand *MMO = M->getMemOperand();
+  SDNode *Selected = CurDAG->SelectNodeTo(N, Opc, N->getVTList(), Ops);
+  CurDAG->setNodeMemRefs(cast<MachineSDNode>(Selected), {MMO});
+}
+
 static unsigned gwsIntrinToOpcode(unsigned IntrID) {
   switch (IntrID) {
   case Intrinsic::amdgcn_ds_gws_init:
     return AMDGPU::DS_GWS_INIT;
   case Intrinsic::amdgcn_ds_gws_barrier:
     return AMDGPU::DS_GWS_BARRIER;
   case Intrinsic::amdgcn_ds_gws_sema_v:
     return AMDGPU::DS_GWS_SEMA_V;
void AMDGPUDAGToDAGISel::SelectINTRINSIC_W_CHAIN(SDNode *N) {
   switch (IntrID) {
   case Intrinsic::amdgcn_ds_append:
   case Intrinsic::amdgcn_ds_consume: {
     if (N->getValueType(0) != MVT::i32)
       break;
     SelectDSAppendConsume(N, IntrID);
     return;
   }
+  case Intrinsic::amdgcn_ds_bvh_stack_rtn:
+    SelectDSBvhStackIntrinsic(N);
+    return;
   }
 
   SelectCode(N);
 }
 
 void AMDGPUDAGToDAGISel::SelectINTRINSIC_WO_CHAIN(SDNode *N) {
   unsigned IntrID = cast<ConstantSDNode>(N->getOperand(0))->getZExtValue();
   unsigned Opcode;

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.h

private:
   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;
   bool selectDSAppendConsume(MachineInstr &MI, bool IsAppend) const;
   bool selectSBarrier(MachineInstr &MI) const;
+  bool selectDSBvhStackIntrinsic(MachineInstr &MI) const;
 
   bool selectImageIntrinsic(MachineInstr &MI,
                             const AMDGPU::ImageDimIntrinsicInfo *Intr) const;
   bool selectG_INTRINSIC_W_SIDE_EFFECTS(MachineInstr &I) const;
   int getS_CMPOpcode(CmpInst::Predicate P, unsigned Size) const;
   bool selectG_ICMP(MachineInstr &I) const;
   bool hasVgprParts(ArrayRef<GEPInfo> AddrInfo) const;
   void getAddrModeInfo(const MachineInstr &Load, const MachineRegisterInfo &MRI,

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

bool AMDGPUInstructionSelector::selectImageIntrinsic(
   }
 
   MI.eraseFromParent();
   constrainSelectedInstRegOperands(*MIB, TII, TRI, RBI);
   TII.enforceOperandRCAlignment(*MIB, AMDGPU::OpName::vaddr);
   return true;
 }
 
+// We need to handle this here because tablegen doesn't support matching
+// instructions with multiple outputs.
+bool AMDGPUInstructionSelector::selectDSBvhStackIntrinsic(
+    MachineInstr &MI) const {
+  Register Dst0 = MI.getOperand(0).getReg();
+  Register Dst1 = MI.getOperand(1).getReg();
+
+  const DebugLoc &DL = MI.getDebugLoc();
+  MachineBasicBlock *MBB = MI.getParent();
+
+  Register Addr = MI.getOperand(3).getReg();
+  Register Data0 = MI.getOperand(4).getReg();
+  Register Data1 = MI.getOperand(5).getReg();
+  unsigned Offset = MI.getOperand(6).getImm();
+
+  auto MIB = BuildMI(*MBB, &MI, DL, TII.get(AMDGPU::DS_BVH_STACK_RTN_B32), Dst0)
+                 .addDef(Dst1)
+                 .addUse(Addr)
+                 .addUse(Data0)
+                 .addUse(Data1)
+                 .addImm(Offset)
+                 .cloneMemRefs(MI);
+
+  MI.eraseFromParent();
+  return constrainSelectedInstRegOperands(*MIB, TII, TRI, RBI);
+}
+
 bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS(
     MachineInstr &I) const {
   unsigned IntrinsicID = I.getIntrinsicID();
   switch (IntrinsicID) {
   case Intrinsic::amdgcn_end_cf:
     return selectEndCfIntrinsic(I);
   case Intrinsic::amdgcn_ds_ordered_add:
   case Intrinsic::amdgcn_ds_ordered_swap:
case Intrinsic::amdgcn_exp_compr:
     if (!STI.hasCompressedExport()) {
       Function &F = I.getMF()->getFunction();
       DiagnosticInfoUnsupported NoFpRet(
           F, "intrinsic not supported on subtarget", I.getDebugLoc(), DS_Error);
       F.getContext().diagnose(NoFpRet);
       return false;
     }
     break;
+  case Intrinsic::amdgcn_ds_bvh_stack_rtn:
+    return selectDSBvhStackIntrinsic(I);
   }
   return selectImpl(I, *CoverageInfo);
 }
 
 bool AMDGPUInstructionSelector::selectG_SELECT(MachineInstr &I) const {
   if (selectImpl(I, *CoverageInfo))
     return true;
 

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

case Intrinsic::amdgcn_lds_direct_load: {
       OpdsMapping[M0Idx] = AMDGPU::getValueMapping(M0Bank, 32);
       break;
     }
     case Intrinsic::amdgcn_ds_add_gs_reg_rtn:
     case Intrinsic::amdgcn_ds_sub_gs_reg_rtn:
       OpdsMapping[0] = getVGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI);
       OpdsMapping[2] = getVGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
       break;
+    case Intrinsic::amdgcn_ds_bvh_stack_rtn: {
+      OpdsMapping[0] =
+          getVGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI); // %vdst
+      OpdsMapping[1] =
+          getVGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI); // %addr
+      OpdsMapping[3] =
+          getVGPROpMapping(MI.getOperand(3).getReg(), MRI, *TRI); // %addr
+      OpdsMapping[4] =
+          getVGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI); // %data0
+      OpdsMapping[5] =
+          getVGPROpMapping(MI.getOperand(5).getReg(), MRI, *TRI); // %data1
+      break;
+    }
+
     default:
       return getInvalidInstructionMapping();
     }
     break;
   }
   case AMDGPU::G_SELECT: {
     unsigned Size = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
     unsigned Op2Bank = getRegBankID(MI.getOperand(2).getReg(), MRI,

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

void AMDGPUAsmParser::cvtDSOffset01(MCInst &Inst,
   addOptionalImmOperand(Inst, Operands, OptionalIdx, AMDGPUOperand::ImmTyGDS);
 
   Inst.addOperand(MCOperand::createReg(AMDGPU::M0)); // m0
 }
 
 void AMDGPUAsmParser::cvtDSImpl(MCInst &Inst, const OperandVector &Operands,
                                 bool IsGdsHardcoded) {
   OptionalImmIndexMap OptionalIdx;
+  const MCInstrDesc &Desc = MII.get(Inst.getOpcode());
   AMDGPUOperand::ImmTy OffsetType = AMDGPUOperand::ImmTyOffset;
 
   for (unsigned i = 1, e = Operands.size(); i != e; ++i) {
     AMDGPUOperand &Op = ((AMDGPUOperand &)*Operands[i]);
 
+    auto TiedTo =
+        Desc.getOperandConstraint(Inst.getNumOperands(), MCOI::TIED_TO);
+
+    if (TiedTo != -1) {
+      assert((unsigned)TiedTo < Inst.getNumOperands());
+      Inst.addOperand(Inst.getOperand(TiedTo));
+    }
+
     // Add the register arguments
     if (Op.isReg()) {
       Op.addRegOperands(Inst, 1);
       continue;
     }
 
     if (Op.isToken() && Op.getToken() == "gds") {
       IsGdsHardcoded = true;

llvm/lib/Target/AMDGPU/DSInstructions.td

multiclass DS_1A2D_Off8_RET_mc<string opName,
                                RegisterClass src = rc> {
   def "" : DS_1A2D_Off8_RET<opName, rc, src>;
 
   let has_m0_read = 0 in {
     def _gfx9 : DS_1A2D_Off8_RET<opName, rc, src>;
   }
 }
 
+class DS_BVH_STACK<string opName>
+: DS_Pseudo<opName,
+  (outs getLdStRegisterOperand<VGPR_32>.ret:$vdst, VGPR_32:$addr),
+  (ins VGPR_32:$addr_in, getLdStRegisterOperand<VGPR_32>.ret:$data0, VReg_128:$data1, offset:$offset),
+  " $vdst, $addr, $data0, $data1$offset"> {
+  let Constraints = "$addr = $addr_in";
+  let DisableEncoding = "$addr_in";
+  let has_gds = 0;
+  let gdsValue = 0;
+  // TODO: Use MMOs in the LDS address space instead of hasSideEffects = 1.
+  let hasSideEffects = 1;
+  let SchedRW = [WriteLDS, WriteLDS];
+}
 
 class DS_1A_RET<string opName, RegisterClass rc = VGPR_32, bit HasTiedOutput = 0, Operand ofs = offset,
                 RegisterOperand data_op = getLdStRegisterOperand<rc>.ret>
 : DS_Pseudo<opName,
   (outs data_op:$vdst),
   !if(HasTiedOutput,
     (ins VGPR_32:$addr, ofs:$offset, gds:$gds, data_op:$vdst_in),
     (ins VGPR_32:$addr, ofs:$offset, gds:$gds)),

 //===----------------------------------------------------------------------===//
 // Instruction definitions for GFX11 and newer.
 //===----------------------------------------------------------------------===//
 
 let SubtargetPredicate = isGFX11Plus in {
 
 def DS_ADD_GS_REG_RTN : DS_0A1D_RET_GDS<"ds_add_gs_reg_rtn", VReg_64, VGPR_32>;
 def DS_SUB_GS_REG_RTN : DS_0A1D_RET_GDS<"ds_sub_gs_reg_rtn", VReg_64, VGPR_32>;
+def DS_BVH_STACK_RTN_B32 : DS_BVH_STACK<"ds_bvh_stack_rtn_b32">;
 
 } // let SubtargetPredicate = isGFX11Plus
 
 //===----------------------------------------------------------------------===//
 // DS Patterns
 //===----------------------------------------------------------------------===//
 
 def : GCNPat <

 defm DS_CMPSTORE_B64                     : DS_Real_gfx11<0x050>;
 defm DS_CMPSTORE_F64                     : DS_Real_gfx11<0x051>;
 defm DS_CMPSTORE_RTN_B64                 : DS_Real_gfx11<0x070>;
 defm DS_CMPSTORE_RTN_F64                 : DS_Real_gfx11<0x071>;
 
 defm DS_ADD_RTN_F32                      : DS_Real_gfx11<0x079>;
 defm DS_ADD_GS_REG_RTN                   : DS_Real_gfx11<0x07a>;
 defm DS_SUB_GS_REG_RTN                   : DS_Real_gfx11<0x07b>;
+defm DS_BVH_STACK_RTN_B32                : DS_Real_gfx11<0x0ad>;
 
 //===----------------------------------------------------------------------===//
 // GFX10.
 //===----------------------------------------------------------------------===//
 
 let AssemblerPredicate = isGFX10Only, DecoderNamespace = "GFX10" in {
   multiclass DS_Real_gfx10<bits<8> op>  {
     def _gfx10 : Base_DS_Real_gfx6_gfx7_gfx10_gfx11<op, !cast<DS_Pseudo>(NAME),

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

bool SITargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info,
   case Intrinsic::amdgcn_global_load_lds: {
     Info.opc = ISD::INTRINSIC_VOID;
     unsigned Width = cast<ConstantInt>(CI.getArgOperand(2))->getZExtValue();
     Info.memVT = EVT::getIntegerVT(CI.getContext(), Width * 8);
     Info.flags |= MachineMemOperand::MOLoad | MachineMemOperand::MOStore |
                   MachineMemOperand::MOVolatile;
     return true;
   }
+  case Intrinsic::amdgcn_ds_bvh_stack_rtn: {
+    Info.opc = ISD::INTRINSIC_W_CHAIN;
+
+    const GCNTargetMachine &TM =
+        static_cast<const GCNTargetMachine &>(getTargetMachine());
+
+    SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
+    Info.ptrVal = MFI->getGWSPSV(TM);
+
+    // This is an abstract access, but we need to specify a type and size.
+    Info.memVT = MVT::i32;
+    Info.size = 4;
+    Info.align = Align(4);
+
+    Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore;
+    return true;
+  }
   default:
     return false;
   }
 }
 
 bool SITargetLowering::getAddrModeArguments(IntrinsicInst *II,
                                             SmallVectorImpl<Value*> &Ops,
                                             Type *&AccessTy) const {

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bvh.stack.rtn.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck %s
+
+declare { i32, i32 } @llvm.amdgcn.ds.bvh.stack.rtn(i32, i32, <4 x i32>, i32 immarg)
+
+define amdgpu_gs void @test_ds_bvh_stack(i32 %addr, i32 %data0, <4 x i32> %data1, i32 addrspace(1)* %out) {
+; CHECK-LABEL: test_ds_bvh_stack:
+; CHECK:       ; %bb.0:
+; CHECK-NEXT:    ds_bvh_stack_rtn_b32 v1, v0, v1, v[2:5]
+; CHECK-NEXT:    s_waitcnt lgkmcnt(0)
+; CHECK-NEXT:    v_add_nc_u32_e32 v0, v1, v0
+; CHECK-NEXT:    global_store_b32 v[6:7], v0, off
+; CHECK-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; CHECK-NEXT:    s_endpgm
+  %pair = call { i32, i32 } @llvm.amdgcn.ds.bvh.stack.rtn(i32 %addr, i32 %data0, <4 x i32> %data1, i32 0)
+  %vdst = extractvalue { i32, i32 } %pair, 0
+  %newaddr = extractvalue { i32, i32 } %pair, 1
+  %res = add i32 %vdst, %newaddr
+  store i32 %res, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+define amdgpu_gs void @test_ds_bvh_stack_1(i32 %addr, i32 %data0, <4 x i32> %data1, i32 addrspace(1)* %out) {
+; CHECK-LABEL: test_ds_bvh_stack_1:
+; CHECK:       ; %bb.0:
+; CHECK-NEXT:    ds_bvh_stack_rtn_b32 v1, v0, v1, v[2:5] offset:1
+; CHECK-NEXT:    s_waitcnt lgkmcnt(0)
+; CHECK-NEXT:    v_add_nc_u32_e32 v0, v1, v0
+; CHECK-NEXT:    global_store_b32 v[6:7], v0, off
+; CHECK-NEXT:    s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; CHECK-NEXT:    s_endpgm
+  %pair = call { i32, i32 } @llvm.amdgcn.ds.bvh.stack.rtn(i32 %addr, i32 %data0, <4 x i32> %data1, i32 1)
+  %vdst = extractvalue { i32, i32 } %pair, 0
+  %newaddr = extractvalue { i32, i32 } %pair, 1
+  %res = add i32 %vdst, %newaddr
+  store i32 %res, i32 addrspace(1)* %out, align 4
+  ret void
+}

llvm/test/MC/AMDGPU/gfx11_asm_ds.s

 ds_swizzle_b32 v8, v2 offset:swizzle(BROADCAST,4,1)
 // GFX11: [0x3c,0x00,0xd4,0xd8,0x02,0x00,0x00,0x08]
 
 ds_swizzle_b32 v8, v2 offset:swizzle(BROADCAST,8,7)
 // GFX11: [0xf8,0x00,0xd4,0xd8,0x02,0x00,0x00,0x08]
 
 ds_swizzle_b32 v8, v2 offset:swizzle(BITMASK_PERM, "01pip")
 // GFX11: [0x07,0x09,0xd4,0xd8,0x02,0x00,0x00,0x08]
+
+ds_bvh_stack_rtn_b32 v255, v254, v253, v[249:252]
+// GFX11: [0x00,0x00,0xb4,0xda,0xfe,0xfd,0xf9,0xff]
+
+ds_bvh_stack_rtn_b32 v255, v254, v253, v[249:252] offset:127
+// GFX11: [0x7f,0x00,0xb4,0xda,0xfe,0xfd,0xf9,0xff]

llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_ds.txt

 # GFX11: ds_sub_gs_reg_rtn v[5:6], v1 gds     ; encoding: [0x00,0x00,0xee,0xd9,0x00,0x01,0x00,0x05]
 0x00,0x00,0xee,0xd9,0x00,0x01,0x00,0x05
 
 # GFX11: ds_sub_gs_reg_rtn v[1:2], v255 offset:65535 gds     ; encoding: [0xff,0xff,0xee,0xd9,0x00,0xff,0x00,0x01]
 0xff,0xff,0xee,0xd9,0x00,0xff,0x00,0x01
 
 # GFX11: ds_sub_gs_reg_rtn v[1:2], v255 gds     ; encoding: [0x00,0x00,0xee,0xd9,0x00,0xff,0x00,0x01]
 0x00,0x00,0xee,0xd9,0x00,0xff,0x00,0x01
+
+# GFX11: ds_bvh_stack_rtn_b32 v1, v2, v3, v[4:7] offset:127     ; encoding: [0x7f,0x00,0xb4,0xda,0x02,0x03,0x04,0x01]
+0x7f,0x00,0xb4,0xda,0x02,0x03,0x04,0x01
+
+# GFX11: ds_bvh_stack_rtn_b32 v1, v2, v3, v[4:7] ; encoding: [0x00,0x00,0xb4,0xda,0x02,0x03,0x04,0x01]
+0x00,0x00,0xb4,0xda,0x02,0x03,0x04,0x01
+
+# GFX11: ds_bvh_stack_rtn_b32 v254, v255, v253, v[5:8] offset:127     ; encoding: [0x7f,0x00,0xb4,0xda,0xff,0xfd,0x05,0xfe]
+0x7f,0x00,0xb4,0xda,0xff,0xfd,0x05,0xfe
+
+# GFX11: ds_bvh_stack_rtn_b32 v254, v255, v253, v[5:8] ; encoding: [0x00,0x00,0xb4,0xda,0xff,0xfd,0x05,0xfe]
+0x00,0x00,0xb4,0xda,0xff,0xfd,0x05,0xfe