From 1a3707813245a3e1d846257e03185b1655c2490c Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Thu, 21 Jan 2021 04:41:21 -0800 Subject: [PATCH] [MLIR][KernelGen] Add chlo.erfc lowerings for f64 PiperOrigin-RevId: 352993223 --- .../mhlo/transforms/chlo_legalize_to_hlo.cc | 56 +++++- tests/chlo_legalize_to_mhlo.mlir | 167 ++++++++++++++++++ 2 files changed, 219 insertions(+), 4 deletions(-) diff --git a/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo.cc b/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo.cc index cf5d85b..2464299 100644 --- a/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo.cc +++ b/lib/Dialect/mhlo/transforms/chlo_legalize_to_hlo.cc @@ -96,6 +96,8 @@ Value MaterializePolynomialApproximation(ConversionPatternRewriter &rewriter, // This implementation is based on Cephes. Value MaterializeErfcApproximationF64ForMagnituteGEOne( ConversionPatternRewriter &rewriter, Location loc, Value x) { + assert(x.getType().cast().getElementType().isF64() && + "expect f64 element type"); const double kMaxlog = 7.09782712893383996843E2; const std::vector kErfcPCoefficients{ 2.46196981473530512524E-10, 5.64189564831068821977E-1, @@ -177,6 +179,8 @@ Value MaterializeErfcApproximationF64ForMagnituteGEOne( // This implementation is based on Cephes. Value MaterializeErfApproximationF64ForMagnituteLEOne( ConversionPatternRewriter &rewriter, Location loc, Value x) { + assert(x.getType().cast().getElementType().isF64() && + "expect f64 element type"); const std::vector kErfTCoefficients{ 9.60497373987051638749E0, 9.00260197203842689217E1, 2.23200534594684319226E3, 7.00332514112805075473E3, @@ -200,7 +204,7 @@ Value MaterializeErfApproximationF64ForMagnituteLEOne( // This implementation is based on Cephes. Value MaterializeErfApproximationF64(ConversionPatternRewriter &rewriter, Location loc, Value x) { - assert(x.getType().cast().getElementType().isF64() && + assert(x.getType().cast().getElementType().isF64() && "expect f64 element type"); // Rely on erf approximation for |x| < 1 @@ -224,10 +228,36 @@ Value MaterializeErfApproximationF64(ConversionPatternRewriter &rewriter, erfc_based_approx); } +Value MaterializeErfcApproximationF64(ConversionPatternRewriter &rewriter, + Location loc, Value x) { + assert(x.getType().cast().getElementType().isF64() && + "expect f64 element type"); + + // Rely on erfc approximation for |x| >= 1 + // erfc(x) = erfc_approx(x) + Value erfc_approx = + MaterializeErfcApproximationF64ForMagnituteGEOne(rewriter, loc, x); + + // Rely on erf approximation for |x| < 1 and materialize erfc as + // erfc(x) = 1 - erf_approx(x) + Value one = chlo::getConstantLike(rewriter, loc, 1.0, x); + Value erf_approx = + MaterializeErfApproximationF64ForMagnituteLEOne(rewriter, loc, x); + Value erf_based_approx = rewriter.create(loc, one, erf_approx); + + // Materialize approximation selection based on argument. + Value abs_x = rewriter.create(loc, x); + const StringAttr kLT = rewriter.getStringAttr( + mhlo::stringifyComparisonDirection(mhlo::ComparisonDirection::LT)); + Value abs_x_lt_one = rewriter.create(loc, abs_x, one, kLT); + return rewriter.create(loc, abs_x_lt_one, erf_based_approx, + erfc_approx); +} + // This is the same approximation as used in Eigen. Value MaterializeErfApproximationF32(ConversionPatternRewriter &rewriter, Location loc, Value operand) { - assert(operand.getType().cast().getElementType().isF32() && + assert(operand.getType().cast().getElementType().isF32() && "expect f32 element type"); const std::vector kAlpha{ -2.72614225801306e-10f, 2.77068142495902e-08f, -2.10102402082508e-06f, @@ -264,7 +294,7 @@ struct ConvertErfOp : public OpConversionPattern { Location loc = op.getLoc(); ErfOp::Adaptor transformed(operands); Value x = transformed.operand(); - Type ty = x.getType().cast().getElementType(); + Type ty = x.getType().cast().getElementType(); // For now, we support only f64, f32, and f16. if (!ty.isF64() && !ty.isF32() && !ty.isF16()) return failure(); @@ -292,6 +322,24 @@ struct ConvertErfOp : public OpConversionPattern { } }; +struct ConvertErfcOp : public OpConversionPattern { + using OpConversionPattern::OpConversionPattern; + LogicalResult matchAndRewrite( + ErfcOp op, ArrayRef operands, + ConversionPatternRewriter &rewriter) const override { + Location loc = op.getLoc(); + ErfcOp::Adaptor transformed(operands); + Value x = transformed.operand(); + Type ty = x.getType().cast().getElementType(); + + // For now, we support only f64. + if (!ty.isF64()) return failure(); + + rewriter.replaceOp(op, MaterializeErfcApproximationF64(rewriter, loc, x)); + return success(); + } +}; + // Converts binary ops that statically are determined to not broadcast directly // to the corresponding mhlo non-broadcasting op. template @@ -443,7 +491,7 @@ void PopulateLegalizeChloToHloPatterns(MLIRContext *context, context, patterns, 5); // Other patterns. - patterns->insert(context); + patterns->insert(context); } } // namespace chlo diff --git a/tests/chlo_legalize_to_mhlo.mlir b/tests/chlo_legalize_to_mhlo.mlir index d0a2640..d46827c 100644 --- a/tests/chlo_legalize_to_mhlo.mlir +++ b/tests/chlo_legalize_to_mhlo.mlir @@ -366,3 +366,170 @@ func @acosh(%arg: tensor) -> tensor { %1 = "chlo.acosh"(%arg) : (tensor) -> tensor return %1 : tensor } + +// CHECK-LABEL: @erfc_f64 +// CHECK-SAME: %[[ARG:.*]]: tensor +func @erfc_f64(%arg : tensor) -> tensor { + // CHECK-NEXT: %[[TMP_0:.*]] = mhlo.multiply %[[ARG]], %[[ARG]] + // CHECK-NEXT: %[[TMP_1:.*]] = "mhlo.negate"(%[[TMP_0]]) + // CHECK-NEXT: %[[TMP_2:.*]] = "mhlo.exponential"(%[[TMP_1]]) + // CHECK-NEXT: %[[TMP_3:.*]] = "mhlo.abs"(%[[ARG]]) + // CHECK-NEXT: %[[TMP_4:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_5:.*]] = mhlo.multiply %[[TMP_4]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_6:.*]] = mhlo.constant dense<2.4619698147353052E-10> + // CHECK-NEXT: %[[TMP_7:.*]] = mhlo.add %[[TMP_5]], %[[TMP_6]] + // CHECK-NEXT: %[[TMP_8:.*]] = mhlo.multiply %[[TMP_7]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_9:.*]] = mhlo.constant dense<0.56418956483106886> + // CHECK-NEXT: %[[TMP_10:.*]] = mhlo.add %[[TMP_8]], %[[TMP_9]] + // CHECK-NEXT: %[[TMP_11:.*]] = mhlo.multiply %[[TMP_10]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_12:.*]] = mhlo.constant dense<7.4632105644226989> + // CHECK-NEXT: %[[TMP_13:.*]] = mhlo.add %[[TMP_11]], %[[TMP_12]] + // CHECK-NEXT: %[[TMP_14:.*]] = mhlo.multiply %[[TMP_13]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_15:.*]] = mhlo.constant dense<48.637197098568137> + // CHECK-NEXT: %[[TMP_16:.*]] = mhlo.add %[[TMP_14]], %[[TMP_15]] + // CHECK-NEXT: %[[TMP_17:.*]] = mhlo.multiply %[[TMP_16]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_18:.*]] = mhlo.constant dense<196.5208329560771> + // CHECK-NEXT: %[[TMP_19:.*]] = mhlo.add %[[TMP_17]], %[[TMP_18]] + // CHECK-NEXT: %[[TMP_20:.*]] = mhlo.multiply %[[TMP_19]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_21:.*]] = mhlo.constant dense<526.44519499547732> + // CHECK-NEXT: %[[TMP_22:.*]] = mhlo.add %[[TMP_20]], %[[TMP_21]] + // CHECK-NEXT: %[[TMP_23:.*]] = mhlo.multiply %[[TMP_22]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_24:.*]] = mhlo.constant dense<934.52852717195765> + // CHECK-NEXT: %[[TMP_25:.*]] = mhlo.add %[[TMP_23]], %[[TMP_24]] + // CHECK-NEXT: %[[TMP_26:.*]] = mhlo.multiply %[[TMP_25]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_27:.*]] = mhlo.constant dense<1027.5518868951572> + // CHECK-NEXT: %[[TMP_28:.*]] = mhlo.add %[[TMP_26]], %[[TMP_27]] + // CHECK-NEXT: %[[TMP_29:.*]] = mhlo.multiply %[[TMP_28]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_30:.*]] = mhlo.constant dense<557.53533536939938> + // CHECK-NEXT: %[[TMP_31:.*]] = mhlo.add %[[TMP_29]], %[[TMP_30]] + // CHECK-NEXT: %[[TMP_32:.*]] = mhlo.multiply %[[TMP_2]], %[[TMP_31]] + // CHECK-NEXT: %[[TMP_33:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_34:.*]] = mhlo.multiply %[[TMP_33]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_35:.*]] = mhlo.constant dense<1.000000e+00> + // CHECK-NEXT: %[[TMP_36:.*]] = mhlo.add %[[TMP_34]], %[[TMP_35]] + // CHECK-NEXT: %[[TMP_37:.*]] = mhlo.multiply %[[TMP_36]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_38:.*]] = mhlo.constant dense<13.228195115474499> + // CHECK-NEXT: %[[TMP_39:.*]] = mhlo.add %[[TMP_37]], %[[TMP_38]] + // CHECK-NEXT: %[[TMP_40:.*]] = mhlo.multiply %[[TMP_39]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_41:.*]] = mhlo.constant dense<86.707214088598973> + // CHECK-NEXT: %[[TMP_42:.*]] = mhlo.add %[[TMP_40]], %[[TMP_41]] + // CHECK-NEXT: %[[TMP_43:.*]] = mhlo.multiply %[[TMP_42]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_44:.*]] = mhlo.constant dense<354.93777888781989> + // CHECK-NEXT: %[[TMP_45:.*]] = mhlo.add %[[TMP_43]], %[[TMP_44]] + // CHECK-NEXT: %[[TMP_46:.*]] = mhlo.multiply %[[TMP_45]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_47:.*]] = mhlo.constant dense<975.70850174320549> + // CHECK-NEXT: %[[TMP_48:.*]] = mhlo.add %[[TMP_46]], %[[TMP_47]] + // CHECK-NEXT: %[[TMP_49:.*]] = mhlo.multiply %[[TMP_48]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_50:.*]] = mhlo.constant dense<1823.9091668790973> + // CHECK-NEXT: %[[TMP_51:.*]] = mhlo.add %[[TMP_49]], %[[TMP_50]] + // CHECK-NEXT: %[[TMP_52:.*]] = mhlo.multiply %[[TMP_51]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_53:.*]] = mhlo.constant dense<2246.3376081871097> + // CHECK-NEXT: %[[TMP_54:.*]] = mhlo.add %[[TMP_52]], %[[TMP_53]] + // CHECK-NEXT: %[[TMP_55:.*]] = mhlo.multiply %[[TMP_54]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_56:.*]] = mhlo.constant dense<1656.6630919416134> + // CHECK-NEXT: %[[TMP_57:.*]] = mhlo.add %[[TMP_55]], %[[TMP_56]] + // CHECK-NEXT: %[[TMP_58:.*]] = mhlo.multiply %[[TMP_57]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_59:.*]] = mhlo.constant dense<557.53534081772773> + // CHECK-NEXT: %[[TMP_60:.*]] = mhlo.add %[[TMP_58]], %[[TMP_59]] + // CHECK-NEXT: %[[TMP_61:.*]] = mhlo.divide %[[TMP_32]], %[[TMP_60]] + // CHECK-NEXT: %[[TMP_62:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_63:.*]] = mhlo.multiply %[[TMP_62]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_64:.*]] = mhlo.constant dense<0.56418958354775506> + // CHECK-NEXT: %[[TMP_65:.*]] = mhlo.add %[[TMP_63]], %[[TMP_64]] + // CHECK-NEXT: %[[TMP_66:.*]] = mhlo.multiply %[[TMP_65]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_67:.*]] = mhlo.constant dense<1.275366707599781> + // CHECK-NEXT: %[[TMP_68:.*]] = mhlo.add %[[TMP_66]], %[[TMP_67]] + // CHECK-NEXT: %[[TMP_69:.*]] = mhlo.multiply %[[TMP_68]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_70:.*]] = mhlo.constant dense<5.0190504225118051> + // CHECK-NEXT: %[[TMP_71:.*]] = mhlo.add %[[TMP_69]], %[[TMP_70]] + // CHECK-NEXT: %[[TMP_72:.*]] = mhlo.multiply %[[TMP_71]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_73:.*]] = mhlo.constant dense<6.160210979930536> + // CHECK-NEXT: %[[TMP_74:.*]] = mhlo.add %[[TMP_72]], %[[TMP_73]] + // CHECK-NEXT: %[[TMP_75:.*]] = mhlo.multiply %[[TMP_74]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_76:.*]] = mhlo.constant dense<7.4097426995044895> + // CHECK-NEXT: %[[TMP_77:.*]] = mhlo.add %[[TMP_75]], %[[TMP_76]] + // CHECK-NEXT: %[[TMP_78:.*]] = mhlo.multiply %[[TMP_77]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_79:.*]] = mhlo.constant dense<2.9788666537210022> + // CHECK-NEXT: %[[TMP_80:.*]] = mhlo.add %[[TMP_78]], %[[TMP_79]] + // CHECK-NEXT: %[[TMP_81:.*]] = mhlo.multiply %[[TMP_2]], %[[TMP_80]] + // CHECK-NEXT: %[[TMP_82:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_83:.*]] = mhlo.multiply %[[TMP_82]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_84:.*]] = mhlo.constant dense<1.000000e+00> + // CHECK-NEXT: %[[TMP_85:.*]] = mhlo.add %[[TMP_83]], %[[TMP_84]] + // CHECK-NEXT: %[[TMP_86:.*]] = mhlo.multiply %[[TMP_85]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_87:.*]] = mhlo.constant dense<2.2605286322011726> + // CHECK-NEXT: %[[TMP_88:.*]] = mhlo.add %[[TMP_86]], %[[TMP_87]] + // CHECK-NEXT: %[[TMP_89:.*]] = mhlo.multiply %[[TMP_88]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_90:.*]] = mhlo.constant dense<9.3960352493800147> + // CHECK-NEXT: %[[TMP_91:.*]] = mhlo.add %[[TMP_89]], %[[TMP_90]] + // CHECK-NEXT: %[[TMP_92:.*]] = mhlo.multiply %[[TMP_91]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_93:.*]] = mhlo.constant dense<12.048953980809666> + // CHECK-NEXT: %[[TMP_94:.*]] = mhlo.add %[[TMP_92]], %[[TMP_93]] + // CHECK-NEXT: %[[TMP_95:.*]] = mhlo.multiply %[[TMP_94]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_96:.*]] = mhlo.constant dense<17.081445074756591> + // CHECK-NEXT: %[[TMP_97:.*]] = mhlo.add %[[TMP_95]], %[[TMP_96]] + // CHECK-NEXT: %[[TMP_98:.*]] = mhlo.multiply %[[TMP_97]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_99:.*]] = mhlo.constant dense<9.6089680906328585> + // CHECK-NEXT: %[[TMP_100:.*]] = mhlo.add %[[TMP_98]], %[[TMP_99]] + // CHECK-NEXT: %[[TMP_101:.*]] = mhlo.multiply %[[TMP_100]], %[[TMP_3]] + // CHECK-NEXT: %[[TMP_102:.*]] = mhlo.constant dense<3.3690764510008151> + // CHECK-NEXT: %[[TMP_103:.*]] = mhlo.add %[[TMP_101]], %[[TMP_102]] + // CHECK-NEXT: %[[TMP_104:.*]] = mhlo.divide %[[TMP_81]], %[[TMP_103]] + // CHECK-NEXT: %[[TMP_105:.*]] = mhlo.constant dense<8.000000e+00> + // CHECK-NEXT: %[[TMP_106:.*]] = "mhlo.compare"(%[[TMP_3]], %[[TMP_105]]) {comparison_direction = "LT"} + // CHECK-NEXT: %[[TMP_107:.*]] = "mhlo.select"(%[[TMP_106]], %[[TMP_61]], %[[TMP_104]]) + // CHECK-NEXT: %[[TMP_108:.*]] = mhlo.constant dense<-709.78271289338397> + // CHECK-NEXT: %[[TMP_109:.*]] = "mhlo.compare"(%[[TMP_1]], %[[TMP_108]]) {comparison_direction = "LT"} + // CHECK-NEXT: %[[TMP_110:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_111:.*]] = "mhlo.select"(%[[TMP_109]], %[[TMP_110]], %[[TMP_107]]) + // CHECK-NEXT: %[[TMP_113:.*]] = "mhlo.compare"(%[[ARG]], %[[TMP_110]]) {comparison_direction = "LT"} + // CHECK-NEXT: %[[TMP_114:.*]] = mhlo.constant dense<2.000000e+00> + // CHECK-NEXT: %[[TMP_115:.*]] = mhlo.subtract %[[TMP_114]], %[[TMP_111]] + // CHECK-NEXT: %[[TMP_116:.*]] = "mhlo.select"(%[[TMP_113]], %[[TMP_115]], %[[TMP_111]]) + // CHECK-NEXT: %[[TMP_117:.*]] = mhlo.constant dense<1.000000e+00> + // CHECK-NEXT: %[[TMP_118:.*]] = mhlo.multiply %[[ARG]], %[[ARG]] + // CHECK-NEXT: %[[TMP_119:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_120:.*]] = mhlo.multiply %[[TMP_119]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_121:.*]] = mhlo.constant dense<9.6049737398705161> + // CHECK-NEXT: %[[TMP_122:.*]] = mhlo.add %[[TMP_120]], %[[TMP_121]] + // CHECK-NEXT: %[[TMP_123:.*]] = mhlo.multiply %[[TMP_122]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_124:.*]] = mhlo.constant dense<90.026019720384269> + // CHECK-NEXT: %[[TMP_125:.*]] = mhlo.add %[[TMP_123]], %[[TMP_124]] + // CHECK-NEXT: %[[TMP_126:.*]] = mhlo.multiply %[[TMP_125]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_127:.*]] = mhlo.constant dense<2232.0053459468431> + // CHECK-NEXT: %[[TMP_128:.*]] = mhlo.add %[[TMP_126]], %[[TMP_127]] + // CHECK-NEXT: %[[TMP_129:.*]] = mhlo.multiply %[[TMP_128]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_130:.*]] = mhlo.constant dense<7003.3251411280507> + // CHECK-NEXT: %[[TMP_131:.*]] = mhlo.add %[[TMP_129]], %[[TMP_130]] + // CHECK-NEXT: %[[TMP_132:.*]] = mhlo.multiply %[[TMP_131]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_133:.*]] = mhlo.constant dense<55592.301301039493> + // CHECK-NEXT: %[[TMP_134:.*]] = mhlo.add %[[TMP_132]], %[[TMP_133]] + // CHECK-NEXT: %[[TMP_135:.*]] = mhlo.multiply %[[ARG]], %[[TMP_134]] + // CHECK-NEXT: %[[TMP_136:.*]] = mhlo.constant dense<0.000000e+00> + // CHECK-NEXT: %[[TMP_137:.*]] = mhlo.multiply %[[TMP_136]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_138:.*]] = mhlo.constant dense<1.000000e+00> + // CHECK-NEXT: %[[TMP_139:.*]] = mhlo.add %[[TMP_137]], %[[TMP_138]] + // CHECK-NEXT: %[[TMP_140:.*]] = mhlo.multiply %[[TMP_139]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_141:.*]] = mhlo.constant dense<33.561714164750313> + // CHECK-NEXT: %[[TMP_142:.*]] = mhlo.add %[[TMP_140]], %[[TMP_141]] + // CHECK-NEXT: %[[TMP_143:.*]] = mhlo.multiply %[[TMP_142]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_144:.*]] = mhlo.constant dense<521.35794978015269> + // CHECK-NEXT: %[[TMP_145:.*]] = mhlo.add %[[TMP_143]], %[[TMP_144]] + // CHECK-NEXT: %[[TMP_146:.*]] = mhlo.multiply %[[TMP_145]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_147:.*]] = mhlo.constant dense<4594.3238297098014> + // CHECK-NEXT: %[[TMP_148:.*]] = mhlo.add %[[TMP_146]], %[[TMP_147]] + // CHECK-NEXT: %[[TMP_149:.*]] = mhlo.multiply %[[TMP_148]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_150:.*]] = mhlo.constant dense<22629.000061389095> + // CHECK-NEXT: %[[TMP_151:.*]] = mhlo.add %[[TMP_149]], %[[TMP_150]] + // CHECK-NEXT: %[[TMP_152:.*]] = mhlo.multiply %[[TMP_151]], %[[TMP_118]] + // CHECK-NEXT: %[[TMP_153:.*]] = mhlo.constant dense<49267.394260863592> + // CHECK-NEXT: %[[TMP_154:.*]] = mhlo.add %[[TMP_152]], %[[TMP_153]] + // CHECK-NEXT: %[[TMP_155:.*]] = mhlo.divide %[[TMP_135]], %[[TMP_154]] + // CHECK-NEXT: %[[TMP_156:.*]] = mhlo.subtract %[[TMP_117]], %[[TMP_155]] + // CHECK-NEXT: %[[TMP_157:.*]] = "mhlo.abs"(%[[ARG]]) + // CHECK-NEXT: %[[TMP_159:.*]] = "mhlo.compare"(%[[TMP_157]], %[[TMP_117]]) {comparison_direction = "LT"} + // CHECK-NEXT: %[[RESULT:.*]] = "mhlo.select"(%[[TMP_159]], %[[TMP_156]], %[[TMP_116]]) + // CHECK-NEXT: return %[[RESULT]] + %1 = "chlo.erfc"(%arg) : (tensor) -> tensor + return %1 : tensor +}