================
@@ -222,10 +211,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_rcpf:
case AMDGPU::BI__builtin_amdgcn_rcph:
case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ mlir::Value src = emitScalarExpr(expr->getArg(0));
+ return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
"amdgcn.rcp",
----------------
skc7 wrote:
Prefer `builder.emitIntrinsicCallOp(...)`. It handles the boilerplate
(intrinsic name StringAttr, perfect-forwarded operands, `.getResult()` so you
directly get an `mlir::Value`.
Use `cir::LLVMIntrinsicCallOp::create(...)` only when you need something the
wrapper doesn't give you.
https://github.com/llvm/llvm-project/pull/197447
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
