[clang] [CIR][AMDGPU] Add support for AMDGCN trig_preop builtins (PR #197399)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 17 13:58:20 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/197399
>From dfbcc968635e4e8e45f484609ee5464349a36dc6 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 05:19:23 -0400
Subject: [PATCH 1/4] 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 <ayokunle321 at gmail.com>
Date: Wed, 13 May 2026 05:44:02 -0400
Subject: [PATCH 2/4] 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>
>From 42d9603ec564a794e34943ef2921e6e69ab3db4e Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 14 May 2026 16:42:33 -0400
Subject: [PATCH 3/4] switchout create instruction
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 9 +++------
1 file changed, 3 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index dfedab81b5837..57a37c67ce9e6 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -96,12 +96,9 @@ 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));
- mlir::Value result =
- LLVMIntrinsicCallOp::create(cgf.getBuilder(), cgf.getLoc(e->getExprLoc()),
- cgf.getBuilder().getStringAttr(intrinsicName),
- src0.getType(), {src0, src1})
- .getResult();
- return result;
+ return cgf.getBuilder().emitIntrinsicCallOp(
+ cgf.getLoc(e->getExprLoc()), "amdgcn.ds.swizzle", src0.getType(),
+ mlir::ValueRange{src0, src1});
}
std::optional<mlir::Value>
>From d3b5d77f39151d713f1c260d471093316c69808a Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 17 Jul 2026 12:02:17 -0400
Subject: [PATCH 4/4] fix incorrect create call
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 57a37c67ce9e6..6838ce1b1d0ec 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -92,13 +92,13 @@ static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e,
}
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
-static mlir::Value emitFPIntBuiltin(CIRGenFunction &cgf, const CallExpr *e,
+static mlir::Value emitFPIntBuiltin(CIRGenFunction &cgf, const CallExpr *expr,
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()), "amdgcn.ds.swizzle", src0.getType(),
- mlir::ValueRange{src0, src1});
+ mlir::Value src0 = cgf.emitScalarExpr(expr->getArg(0));
+ mlir::Value src1 = cgf.emitScalarExpr(expr->getArg(1));
+ return cgf.getBuilder().emitIntrinsicCallOp(cgf.getLoc(e->getExprLoc()),
+ intrinsicName, src0.getType(),
+ mlir::ValueRange{src0, src1});
}
std::optional<mlir::Value>
More information about the cfe-commits
mailing list