[AMDGPU] Change llvm.amdgcn.image.bvh.intersect.ray to take vec3 args

[foad]

Summary

The ray_origin, ray_dir and ray_inv_dir arguments should all be vec3 to
match how the hardware instruction works.

Also update the API of the corresponding OpenCL builtins.

[foad]

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

[yaxunl]

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

[arsenm]

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I do not think it's worth introducing a macro for this. Are there actually C users of these builtins?

[yaxunl]

In D114957#3166861, @arsenm wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I do not think it's worth introducing a macro for this. Are there actually C users of these builtins?

Yes we have users who use these clang builtins. We have received quite a few complaints about making breaking API changes without a way to detect them in the program.

[b-sumner]

In D114957#3166882, @yaxunl wrote:

In D114957#3166861, @arsenm wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I do not think it's worth introducing a macro for this. Are there actually C users of these builtins?

Yes we have users who use these clang builtins. We have received quite a few complaints about making breaking API changes without a way to detect them in the program.

But builtins are not part of the documented API and we have advised developers using them that they are subject to change.

[foad]

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I don't know anything about OpenCL macros. Is it good enough to put this in AMDGPUTargetInfo::getTargetDefines:

if (Opts.OpenCL)
  Builder.defineMacro("__amdgcn_bvh_use_vec3__");

Does it need tests, documentation, etc?

[b-sumner]

In D114957#3166936, @foad wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I don't know anything about OpenCL macros. Is it good enough to put this in AMDGPUTargetInfo::getTargetDefines:

if (Opts.OpenCL)
  Builder.defineMacro("__amdgcn_bvh_use_vec3__");

Does it need tests, documentation, etc?

But how long would that be carried? And then deprecated?

[foad]

In D114957#3166948, @b-sumner wrote:

In D114957#3166936, @foad wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I don't know anything about OpenCL macros. Is it good enough to put this in AMDGPUTargetInfo::getTargetDefines:

if (Opts.OpenCL)
  Builder.defineMacro("__amdgcn_bvh_use_vec3__");

Does it need tests, documentation, etc?

But how long would that be carried? And then deprecated?

Then do you think the patch is OK as-is?

[yaxunl]

In D114957#3166974, @foad wrote:

In D114957#3166948, @b-sumner wrote:

In D114957#3166936, @foad wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I don't know anything about OpenCL macros. Is it good enough to put this in AMDGPUTargetInfo::getTargetDefines:

if (Opts.OpenCL)
  Builder.defineMacro("__amdgcn_bvh_use_vec3__");

Does it need tests, documentation, etc?

But how long would that be carried? And then deprecated?

Then do you think the patch is OK as-is?

Let's discuss with the users and see whether the macro is needed or not.

[yaxunl]

In D114957#3166936, @foad wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I don't know anything about OpenCL macros. Is it good enough to put this in AMDGPUTargetInfo::getTargetDefines:

if (Opts.OpenCL)
  Builder.defineMacro("__amdgcn_bvh_use_vec3__");

Does it need tests, documentation, etc?

Yes you can add it to AMDGPUTargetInfo::getTargetDefines but without if (Opts.OpenCL) since HIP uses it too.

You can add a test to https://github.com/llvm/llvm-project/blob/b2a2c38349a18b89b04d342632d5ea02f86dfdd6/clang/test/Preprocessor/predefined-arch-macros.c#L3731

[b-sumner]

Can we agree to drop the macro in LLVM 15 and note that in a comment or elsewhere?

[foad]

Define amdgcn_bvh_use_vec3.

[yaxunl]

LGTM from clang side. Thanks.

[arsenm]

I think this macro is purely terrible and should not be added (and at least should be all caps?). If we can't just hard break users, I would rather just leave the builtin signatures broken

[arsenm]

In D114957#3167700, @arsenm wrote:

I think this macro is purely terrible and should not be added (and at least should be all caps?). If we can't just hard break users, I would rather just leave the builtin signatures broken

Rather than adding an ad-hoc named macro, could they just directly check the clang version?

[yaxunl]

In D114957#3167703, @arsenm wrote:

In D114957#3167700, @arsenm wrote:

I think this macro is purely terrible and should not be added (and at least should be all caps?). If we can't just hard break users, I would rather just leave the builtin signatures broken

Rather than adding an ad-hoc named macro, could they just directly check the clang version?

For the previous and the next clang release, it can be determined by checking clang version. Problem is that ROCm release is not synch'ed with clang release. So for ROCm release that uses the current clang release, it could have this change or could not, which cannot be determined by checking clang version.

[foad]

In D114957#3167700, @arsenm wrote:

I think this macro is purely terrible and should not be added (and at least should be all caps?). If we can't just hard break users, I would rather just leave the builtin signatures broken

OK, how about D115032?

Personally I have no opinion about what's best to do with the OpenCL builtins, but I would like to make progress with changing the intrinsics. So I have a slight preference for D115032 because it gives me a way forward without changing OpenCL behaviour. The OpenCL team can then decide whether or not to update the builtins at their leisure.

[foad]

Abandoned in favour of D115032.

clang/include/clang/Basic/BuiltinsAMDGPU.def

 // GFX10+ only builtins.
 //===----------------------------------------------------------------------===//
 TARGET_BUILTIN(__builtin_amdgcn_permlane16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts")
 TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts")
 TARGET_BUILTIN(__builtin_amdgcn_mov_dpp8, "UiUiIUi", "nc", "gfx10-insts")
 
 //===----------------------------------------------------------------------===//
 // Raytracing builtins.
-// By default the 1st argument is i32 and the 4/5-th arguments are float4.
+// By default the 1st argument is i32 and the 4/5-th arguments are float3.
 // Postfix l indicates the 1st argument is i64.
-// Postfix h indicates the 4/5-th arguments are half4.
+// Postfix h indicates the 4/5-th arguments are half3.
 //===----------------------------------------------------------------------===//
-TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray, "V4UiUifV4fV4fV4fV4Ui", "nc", "gfx10-insts")
-TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_h, "V4UiUifV4fV4hV4hV4Ui", "nc", "gfx10-insts")
-TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_l, "V4UiWUifV4fV4fV4fV4Ui", "nc", "gfx10-insts")
-TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_lh, "V4UiWUifV4fV4hV4hV4Ui", "nc", "gfx10-insts")
+TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray, "V4UiUifV3fV3fV3fV4Ui", "nc", "gfx10-insts")
+TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_h, "V4UiUifV3fV3hV3hV4Ui", "nc", "gfx10-insts")
+TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_l, "V4UiWUifV3fV3fV3fV4Ui", "nc", "gfx10-insts")
+TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_lh, "V4UiWUifV3fV3hV3hV4Ui", "nc", "gfx10-insts")
 
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_read_exec, "WUi", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
 

clang/lib/Basic/Targets/AMDGPU.cpp

void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
   if (hasLDEXPF())
     Builder.defineMacro("__HAS_LDEXPF__");
   if (hasFP64())
     Builder.defineMacro("__HAS_FP64__");
   if (hasFastFMA())
     Builder.defineMacro("FP_FAST_FMA");
 
   Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE", Twine(WavefrontSize));
+
+  // Temporary macro to help clients transition to the current definition of
+  // __builtin_amdgcn_image_bvh_intersect_ray* from the previous one which used
+  // vec4 arguments.
+  // TODO: Remove this in LLVM 15.
+  Builder.defineMacro("__amdgcn_bvh_use_vec3__");
 }
 
 void AMDGPUTargetInfo::setAuxTarget(const TargetInfo *Aux) {
   assert(HalfFormat == Aux->HalfFormat);
   assert(FloatFormat == Aux->FloatFormat);
   assert(DoubleFormat == Aux->DoubleFormat);
 
   // On x86_64 long double is 80-bit extended precision format, which is

clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S \
 // RUN:   -emit-llvm -cl-std=CL2.0 -o - %s | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S \
 // RUN:   -cl-std=CL2.0 -o - %s | FileCheck -check-prefix=ISA %s
 
 // Test llvm.amdgcn.image.bvh.intersect.ray intrinsic.
 
 // The clang builtin functions __builtin_amdgcn_image_bvh_intersect_ray* use
 // postfixes to indicate the types of the 1st, 4th, and 5th arguments.
-// By default, the 1st argument is i32, the 4/5-th arguments are float4.
+// By default, the 1st argument is i32, the 4/5-th arguments are float3.
 // Postfix l indicates the 1st argument is i64 and postfix h indicates
-// the 4/5-th arguments are half4.
+// the 4/5-th arguments are half3.
 
 typedef unsigned int uint;
 typedef unsigned long ulong;
-typedef float float4 __attribute__((ext_vector_type(4)));
-typedef double double4 __attribute__((ext_vector_type(4)));
-typedef half half4 __attribute__((ext_vector_type(4)));
+typedef float float3 __attribute__((ext_vector_type(3)));
+typedef half half3 __attribute__((ext_vector_type(3)));
 typedef uint uint4 __attribute__((ext_vector_type(4)));
 
-// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32
+// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v3f32
 // ISA: image_bvh_intersect_ray
 void test_image_bvh_intersect_ray(global uint4* out, uint node_ptr,
-  float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir,
+  float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir,
   uint4 texture_descr)
 {
   *out = __builtin_amdgcn_image_bvh_intersect_ray(node_ptr, ray_extent,
            ray_origin, ray_dir, ray_inv_dir, texture_descr);
 }
 
-// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16
+// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v3f16
 // ISA: image_bvh_intersect_ray
 void test_image_bvh_intersect_ray_h(global uint4* out, uint node_ptr,
-  float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir,
+  float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir,
   uint4 texture_descr)
 {
   *out = __builtin_amdgcn_image_bvh_intersect_ray_h(node_ptr, ray_extent,
            ray_origin, ray_dir, ray_inv_dir, texture_descr);
 }
 
-// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32
+// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v3f32
 // ISA: image_bvh_intersect_ray
 void test_image_bvh_intersect_ray_l(global uint4* out, ulong node_ptr,
-  float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir,
+  float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir,
   uint4 texture_descr)
 {
   *out = __builtin_amdgcn_image_bvh_intersect_ray_l(node_ptr, ray_extent,
            ray_origin, ray_dir, ray_inv_dir, texture_descr);
 }
 
-// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16
+// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v3f16
 // ISA: image_bvh_intersect_ray
 void test_image_bvh_intersect_ray_lh(global uint4* out, ulong node_ptr,
-  float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir,
+  float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir,
   uint4 texture_descr)
 {
   *out = __builtin_amdgcn_image_bvh_intersect_ray_lh(node_ptr, ray_extent,
            ray_origin, ray_dir, ray_inv_dir, texture_descr);
 }
 

clang/test/Preprocessor/predefined-macros.c

 // CHECK-SPIRV64-NOT: #define __SPIRV32__ 1
 
 // RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa      \
 // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP
 // CHECK-HIP-NOT: #define __CUDA_ARCH__
 // CHECK-HIP: #define __HIPCC__ 1
 // CHECK-HIP-NOT: #define __HIP_DEVICE_COMPILE__ 1
 // CHECK-HIP: #define __HIP__ 1
+// CHECK-HIP: #define __amdgcn_bvh_use_vec3__ 1
 
 // RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device \
 // RUN:   | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP-DEV
 // CHECK-HIP-DEV-NOT: #define __CUDA_ARCH__
 // CHECK-HIP-DEV: #define __HIPCC__ 1
 // CHECK-HIP-DEV: #define __HIP_DEVICE_COMPILE__ 1
 // CHECK-HIP-DEV: #define __HIP__ 1
+// CHECK-HIP-DEV: #define __amdgcn_bvh_use_vec3__ 1

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

class AMDGPUGlobalAtomicRtn<LLVMType vt> : Intrinsic <
    vt],               // vdata(VGPR)
   [IntrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>], "",
   [SDNPMemOperand]>;
 
 def int_amdgcn_global_atomic_csub : AMDGPUGlobalAtomicRtn<llvm_i32_ty>;
 
 // uint4 llvm.amdgcn.image.bvh.intersect.ray <node_ptr>, <ray_extent>, <ray_origin>,
 //                                           <ray_dir>, <ray_inv_dir>, <texture_descr>
+// <node_ptr> is i32 or i64.
+// <ray_dir> and <ray_inv_dir> are both v3f16 or both v3f32.
 def int_amdgcn_image_bvh_intersect_ray :
   Intrinsic<[llvm_v4i32_ty],
-            [llvm_anyint_ty, llvm_float_ty, llvm_v4f32_ty, llvm_anyvector_ty,
+            [llvm_anyint_ty, llvm_float_ty, llvm_v3f32_ty, llvm_anyvector_ty,
              LLVMMatchType<1>, llvm_v4i32_ty],
             [IntrReadMem, IntrWillReturn]>;
 
 //===----------------------------------------------------------------------===//
 // Deep learning intrinsics.
 //===----------------------------------------------------------------------===//
 
 // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

if (Is64) {
     auto Unmerge = B.buildUnmerge({S32, S32}, NodePtr);
     Ops.push_back(Unmerge.getReg(0));
     Ops.push_back(Unmerge.getReg(1));
   } else {
     Ops.push_back(NodePtr);
   }
   Ops.push_back(RayExtent);
 
-  auto packLanes = [&Ops, &S32, &B] (Register Src) {
-    auto Unmerge = B.buildUnmerge({S32, S32, S32, S32}, Src);
+  auto packLanes = [&Ops, &S32, &B](Register Src) {
+    auto Unmerge = B.buildUnmerge({S32, S32, S32}, Src);
     Ops.push_back(Unmerge.getReg(0));
     Ops.push_back(Unmerge.getReg(1));
     Ops.push_back(Unmerge.getReg(2));
   };
 
   packLanes(RayOrigin);
   if (IsA16) {
-    auto UnmergeRayDir = B.buildUnmerge({S16, S16, S16, S16}, RayDir);
-    auto UnmergeRayInvDir = B.buildUnmerge({S16, S16, S16, S16}, RayInvDir);
+    auto UnmergeRayDir = B.buildUnmerge({S16, S16, S16}, RayDir);
+    auto UnmergeRayInvDir = B.buildUnmerge({S16, S16, S16}, RayInvDir);
     Register R1 = MRI.createGenericVirtualRegister(S32);
     Register R2 = MRI.createGenericVirtualRegister(S32);
     Register R3 = MRI.createGenericVirtualRegister(S32);
     B.buildMerge(R1, {UnmergeRayDir.getReg(0), UnmergeRayDir.getReg(1)});
     B.buildMerge(R2, {UnmergeRayDir.getReg(2), UnmergeRayInvDir.getReg(0)});
     B.buildMerge(R3, {UnmergeRayInvDir.getReg(1), UnmergeRayInvDir.getReg(2)});
     Ops.push_back(R1);
     Ops.push_back(R2);

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

case Intrinsic::amdgcn_image_bvh_intersect_ray: {
     SDValue RayExtent = M->getOperand(3);
     SDValue RayOrigin = M->getOperand(4);
     SDValue RayDir = M->getOperand(5);
     SDValue RayInvDir = M->getOperand(6);
     SDValue TDescr = M->getOperand(7);
 
     assert(NodePtr.getValueType() == MVT::i32 ||
            NodePtr.getValueType() == MVT::i64);
-    assert(RayDir.getValueType() == MVT::v4f16 ||
-           RayDir.getValueType() == MVT::v4f32);
+    assert(RayDir.getValueType() == MVT::v3f16 ||
+           RayDir.getValueType() == MVT::v3f32);
 
     if (!Subtarget->hasGFX10_AEncoding()) {
       emitRemovedIntrinsicError(DAG, DL, Op.getValueType());
       return SDValue();
     }
 
     const bool IsA16 = RayDir.getValueType().getVectorElementType() == MVT::f16;
     const bool Is64 = NodePtr.getValueType() == MVT::i64;

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX1030 %s
 ; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1013 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX1013 %s
 ; RUN: not --crash llc -global-isel -march=amdgcn -mcpu=gfx1012 -verify-machineinstrs < %s -o /dev/null 2>&1 | FileCheck -check-prefix=ERR %s
 
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, uint4 texture_descr)
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, uint4 texture_descr)
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(ulong node_ptr, float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, uint4 texture_descr)
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(ulong node_ptr, float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, uint4 texture_descr)
-
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32, float, <4 x float>, <4 x float>, <4 x float>, <4 x i32>)
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32, float, <4 x float>, <4 x half>, <4 x half>, <4 x i32>)
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64, float, <4 x float>, <4 x float>, <4 x float>, <4 x i32>)
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64, float, <4 x float>, <4 x half>, <4 x half>, <4 x i32>)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(ulong node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(ulong node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr)
+
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32, float, <3 x float>, <3 x float>, <3 x float>, <4 x i32>)
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32, float, <3 x float>, <3 x half>, <3 x half>, <4 x i32>)
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64, float, <3 x float>, <3 x float>, <3 x float>, <4 x i32>)
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64, float, <3 x float>, <3 x half>, <3 x half>, <4 x i32>)
 declare i32 @llvm.amdgcn.workitem.id.x()
 
-define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> inreg %tdescr) {
-; GFX1030-LABEL: image_bvh_intersect_ray:
-; GFX1030:       ; %bb.0:
-; GFX1030-NEXT:    image_bvh_intersect_ray v[0:3], [v0, v1, v2, v3, v4, v6, v7, v8, v10, v11, v12], s[0:3]
-; GFX1030-NEXT:    s_waitcnt vmcnt(0)
-; GFX1030-NEXT:    ; return to shader part epilog
-;
-; GFX1013-LABEL: image_bvh_intersect_ray:
-; GFX1013:       ; %bb.0:
-; GFX1013-NEXT:    v_mov_b32_e32 v5, v6
-; GFX1013-NEXT:    v_mov_b32_e32 v6, v7
-; GFX1013-NEXT:    v_mov_b32_e32 v7, v8
-; GFX1013-NEXT:    v_mov_b32_e32 v8, v10
-; GFX1013-NEXT:    v_mov_b32_e32 v9, v11
-; GFX1013-NEXT:    v_mov_b32_e32 v10, v12
-; GFX1013-NEXT:    image_bvh_intersect_ray v[0:3], v[0:15], s[0:3]
-; GFX1013-NEXT:    s_waitcnt vmcnt(0)
-; GFX1013-NEXT:    ; return to shader part epilog
+define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> inreg %tdescr) {
+; GCN-LABEL: image_bvh_intersect_ray:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    image_bvh_intersect_ray v[0:3], v[0:15], s[0:3]
+; GCN-NEXT:    s_waitcnt vmcnt(0)
+; GCN-NEXT:    ; return to shader part epilog
 ; ERR: in function image_bvh_intersect_ray{{.*}}intrinsic not supported on subtarget
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
 define amdgpu_ps <4 x float> @image_bvh_intersect_ray_flat(i32 %node_ptr, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, float %ray_inv_dir_x, float %ray_inv_dir_y, float %ray_inv_dir_z, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh_intersect_ray_flat:
 ; GCN:       ; %bb.0:
 ; GCN-NEXT:    image_bvh_intersect_ray v[0:3], v[0:15], s[0:3]
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
-  %ray_origin0 = insertelement <4 x float> undef, float %ray_origin_x, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float %ray_origin_y, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float %ray_origin_z, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float %ray_dir_x, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float %ray_dir_y, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float %ray_dir_z, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float %ray_inv_dir_x, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float %ray_origin_x, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float %ray_origin_y, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float %ray_origin_z, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float %ray_dir_x, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float %ray_dir_y, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float %ray_dir_z, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float %ray_inv_dir_x, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
  %r = bitcast <4 x i32> %v to <4 x float>
  ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh_intersect_ray_a16(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> inreg %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh_intersect_ray_a16(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh_intersect_ray_a16:
 ; GCN:       ; %bb.0:
 ; GCN-NEXT:    s_mov_b32 s4, 0xffff
-; GCN-NEXT:    v_lshrrev_b32_e32 v5, 16, v6
-; GCN-NEXT:    v_and_b32_e32 v10, s4, v8
-; GCN-NEXT:    v_lshrrev_b32_e32 v8, 16, v8
-; GCN-NEXT:    v_and_b32_e32 v9, s4, v9
-; GCN-NEXT:    v_lshlrev_b32_e32 v5, 16, v5
+; GCN-NEXT:    v_lshrrev_b32_e32 v9, 16, v5
+; GCN-NEXT:    v_and_b32_e32 v10, s4, v7
+; GCN-NEXT:    v_lshrrev_b32_e32 v7, 16, v7
+; GCN-NEXT:    v_and_b32_e32 v8, s4, v8
+; GCN-NEXT:    v_lshlrev_b32_e32 v9, 16, v9
 ; GCN-NEXT:    v_lshlrev_b32_e32 v10, 16, v10
-; GCN-NEXT:    v_and_or_b32 v5, v6, s4, v5
-; GCN-NEXT:    v_and_or_b32 v6, v7, s4, v10
-; GCN-NEXT:    v_lshl_or_b32 v7, v9, 16, v8
+; GCN-NEXT:    v_lshl_or_b32 v7, v8, 16, v7
+; GCN-NEXT:    v_and_or_b32 v5, v5, s4, v9
+; GCN-NEXT:    v_and_or_b32 v6, v6, s4, v10
 ; GCN-NEXT:    image_bvh_intersect_ray v[0:3], v[0:7], s[0:3] a16
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh64_intersect_ray(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> inreg %tdescr) {
-; GFX1030-LABEL: image_bvh64_intersect_ray:
-; GFX1030:       ; %bb.0:
-; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], [v0, v1, v2, v3, v4, v5, v7, v8, v9, v11, v12, v13], s[0:3]
-; GFX1030-NEXT:    s_waitcnt vmcnt(0)
-; GFX1030-NEXT:    ; return to shader part epilog
-;
-; GFX1013-LABEL: image_bvh64_intersect_ray:
-; GFX1013:       ; %bb.0:
-; GFX1013-NEXT:    v_mov_b32_e32 v6, v7
-; GFX1013-NEXT:    v_mov_b32_e32 v7, v8
-; GFX1013-NEXT:    v_mov_b32_e32 v8, v9
-; GFX1013-NEXT:    v_mov_b32_e32 v9, v11
-; GFX1013-NEXT:    v_mov_b32_e32 v10, v12
-; GFX1013-NEXT:    v_mov_b32_e32 v11, v13
-; GFX1013-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3]
-; GFX1013-NEXT:    s_waitcnt vmcnt(0)
-; GFX1013-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+define amdgpu_ps <4 x float> @image_bvh64_intersect_ray(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> inreg %tdescr) {
+; GCN-LABEL: image_bvh64_intersect_ray:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3]
+; GCN-NEXT:    s_waitcnt vmcnt(0)
+; GCN-NEXT:    ; return to shader part epilog
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
 define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_flat(<2 x i32> %node_ptr_vec, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, float %ray_inv_dir_x, float %ray_inv_dir_y, float %ray_inv_dir_z, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh64_intersect_ray_flat:
 ; GCN:       ; %bb.0:
 ; GCN-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3]
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
   %node_ptr = bitcast <2 x i32> %node_ptr_vec to i64
-  %ray_origin0 = insertelement <4 x float> undef, float %ray_origin_x, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float %ray_origin_y, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float %ray_origin_z, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float %ray_dir_x, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float %ray_dir_y, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float %ray_dir_z, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float %ray_inv_dir_x, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float %ray_origin_x, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float %ray_origin_y, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float %ray_origin_z, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float %ray_dir_x, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float %ray_dir_y, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float %ray_dir_z, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float %ray_inv_dir_x, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
  %r = bitcast <4 x i32> %v to <4 x float>
  ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_a16(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> inreg %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_a16(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh64_intersect_ray_a16:
 ; GCN:       ; %bb.0:
 ; GCN-NEXT:    s_mov_b32 s4, 0xffff
-; GCN-NEXT:    v_lshrrev_b32_e32 v6, 16, v7
-; GCN-NEXT:    v_and_b32_e32 v11, s4, v9
-; GCN-NEXT:    v_lshrrev_b32_e32 v9, 16, v9
-; GCN-NEXT:    v_and_b32_e32 v10, s4, v10
-; GCN-NEXT:    v_lshlrev_b32_e32 v6, 16, v6
+; GCN-NEXT:    v_lshrrev_b32_e32 v10, 16, v6
+; GCN-NEXT:    v_and_b32_e32 v11, s4, v8
+; GCN-NEXT:    v_lshrrev_b32_e32 v8, 16, v8
+; GCN-NEXT:    v_and_b32_e32 v9, s4, v9
+; GCN-NEXT:    v_lshlrev_b32_e32 v10, 16, v10
 ; GCN-NEXT:    v_lshlrev_b32_e32 v11, 16, v11
-; GCN-NEXT:    v_and_or_b32 v6, v7, s4, v6
-; GCN-NEXT:    v_and_or_b32 v7, v8, s4, v11
-; GCN-NEXT:    v_lshl_or_b32 v8, v10, 16, v9
+; GCN-NEXT:    v_lshl_or_b32 v8, v9, 16, v8
+; GCN-NEXT:    v_and_or_b32 v6, v6, s4, v10
+; GCN-NEXT:    v_and_or_b32 v7, v7, s4, v11
 ; GCN-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3] a16
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh_intersect_ray_vgpr_descr(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh_intersect_ray_vgpr_descr(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr) {
 ; GFX1030-LABEL: image_bvh_intersect_ray_vgpr_descr:
 ; GFX1030:       ; %bb.0:
-; GFX1030-NEXT:    v_mov_b32_e32 v18, v0
-; GFX1030-NEXT:    v_mov_b32_e32 v19, v1
-; GFX1030-NEXT:    v_mov_b32_e32 v20, v2
-; GFX1030-NEXT:    v_mov_b32_e32 v21, v3
-; GFX1030-NEXT:    v_mov_b32_e32 v22, v4
-; GFX1030-NEXT:    v_mov_b32_e32 v23, v6
-; GFX1030-NEXT:    v_mov_b32_e32 v24, v7
-; GFX1030-NEXT:    v_mov_b32_e32 v25, v8
-; GFX1030-NEXT:    v_mov_b32_e32 v26, v10
-; GFX1030-NEXT:    v_mov_b32_e32 v27, v11
-; GFX1030-NEXT:    v_mov_b32_e32 v28, v12
+; GFX1030-NEXT:    v_mov_b32_e32 v15, v0
+; GFX1030-NEXT:    v_mov_b32_e32 v16, v1
+; GFX1030-NEXT:    v_mov_b32_e32 v17, v2
+; GFX1030-NEXT:    v_mov_b32_e32 v18, v3
+; GFX1030-NEXT:    v_mov_b32_e32 v19, v4
+; GFX1030-NEXT:    v_mov_b32_e32 v20, v5
+; GFX1030-NEXT:    v_mov_b32_e32 v21, v6
+; GFX1030-NEXT:    v_mov_b32_e32 v22, v7
+; GFX1030-NEXT:    v_mov_b32_e32 v23, v8
+; GFX1030-NEXT:    v_mov_b32_e32 v24, v9
+; GFX1030-NEXT:    v_mov_b32_e32 v25, v10
 ; GFX1030-NEXT:    s_mov_b32 s1, exec_lo
 ; GFX1030-NEXT:  .LBB6_1: ; =>This Inner Loop Header: Depth=1
-; GFX1030-NEXT:    v_readfirstlane_b32 s4, v14
-; GFX1030-NEXT:    v_readfirstlane_b32 s5, v15
-; GFX1030-NEXT:    v_readfirstlane_b32 s6, v16
-; GFX1030-NEXT:    v_readfirstlane_b32 s7, v17
-; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[14:15]
-; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[16:17]
+; GFX1030-NEXT:    v_readfirstlane_b32 s4, v11
+; GFX1030-NEXT:    v_readfirstlane_b32 s5, v12
+; GFX1030-NEXT:    v_readfirstlane_b32 s6, v13
+; GFX1030-NEXT:    v_readfirstlane_b32 s7, v14
+; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[11:12]
+; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[13:14]
 ; GFX1030-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1030-NEXT:    s_and_saveexec_b32 s0, s0
-; GFX1030-NEXT:    image_bvh_intersect_ray v[0:3], v[18:33], s[4:7]
-; GFX1030-NEXT:    ; implicit-def: $vgpr14_vgpr15
+; GFX1030-NEXT:    image_bvh_intersect_ray v[0:3], v[15:30], s[4:7]
+; GFX1030-NEXT:    ; implicit-def: $vgpr11_vgpr12
+; GFX1030-NEXT:    ; implicit-def: $vgpr15
+; GFX1030-NEXT:    ; implicit-def: $vgpr16
+; GFX1030-NEXT:    ; implicit-def: $vgpr17
 ; GFX1030-NEXT:    ; implicit-def: $vgpr18
 ; GFX1030-NEXT:    ; implicit-def: $vgpr19
 ; GFX1030-NEXT:    ; implicit-def: $vgpr20
 ; GFX1030-NEXT:    ; implicit-def: $vgpr21
 ; GFX1030-NEXT:    ; implicit-def: $vgpr22
 ; GFX1030-NEXT:    ; implicit-def: $vgpr23
 ; GFX1030-NEXT:    ; implicit-def: $vgpr24
 ; GFX1030-NEXT:    ; implicit-def: $vgpr25
-; GFX1030-NEXT:    ; implicit-def: $vgpr26
-; GFX1030-NEXT:    ; implicit-def: $vgpr27
-; GFX1030-NEXT:    ; implicit-def: $vgpr28
-; GFX1030-NEXT:    ; implicit-def: $vgpr14_vgpr15_vgpr16_vgpr17
+; GFX1030-NEXT:    ; implicit-def: $vgpr11_vgpr12_vgpr13_vgpr14
 ; GFX1030-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1030-NEXT:    s_cbranch_execnz .LBB6_1
 ; GFX1030-NEXT:  ; %bb.2:
 ; GFX1030-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1030-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1030-NEXT:    ; return to shader part epilog
 ;
 ; GFX1013-LABEL: image_bvh_intersect_ray_vgpr_descr:
 ; GFX1013:       ; %bb.0:
-; GFX1013-NEXT:    v_mov_b32_e32 v5, v6
-; GFX1013-NEXT:    v_mov_b32_e32 v6, v7
-; GFX1013-NEXT:    v_mov_b32_e32 v7, v8
-; GFX1013-NEXT:    v_mov_b32_e32 v8, v10
-; GFX1013-NEXT:    v_mov_b32_e32 v9, v11
-; GFX1013-NEXT:    v_mov_b32_e32 v10, v12
-; GFX1013-NEXT:    v_mov_b32_e32 v18, v14
-; GFX1013-NEXT:    v_mov_b32_e32 v19, v15
+; GFX1013-NEXT:    v_mov_b32_e32 v16, v11
+; GFX1013-NEXT:    v_mov_b32_e32 v17, v12
+; GFX1013-NEXT:    v_mov_b32_e32 v18, v13
+; GFX1013-NEXT:    v_mov_b32_e32 v19, v14
 ; GFX1013-NEXT:    s_mov_b32 s1, exec_lo
 ; GFX1013-NEXT:  .LBB6_1: ; =>This Inner Loop Header: Depth=1
-; GFX1013-NEXT:    v_readfirstlane_b32 s4, v18
-; GFX1013-NEXT:    v_readfirstlane_b32 s5, v19
-; GFX1013-NEXT:    v_readfirstlane_b32 s6, v16
-; GFX1013-NEXT:    v_readfirstlane_b32 s7, v17
-; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[18:19]
-; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[16:17]
+; GFX1013-NEXT:    v_readfirstlane_b32 s4, v16
+; GFX1013-NEXT:    v_readfirstlane_b32 s5, v17
+; GFX1013-NEXT:    v_readfirstlane_b32 s6, v18
+; GFX1013-NEXT:    v_readfirstlane_b32 s7, v19
+; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[16:17]
+; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[18:19]
 ; GFX1013-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1013-NEXT:    s_and_saveexec_b32 s0, s0
 ; GFX1013-NEXT:    image_bvh_intersect_ray v[20:23], v[0:15], s[4:7]
+; GFX1013-NEXT:    ; implicit-def: $vgpr16_vgpr17
 ; GFX1013-NEXT:    ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15
-; GFX1013-NEXT:    ; implicit-def: $vgpr18_vgpr19
-; GFX1013-NEXT:    ; implicit-def: $vgpr14_vgpr15_vgpr16_vgpr17
+; GFX1013-NEXT:    ; implicit-def: $vgpr16_vgpr17_vgpr18_vgpr19
 ; GFX1013-NEXT:    s_waitcnt_depctr 0xffe3
 ; GFX1013-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1013-NEXT:    s_cbranch_execnz .LBB6_1
 ; GFX1013-NEXT:  ; %bb.2:
 ; GFX1013-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1013-NEXT:    v_mov_b32_e32 v0, v20
 ; GFX1013-NEXT:    v_mov_b32_e32 v1, v21
 ; GFX1013-NEXT:    v_mov_b32_e32 v2, v22
 ; GFX1013-NEXT:    v_mov_b32_e32 v3, v23
 ; GFX1013-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh_intersect_ray_a16_vgpr_descr(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh_intersect_ray_a16_vgpr_descr(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr) {
 ; GFX1030-LABEL: image_bvh_intersect_ray_a16_vgpr_descr:
 ; GFX1030:       ; %bb.0:
 ; GFX1030-NEXT:    s_mov_b32 s0, 0xffff
-; GFX1030-NEXT:    v_mov_b32_e32 v14, v0
-; GFX1030-NEXT:    v_mov_b32_e32 v15, v1
-; GFX1030-NEXT:    v_lshrrev_b32_e32 v0, 16, v6
-; GFX1030-NEXT:    v_and_b32_e32 v1, s0, v8
-; GFX1030-NEXT:    v_mov_b32_e32 v16, v2
-; GFX1030-NEXT:    v_mov_b32_e32 v17, v3
-; GFX1030-NEXT:    v_lshrrev_b32_e32 v2, 16, v8
+; GFX1030-NEXT:    v_mov_b32_e32 v13, v0
+; GFX1030-NEXT:    v_mov_b32_e32 v14, v1
+; GFX1030-NEXT:    v_lshrrev_b32_e32 v0, 16, v5
+; GFX1030-NEXT:    v_and_b32_e32 v1, s0, v7
+; GFX1030-NEXT:    v_mov_b32_e32 v15, v2
+; GFX1030-NEXT:    v_mov_b32_e32 v16, v3
+; GFX1030-NEXT:    v_lshrrev_b32_e32 v2, 16, v7
 ; GFX1030-NEXT:    v_lshlrev_b32_e32 v0, 16, v0
 ; GFX1030-NEXT:    v_lshlrev_b32_e32 v1, 16, v1
-; GFX1030-NEXT:    v_and_b32_e32 v3, s0, v9
-; GFX1030-NEXT:    v_mov_b32_e32 v18, v4
+; GFX1030-NEXT:    v_and_b32_e32 v3, s0, v8
+; GFX1030-NEXT:    v_mov_b32_e32 v17, v4
 ; GFX1030-NEXT:    s_mov_b32 s1, exec_lo
-; GFX1030-NEXT:    v_and_or_b32 v19, v6, s0, v0
-; GFX1030-NEXT:    v_and_or_b32 v20, v7, s0, v1
-; GFX1030-NEXT:    v_lshl_or_b32 v21, v3, 16, v2
+; GFX1030-NEXT:    v_and_or_b32 v18, v5, s0, v0
+; GFX1030-NEXT:    v_and_or_b32 v19, v6, s0, v1
+; GFX1030-NEXT:    v_lshl_or_b32 v20, v3, 16, v2
 ; GFX1030-NEXT:  .LBB7_1: ; =>This Inner Loop Header: Depth=1
-; GFX1030-NEXT:    v_readfirstlane_b32 s4, v10
-; GFX1030-NEXT:    v_readfirstlane_b32 s5, v11
-; GFX1030-NEXT:    v_readfirstlane_b32 s6, v12
-; GFX1030-NEXT:    v_readfirstlane_b32 s7, v13
-; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[10:11]
-; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[12:13]
+; GFX1030-NEXT:    v_readfirstlane_b32 s4, v9
+; GFX1030-NEXT:    v_readfirstlane_b32 s5, v10
+; GFX1030-NEXT:    v_readfirstlane_b32 s6, v11
+; GFX1030-NEXT:    v_readfirstlane_b32 s7, v12
+; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[9:10]
+; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[11:12]
 ; GFX1030-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1030-NEXT:    s_and_saveexec_b32 s0, s0
-; GFX1030-NEXT:    image_bvh_intersect_ray v[0:3], v[14:21], s[4:7] a16
-; GFX1030-NEXT:    ; implicit-def: $vgpr10_vgpr11
+; GFX1030-NEXT:    image_bvh_intersect_ray v[0:3], v[13:20], s[4:7] a16
+; GFX1030-NEXT:    ; implicit-def: $vgpr9_vgpr10
+; GFX1030-NEXT:    ; implicit-def: $vgpr13
 ; GFX1030-NEXT:    ; implicit-def: $vgpr14
 ; GFX1030-NEXT:    ; implicit-def: $vgpr15
 ; GFX1030-NEXT:    ; implicit-def: $vgpr16
 ; GFX1030-NEXT:    ; implicit-def: $vgpr17
 ; GFX1030-NEXT:    ; implicit-def: $vgpr18
 ; GFX1030-NEXT:    ; implicit-def: $vgpr19
 ; GFX1030-NEXT:    ; implicit-def: $vgpr20
-; GFX1030-NEXT:    ; implicit-def: $vgpr21
-; GFX1030-NEXT:    ; implicit-def: $vgpr10_vgpr11_vgpr12_vgpr13
+; GFX1030-NEXT:    ; implicit-def: $vgpr9_vgpr10_vgpr11_vgpr12
 ; GFX1030-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1030-NEXT:    s_cbranch_execnz .LBB7_1
 ; GFX1030-NEXT:  ; %bb.2:
 ; GFX1030-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1030-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1030-NEXT:    ; return to shader part epilog
 ;
 ; GFX1013-LABEL: image_bvh_intersect_ray_a16_vgpr_descr:
 ; GFX1013:       ; %bb.0:
 ; GFX1013-NEXT:    s_mov_b32 s0, 0xffff
-; GFX1013-NEXT:    v_lshrrev_b32_e32 v5, 16, v6
-; GFX1013-NEXT:    v_and_b32_e32 v14, s0, v8
-; GFX1013-NEXT:    v_lshrrev_b32_e32 v8, 16, v8
-; GFX1013-NEXT:    v_and_b32_e32 v9, s0, v9
+; GFX1013-NEXT:    v_lshrrev_b32_e32 v13, 16, v5
+; GFX1013-NEXT:    v_and_b32_e32 v14, s0, v7
+; GFX1013-NEXT:    v_lshrrev_b32_e32 v7, 16, v7
+; GFX1013-NEXT:    v_and_b32_e32 v8, s0, v8
 ; GFX1013-NEXT:    s_mov_b32 s1, exec_lo
-; GFX1013-NEXT:    v_lshlrev_b32_e32 v5, 16, v5
+; GFX1013-NEXT:    v_lshlrev_b32_e32 v13, 16, v13
 ; GFX1013-NEXT:    v_lshlrev_b32_e32 v14, 16, v14
-; GFX1013-NEXT:    v_and_or_b32 v5, v6, s0, v5
-; GFX1013-NEXT:    v_and_or_b32 v6, v7, s0, v14
-; GFX1013-NEXT:    v_lshl_or_b32 v7, v9, 16, v8
+; GFX1013-NEXT:    v_lshl_or_b32 v7, v8, 16, v7
+; GFX1013-NEXT:    v_and_or_b32 v5, v5, s0, v13
+; GFX1013-NEXT:    v_and_or_b32 v6, v6, s0, v14
 ; GFX1013-NEXT:  .LBB7_1: ; =>This Inner Loop Header: Depth=1
-; GFX1013-NEXT:    v_readfirstlane_b32 s4, v10
-; GFX1013-NEXT:    v_readfirstlane_b32 s5, v11
-; GFX1013-NEXT:    v_readfirstlane_b32 s6, v12
-; GFX1013-NEXT:    v_readfirstlane_b32 s7, v13
-; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[10:11]
-; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[12:13]
+; GFX1013-NEXT:    v_readfirstlane_b32 s4, v9
+; GFX1013-NEXT:    v_readfirstlane_b32 s5, v10
+; GFX1013-NEXT:    v_readfirstlane_b32 s6, v11
+; GFX1013-NEXT:    v_readfirstlane_b32 s7, v12
+; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[9:10]
+; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[11:12]
 ; GFX1013-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1013-NEXT:    s_and_saveexec_b32 s0, s0
-; GFX1013-NEXT:    image_bvh_intersect_ray v[14:17], v[0:7], s[4:7] a16
-; GFX1013-NEXT:    ; implicit-def: $vgpr10_vgpr11
+; GFX1013-NEXT:    image_bvh_intersect_ray v[13:16], v[0:7], s[4:7] a16
+; GFX1013-NEXT:    ; implicit-def: $vgpr9_vgpr10
 ; GFX1013-NEXT:    ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7
-; GFX1013-NEXT:    ; implicit-def: $vgpr10_vgpr11_vgpr12_vgpr13
+; GFX1013-NEXT:    ; implicit-def: $vgpr9_vgpr10_vgpr11_vgpr12
 ; GFX1013-NEXT:    s_waitcnt_depctr 0xffe3
 ; GFX1013-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1013-NEXT:    s_cbranch_execnz .LBB7_1
 ; GFX1013-NEXT:  ; %bb.2:
 ; GFX1013-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
-; GFX1013-NEXT:    v_mov_b32_e32 v0, v14
-; GFX1013-NEXT:    v_mov_b32_e32 v1, v15
-; GFX1013-NEXT:    v_mov_b32_e32 v2, v16
-; GFX1013-NEXT:    v_mov_b32_e32 v3, v17
+; GFX1013-NEXT:    v_mov_b32_e32 v0, v13
+; GFX1013-NEXT:    v_mov_b32_e32 v1, v14
+; GFX1013-NEXT:    v_mov_b32_e32 v2, v15
+; GFX1013-NEXT:    v_mov_b32_e32 v3, v16
 ; GFX1013-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_vgpr_descr(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_vgpr_descr(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr) {
 ; GFX1030-LABEL: image_bvh64_intersect_ray_vgpr_descr:
 ; GFX1030:       ; %bb.0:
-; GFX1030-NEXT:    v_mov_b32_e32 v19, v0
-; GFX1030-NEXT:    v_mov_b32_e32 v20, v1
-; GFX1030-NEXT:    v_mov_b32_e32 v21, v2
-; GFX1030-NEXT:    v_mov_b32_e32 v22, v3
-; GFX1030-NEXT:    v_mov_b32_e32 v23, v4
-; GFX1030-NEXT:    v_mov_b32_e32 v24, v5
-; GFX1030-NEXT:    v_mov_b32_e32 v25, v7
-; GFX1030-NEXT:    v_mov_b32_e32 v26, v8
-; GFX1030-NEXT:    v_mov_b32_e32 v27, v9
-; GFX1030-NEXT:    v_mov_b32_e32 v28, v11
-; GFX1030-NEXT:    v_mov_b32_e32 v29, v12
-; GFX1030-NEXT:    v_mov_b32_e32 v30, v13
+; GFX1030-NEXT:    v_mov_b32_e32 v16, v0
+; GFX1030-NEXT:    v_mov_b32_e32 v17, v1
+; GFX1030-NEXT:    v_mov_b32_e32 v18, v2
+; GFX1030-NEXT:    v_mov_b32_e32 v19, v3
+; GFX1030-NEXT:    v_mov_b32_e32 v20, v4
+; GFX1030-NEXT:    v_mov_b32_e32 v21, v5
+; GFX1030-NEXT:    v_mov_b32_e32 v22, v6
+; GFX1030-NEXT:    v_mov_b32_e32 v23, v7
+; GFX1030-NEXT:    v_mov_b32_e32 v24, v8
+; GFX1030-NEXT:    v_mov_b32_e32 v25, v9
+; GFX1030-NEXT:    v_mov_b32_e32 v26, v10
+; GFX1030-NEXT:    v_mov_b32_e32 v27, v11
 ; GFX1030-NEXT:    s_mov_b32 s1, exec_lo
 ; GFX1030-NEXT:  .LBB8_1: ; =>This Inner Loop Header: Depth=1
-; GFX1030-NEXT:    v_readfirstlane_b32 s4, v15
-; GFX1030-NEXT:    v_readfirstlane_b32 s5, v16
-; GFX1030-NEXT:    v_readfirstlane_b32 s6, v17
-; GFX1030-NEXT:    v_readfirstlane_b32 s7, v18
-; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[15:16]
-; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[17:18]
+; GFX1030-NEXT:    v_readfirstlane_b32 s4, v12
+; GFX1030-NEXT:    v_readfirstlane_b32 s5, v13
+; GFX1030-NEXT:    v_readfirstlane_b32 s6, v14
+; GFX1030-NEXT:    v_readfirstlane_b32 s7, v15
+; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[12:13]
+; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[14:15]
 ; GFX1030-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1030-NEXT:    s_and_saveexec_b32 s0, s0
-; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], v[19:34], s[4:7]
-; GFX1030-NEXT:    ; implicit-def: $vgpr15_vgpr16
+; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], v[16:31], s[4:7]
+; GFX1030-NEXT:    ; implicit-def: $vgpr12_vgpr13
+; GFX1030-NEXT:    ; implicit-def: $vgpr16
+; GFX1030-NEXT:    ; implicit-def: $vgpr17
+; GFX1030-NEXT:    ; implicit-def: $vgpr18
 ; GFX1030-NEXT:    ; implicit-def: $vgpr19
 ; GFX1030-NEXT:    ; implicit-def: $vgpr20
 ; GFX1030-NEXT:    ; implicit-def: $vgpr21
 ; GFX1030-NEXT:    ; implicit-def: $vgpr22
 ; GFX1030-NEXT:    ; implicit-def: $vgpr23
 ; GFX1030-NEXT:    ; implicit-def: $vgpr24
 ; GFX1030-NEXT:    ; implicit-def: $vgpr25
 ; GFX1030-NEXT:    ; implicit-def: $vgpr26
 ; GFX1030-NEXT:    ; implicit-def: $vgpr27
-; GFX1030-NEXT:    ; implicit-def: $vgpr28
-; GFX1030-NEXT:    ; implicit-def: $vgpr29
-; GFX1030-NEXT:    ; implicit-def: $vgpr30
-; GFX1030-NEXT:    ; implicit-def: $vgpr15_vgpr16_vgpr17_vgpr18
+; GFX1030-NEXT:    ; implicit-def: $vgpr12_vgpr13_vgpr14_vgpr15
 ; GFX1030-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1030-NEXT:    s_cbranch_execnz .LBB8_1
 ; GFX1030-NEXT:  ; %bb.2:
 ; GFX1030-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1030-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1030-NEXT:    ; return to shader part epilog
 ;
 ; GFX1013-LABEL: image_bvh64_intersect_ray_vgpr_descr:
 ; GFX1013:       ; %bb.0:
-; GFX1013-NEXT:    v_mov_b32_e32 v6, v7
-; GFX1013-NEXT:    v_mov_b32_e32 v7, v8
-; GFX1013-NEXT:    v_mov_b32_e32 v8, v9
-; GFX1013-NEXT:    v_mov_b32_e32 v9, v11
-; GFX1013-NEXT:    v_mov_b32_e32 v10, v12
-; GFX1013-NEXT:    v_mov_b32_e32 v11, v13
+; GFX1013-NEXT:    v_mov_b32_e32 v16, v12
+; GFX1013-NEXT:    v_mov_b32_e32 v17, v13
+; GFX1013-NEXT:    v_mov_b32_e32 v18, v14
 ; GFX1013-NEXT:    v_mov_b32_e32 v19, v15
-; GFX1013-NEXT:    v_mov_b32_e32 v20, v16
 ; GFX1013-NEXT:    s_mov_b32 s1, exec_lo
 ; GFX1013-NEXT:  .LBB8_1: ; =>This Inner Loop Header: Depth=1
-; GFX1013-NEXT:    v_readfirstlane_b32 s4, v19
-; GFX1013-NEXT:    v_readfirstlane_b32 s5, v20
-; GFX1013-NEXT:    v_readfirstlane_b32 s6, v17
-; GFX1013-NEXT:    v_readfirstlane_b32 s7, v18
-; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[19:20]
-; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[17:18]
+; GFX1013-NEXT:    v_readfirstlane_b32 s4, v16
+; GFX1013-NEXT:    v_readfirstlane_b32 s5, v17
+; GFX1013-NEXT:    v_readfirstlane_b32 s6, v18
+; GFX1013-NEXT:    v_readfirstlane_b32 s7, v19
+; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[16:17]
+; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[18:19]
 ; GFX1013-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1013-NEXT:    s_and_saveexec_b32 s0, s0
-; GFX1013-NEXT:    image_bvh64_intersect_ray v[21:24], v[0:15], s[4:7]
+; GFX1013-NEXT:    image_bvh64_intersect_ray v[20:23], v[0:15], s[4:7]
+; GFX1013-NEXT:    ; implicit-def: $vgpr16_vgpr17
 ; GFX1013-NEXT:    ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15
-; GFX1013-NEXT:    ; implicit-def: $vgpr19_vgpr20
-; GFX1013-NEXT:    ; implicit-def: $vgpr15_vgpr16_vgpr17_vgpr18
+; GFX1013-NEXT:    ; implicit-def: $vgpr16_vgpr17_vgpr18_vgpr19
 ; GFX1013-NEXT:    s_waitcnt_depctr 0xffe3
 ; GFX1013-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1013-NEXT:    s_cbranch_execnz .LBB8_1
 ; GFX1013-NEXT:  ; %bb.2:
 ; GFX1013-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
-; GFX1013-NEXT:    v_mov_b32_e32 v0, v21
-; GFX1013-NEXT:    v_mov_b32_e32 v1, v22
-; GFX1013-NEXT:    v_mov_b32_e32 v2, v23
-; GFX1013-NEXT:    v_mov_b32_e32 v3, v24
+; GFX1013-NEXT:    v_mov_b32_e32 v0, v20
+; GFX1013-NEXT:    v_mov_b32_e32 v1, v21
+; GFX1013-NEXT:    v_mov_b32_e32 v2, v22
+; GFX1013-NEXT:    v_mov_b32_e32 v3, v23
 ; GFX1013-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_a16_vgpr_descr(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_a16_vgpr_descr(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr) {
 ; GFX1030-LABEL: image_bvh64_intersect_ray_a16_vgpr_descr:
 ; GFX1030:       ; %bb.0:
 ; GFX1030-NEXT:    s_mov_b32 s0, 0xffff
-; GFX1030-NEXT:    v_mov_b32_e32 v15, v0
-; GFX1030-NEXT:    v_mov_b32_e32 v16, v1
-; GFX1030-NEXT:    v_lshrrev_b32_e32 v0, 16, v7
-; GFX1030-NEXT:    v_and_b32_e32 v1, s0, v9
-; GFX1030-NEXT:    v_mov_b32_e32 v17, v2
-; GFX1030-NEXT:    v_mov_b32_e32 v18, v3
-; GFX1030-NEXT:    v_lshrrev_b32_e32 v2, 16, v9
+; GFX1030-NEXT:    v_mov_b32_e32 v14, v0
+; GFX1030-NEXT:    v_mov_b32_e32 v15, v1
+; GFX1030-NEXT:    v_lshrrev_b32_e32 v0, 16, v6
+; GFX1030-NEXT:    v_and_b32_e32 v1, s0, v8
+; GFX1030-NEXT:    v_mov_b32_e32 v16, v2
+; GFX1030-NEXT:    v_mov_b32_e32 v17, v3
+; GFX1030-NEXT:    v_lshrrev_b32_e32 v2, 16, v8
 ; GFX1030-NEXT:    v_lshlrev_b32_e32 v0, 16, v0
 ; GFX1030-NEXT:    v_lshlrev_b32_e32 v1, 16, v1
-; GFX1030-NEXT:    v_and_b32_e32 v3, s0, v10
-; GFX1030-NEXT:    v_mov_b32_e32 v19, v4
-; GFX1030-NEXT:    v_mov_b32_e32 v20, v5
-; GFX1030-NEXT:    v_and_or_b32 v21, v7, s0, v0
-; GFX1030-NEXT:    v_and_or_b32 v22, v8, s0, v1
-; GFX1030-NEXT:    v_lshl_or_b32 v23, v3, 16, v2
+; GFX1030-NEXT:    v_and_b32_e32 v3, s0, v9
+; GFX1030-NEXT:    v_mov_b32_e32 v18, v4
+; GFX1030-NEXT:    v_mov_b32_e32 v19, v5
+; GFX1030-NEXT:    v_and_or_b32 v20, v6, s0, v0
+; GFX1030-NEXT:    v_and_or_b32 v21, v7, s0, v1
+; GFX1030-NEXT:    v_lshl_or_b32 v22, v3, 16, v2
 ; GFX1030-NEXT:    s_mov_b32 s1, exec_lo
 ; GFX1030-NEXT:  .LBB9_1: ; =>This Inner Loop Header: Depth=1
-; GFX1030-NEXT:    v_readfirstlane_b32 s4, v11
-; GFX1030-NEXT:    v_readfirstlane_b32 s5, v12
-; GFX1030-NEXT:    v_readfirstlane_b32 s6, v13
-; GFX1030-NEXT:    v_readfirstlane_b32 s7, v14
-; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[11:12]
-; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[13:14]
+; GFX1030-NEXT:    v_readfirstlane_b32 s4, v10
+; GFX1030-NEXT:    v_readfirstlane_b32 s5, v11
+; GFX1030-NEXT:    v_readfirstlane_b32 s6, v12
+; GFX1030-NEXT:    v_readfirstlane_b32 s7, v13
+; GFX1030-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[10:11]
+; GFX1030-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[12:13]
 ; GFX1030-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1030-NEXT:    s_and_saveexec_b32 s0, s0
-; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], v[15:30], s[4:7] a16
-; GFX1030-NEXT:    ; implicit-def: $vgpr11_vgpr12
+; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], v[14:29], s[4:7] a16
+; GFX1030-NEXT:    ; implicit-def: $vgpr10_vgpr11
+; GFX1030-NEXT:    ; implicit-def: $vgpr14
 ; GFX1030-NEXT:    ; implicit-def: $vgpr15
 ; GFX1030-NEXT:    ; implicit-def: $vgpr16
 ; GFX1030-NEXT:    ; implicit-def: $vgpr17
 ; GFX1030-NEXT:    ; implicit-def: $vgpr18
 ; GFX1030-NEXT:    ; implicit-def: $vgpr19
 ; GFX1030-NEXT:    ; implicit-def: $vgpr20
 ; GFX1030-NEXT:    ; implicit-def: $vgpr21
 ; GFX1030-NEXT:    ; implicit-def: $vgpr22
-; GFX1030-NEXT:    ; implicit-def: $vgpr23
-; GFX1030-NEXT:    ; implicit-def: $vgpr11_vgpr12_vgpr13_vgpr14
+; GFX1030-NEXT:    ; implicit-def: $vgpr10_vgpr11_vgpr12_vgpr13
 ; GFX1030-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1030-NEXT:    s_cbranch_execnz .LBB9_1
 ; GFX1030-NEXT:  ; %bb.2:
 ; GFX1030-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1030-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1030-NEXT:    ; return to shader part epilog
 ;
 ; GFX1013-LABEL: image_bvh64_intersect_ray_a16_vgpr_descr:
 ; GFX1013:       ; %bb.0:
 ; GFX1013-NEXT:    s_mov_b32 s0, 0xffff
-; GFX1013-NEXT:    v_mov_b32_e32 v16, v11
-; GFX1013-NEXT:    v_lshrrev_b32_e32 v6, 16, v7
-; GFX1013-NEXT:    v_and_b32_e32 v11, s0, v9
-; GFX1013-NEXT:    v_lshrrev_b32_e32 v9, 16, v9
-; GFX1013-NEXT:    v_and_b32_e32 v10, s0, v10
-; GFX1013-NEXT:    v_mov_b32_e32 v17, v12
-; GFX1013-NEXT:    v_lshlrev_b32_e32 v6, 16, v6
+; GFX1013-NEXT:    v_mov_b32_e32 v16, v10
+; GFX1013-NEXT:    v_mov_b32_e32 v17, v11
+; GFX1013-NEXT:    v_lshrrev_b32_e32 v10, 16, v6
+; GFX1013-NEXT:    v_and_b32_e32 v11, s0, v8
+; GFX1013-NEXT:    v_lshrrev_b32_e32 v8, 16, v8
+; GFX1013-NEXT:    v_and_b32_e32 v9, s0, v9
+; GFX1013-NEXT:    v_mov_b32_e32 v18, v12
+; GFX1013-NEXT:    v_lshlrev_b32_e32 v10, 16, v10
 ; GFX1013-NEXT:    v_lshlrev_b32_e32 v11, 16, v11
-; GFX1013-NEXT:    v_mov_b32_e32 v18, v13
-; GFX1013-NEXT:    v_mov_b32_e32 v19, v14
+; GFX1013-NEXT:    v_mov_b32_e32 v19, v13
+; GFX1013-NEXT:    v_lshl_or_b32 v8, v9, 16, v8
 ; GFX1013-NEXT:    s_mov_b32 s1, exec_lo
-; GFX1013-NEXT:    v_and_or_b32 v6, v7, s0, v6
-; GFX1013-NEXT:    v_and_or_b32 v7, v8, s0, v11
-; GFX1013-NEXT:    v_lshl_or_b32 v8, v10, 16, v9
+; GFX1013-NEXT:    v_and_or_b32 v6, v6, s0, v10
+; GFX1013-NEXT:    v_and_or_b32 v7, v7, s0, v11
 ; GFX1013-NEXT:  .LBB9_1: ; =>This Inner Loop Header: Depth=1
 ; GFX1013-NEXT:    v_readfirstlane_b32 s4, v16
 ; GFX1013-NEXT:    v_readfirstlane_b32 s5, v17
 ; GFX1013-NEXT:    v_readfirstlane_b32 s6, v18
 ; GFX1013-NEXT:    v_readfirstlane_b32 s7, v19
 ; GFX1013-NEXT:    v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[16:17]
 ; GFX1013-NEXT:    v_cmp_eq_u64_e64 s0, s[6:7], v[18:19]
 ; GFX1013-NEXT:    s_and_b32 s0, s0, vcc_lo
 ; GFX1013-NEXT:    s_and_saveexec_b32 s0, s0
 ; GFX1013-NEXT:    image_bvh64_intersect_ray v[20:23], v[0:15], s[4:7] a16
 ; GFX1013-NEXT:    ; implicit-def: $vgpr16_vgpr17
 ; GFX1013-NEXT:    ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15
 ; GFX1013-NEXT:    ; implicit-def: $vgpr16_vgpr17_vgpr18_vgpr19
 ; GFX1013-NEXT:    s_waitcnt_depctr 0xffe3
 ; GFX1013-NEXT:    s_xor_b32 exec_lo, exec_lo, s0
 ; GFX1013-NEXT:    s_cbranch_execnz .LBB9_1
 ; GFX1013-NEXT:  ; %bb.2:
 ; GFX1013-NEXT:    s_mov_b32 exec_lo, s1
 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1013-NEXT:    v_mov_b32_e32 v0, v20
 ; GFX1013-NEXT:    v_mov_b32_e32 v1, v21
 ; GFX1013-NEXT:    v_mov_b32_e32 v2, v22
 ; GFX1013-NEXT:    v_mov_b32_e32 v3, v23
 ; GFX1013-NEXT:    ; return to shader part epilog
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
 define amdgpu_kernel void @image_bvh_intersect_ray_nsa_reassign(i32* %p_node_ptr, float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1030-LABEL: image_bvh_intersect_ray_nsa_reassign:
 ; GFX1030:       ; %bb.0:
 ; GFX1030-NEXT:    s_clause 0x1

 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1013-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1013-NEXT:    s_endpgm
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_node_ptr = getelementptr inbounds i32, i32* %p_node_ptr, i32 %lid
   %node_ptr = load i32, i32* %gep_node_ptr, align 4
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float 3.0, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float 4.0, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float 3.0, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float 4.0, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 define amdgpu_kernel void @image_bvh_intersect_ray_a16_nsa_reassign(i32* %p_node_ptr, float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1030-LABEL: image_bvh_intersect_ray_a16_nsa_reassign:
 ; GFX1030:       ; %bb.0:
 ; GFX1030-NEXT:    s_clause 0x1

 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1013-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1013-NEXT:    s_endpgm
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_node_ptr = getelementptr inbounds i32, i32* %p_node_ptr, i32 %lid
   %node_ptr = load i32, i32* %gep_node_ptr, align 4
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x half> undef, half 3.0, i32 0
-  %ray_dir1 = insertelement <4 x half> %ray_dir0, half 4.0, i32 1
-  %ray_dir = insertelement <4 x half> %ray_dir1, half 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x half> undef, half 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x half> %ray_inv_dir0, half 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x half> %ray_inv_dir1, half 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x half> undef, half 3.0, i32 0
+  %ray_dir1 = insertelement <3 x half> %ray_dir0, half 4.0, i32 1
+  %ray_dir = insertelement <3 x half> %ray_dir1, half 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x half> undef, half 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x half> %ray_inv_dir0, half 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x half> %ray_inv_dir1, half 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 define amdgpu_kernel void @image_bvh64_intersect_ray_nsa_reassign(float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1030-LABEL: image_bvh64_intersect_ray_nsa_reassign:
 ; GFX1030:       ; %bb.0:
 ; GFX1030-NEXT:    s_clause 0x1

 ; GFX1013-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
 ; GFX1013-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[4:7]
 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1013-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1013-NEXT:    s_endpgm
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float 3.0, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float 4.0, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 1111111111111, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float 3.0, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float 4.0, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 1111111111111, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 define amdgpu_kernel void @image_bvh64_intersect_ray_a16_nsa_reassign(float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1030-LABEL: image_bvh64_intersect_ray_a16_nsa_reassign:
 ; GFX1030:       ; %bb.0:
 ; GFX1030-NEXT:    s_clause 0x1

 ; GFX1013-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
 ; GFX1013-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[4:7] a16
 ; GFX1013-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1013-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1013-NEXT:    s_endpgm
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x half> undef, half 3.0, i32 0
-  %ray_dir1 = insertelement <4 x half> %ray_dir0, half 4.0, i32 1
-  %ray_dir = insertelement <4 x half> %ray_dir1, half 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x half> undef, half 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x half> %ray_inv_dir0, half 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x half> %ray_inv_dir1, half 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 1111111111110, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x half> undef, half 3.0, i32 0
+  %ray_dir1 = insertelement <3 x half> %ray_dir0, half 4.0, i32 1
+  %ray_dir = insertelement <3 x half> %ray_dir1, half 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x half> undef, half 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x half> %ray_inv_dir0, half 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x half> %ray_inv_dir1, half 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 1111111111110, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }

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

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -march=amdgcn -mcpu=gfx1013 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX1013 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1030 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX1030 %s
 ; RUN: not --crash llc -march=amdgcn -mcpu=gfx1012 -verify-machineinstrs < %s 2>&1 | FileCheck -check-prefix=ERR %s
 
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, uint4 texture_descr)
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, uint4 texture_descr)
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(ulong node_ptr, float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, uint4 texture_descr)
-; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(ulong node_ptr, float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, uint4 texture_descr)
-
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32, float, <4 x float>, <4 x float>, <4 x float>, <4 x i32>)
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32, float, <4 x float>, <4 x half>, <4 x half>, <4 x i32>)
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64, float, <4 x float>, <4 x float>, <4 x float>, <4 x i32>)
-declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64, float, <4 x float>, <4 x half>, <4 x half>, <4 x i32>)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(ulong node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr)
+; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(ulong node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr)
+
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32, float, <3 x float>, <3 x float>, <3 x float>, <4 x i32>)
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32, float, <3 x float>, <3 x half>, <3 x half>, <4 x i32>)
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64, float, <3 x float>, <3 x float>, <3 x float>, <4 x i32>)
+declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64, float, <3 x float>, <3 x half>, <3 x half>, <4 x i32>)
 
 ; ERR: in function image_bvh_intersect_ray{{.*}}intrinsic not supported on subtarget
 ; Arguments are flattened to represent the actual VGPR_A layout, so we have no
 ; extra moves in the generated kernel.
 define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, float %ray_inv_dir_x, float %ray_inv_dir_y, float %ray_inv_dir_z, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh_intersect_ray:
 ; GCN:       ; %bb.0: ; %main_body
 ; GCN-NEXT:    image_bvh_intersect_ray v[0:3], v[0:15], s[0:3]
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
 main_body:
-  %ray_origin0 = insertelement <4 x float> undef, float %ray_origin_x, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float %ray_origin_y, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float %ray_origin_z, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float %ray_dir_x, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float %ray_dir_y, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float %ray_dir_z, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float %ray_inv_dir_x, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float %ray_origin_x, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float %ray_origin_y, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float %ray_origin_z, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float %ray_dir_x, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float %ray_dir_y, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float %ray_dir_z, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float %ray_inv_dir_x, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
  %r = bitcast <4 x i32> %v to <4 x float>
  ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh_intersect_ray_a16(i32 inreg %node_ptr, float inreg %ray_extent, <4 x float> inreg %ray_origin, <4 x half> inreg %ray_dir, <4 x half> inreg %ray_inv_dir, <4 x i32> inreg %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh_intersect_ray_a16(i32 inreg %node_ptr, float inreg %ray_extent, <3 x float> inreg %ray_origin, <3 x half> inreg %ray_dir, <3 x half> inreg %ray_inv_dir, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh_intersect_ray_a16:
 ; GCN:       ; %bb.0: ; %main_body
-; GCN-NEXT:    s_lshr_b32 s5, s8, 16
-; GCN-NEXT:    s_pack_ll_b32_b16 s7, s7, s8
-; GCN-NEXT:    s_pack_ll_b32_b16 s5, s5, s9
+; GCN-NEXT:    s_mov_b32 s15, s12
+; GCN-NEXT:    s_mov_b32 s12, s9
+; GCN-NEXT:    s_lshr_b32 s9, s7, 16
+; GCN-NEXT:    s_pack_ll_b32_b16 s6, s6, s7
+; GCN-NEXT:    s_pack_ll_b32_b16 s7, s9, s8
 ; GCN-NEXT:    v_mov_b32_e32 v0, s0
 ; GCN-NEXT:    v_mov_b32_e32 v1, s1
 ; GCN-NEXT:    v_mov_b32_e32 v2, s2
 ; GCN-NEXT:    v_mov_b32_e32 v3, s3
 ; GCN-NEXT:    v_mov_b32_e32 v4, s4
-; GCN-NEXT:    v_mov_b32_e32 v5, s6
-; GCN-NEXT:    v_mov_b32_e32 v6, s7
-; GCN-NEXT:    v_mov_b32_e32 v7, s5
-; GCN-NEXT:    s_mov_b32 s15, s13
-; GCN-NEXT:    s_mov_b32 s14, s12
-; GCN-NEXT:    s_mov_b32 s13, s11
-; GCN-NEXT:    s_mov_b32 s12, s10
+; GCN-NEXT:    v_mov_b32_e32 v5, s5
+; GCN-NEXT:    v_mov_b32_e32 v6, s6
+; GCN-NEXT:    v_mov_b32_e32 v7, s7
+; GCN-NEXT:    s_mov_b32 s14, s11
+; GCN-NEXT:    s_mov_b32 s13, s10
 ; GCN-NEXT:    image_bvh_intersect_ray v[0:3], v[0:7], s[12:15] a16
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
 main_body:
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
 ; Arguments are flattened to represent the actual VGPR_A layout, so we have no
 ; extra moves in the generated kernel.
 define amdgpu_ps <4 x float> @image_bvh64_intersect_ray(<2 x i32> %node_ptr_vec, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, float %ray_inv_dir_x, float %ray_inv_dir_y, float %ray_inv_dir_z, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh64_intersect_ray:
 ; GCN:       ; %bb.0: ; %main_body
 ; GCN-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3]
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
 main_body:
   %node_ptr = bitcast <2 x i32> %node_ptr_vec to i64
-  %ray_origin0 = insertelement <4 x float> undef, float %ray_origin_x, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float %ray_origin_y, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float %ray_origin_z, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float %ray_dir_x, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float %ray_dir_y, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float %ray_dir_z, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float %ray_inv_dir_x, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float %ray_origin_x, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float %ray_origin_y, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float %ray_origin_z, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float %ray_dir_x, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float %ray_dir_y, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float %ray_dir_z, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float %ray_inv_dir_x, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float %ray_inv_dir_y, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float %ray_inv_dir_z, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
  %r = bitcast <4 x i32> %v to <4 x float>
  ret <4 x float> %r
 }
 
-define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_a16(i64 inreg %node_ptr, float inreg %ray_extent, <4 x float> inreg %ray_origin, <4 x half> inreg %ray_dir, <4 x half> inreg %ray_inv_dir, <4 x i32> inreg %tdescr) {
+define amdgpu_ps <4 x float> @image_bvh64_intersect_ray_a16(i64 inreg %node_ptr, float inreg %ray_extent, <3 x float> inreg %ray_origin, <3 x half> inreg %ray_dir, <3 x half> inreg %ray_inv_dir, <4 x i32> inreg %tdescr) {
 ; GCN-LABEL: image_bvh64_intersect_ray_a16:
 ; GCN:       ; %bb.0: ; %main_body
-; GCN-NEXT:    s_lshr_b32 s6, s9, 16
-; GCN-NEXT:    s_pack_ll_b32_b16 s8, s8, s9
-; GCN-NEXT:    s_pack_ll_b32_b16 s6, s6, s10
+; GCN-NEXT:    s_mov_b32 s14, s12
+; GCN-NEXT:    s_mov_b32 s12, s10
+; GCN-NEXT:    s_lshr_b32 s10, s8, 16
+; GCN-NEXT:    s_pack_ll_b32_b16 s7, s7, s8
+; GCN-NEXT:    s_pack_ll_b32_b16 s8, s10, s9
 ; GCN-NEXT:    v_mov_b32_e32 v0, s0
 ; GCN-NEXT:    v_mov_b32_e32 v1, s1
 ; GCN-NEXT:    v_mov_b32_e32 v2, s2
 ; GCN-NEXT:    v_mov_b32_e32 v3, s3
 ; GCN-NEXT:    v_mov_b32_e32 v4, s4
 ; GCN-NEXT:    v_mov_b32_e32 v5, s5
-; GCN-NEXT:    v_mov_b32_e32 v6, s7
-; GCN-NEXT:    v_mov_b32_e32 v7, s8
-; GCN-NEXT:    v_mov_b32_e32 v8, s6
-; GCN-NEXT:    s_mov_b32 s15, s14
-; GCN-NEXT:    s_mov_b32 s14, s13
-; GCN-NEXT:    s_mov_b32 s13, s12
-; GCN-NEXT:    s_mov_b32 s12, s11
+; GCN-NEXT:    v_mov_b32_e32 v6, s6
+; GCN-NEXT:    v_mov_b32_e32 v7, s7
+; GCN-NEXT:    v_mov_b32_e32 v8, s8
+; GCN-NEXT:    s_mov_b32 s15, s13
+; GCN-NEXT:    s_mov_b32 s13, s11
 ; GCN-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[12:15] a16
 ; GCN-NEXT:    s_waitcnt vmcnt(0)
 ; GCN-NEXT:    ; return to shader part epilog
 main_body:
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   %r = bitcast <4 x i32> %v to <4 x float>
   ret <4 x float> %r
 }
 
 ; TODO: NSA reassign is very limited and cannot work with VGPR tuples and subregs.
 
 define amdgpu_kernel void @image_bvh_intersect_ray_nsa_reassign(i32* %p_node_ptr, float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1013-LABEL: image_bvh_intersect_ray_nsa_reassign:

 ; GFX1030-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1030-NEXT:    s_endpgm
 main_body:
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_node_ptr = getelementptr inbounds i32, i32* %p_node_ptr, i32 %lid
   %node_ptr = load i32, i32* %gep_node_ptr, align 4
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float 3.0, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float 4.0, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float 3.0, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float 4.0, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 define amdgpu_kernel void @image_bvh_intersect_ray_a16_nsa_reassign(i32* %p_node_ptr, float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1013-LABEL: image_bvh_intersect_ray_a16_nsa_reassign:
 ; GFX1013:       ; %bb.0: ; %main_body
 ; GFX1013-NEXT:    s_load_dwordx4 s[4:7], s[0:1], 0x24

 ; GFX1030-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1030-NEXT:    s_endpgm
 main_body:
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_node_ptr = getelementptr inbounds i32, i32* %p_node_ptr, i32 %lid
   %node_ptr = load i32, i32* %gep_node_ptr, align 4
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x half> undef, half 3.0, i32 0
-  %ray_dir1 = insertelement <4 x half> %ray_dir0, half 4.0, i32 1
-  %ray_dir = insertelement <4 x half> %ray_dir1, half 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x half> undef, half 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x half> %ray_inv_dir0, half 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x half> %ray_inv_dir1, half 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x half> undef, half 3.0, i32 0
+  %ray_dir1 = insertelement <3 x half> %ray_dir0, half 4.0, i32 1
+  %ray_dir = insertelement <3 x half> %ray_dir1, half 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x half> undef, half 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x half> %ray_inv_dir0, half 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x half> %ray_inv_dir1, half 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 define amdgpu_kernel void @image_bvh64_intersect_ray_nsa_reassign(float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1013-LABEL: image_bvh64_intersect_ray_nsa_reassign:
 ; GFX1013:       ; %bb.0: ; %main_body
 ; GFX1013-NEXT:    s_load_dwordx2 s[4:5], s[0:1], 0x24

 ; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3]
 ; GFX1030-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1030-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1030-NEXT:    s_endpgm
 main_body:
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x float> undef, float 3.0, i32 0
-  %ray_dir1 = insertelement <4 x float> %ray_dir0, float 4.0, i32 1
-  %ray_dir = insertelement <4 x float> %ray_dir1, float 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x float> undef, float 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x float> %ray_inv_dir0, float 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x float> %ray_inv_dir1, float 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 1111111111111, float %ray_extent, <4 x float> %ray_origin, <4 x float> %ray_dir, <4 x float> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x float> undef, float 3.0, i32 0
+  %ray_dir1 = insertelement <3 x float> %ray_dir0, float 4.0, i32 1
+  %ray_dir = insertelement <3 x float> %ray_dir1, float 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x float> undef, float 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x float> %ray_inv_dir0, float 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x float> %ray_inv_dir1, float 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64 1111111111111, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 define amdgpu_kernel void @image_bvh64_intersect_ray_a16_nsa_reassign(float* %p_ray, <4 x i32> inreg %tdescr) {
 ; GFX1013-LABEL: image_bvh64_intersect_ray_a16_nsa_reassign:
 ; GFX1013:       ; %bb.0: ; %main_body
 ; GFX1013-NEXT:    s_load_dwordx2 s[4:5], s[0:1], 0x24

 ; GFX1030-NEXT:    image_bvh64_intersect_ray v[0:3], v[0:15], s[0:3] a16
 ; GFX1030-NEXT:    s_waitcnt vmcnt(0)
 ; GFX1030-NEXT:    flat_store_dwordx4 v[0:1], v[0:3]
 ; GFX1030-NEXT:    s_endpgm
 main_body:
   %lid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %gep_ray = getelementptr inbounds float, float* %p_ray, i32 %lid
   %ray_extent = load float, float* %gep_ray, align 4
-  %ray_origin0 = insertelement <4 x float> undef, float 0.0, i32 0
-  %ray_origin1 = insertelement <4 x float> %ray_origin0, float 1.0, i32 1
-  %ray_origin = insertelement <4 x float> %ray_origin1, float 2.0, i32 2
-  %ray_dir0 = insertelement <4 x half> undef, half 3.0, i32 0
-  %ray_dir1 = insertelement <4 x half> %ray_dir0, half 4.0, i32 1
-  %ray_dir = insertelement <4 x half> %ray_dir1, half 5.0, i32 2
-  %ray_inv_dir0 = insertelement <4 x half> undef, half 6.0, i32 0
-  %ray_inv_dir1 = insertelement <4 x half> %ray_inv_dir0, half 7.0, i32 1
-  %ray_inv_dir = insertelement <4 x half> %ray_inv_dir1, half 8.0, i32 2
-  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 1111111111110, float %ray_extent, <4 x float> %ray_origin, <4 x half> %ray_dir, <4 x half> %ray_inv_dir, <4 x i32> %tdescr)
+  %ray_origin0 = insertelement <3 x float> undef, float 0.0, i32 0
+  %ray_origin1 = insertelement <3 x float> %ray_origin0, float 1.0, i32 1
+  %ray_origin = insertelement <3 x float> %ray_origin1, float 2.0, i32 2
+  %ray_dir0 = insertelement <3 x half> undef, half 3.0, i32 0
+  %ray_dir1 = insertelement <3 x half> %ray_dir0, half 4.0, i32 1
+  %ray_dir = insertelement <3 x half> %ray_dir1, half 5.0, i32 2
+  %ray_inv_dir0 = insertelement <3 x half> undef, half 6.0, i32 0
+  %ray_inv_dir1 = insertelement <3 x half> %ray_inv_dir0, half 7.0, i32 1
+  %ray_inv_dir = insertelement <3 x half> %ray_inv_dir1, half 8.0, i32 2
+  %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64 1111111111110, float %ray_extent, <3 x float> %ray_origin, <3 x half> %ray_dir, <3 x half> %ray_inv_dir, <4 x i32> %tdescr)
   store <4 x i32> %v, <4 x i32>* undef
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x()