yaxunl updated this revision to Diff 355309.
yaxunl added a comment.
Add check for ISA. Add comments about argument types.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D104946/new/
https://reviews.llvm.org/D104946
Files:
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/lib/CodeGen/CGBuiltin.cpp
clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl
Index: clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl
===================================================================
--- /dev/null
+++ 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);
+}
+
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -15745,6 +15745,23 @@
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);
Index: clang/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -215,6 +215,17 @@
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.
//===----------------------------------------------------------------------===//
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
