| 276 | } |
| 277 | |
| 278 | XlaOp Erfc(XlaOp x) { |
| 279 | auto& b = *x.builder(); |
| 280 | return b.ReportErrorOrReturn([&]() -> StatusOr<XlaOp> { |
| 281 | TF_RETURN_IF_ERROR(EnsureOperandIsRealFp("Erfc", x)); |
| 282 | TF_ASSIGN_OR_RETURN(auto shape, b.GetShape(x)); |
| 283 | // erfc(x) = |
| 284 | // erfc_impl(x) if x > 1 |
| 285 | // 1 - erf_impl(x) otherwise |
| 286 | if (shape.element_type() == F64) { |
| 287 | return Select(Gt(Abs(x), ScalarLike(x, 1)), ErfcImpl64(x), |
| 288 | ScalarLike(x, 1) - ErfImpl64(x)); |
| 289 | } |
| 290 | // Erf(c)Impl don't have enough precision when run with bf16 intermediates |
| 291 | // (not surprising!), so upcast to f32 in this case. |
| 292 | return DoWithUpcastToF32(x, {BF16, F16}, [](XlaOp x) { |
| 293 | return Select(Gt(Abs(x), ScalarLike(x, 1)), ErfcImpl32(x), |
| 294 | ScalarLike(x, 1) - ErfImpl32(x)); |
| 295 | }); |
| 296 | }); |
| 297 | } |
| 298 | |
| 299 | XlaOp Erf(XlaOp x) { |
| 300 | auto& b = *x.builder(); |