[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