[all-commits] [llvm/llvm-project] d6d0f7: [AMDGPU][Clang] Add builtins for gfx12 ray tracing...

Shilei Tian via All-commits all-commits at lists.llvm.org
Thu Apr 10 10:52:14 PDT 2025


  Branch: refs/heads/users/shiltian/missing-ray-tracing-builtins-gfx12
  Home:   https://github.com/llvm/llvm-project
  Commit: d6d0f7ff400d1540db414731c0a368e309f24818
      https://github.com/llvm/llvm-project/commit/d6d0f7ff400d1540db414731c0a368e309f24818
  Author: Shilei Tian <i at tianshilei.me>
  Date:   2025-04-10 (Thu, 10 Apr 2025)

  Changed paths:
    M clang/include/clang/Basic/BuiltinsAMDGPU.def
    M clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
    M clang/test/CodeGenOpenCL/builtins-amdgcn-raytracing.cl

  Log Message:
  -----------
  [AMDGPU][Clang] Add builtins for gfx12 ray tracing intrinsics

__builtin_amdgcn_image_bvh8_intersect_ray
__builtin_amdgcn_image_bvh_dual_intersect_ray

For the above two builtins, the second and third return values of the intrinsics
are returned through pointer-type function arguments.

__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn
__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn
__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn

For the last builtin, the intrinsic returns `{i64, i32}`, the builtin returns
`<2 x i64>`. The second return value of the intrinsic is zext'ed.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications


More information about the All-commits mailing list