1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \ 3// RUN: -emit-llvm -cl-std=CL2.0 -o - %s | FileCheck %s 4// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S \ 5// RUN: -cl-std=CL2.0 -o - %s | FileCheck -check-prefix=ISA %s 6 7// Test llvm.amdgcn.image.bvh.intersect.ray intrinsic. 8 9// The clang builtin functions __builtin_amdgcn_image_bvh_intersect_ray* use 10// postfixes to indicate the types of the 1st, 4th, and 5th arguments. 11// By default, the 1st argument is i32, the 4/5-th arguments are float4. 12// Postfix l indicates the 1st argument is i64 and postfix h indicates 13// the 4/5-th arguments are half4. 14 15typedef unsigned int uint; 16typedef unsigned long ulong; 17typedef float float4 __attribute__((ext_vector_type(4))); 18typedef double double4 __attribute__((ext_vector_type(4))); 19typedef half half4 __attribute__((ext_vector_type(4))); 20typedef uint uint4 __attribute__((ext_vector_type(4))); 21 22// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v3f32 23// ISA: image_bvh_intersect_ray 24void test_image_bvh_intersect_ray(global uint4* out, uint node_ptr, 25 float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, 26 uint4 texture_descr) 27{ 28 *out = __builtin_amdgcn_image_bvh_intersect_ray(node_ptr, ray_extent, 29 ray_origin, ray_dir, ray_inv_dir, texture_descr); 30} 31 32// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v3f16 33// ISA: image_bvh_intersect_ray 34void test_image_bvh_intersect_ray_h(global uint4* out, uint node_ptr, 35 float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, 36 uint4 texture_descr) 37{ 38 *out = __builtin_amdgcn_image_bvh_intersect_ray_h(node_ptr, ray_extent, 39 ray_origin, ray_dir, ray_inv_dir, texture_descr); 40} 41 42// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v3f32 43// ISA: image_bvh_intersect_ray 44void test_image_bvh_intersect_ray_l(global uint4* out, ulong node_ptr, 45 float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, 46 uint4 texture_descr) 47{ 48 *out = __builtin_amdgcn_image_bvh_intersect_ray_l(node_ptr, ray_extent, 49 ray_origin, ray_dir, ray_inv_dir, texture_descr); 50} 51 52// CHECK: call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v3f16 53// ISA: image_bvh_intersect_ray 54void test_image_bvh_intersect_ray_lh(global uint4* out, ulong node_ptr, 55 float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, 56 uint4 texture_descr) 57{ 58 *out = __builtin_amdgcn_image_bvh_intersect_ray_lh(node_ptr, ray_extent, 59 ray_origin, ray_dir, ray_inv_dir, texture_descr); 60} 61 62