https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/199518
>From f34781310e3de990237ff2fe86fa4706e5d66017 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Mon, 25 May 2026 17:11:23 +0530 Subject: [PATCH 1/2] [CIR] Add emitBuiltinWithOneOverloadedType helper --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 44 +++++-------------- clang/lib/CIR/CodeGen/CIRGenFunction.h | 19 ++++++++ 2 files changed, 31 insertions(+), 32 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 7e6e0f1a06046..432e01a696a4e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -149,25 +149,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return result; } case AMDGPU::BI__builtin_amdgcn_div_fmas: - case AMDGPU::BI__builtin_amdgcn_div_fmasf: { - mlir::Value src0 = emitScalarExpr(expr->getArg(0)); - mlir::Value src1 = emitScalarExpr(expr->getArg(1)); - mlir::Value src2 = emitScalarExpr(expr->getArg(2)); - mlir::Value src3 = emitScalarExpr(expr->getArg(3)); - mlir::Value result = cir::LLVMIntrinsicCallOp::create( - builder, getLoc(expr->getExprLoc()), - builder.getStringAttr("amdgcn.div.fmas"), - src0.getType(), {src0, src1, src2, src3}) - .getResult(); - return result; - } - case AMDGPU::BI__builtin_amdgcn_ds_swizzle: { - mlir::Value src0 = emitScalarExpr(expr->getArg(0)); - mlir::Value src1 = emitScalarExpr(expr->getArg(1)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.ds.swizzle", src0.getType(), - mlir::ValueRange{src0, src1}); - } + case AMDGPU::BI__builtin_amdgcn_div_fmasf: + return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.div.fmas") + .getValue(); + case AMDGPU::BI__builtin_amdgcn_ds_swizzle: + return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ds.swizzle") + .getValue(); case AMDGPU::BI__builtin_amdgcn_mov_dpp8: case AMDGPU::BI__builtin_amdgcn_mov_dpp: case AMDGPU::BI__builtin_amdgcn_update_dpp: { @@ -184,19 +171,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } - case AMDGPU::BI__builtin_amdgcn_readlane: { - mlir::Value src0 = emitScalarExpr(expr->getArg(0)); - mlir::Value src1 = emitScalarExpr(expr->getArg(1)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.readlane", src0.getType(), - mlir::ValueRange{src0, src1}); - } - case AMDGPU::BI__builtin_amdgcn_readfirstlane: { - mlir::Value src0 = emitScalarExpr(expr->getArg(0)); - return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), - "amdgcn.readfirstlane", src0.getType(), - mlir::ValueRange{src0}); - } + case AMDGPU::BI__builtin_amdgcn_readlane: + return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") + .getValue(); + case AMDGPU::BI__builtin_amdgcn_readfirstlane: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane") + .getValue(); case AMDGPU::BI__builtin_amdgcn_wave_shuffle: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 9f2facd12f417..3b04398c2c1cd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -1625,6 +1625,25 @@ class CIRGenFunction : public CIRGenTypeCache { void instantiateIndirectGotoBlock(); + /// Emit a simple LLVM intrinsic that takes N scalar arguments and whose + /// return type matches the type of the first argument. The intrinsic name is + /// used verbatim; any overload mangling (e.g. `.f32`, `.p1`) must be baked + /// into \p Name by the caller. + template <uint32_t N> + [[maybe_unused]] RValue + emitBuiltinWithOneOverloadedType(const CallExpr *E, llvm::StringRef Name) { + static_assert(N, "expect non-empty argument"); + mlir::Type cirTy = convertType(E->getArg(0)->getType()); + SmallVector<mlir::Value, N> args; + for (uint32_t i = 0; i < N; ++i) { + args.push_back(emitScalarExpr(E->getArg(i))); + } + const auto call = cir::LLVMIntrinsicCallOp::create( + builder, getLoc(E->getExprLoc()), builder.getStringAttr(Name), cirTy, + args); + return RValue::get(call->getResult(0)); + } + RValue emitCall(const CIRGenFunctionInfo &funcInfo, const CIRGenCallee &callee, ReturnValueSlot returnValue, const CallArgList &args, cir::CIRCallOpInterface *callOp, >From f7d5de2d23281a996de2e85d576670d0a1c7f1b8 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Wed, 27 May 2026 10:49:31 +0530 Subject: [PATCH 2/2] code format update --- clang/lib/CIR/CodeGen/CIRGenFunction.h | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index 3b04398c2c1cd..667d2292799fb 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -1628,19 +1628,19 @@ class CIRGenFunction : public CIRGenTypeCache { /// Emit a simple LLVM intrinsic that takes N scalar arguments and whose /// return type matches the type of the first argument. The intrinsic name is /// used verbatim; any overload mangling (e.g. `.f32`, `.p1`) must be baked - /// into \p Name by the caller. - template <uint32_t N> + /// into \p intrinName by the caller. + template <unsigned N> [[maybe_unused]] RValue - emitBuiltinWithOneOverloadedType(const CallExpr *E, llvm::StringRef Name) { + emitBuiltinWithOneOverloadedType(const CallExpr *e, + llvm::StringRef intrinName) { static_assert(N, "expect non-empty argument"); - mlir::Type cirTy = convertType(E->getArg(0)->getType()); + mlir::Type cirTy = convertType(e->getArg(0)->getType()); SmallVector<mlir::Value, N> args; - for (uint32_t i = 0; i < N; ++i) { - args.push_back(emitScalarExpr(E->getArg(i))); - } + for (unsigned i = 0; i < N; ++i) + args.push_back(emitScalarExpr(e->getArg(i))); const auto call = cir::LLVMIntrinsicCallOp::create( - builder, getLoc(E->getExprLoc()), builder.getStringAttr(Name), cirTy, - args); + builder, getLoc(e->getExprLoc()), builder.getStringAttr(intrinName), + cirTy, args); return RValue::get(call->getResult(0)); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
