================
@@ -317,12 +317,24 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
return mlir::Value{};
}
case AMDGPU::BI__builtin_amdgcn_ldexp:
- case AMDGPU::BI__builtin_amdgcn_ldexpf:
+ case AMDGPU::BI__builtin_amdgcn_ldexpf: {
+ mlir::Value src0 = emitScalarExpr(expr->getArg(0));
+ mlir::Value src1 = emitScalarExpr(expr->getArg(1));
+ return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp",
+ src0.getType(),
+ mlir::ValueRange{src0, src1});
+ }
case AMDGPU::BI__builtin_amdgcn_ldexph: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ // The raw instruction has a different behavior for out of bounds exponent
+ // values (implicit truncation instead of saturate to short_min/short_max).
+ mlir::Value src0 = emitScalarExpr(expr->getArg(0));
+ mlir::Value src1 = emitScalarExpr(expr->getArg(1));
+ mlir::Value truncSrc1 = cir::CastOp::create(
+ builder, getLoc(expr->getExprLoc()), builder.getSInt16Ty(),
+ cir::CastKind::integral, src1);
+ return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp",
+ src0.getType(),
+ mlir::ValueRange{src0, truncSrc1});
----------------
ayokunle321 wrote:
Done!
https://github.com/llvm/llvm-project/pull/198087
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
