[AMDGPU] Add builtin functions image_bvh_intersect_ray

Reviewed by: Stanislav Mekhanoshin, Matt Arsenault

Differential Revision: https://reviews.llvm.org/D104946

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -215,6 +215,17 @@ 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.
+// Postfix l indicates the 1st argument is i64.
+// Postfix h indicates the 4/5-th arguments are half4.
+//===----------------------------------------------------------------------===//
+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")
+
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//

clang/lib/CodeGen/CGBuiltin.cpp

@@ -15850,6 +15850,23 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     CI->setConvergent();
     return CI;
   }
+  case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
+  case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
+  case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
+  case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
+    llvm::Value *NodePtr = EmitScalarExpr(E->getArg(0));
+    llvm::Value *RayExtent = EmitScalarExpr(E->getArg(1));
+    llvm::Value *RayOrigin = EmitScalarExpr(E->getArg(2));
+    llvm::Value *RayDir = EmitScalarExpr(E->getArg(3));
+    llvm::Value *RayInverseDir = EmitScalarExpr(E->getArg(4));
+    llvm::Value *TextureDescr = EmitScalarExpr(E->getArg(5));
+
+    Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_image_bvh_intersect_ray,
+                                   {NodePtr->getType(), RayDir->getType()});
+    return Builder.CreateCall(F, {NodePtr, RayExtent, RayOrigin, RayDir,
+                                  RayInverseDir, TextureDescr});
+  }
+
   // amdgcn workitem
   case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0, 1024);

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

@@ -0,0 +1,61 @@
+// 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.
+// Postfix l indicates the 1st argument is i64 and postfix h indicates
+// the 4/5-th arguments are half4.
+
+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 uint uint4 __attribute__((ext_vector_type(4)));
+
+// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32
+// 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,
+  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
+// 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,
+  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
+// 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,
+  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
+// 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,
+  uint4 texture_descr)
+{
+  *out = __builtin_amdgcn_image_bvh_intersect_ray_lh(node_ptr, ray_extent,
+           ray_origin, ray_dir, ray_inv_dir, texture_descr);
+}
+