Changes after review

Author: mariusz-sikora-at-amd

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

@@ -2810,7 +2810,7 @@ def int_amdgcn_image_bvh_dual_intersect_ray :
   Intrinsic<[llvm_v10i32_ty, llvm_v3f32_ty, llvm_v3f32_ty],
             [llvm_i64_ty, llvm_float_ty, llvm_i8_ty, llvm_v3f32_ty,
              llvm_v3f32_ty, llvm_v2i32_ty, llvm_v4i32_ty],
-            [IntrReadMem, IntrWillReturn]>;
+            [IntrReadMem, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 // llvm.amdgcn.permlane16.var <old> <src0> <src1> <fi> <bound_control>
 def int_amdgcn_permlane16_var : ClangBuiltin<"__builtin_amdgcn_permlane16_var">,

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

@@ -7214,25 +7214,21 @@ bool AMDGPULegalizerInfo::legalizeBVHDualIntrinsic(MachineInstr &MI,
                                      NumVAddrDwords);
   assert(Opcode != -1);
 
-  SmallVector<Register, 12> Ops;
-  Ops.push_back(NodePtr);
-  Ops.push_back(B.buildMergeLikeInstr(
-                     V2S32, {RayExtent, B.buildAnyExt(S32, InstanceMask)})
-                    .getReg(0));
-  Ops.push_back(RayOrigin);
-  Ops.push_back(RayDir);
-  Ops.push_back(Offsets);
-
-  auto MIB = B.buildInstr(AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY)
-                 .addDef(DstReg)
-                 .addDef(DstOrigin)
-                 .addDef(DstDir)
-                 .addImm(Opcode);
-
-  for (Register R : Ops)
-    MIB.addUse(R);
-
-  MIB.addUse(TDescr).cloneMemRefs(MI);
+  auto RayExtentInstanceMaskVec = B.buildMergeLikeInstr(
+      V2S32, {RayExtent, B.buildAnyExt(S32, InstanceMask)});
+
+  B.buildInstr(AMDGPU::G_AMDGPU_BVH_DUAL_INTERSECT_RAY)
+      .addDef(DstReg)
+      .addDef(DstOrigin)
+      .addDef(DstDir)
+      .addImm(Opcode)
+      .addUse(NodePtr)
+      .addUse(RayExtentInstanceMaskVec.getReg(0))
+      .addUse(RayOrigin)
+      .addUse(RayDir)
+      .addUse(Offsets)
+      .addUse(TDescr)
+      .cloneMemRefs(MI);
 
   MI.eraseFromParent();
   return true;

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

@@ -9465,7 +9465,7 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op,
                                        NumVAddrDwords);
     assert(Opcode != -1);
 
-    SmallVector<SDValue, 16> Ops;
+    SmallVector<SDValue, 7> Ops;
     Ops.push_back(NodePtr);
     Ops.push_back(DAG.getBuildVector(
         MVT::v2i32, DL,