[PATCH] D104946: [AMDGPU] Add builtin functions image_bvh_intersect_ray

Yaxun Liu via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Mon Jun 28 09:50:38 PDT 2021


yaxunl marked an inline comment as done.
yaxunl added inline comments.


================
Comment at: clang/include/clang/Basic/BuiltinsAMDGPU.def:221-224
+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")
----------------
arsenm wrote:
> The intrinsic signature suggests the 1st and 4th/5th arguments are overloadable. How does this handle the various supported types?
By convention, we do not define overloaded target specific clang builtins. In stead we use postfixes to indicate the types. In this case, we use the following naming convention:

h - 4/5-th args are half
d - 4/5-th args are double
l - first arg is i64

by default - 1st arg is i32, 4/5-th args are float.

Currently we are missing functions for 4/5-th args are double. I will add them.



CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D104946/new/

https://reviews.llvm.org/D104946



More information about the cfe-commits mailing list