| 297 | } |
| 298 | |
| 299 | XlaOp Erf(XlaOp x) { |
| 300 | auto& b = *x.builder(); |
| 301 | return b.ReportErrorOrReturn([&]() -> StatusOr<XlaOp> { |
| 302 | TF_RETURN_IF_ERROR(EnsureOperandIsRealFp("Erf", x)); |
| 303 | TF_ASSIGN_OR_RETURN(auto shape, b.GetShape(x)); |
| 304 | // erf(x) = |
| 305 | // erf_impl(x) if x < 1 |
| 306 | // 1 - erfc_impl(x) otherwise |
| 307 | if (shape.element_type() == F64) { |
| 308 | return Select(Lt(Abs(x), ScalarLike(x, 1)), ErfImpl64(x), |
| 309 | ScalarLike(x, 1) - ErfcImpl64(x)); |
| 310 | } |
| 311 | // Erf(c)Impl don't have enough precision when run with bf16 intermediates |
| 312 | // (not surprising!), so upcast to f32 in this case. |
| 313 | return DoWithUpcastToF32(x, {BF16, F16}, [](XlaOp x) { |
| 314 | return Select(Lt(Abs(x), ScalarLike(x, 1)), ErfImpl32(x), |
| 315 | ScalarLike(x, 1) - ErfcImpl32(x)); |
| 316 | }); |
| 317 | }); |
| 318 | } |
| 319 | |
| 320 | namespace { |
| 321 | |