Author: Ayokunle Amodu Date: 2026-07-31T14:55:38-04:00 New Revision: 3ece94d8f394f9142f0b6022761aa25a89946b20
URL: https://github.com/llvm/llvm-project/commit/3ece94d8f394f9142f0b6022761aa25a89946b20 DIFF: https://github.com/llvm/llvm-project/commit/3ece94d8f394f9142f0b6022761aa25a89946b20.diff LOG: [CIR][AMDGPU] Add support for AMDGCN trig_preop builtins (#197399) Adds codegen for the following AMDGCN trigonometric pre-operation builtins: - __builtin_amdgcn_trig_preop (double) - __builtin_amdgcn_trig_preopf (float) These are lowered to the corresponding `llvm.amdgcn.trig.preop` intrinsic. Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index b9ae5948ba013..b681455c4e4b0 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -241,12 +241,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, mlir::ValueRange{src0, src1, src2}); } case AMDGPU::BI__builtin_amdgcn_trig_preop: - case AMDGPU::BI__builtin_amdgcn_trig_preopf: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_trig_preopf: + return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.trig.preop") + .getValue(); case AMDGPU::BI__builtin_amdgcn_rcp: case AMDGPU::BI__builtin_amdgcn_rcpf: case AMDGPU::BI__builtin_amdgcn_rcph: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 4a129a7e534ac..759f981f3fa12 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -223,3 +223,19 @@ __device__ void test_frexp_mant_f32(float* out, float a) { __device__ void test_frexp_mant_f64(double* out, double a) { *out = __builtin_amdgcn_frexp_mant(a); } + +// CIR-LABEL: @_Z19test_trig_preop_f32Pffi +// CIR: cir.call_llvm_intrinsic "amdgcn.trig.preop" {{.*}} : (!cir.float, !s32i) -> !cir.float +// LLVM: define{{.*}} void @_Z19test_trig_preop_f32Pffi +// LLVM: call{{.*}} float @llvm.amdgcn.trig.preop.f32(float %{{.+}}, i32 %{{.+}}) +__device__ void test_trig_preop_f32(float* out, float a, int b) { + *out = __builtin_amdgcn_trig_preopf(a, b); +} + +// CIR-LABEL: @_Z19test_trig_preop_f64Pddi +// CIR: cir.call_llvm_intrinsic "amdgcn.trig.preop" {{.*}} : (!cir.double, !s32i) -> !cir.double +// LLVM: define{{.*}} void @_Z19test_trig_preop_f64Pddi +// LLVM: call{{.*}} double @llvm.amdgcn.trig.preop.f64(double %{{.+}}, i32 %{{.+}}) +__device__ void test_trig_preop_f64(double* out, double a, int b) { + *out = __builtin_amdgcn_trig_preop(a, b); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
