File: builtins-amdgcn-raytracing.cl

package info (click to toggle)
llvm-toolchain-19 1%3A19.1.7-3
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid, trixie
  • size: 1,998,520 kB
  • sloc: cpp: 6,951,680; ansic: 1,486,157; asm: 913,598; python: 232,024; f90: 80,126; objc: 75,281; lisp: 37,276; pascal: 16,990; sh: 10,009; ml: 5,058; perl: 4,724; awk: 3,523; makefile: 3,167; javascript: 2,504; xml: 892; fortran: 664; cs: 573
file content (61 lines) | stat: -rw-r--r-- 2,624 bytes parent folder | download | duplicates (5)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 \
// 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.v3f32
// 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.v3f16
// 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.v3f32
// 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.v3f16
// 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);
}