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

Yaxun Liu via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Tue Jun 29 11:04:51 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")
----------------
yaxunl wrote:
> 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.
> 
It seems we do not support 4/50th args to be double. So the current definitions are all that we support.


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

https://reviews.llvm.org/D104946



More information about the cfe-commits mailing list