llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Chaitanya (skc7) <details> <summary>Changes</summary> Port `emitBuiltinWithOneOverloadedType<N>` from classic CodeGen(`CGBuiltin.h`) to `CIRGenFunction` so the common pattern "emit N scalar operands and call an LLVM intrinsic whose return type matches the first operand" can be written as a one-liner across all targets. Used it to simplify a few existing AMDGPU builtin lowerings (`div_fmas`, `ds_swizzle`, `readlane`, `readfirstlane`). --- Full diff: https://github.com/llvm/llvm-project/pull/199518.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+12-32) - (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+19) ``````````diff 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, `````````` </details> https://github.com/llvm/llvm-project/pull/199518 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
