https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/197399
Adds codegen support for the following AMDGPU trig preop builtins: `__builtin_amdgcn_trig_preop` (double) `__builtin_amdgcn_trig_preopf` (float) These are lowered to the corresponding `llvm.amdgcn.trig.preop` intrinsic. >From dfbcc968635e4e8e45f484609ee5464349a36dc6 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 13 May 2026 05:19:23 -0400 Subject: [PATCH 1/2] add amdgcn trig_preop builtin --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 15 +++++++++++---- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 16 ++++++++++++++++ 2 files changed, 27 insertions(+), 4 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 04ab1c29b0d63..476a88c67cd18 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -91,6 +91,16 @@ static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, return res; } +// Emit an intrinsic that has 1 float or double operand, and 1 integer. +static mlir::Value emitFPIntBuiltin(CIRGenFunction &cgf, const CallExpr *e, + llvm::StringRef intrinsicName) { + mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0)); + mlir::Value src1 = cgf.emitScalarExpr(e->getArg(1)); + return cgf.getBuilder().emitIntrinsicCallOp(cgf.getLoc(e->getExprLoc()), + intrinsicName, src0.getType(), + mlir::ValueRange{src0, src1}); +} + std::optional<mlir::Value> CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, const CallExpr *expr) { @@ -202,10 +212,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, } 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{}; + return emitFPIntBuiltin(*this, expr, "amdgcn.trig.preop"); } case AMDGPU::BI__builtin_amdgcn_rcp: case AMDGPU::BI__builtin_amdgcn_rcpf: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 4a61fde7aa90c..58504c9bfdf40 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -71,3 +71,19 @@ __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int __device__ void test_ds_swizzle(int* out, int a) { *out = __builtin_amdgcn_ds_swizzle(a, 32); } + +// 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); +} >From 4d800b18ef05485c365a8f875673ffdd5b0b57ea Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 13 May 2026 05:44:02 -0400 Subject: [PATCH 2/2] switch out intrinsic call function --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 476a88c67cd18..dfedab81b5837 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -96,9 +96,12 @@ static mlir::Value emitFPIntBuiltin(CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName) { mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0)); mlir::Value src1 = cgf.emitScalarExpr(e->getArg(1)); - return cgf.getBuilder().emitIntrinsicCallOp(cgf.getLoc(e->getExprLoc()), - intrinsicName, src0.getType(), - mlir::ValueRange{src0, src1}); + mlir::Value result = + LLVMIntrinsicCallOp::create(cgf.getBuilder(), cgf.getLoc(e->getExprLoc()), + cgf.getBuilder().getStringAttr(intrinsicName), + src0.getType(), {src0, src1}) + .getResult(); + return result; } std::optional<mlir::Value> _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
