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

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S \
   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
  15 typedef unsigned int uint;
  16 typedef unsigned long ulong;
  17 typedef float float4 __attribute__((ext_vector_type(4)));
  18 typedef double double4 __attribute__((ext_vector_type(4)));
  19 typedef half half4 __attribute__((ext_vector_type(4)));
  20 typedef 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
  24 void 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
  34 void 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
  44 void 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
  54 void 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