Skip to content

Commit 575fde0

Browse files
mariusz-sikora-at-amdmatejaMarjanovic
mariusz-sikora-at-amd
and
matejaMarjanovic
authored
[AMDGPU] Add intrinsic and MI for image_bvh_dual_intersect_ray (#130038)
- Add llvm.amdgcn.image.bvh.dual.intersect.ray intrinsic and image_bvh_dual_intersect_ray machine instruction. - Add llvm_v10i32_ty and llvm_v10f32_ty --------- Co-authored-by: Mateja Marjanovic <[email protected]>
1 parent f017073 commit 575fde0

16 files changed

+302
-30
lines changed

‎llvm/include/llvm/IR/Intrinsics.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -561,6 +561,7 @@ def llvm_v3i32_ty : LLVMType<v3i32>; // 3 x i32
561561
def llvm_v4i32_ty : LLVMType<v4i32>; // 4 x i32
562562
def llvm_v6i32_ty : LLVMType<v6i32>; // 6 x i32
563563
def llvm_v8i32_ty : LLVMType<v8i32>; // 8 x i32
564+
def llvm_v10i32_ty : LLVMType<v10i32>; // 10 x i32
564565
def llvm_v16i32_ty : LLVMType<v16i32>; // 16 x i32
565566
def llvm_v32i32_ty : LLVMType<v32i32>; // 32 x i32
566567
def llvm_v64i32_ty : LLVMType<v64i32>; // 64 x i32
@@ -591,6 +592,7 @@ def llvm_v2f32_ty : LLVMType<v2f32>; // 2 x float
591592
def llvm_v3f32_ty : LLVMType<v3f32>; // 3 x float
592593
def llvm_v4f32_ty : LLVMType<v4f32>; // 4 x float
593594
def llvm_v8f32_ty : LLVMType<v8f32>; // 8 x float
595+
def llvm_v10f32_ty : LLVMType<v10f32>; // 10 x float
594596
def llvm_v16f32_ty : LLVMType<v16f32>; // 16 x float
595597
def llvm_v32f32_ty : LLVMType<v32f32>; // 32 x float
596598
def llvm_v1f64_ty : LLVMType<v1f64>; // 1 x double

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2813,6 +2813,17 @@ def int_amdgcn_ds_bvh_stack_push8_pop1_rtn : IntDSBVHStackRtn<vdst = llvm_i32_t
28132813
def int_amdgcn_ds_bvh_stack_push8_pop2_rtn : IntDSBVHStackRtn<vdst = llvm_i64_ty,
28142814
data1 = llvm_v8i32_ty>;
28152815

2816+
// <vdata>, <ray_origin>, <ray_dir>
2817+
// llvm.amdgcn.image.bvh.dual.intersect.ray <node_ptr>, <ray_extent>,
2818+
// <instance_mask>, <ray_origin>,
2819+
// <ray_dir>, <offsets>,
2820+
// <texture_descr>
2821+
def int_amdgcn_image_bvh_dual_intersect_ray :
2822+
Intrinsic<[llvm_v10i32_ty, llvm_v3f32_ty, llvm_v3f32_ty],
2823+
[llvm_i64_ty, llvm_float_ty, llvm_i8_ty, llvm_v3f32_ty,
2824+
llvm_v3f32_ty, llvm_v2i32_ty, llvm_v4i32_ty],
2825+
[IntrReadMem, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
2826+
28162827
// llvm.amdgcn.permlane16.var <old> <src0> <src1> <fi> <bound_control>
28172828
def int_amdgcn_permlane16_var : ClangBuiltin<"__builtin_amdgcn_permlane16_var">,
28182829
Intrinsic<[llvm_i32_ty],

‎llvm/lib/Target/AMDGPU/AMDGPU.td

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1094,6 +1094,12 @@ def FeaturePrngInst : SubtargetFeature<"prng-inst",
10941094
"Has v_prng_b32 instruction"
10951095
>;
10961096

1097+
def FeatureBVHDualInst : SubtargetFeature<"bvh-dual-inst",
1098+
"HasBVHDualInst",
1099+
"true",
1100+
"Has image_bvh_dual_intersect_ray instruction"
1101+
>;
1102+
10971103
//===------------------------------------------------------------===//
10981104
// Subtarget Features (options and debugging)
10991105
//===------------------------------------------------------------===//
@@ -1844,7 +1850,8 @@ def FeatureISAVersion12 : FeatureSet<
18441850
FeatureDPPSrc1SGPR,
18451851
FeatureMaxHardClauseLength32,
18461852
Feature1_5xVGPRs,
1847-
FeatureMemoryAtomicFAddF32DenormalSupport
1853+
FeatureMemoryAtomicFAddF32DenormalSupport,
1854+
FeatureBVHDualInst
18481855
]>;
18491856

18501857
def FeatureISAVersion12_Generic: FeatureSet<
@@ -2500,6 +2507,9 @@ def HasBitOp3Insts : Predicate<"Subtarget->hasBitOp3Insts()">,
25002507
def HasPrngInst : Predicate<"Subtarget->hasPrngInst()">,
25012508
AssemblerPredicate<(all_of FeaturePrngInst)>;
25022509

2510+
def HasBVHDualInst : Predicate<"Subtarget->hasBVHDualInst()">,
2511+
AssemblerPredicate<(all_of FeatureBVHDualInst)>;
2512+
25032513
def HasFP8ConversionScaleInsts : Predicate<"Subtarget->hasFP8ConversionScaleInsts()">,
25042514
AssemblerPredicate<(all_of FeatureFP8ConversionScaleInsts)>;
25052515

‎llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3597,10 +3597,12 @@ bool AMDGPUInstructionSelector::selectGlobalLoadLds(MachineInstr &MI) const{
35973597

35983598
bool AMDGPUInstructionSelector::selectBVHIntersectRayIntrinsic(
35993599
MachineInstr &MI) const {
3600-
MI.setDesc(TII.get(MI.getOperand(1).getImm()));
3601-
MI.removeOperand(1);
3600+
unsigned OpcodeOpIdx =
3601+
MI.getOpcode() == AMDGPU::G_AMDGPU_BVH_INTERSECT_RAY ? 1 : 3;
3602+
MI.setDesc(TII.get(MI.getOperand(OpcodeOpIdx).getImm()));
3603+
MI.removeOperand(OpcodeOpIdx);
36023604
MI.addImplicitDefUseOperands(*MI.getParent()->getParent());
3603-
return true;
3605+
return constrainSelectedInstRegOperands(MI, TII, TRI, RBI);
36043606
}
36053607

36063608
// FIXME: This should be removed and let the patterns select. We just need the
@@ -4114,6 +4116,7 @@ bool AMDGPUInstructionSelector::select(MachineInstr &I) {
41144116
assert(Intr && "not an image intrinsic with image pseudo");
41154117
return selectImageIntrinsic(I, Intr);
41164118
}
4119+
case AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY:
41174120
case AMDGPU::G_AMDGPU_BVH_INTERSECT_RAY:
41184121
return selectBVHIntersectRayIntrinsic(I);
41194122
case AMDGPU::G_SBFX:

‎llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7198,6 +7198,57 @@ bool AMDGPULegalizerInfo::legalizeBVHIntersectRayIntrinsic(
71987198
return true;
71997199
}
72007200

7201+
bool AMDGPULegalizerInfo::legalizeBVHDualIntrinsic(MachineInstr &MI,
7202+
MachineIRBuilder &B) const {
7203+
const LLT S32 = LLT::scalar(32);
7204+
const LLT V2S32 = LLT::fixed_vector(2, 32);
7205+
7206+
Register DstReg = MI.getOperand(0).getReg();
7207+
Register DstOrigin = MI.getOperand(1).getReg();
7208+
Register DstDir = MI.getOperand(2).getReg();
7209+
Register NodePtr = MI.getOperand(4).getReg();
7210+
Register RayExtent = MI.getOperand(5).getReg();
7211+
Register InstanceMask = MI.getOperand(6).getReg();
7212+
Register RayOrigin = MI.getOperand(7).getReg();
7213+
Register RayDir = MI.getOperand(8).getReg();
7214+
Register Offsets = MI.getOperand(9).getReg();
7215+
Register TDescr = MI.getOperand(10).getReg();
7216+
7217+
if (!ST.hasBVHDualInst()) {
7218+
DiagnosticInfoUnsupported BadIntrin(B.getMF().getFunction(),
7219+
"intrinsic not supported on subtarget",
7220+
MI.getDebugLoc());
7221+
B.getMF().getFunction().getContext().diagnose(BadIntrin);
7222+
return false;
7223+
}
7224+
7225+
const unsigned NumVDataDwords = 10;
7226+
const unsigned NumVAddrDwords = 12;
7227+
int Opcode = AMDGPU::getMIMGOpcode(AMDGPU::IMAGE_BVH_DUAL_INTERSECT_RAY,
7228+
AMDGPU::MIMGEncGfx12, NumVDataDwords,
7229+
NumVAddrDwords);
7230+
assert(Opcode != -1);
7231+
7232+
auto RayExtentInstanceMaskVec = B.buildMergeLikeInstr(
7233+
V2S32, {RayExtent, B.buildAnyExt(S32, InstanceMask)});
7234+
7235+
B.buildInstr(AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY)
7236+
.addDef(DstReg)
7237+
.addDef(DstOrigin)
7238+
.addDef(DstDir)
7239+
.addImm(Opcode)
7240+
.addUse(NodePtr)
7241+
.addUse(RayExtentInstanceMaskVec.getReg(0))
7242+
.addUse(RayOrigin)
7243+
.addUse(RayDir)
7244+
.addUse(Offsets)
7245+
.addUse(TDescr)
7246+
.cloneMemRefs(MI);
7247+
7248+
MI.eraseFromParent();
7249+
return true;
7250+
}
7251+
72017252
bool AMDGPULegalizerInfo::legalizeStackSave(MachineInstr &MI,
72027253
MachineIRBuilder &B) const {
72037254
const SITargetLowering *TLI = ST.getTargetLowering();
@@ -7546,6 +7597,8 @@ bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper,
75467597
return legalizeRsqClampIntrinsic(MI, MRI, B);
75477598
case Intrinsic::amdgcn_image_bvh_intersect_ray:
75487599
return legalizeBVHIntersectRayIntrinsic(MI, B);
7600+
case Intrinsic::amdgcn_image_bvh_dual_intersect_ray:
7601+
return legalizeBVHDualIntrinsic(MI, B);
75497602
case Intrinsic::amdgcn_swmmac_f16_16x16x32_f16:
75507603
case Intrinsic::amdgcn_swmmac_bf16_16x16x32_bf16:
75517604
case Intrinsic::amdgcn_swmmac_f32_16x16x32_bf16:

‎llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -208,6 +208,8 @@ class AMDGPULegalizerInfo final : public LegalizerInfo {
208208
bool legalizeBVHIntersectRayIntrinsic(MachineInstr &MI,
209209
MachineIRBuilder &B) const;
210210

211+
bool legalizeBVHDualIntrinsic(MachineInstr &MI, MachineIRBuilder &B) const;
212+
211213
bool legalizeLaneOp(LegalizerHelper &Helper, MachineInstr &MI,
212214
Intrinsic::ID IID) const;
213215

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

Lines changed: 26 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -3239,10 +3239,13 @@ void AMDGPURegisterBankInfo::applyMappingImpl(
32393239
applyMappingImage(B, MI, OpdMapper, RSrcIntrin->RsrcArg);
32403240
return;
32413241
}
3242-
case AMDGPU::G_AMDGPU_BVH_INTERSECT_RAY: {
3243-
unsigned N = MI.getNumExplicitOperands() - 2;
3242+
case AMDGPU::G_AMDGPU_BVH_INTERSECT_RAY:
3243+
case AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY: {
3244+
bool IsDual = MI.getOpcode() == AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY;
3245+
unsigned NumMods = IsDual ? 0 : 1; // Has A16 modifier
3246+
unsigned LastRegOpIdx = MI.getNumExplicitOperands() - 1 - NumMods;
32443247
applyDefaultMapping(OpdMapper);
3245-
executeInWaterfallLoop(B, MI, {N});
3248+
executeInWaterfallLoop(B, MI, {LastRegOpIdx});
32463249
return;
32473250
}
32483251
case AMDGPU::G_INTRINSIC_W_SIDE_EFFECTS:
@@ -5032,19 +5035,33 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
50325035
assert(RSrcIntrin->IsImage);
50335036
return getImageMapping(MRI, MI, RSrcIntrin->RsrcArg);
50345037
}
5035-
case AMDGPU::G_AMDGPU_BVH_INTERSECT_RAY: {
5036-
unsigned N = MI.getNumExplicitOperands() - 2;
5037-
OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 128);
5038-
OpdsMapping[N] = getSGPROpMapping(MI.getOperand(N).getReg(), MRI, *TRI);
5039-
if (N == 3) {
5038+
case AMDGPU::G_AMDGPU_BVH_INTERSECT_RAY:
5039+
case AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY: {
5040+
bool IsDual = MI.getOpcode() == AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY;
5041+
unsigned NumMods = IsDual ? 0 : 1; // Has A16 modifier
5042+
unsigned LastRegOpIdx = MI.getNumExplicitOperands() - 1 - NumMods;
5043+
unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
5044+
OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, DstSize);
5045+
if (IsDual) {
5046+
OpdsMapping[1] = AMDGPU::getValueMapping(
5047+
AMDGPU::VGPRRegBankID,
5048+
MRI.getType(MI.getOperand(1).getReg()).getSizeInBits());
5049+
OpdsMapping[2] = AMDGPU::getValueMapping(
5050+
AMDGPU::VGPRRegBankID,
5051+
MRI.getType(MI.getOperand(2).getReg()).getSizeInBits());
5052+
}
5053+
OpdsMapping[LastRegOpIdx] =
5054+
getSGPROpMapping(MI.getOperand(LastRegOpIdx).getReg(), MRI, *TRI);
5055+
if (LastRegOpIdx == 3) {
50405056
// Sequential form: all operands combined into VGPR256/VGPR512
50415057
unsigned Size = MRI.getType(MI.getOperand(2).getReg()).getSizeInBits();
50425058
if (Size > 256)
50435059
Size = 512;
50445060
OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size);
50455061
} else {
50465062
// NSA form
5047-
for (unsigned I = 2; I < N; ++I) {
5063+
unsigned FirstSrcOpIdx = IsDual ? 4 : 2;
5064+
for (unsigned I = FirstSrcOpIdx; I < LastRegOpIdx; ++I) {
50485065
unsigned Size = MRI.getType(MI.getOperand(I).getReg()).getSizeInBits();
50495066
OpdsMapping[I] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size);
50505067
}

‎llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -268,6 +268,7 @@ DECODE_OPERAND_REG_8(VReg_128)
268268
DECODE_OPERAND_REG_8(VReg_192)
269269
DECODE_OPERAND_REG_8(VReg_256)
270270
DECODE_OPERAND_REG_8(VReg_288)
271+
DECODE_OPERAND_REG_8(VReg_320)
271272
DECODE_OPERAND_REG_8(VReg_352)
272273
DECODE_OPERAND_REG_8(VReg_384)
273274
DECODE_OPERAND_REG_8(VReg_512)

‎llvm/lib/Target/AMDGPU/GCNSubtarget.h

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -228,6 +228,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
228228
bool HasRestrictedSOffset = false;
229229
bool HasBitOp3Insts = false;
230230
bool HasPrngInst = false;
231+
bool HasBVHDualInst = false;
231232
bool HasPermlane16Swap = false;
232233
bool HasPermlane32Swap = false;
233234
bool HasVcmpxPermlaneHazard = false;
@@ -1364,6 +1365,8 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
13641365

13651366
bool hasPrngInst() const { return HasPrngInst; }
13661367

1368+
bool hasBVHDualInst() const { return HasBVHDualInst; }
1369+
13671370
/// Return the maximum number of waves per SIMD for kernels using \p SGPRs
13681371
/// SGPRs
13691372
unsigned getOccupancyWithNumSGPRs(unsigned SGPRs) const;

‎llvm/lib/Target/AMDGPU/MIMGInstructions.td

Lines changed: 31 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -1509,17 +1509,18 @@ multiclass MIMG_Gather <mimgopc op, AMDGPUSampleVariant sample, bit wqm = 0,
15091509
multiclass MIMG_Gather_WQM <mimgopc op, AMDGPUSampleVariant sample>
15101510
: MIMG_Gather<op, sample, 1>;
15111511

1512-
class MIMG_IntersectRay_Helper<bit Is64, bit IsA16> {
1512+
class MIMG_IntersectRay_Helper<bit Is64, bit IsA16, bit isDual> {
15131513
int num_addrs = !if(Is64, !if(IsA16, 9, 12), !if(IsA16, 8, 11));
15141514
RegisterClass RegClass = MIMGAddrSize<num_addrs, 0>.RegClass;
15151515
int VAddrDwords = !srl(RegClass.Size, 5);
15161516

15171517
int GFX11PlusNSAAddrs = !if(IsA16, 4, 5);
15181518
RegisterClass node_ptr_type = !if(Is64, VReg_64, VGPR_32);
15191519
list<RegisterClass> GFX11PlusAddrTypes =
1520-
!if(IsA16,
1520+
!if(isDual, [VReg_64, VReg_64, VReg_96, VReg_96, VReg_64],
1521+
!if(IsA16,
15211522
[node_ptr_type, VGPR_32, VReg_96, VReg_96],
1522-
[node_ptr_type, VGPR_32, VReg_96, VReg_96, VReg_96]);
1523+
[node_ptr_type, VGPR_32, VReg_96, VReg_96, VReg_96]));
15231524
}
15241525

15251526
class MIMG_IntersectRay_gfx10<mimgopc op, string opcode, RegisterClass AddrRC>
@@ -1553,15 +1554,26 @@ class MIMG_IntersectRay_nsa_gfx11<mimgopc op, string opcode, int num_addrs,
15531554
}
15541555

15551556
class VIMAGE_IntersectRay_gfx12<mimgopc op, string opcode, int num_addrs,
1557+
bit isDual,
15561558
list<RegisterClass> addr_types>
1557-
: VIMAGE_gfx12<op.GFX12, (outs VReg_128:$vdata),
1559+
: VIMAGE_gfx12<op.GFX12, !if(isDual,
1560+
(outs VReg_320:$vdata, VReg_96:$ray_origin_out,
1561+
VReg_96:$ray_dir_out),
1562+
(outs VReg_128:$vdata)),
15581563
num_addrs, "GFX12", addr_types> {
1559-
let InOperandList = !con(nsah.AddrIns, (ins SReg_128_XNULL:$rsrc, A16:$a16));
1560-
let AsmString = opcode#" $vdata, "#nsah.AddrAsm#", $rsrc$a16";
1561-
}
1562-
1563-
multiclass MIMG_IntersectRay<mimgopc op, string opcode, bit Is64, bit IsA16> {
1564-
defvar info = MIMG_IntersectRay_Helper<Is64, IsA16>;
1564+
let Constraints = !if(isDual,
1565+
"$ray_origin_out = $vaddr2, $ray_dir_out = $vaddr3", "");
1566+
let InOperandList = !con(nsah.AddrIns, (ins SReg_128_XNULL:$rsrc),
1567+
!if(isDual, (ins), (ins A16:$a16)));
1568+
let AsmString = opcode#" $vdata, "#nsah.AddrAsm#", $rsrc"#
1569+
!if(isDual, "", "$a16");
1570+
let SchedRW = !if(isDual,
1571+
[WriteVMEM, WriteVMEM, WriteVMEM], [WriteVMEM]);
1572+
}
1573+
1574+
multiclass MIMG_IntersectRay<mimgopc op, string opcode, bit Is64, bit IsA16,
1575+
bit isDual> {
1576+
defvar info = MIMG_IntersectRay_Helper<Is64, IsA16, isDual>;
15651577
def "" : MIMGBaseOpcode {
15661578
let BVH = 1;
15671579
let A16 = IsA16;
@@ -1599,7 +1611,8 @@ multiclass MIMG_IntersectRay<mimgopc op, string opcode, bit Is64, bit IsA16> {
15991611
}
16001612
}
16011613
def _gfx12 : VIMAGE_IntersectRay_gfx12<op, opcode, info.GFX11PlusNSAAddrs,
1602-
info.GFX11PlusAddrTypes> {
1614+
isDual, info.GFX11PlusAddrTypes> {
1615+
let VDataDwords = !if(isDual, 10, 4);
16031616
let VAddrDwords = info.num_addrs;
16041617
}
16051618
}
@@ -1771,15 +1784,18 @@ defm IMAGE_MSAA_LOAD_X : MIMG_NoSampler <mimgopc<MIMG.NOP, MIMG.NOP, 0x80>, "ima
17711784
let OtherPredicates = [HasImageInsts, HasGFX10_AEncoding] in {
17721785
defm IMAGE_MSAA_LOAD : MIMG_MSAA_Load <mimgopc<0x18, 0x18, MIMG.NOP>, "image_msaa_load">;
17731786

1774-
defm IMAGE_BVH_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 0>;
1775-
defm IMAGE_BVH_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 1>;
1776-
defm IMAGE_BVH64_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 0>;
1777-
defm IMAGE_BVH64_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 1>;
1787+
defm IMAGE_BVH_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 0, 0>;
1788+
defm IMAGE_BVH_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 1, 0>;
1789+
defm IMAGE_BVH64_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 0, 0>;
1790+
defm IMAGE_BVH64_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 1, 0>;
17781791
} // End OtherPredicates = [HasImageInsts, HasGFX10_AEncoding]
17791792

1793+
defm IMAGE_BVH_DUAL_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<0x80, MIMG.NOP, MIMG.NOP>, "image_bvh_dual_intersect_ray", 1, 0, 1>;
1794+
17801795
let SubtargetPredicate = isGFX12Plus in {
17811796
def : AMDGPUMnemonicAlias<"bvh_intersect_ray", "image_bvh_intersect_ray">;
17821797
def : AMDGPUMnemonicAlias<"bvh64_intersect_ray", "image_bvh64_intersect_ray">;
1798+
def : AMDGPUMnemonicAlias<"bvh_dual_intersect_ray", "image_bvh_dual_intersect_ray">;
17831799
}
17841800

17851801
} // End let OtherPredicates = [HasImageInsts]

0 commit comments

Comments
 (0)