From: Yaxun (Sam) Liu Date: Sat, 21 Mar 2020 21:06:39 +0000 (-0400) Subject: [AMDGPU] Add builtin functions image_bvh_intersect_ray X-Git-Tag: llvmorg-14-init~2588 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=434bd5bf5479615ea42e64a80ff994059f31f5f4;p=platform%2Fupstream%2Fllvm.git [AMDGPU] Add builtin functions image_bvh_intersect_ray Reviewed by: Stanislav Mekhanoshin, Matt Arsenault Differential Revision: https://reviews.llvm.org/D104946 --- diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index f9d079a..3570431 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -216,6 +216,17 @@ TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb", "nc", "gfx10-inst 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. //===----------------------------------------------------------------------===// BUILTIN(__builtin_amdgcn_read_exec, "WUi", "nc") diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 0e13b55..9061abc 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/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); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl new file mode 100644 index 0000000..805d17a --- /dev/null +++ b/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); +} +