Index: clang/include/clang/Basic/BuiltinsAMDGPU.def =================================================================== --- clang/include/clang/Basic/BuiltinsAMDGPU.def +++ clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -215,6 +215,14 @@ TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts") TARGET_BUILTIN(__builtin_amdgcn_mov_dpp8, "UiUiIUi", "nc", "gfx10-insts") +//===----------------------------------------------------------------------===// +// Raytracing builtins. +//===----------------------------------------------------------------------===// +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. //===----------------------------------------------------------------------===// 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/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl =================================================================== --- /dev/null +++ clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl @@ -0,0 +1,46 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S \ +// RUN: -emit-llvm -o - %s | FileCheck %s + +typedef unsigned int uint; +typedef unsigned long ulong; +typedef float float4 __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 +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 +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 +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 +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); +} +