diff --git a/mhlo/transforms/chlo_legalize_to_hlo/chlo_legalize_to_hlo.cc b/mhlo/transforms/chlo_legalize_to_hlo/chlo_legalize_to_hlo.cc index 3bd2c4778..0a9878876 100644 --- a/mhlo/transforms/chlo_legalize_to_hlo/chlo_legalize_to_hlo.cc +++ b/mhlo/transforms/chlo_legalize_to_hlo/chlo_legalize_to_hlo.cc @@ -19,6 +19,7 @@ limitations under the License. // https://docs.microsoft.com/en-us/cpp/c-runtime-library/math-constants #define _USE_MATH_DEFINES #include +#include #include #include #include @@ -633,6 +634,184 @@ struct ConvertErfcOp : public OpConversionPattern { } }; +Value erfInv32(ConversionPatternRewriter &b, Location loc, ValueRange args) { + constexpr int kDegree = 9; + constexpr std::array wLessThan5Constants = { + 2.81022636e-08f, 3.43273939e-07f, -3.5233877e-06f, + -4.39150654e-06f, 0.00021858087f, -0.00125372503f, + -0.00417768164f, 0.246640727f, 1.50140941f}; + constexpr std::array wGreaterThan5Constants = { + -0.000200214257f, 0.000100950558f, 0.00134934322f, + -0.00367342844f, 0.00573950773f, -0.0076224613f, + 0.00943887047f, 1.00167406f, 2.83297682f}; + + Value x = args[0]; + // Compute logarithm of (1+arg) using log1p(arg) which is more precise than + // log(1+arg) when arg is close to zero. For more details, see + // https://en.cppreference.com/w/cpp/numeric/math/log1p + Value minusXSquared = + b.create(loc, x, b.create(loc, x)); + Value w = + b.create(loc, b.create(loc, minusXSquared)); + + Value lt = b.create(loc, w, getConstantLike(b, loc, 5.0, x), + mhlo::ComparisonDirection::LT); + auto coefficient = [&](int i) { + return b.create( + loc, lt, getConstantLike(b, loc, wLessThan5Constants[i], x), + getConstantLike(b, loc, wGreaterThan5Constants[i], x)); + }; + w = b.create( + loc, lt, + b.create(loc, w, getConstantLike(b, loc, 2.5, x)), + b.create(loc, b.create(loc, w), + getConstantLike(b, loc, 3.0, x))); + Value p = coefficient(0); + for (int i = 1; i < kDegree; ++i) { + p = b.create(loc, coefficient(i), + b.create(loc, p, w)); + } + + // Result modulo edge cases. + Value result = b.create(loc, p, x); + + // Handle edge cases, namely erfinv(+/-1) = +/-inf. (The above computation is + // indeterminate, and can give nan or -/+inf.) + return b.create( + loc, + b.create(loc, b.create(loc, x), + getConstantLike(b, loc, 1, x), + mhlo::ComparisonDirection::EQ), + b.create(loc, x, getConstantLikeMaxFiniteValue(b, loc, x)), + result); +} + +Value erfInv64(ConversionPatternRewriter &b, Location loc, ValueRange args) { + constexpr std::array wLessThan625Constants = { + -3.6444120640178196996e-21, -1.685059138182016589e-19, + 1.2858480715256400167e-18, 1.115787767802518096e-17, + -1.333171662854620906e-16, 2.0972767875968561637e-17, + 6.6376381343583238325e-15, -4.0545662729752068639e-14, + -8.1519341976054721522e-14, 2.6335093153082322977e-12, + -1.2975133253453532498e-11, -5.4154120542946279317e-11, + 1.051212273321532285e-09, -4.1126339803469836976e-09, + -2.9070369957882005086e-08, 4.2347877827932403518e-07, + -1.3654692000834678645e-06, -1.3882523362786468719e-05, + 0.0001867342080340571352, -0.00074070253416626697512, + -0.0060336708714301490533, 0.24015818242558961693, + 1.6536545626831027356}; + constexpr std::array wLessThan16Constants = { + 2.2137376921775787049e-09, 9.0756561938885390979e-08, + -2.7517406297064545428e-07, 1.8239629214389227755e-08, + 1.5027403968909827627e-06, -4.013867526981545969e-06, + 2.9234449089955446044e-06, 1.2475304481671778723e-05, + -4.7318229009055733981e-05, 6.8284851459573175448e-05, + 2.4031110387097893999e-05, -0.0003550375203628474796, + 0.00095328937973738049703, -0.0016882755560235047313, + 0.0024914420961078508066, -0.0037512085075692412107, + 0.005370914553590063617, 1.0052589676941592334, + 3.0838856104922207635, + }; + constexpr std::array wGreaterThan16Constants = { + -2.7109920616438573243e-11, -2.5556418169965252055e-10, + 1.5076572693500548083e-09, -3.7894654401267369937e-09, + 7.6157012080783393804e-09, -1.4960026627149240478e-08, + 2.9147953450901080826e-08, -6.7711997758452339498e-08, + 2.2900482228026654717e-07, -9.9298272942317002539e-07, + 4.5260625972231537039e-06, -1.9681778105531670567e-05, + 7.5995277030017761139e-05, -0.00021503011930044477347, + -0.00013871931833623122026, 1.0103004648645343977, + 4.8499064014085844221, + }; + + Value x = args[0]; + // Compute logarithm of (1+arg) using log1p(arg) which is more precise than + // log(1+arg) when arg is close to zero. For more details, see + // https://en.cppreference.com/w/cpp/numeric/math/log1p + Value minusXSquared = + b.create(loc, x, b.create(loc, x)); + Value w = + b.create(loc, b.create(loc, minusXSquared)); + + Value lt625 = b.create( + loc, w, getConstantLike(b, loc, 6.25, x), mhlo::ComparisonDirection::LT); + Value lt16 = b.create(loc, w, getConstantLike(b, loc, 16, x), + mhlo::ComparisonDirection::LT); + + auto coefficient = [&](int i) { + Value c = getConstantLike(b, loc, wLessThan625Constants[i], x); + if (i < 19) { + c = b.create( + loc, lt625, c, getConstantLike(b, loc, wLessThan16Constants[i], x)); + } + if (i < 17) { + c = b.create( + loc, lt16, c, getConstantLike(b, loc, wGreaterThan16Constants[i], x)); + } + return c; + }; + + Value sqrtW = b.create(loc, w); + Value wMinus3125 = + b.create(loc, w, getConstantLike(b, loc, 3.125, x)); + Value select2 = + b.create(loc, lt16, getConstantLike(b, loc, 3.25, w), + getConstantLike(b, loc, 5.0, w)); + Value select2Result = b.create(loc, sqrtW, select2); + w = b.create(loc, lt625, wMinus3125, select2Result); + + Value p = coefficient(0); + for (int i = 1; i < 17; ++i) { + p = b.create(loc, coefficient(i), + b.create(loc, p, w)); + } + for (int i = 17; i < 19; ++i) { + p = b.create( + loc, lt16, + b.create(loc, coefficient(i), + b.create(loc, p, w)), + p); + } + for (int i = 19; i < 23; ++i) { + p = b.create( + loc, lt625, + b.create(loc, coefficient(i), + b.create(loc, p, w)), + p); + } + + // Result modulo edge cases. + Value result = b.create(loc, p, x); + + // Handle edge cases, namely erfinv(+/-1) = +/-inf. (The above computation is + // indeterminate, and can give nan or -/+inf.) + return b.create( + loc, + b.create(loc, b.create(loc, x), + getConstantLike(b, loc, 1, x), + mhlo::ComparisonDirection::EQ), + b.create(loc, x, getConstantLikeMaxFiniteValue(b, loc, x)), + result); +} + +struct ConvertErfInvOp : public OpConversionPattern { + using OpConversionPattern::OpConversionPattern; + LogicalResult matchAndRewrite( + ErfInvOp op, OpAdaptor adaptor, + ConversionPatternRewriter &rewriter) const override { + Location loc = op.getLoc(); + if (op.getResult().getType().getElementType().isF64()) { + rewriter.replaceOp(op, erfInv64(rewriter, loc, adaptor.getOperands())); + return success(); + } + FloatType minPrecisionTy = rewriter.getF32Type(); + rewriter.replaceOp( + op, materializeWithUpcast(rewriter, loc, adaptor.getOperands(), + minPrecisionTy, &erfInv32)); + return success(); + } +}; + // Coefficients for the Lanczos approximation of the gamma function. The // coefficients are uniquely determined by the choice of g and n (kLanczosGamma // and kLanczosCoefficients.size() + 1). The coefficients below correspond to @@ -1705,6 +1884,7 @@ void populateDecomposeChloPatterns(MLIRContext *context, ConvertDigammaOp, ConvertErfOp, ConvertErfcOp, + ConvertErfInvOp, ConvertLgammaOp, ConvertNextAfterOp, ConvertPolygammaOp, diff --git a/stablehlo/stablehlo/dialect/ChloOps.cpp b/stablehlo/stablehlo/dialect/ChloOps.cpp index 00c8d9171..5d2eefd14 100644 --- a/stablehlo/stablehlo/dialect/ChloOps.cpp +++ b/stablehlo/stablehlo/dialect/ChloOps.cpp @@ -66,6 +66,7 @@ INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(CoshOp) INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(DigammaOp) INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(ErfOp) INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(ErfcOp) +INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(ErfInvOp) INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(LgammaOp) INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(NextAfterOp) INFER_RETURN_TYPE_COMPONENTS_FROM_OPERANDS(PolygammaOp) diff --git a/stablehlo/stablehlo/dialect/ChloOps.td b/stablehlo/stablehlo/dialect/ChloOps.td index 7947d7f4d..478e566cc 100644 --- a/stablehlo/stablehlo/dialect/ChloOps.td +++ b/stablehlo/stablehlo/dialect/ChloOps.td @@ -673,6 +673,15 @@ def CHLO_ErfOp : CHLO_UnaryElementwiseOp<"erf", }]; } +def CHLO_ErfInvOp : CHLO_UnaryElementwiseOp<"erf_inv", + [HLO_CompatibleOperandsAndResultType], HLO_FpTensor, HLO_FpTensor> { + let summary = "Inverse Erf"; + let description = [{ + Returns `ErfInv(operand)` element-wise. + }]; +} + + def CHLO_ErfcOp : CHLO_UnaryElementwiseOp<"erfc", [HLO_CompatibleOperandsAndResultType], HLO_FpTensor, HLO_FpTensor> { let summary = "Erfc operator"; diff --git a/stablehlo/stablehlo/tests/ops_chlo.mlir b/stablehlo/stablehlo/tests/ops_chlo.mlir index df29f97ae..191715565 100644 --- a/stablehlo/stablehlo/tests/ops_chlo.mlir +++ b/stablehlo/stablehlo/tests/ops_chlo.mlir @@ -133,3 +133,10 @@ func.func @top_k(%arg0 : tensor<16x16xf32>) { %0:2 = chlo.top_k(%arg0, k=8) : tensor<16x16xf32> -> (tensor<16x8xf32>, tensor<16x8xi32>) return } + +// ----- + +func.func @erf_inv(%arg0 : tensor<16x16xf32>) { + %0 = chlo.erf_inv %arg0 : tensor<16x16xf32> -> tensor<16x16xf32> + return +} diff --git a/stablehlo/stablehlo/tests/ops_chlo_roundtrip.mlir b/stablehlo/stablehlo/tests/ops_chlo_roundtrip.mlir index 9d8f7cf33..a0a1ea930 100644 --- a/stablehlo/stablehlo/tests/ops_chlo_roundtrip.mlir +++ b/stablehlo/stablehlo/tests/ops_chlo_roundtrip.mlir @@ -432,3 +432,11 @@ func.func @chlo_rank_specialization_cluster(%arg0 : tensor<*xf32>, %arg1 : tenso }) : (tensor<*xf32>, tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32> func.return %0 : tensor<*xf32> } + +// CHECK-LABEL: func @chlo_erf_inv +// CHECK-SAME: %[[A0:.*0]]: tensor<16x16xf32>) +// CHECK: chlo.erf_inv %[[A0]] : tensor<16x16xf32> -> tensor<16x16xf32> +func.func @chlo_erf_inv(%arg0 : tensor<16x16xf32>) { + %0 = "chlo.erf_inv"(%arg0) : (tensor<16x16xf32>) -> tensor<16x16xf32> + return +} diff --git a/tests/Dialect/chlo/chlo_legalize_to_mhlo.mlir b/tests/Dialect/chlo/chlo_legalize_to_mhlo.mlir index 1a3197e0e..ea723c818 100644 --- a/tests/Dialect/chlo/chlo_legalize_to_mhlo.mlir +++ b/tests/Dialect/chlo/chlo_legalize_to_mhlo.mlir @@ -3053,3 +3053,246 @@ func.func @bessel_i1e_f64(%arg : tensor<16x16xf64>) -> tensor<16x16xf64> { %0 = chlo.bessel_i1e %arg : tensor<16x16xf64> -> tensor<16x16xf64> func.return %0 : tensor<16x16xf64> } + +// CHECK-LABEL: func @erf_inv +// CHECK: [[ARG_0:%.*]]: tensor<16x16xf32>) { +// CHECK-DAG: [[TMP_0:%.*]] = mhlo.negate [[ARG_0]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_1:%.*]] = mhlo.multiply [[ARG_0]], [[TMP_0]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_2:%.*]] = mhlo.log_plus_one [[TMP_1]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_3:%.*]] = mhlo.negate [[TMP_2]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_4:%.*]] = mhlo.constant dense<5.000000e+00> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_5:%.*]] = mhlo.compare LT, [[TMP_3]], [[TMP_4]], NOTYPE : (tensor<16x16xf32>, tensor<16x16xf32>) -> tensor<16x16xi1> +// CHECK-DAG: [[TMP_6:%.*]] = mhlo.constant dense<2.500000e+00> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_7:%.*]] = mhlo.subtract [[TMP_3]], [[TMP_6]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_8:%.*]] = mhlo.sqrt [[TMP_3]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_9:%.*]] = mhlo.constant dense<3.000000e+00> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_10:%.*]] = mhlo.subtract [[TMP_8]], [[TMP_9]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_11:%.*]] = mhlo.select [[TMP_5]], [[TMP_7]], [[TMP_10]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_12:%.*]] = mhlo.constant dense<2.81022636E-8> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_13:%.*]] = mhlo.constant dense<-2.00214257E-4> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_14:%.*]] = mhlo.select [[TMP_5]], [[TMP_12]], [[TMP_13]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_15:%.*]] = mhlo.constant dense<3.43273939E-7> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_16:%.*]] = mhlo.constant dense<1.00950558E-4> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_17:%.*]] = mhlo.select [[TMP_5]], [[TMP_15]], [[TMP_16]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_18:%.*]] = mhlo.multiply [[TMP_14]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_19:%.*]] = mhlo.add [[TMP_17]], [[TMP_18]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_20:%.*]] = mhlo.constant dense<-3.5233877E-6> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_21:%.*]] = mhlo.constant dense<0.00134934322> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_22:%.*]] = mhlo.select [[TMP_5]], [[TMP_20]], [[TMP_21]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_23:%.*]] = mhlo.multiply [[TMP_19]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_24:%.*]] = mhlo.add [[TMP_22]], [[TMP_23]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_25:%.*]] = mhlo.constant dense<-4.39150654E-6> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_26:%.*]] = mhlo.constant dense<-0.00367342844> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_27:%.*]] = mhlo.select [[TMP_5]], [[TMP_25]], [[TMP_26]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_28:%.*]] = mhlo.multiply [[TMP_24]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_29:%.*]] = mhlo.add [[TMP_27]], [[TMP_28]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_30:%.*]] = mhlo.constant dense<2.1858087E-4> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_31:%.*]] = mhlo.constant dense<0.00573950773> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_32:%.*]] = mhlo.select [[TMP_5]], [[TMP_30]], [[TMP_31]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_33:%.*]] = mhlo.multiply [[TMP_29]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_34:%.*]] = mhlo.add [[TMP_32]], [[TMP_33]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_35:%.*]] = mhlo.constant dense<-0.00125372503> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_36:%.*]] = mhlo.constant dense<-0.0076224613> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_37:%.*]] = mhlo.select [[TMP_5]], [[TMP_35]], [[TMP_36]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_38:%.*]] = mhlo.multiply [[TMP_34]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_39:%.*]] = mhlo.add [[TMP_37]], [[TMP_38]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_40:%.*]] = mhlo.constant dense<-0.00417768164> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_41:%.*]] = mhlo.constant dense<0.00943887047> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_42:%.*]] = mhlo.select [[TMP_5]], [[TMP_40]], [[TMP_41]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_43:%.*]] = mhlo.multiply [[TMP_39]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_44:%.*]] = mhlo.add [[TMP_42]], [[TMP_43]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_45:%.*]] = mhlo.constant dense<0.246640727> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_46:%.*]] = mhlo.constant dense<1.00167406> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_47:%.*]] = mhlo.select [[TMP_5]], [[TMP_45]], [[TMP_46]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_48:%.*]] = mhlo.multiply [[TMP_44]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_49:%.*]] = mhlo.add [[TMP_47]], [[TMP_48]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_50:%.*]] = mhlo.constant dense<1.50140941> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_51:%.*]] = mhlo.constant dense<2.83297682> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_52:%.*]] = mhlo.select [[TMP_5]], [[TMP_50]], [[TMP_51]] : tensor<16x16xi1>, tensor<16x16xf32> +// CHECK-DAG: [[TMP_53:%.*]] = mhlo.multiply [[TMP_49]], [[TMP_11]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_54:%.*]] = mhlo.add [[TMP_52]], [[TMP_53]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_55:%.*]] = mhlo.multiply [[TMP_54]], [[ARG_0]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_56:%.*]] = mhlo.abs [[ARG_0]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_57:%.*]] = mhlo.constant dense<1.000000e+00> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_58:%.*]] = mhlo.compare EQ, [[TMP_56]], [[TMP_57]], NOTYPE : (tensor<16x16xf32>, tensor<16x16xf32>) -> tensor<16x16xi1> +// CHECK-DAG: [[TMP_59:%.*]] = mhlo.constant dense<3.40282347E+38> : tensor<16x16xf32> +// CHECK-DAG: [[TMP_60:%.*]] = mhlo.multiply [[ARG_0]], [[TMP_59]] : tensor<16x16xf32> +// CHECK-DAG: [[TMP_61:%.*]] = mhlo.select [[TMP_58]], [[TMP_60]], [[TMP_55]] : tensor<16x16xi1>, tensor<16x16xf32> +func.func @erf_inv(%arg0 : tensor<16x16xf32>) { + %0 = chlo.erf_inv %arg0 : tensor<16x16xf32> -> tensor<16x16xf32> + return +} + +// CHECK: @erf_inv_wide([[ARG_0:%.*]]: tensor<16x16xf64>) { +// CHECK-DAG: [[VAL_0:%.*]] = mhlo.negate [[ARG_0]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_1:%.*]] = mhlo.multiply [[ARG_0]], [[VAL_0]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_2:%.*]] = mhlo.log_plus_one [[VAL_1]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_3:%.*]] = mhlo.negate [[VAL_2]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_4:%.*]] = mhlo.constant dense<6.250000e+00> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_5:%.*]] = mhlo.compare LT, [[VAL_3]], [[VAL_4]], NOTYPE : (tensor<16x16xf64>, tensor<16x16xf64>) -> tensor<16x16xi1> +// CHECK-DAG: [[VAL_6:%.*]] = mhlo.constant dense<1.600000e+01> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_7:%.*]] = mhlo.compare LT, [[VAL_3]], [[VAL_6]], NOTYPE : (tensor<16x16xf64>, tensor<16x16xf64>) -> tensor<16x16xi1> +// CHECK-DAG: [[VAL_8:%.*]] = mhlo.sqrt [[VAL_3]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_9:%.*]] = mhlo.constant dense<3.125000e+00> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_10:%.*]] = mhlo.subtract [[VAL_3]], [[VAL_9]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_11:%.*]] = mhlo.constant dense<3.250000e+00> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_12:%.*]] = mhlo.constant dense<5.000000e+00> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_13:%.*]] = mhlo.select [[VAL_7]], [[VAL_11]], [[VAL_12]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_14:%.*]] = mhlo.subtract [[VAL_8]], [[VAL_13]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_15:%.*]] = mhlo.select [[VAL_5]], [[VAL_10]], [[VAL_14]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_16:%.*]] = mhlo.constant dense<-3.6444120640178197E-21> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_17:%.*]] = mhlo.constant dense<2.2137376921775787E-9> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_18:%.*]] = mhlo.select [[VAL_5]], [[VAL_16]], [[VAL_17]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_19:%.*]] = mhlo.constant dense<-2.7109920616438573E-11> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_20:%.*]] = mhlo.select [[VAL_7]], [[VAL_18]], [[VAL_19]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_21:%.*]] = mhlo.constant dense<-1.6850591381820166E-19> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_22:%.*]] = mhlo.constant dense<9.075656193888539E-8> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_23:%.*]] = mhlo.select [[VAL_5]], [[VAL_21]], [[VAL_22]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_24:%.*]] = mhlo.constant dense<-2.5556418169965252E-10> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_25:%.*]] = mhlo.select [[VAL_7]], [[VAL_23]], [[VAL_24]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_26:%.*]] = mhlo.multiply [[VAL_20]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_27:%.*]] = mhlo.add [[VAL_25]], [[VAL_26]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_28:%.*]] = mhlo.constant dense<1.28584807152564E-18> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_29:%.*]] = mhlo.constant dense<-2.7517406297064545E-7> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_30:%.*]] = mhlo.select [[VAL_5]], [[VAL_28]], [[VAL_29]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_31:%.*]] = mhlo.constant dense<1.5076572693500548E-9> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_32:%.*]] = mhlo.select [[VAL_7]], [[VAL_30]], [[VAL_31]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_33:%.*]] = mhlo.multiply [[VAL_27]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_34:%.*]] = mhlo.add [[VAL_32]], [[VAL_33]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_35:%.*]] = mhlo.constant dense<1.1157877678025181E-17> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_36:%.*]] = mhlo.constant dense<1.8239629214389228E-8> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_37:%.*]] = mhlo.select [[VAL_5]], [[VAL_35]], [[VAL_36]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_38:%.*]] = mhlo.constant dense<-3.789465440126737E-9> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_39:%.*]] = mhlo.select [[VAL_7]], [[VAL_37]], [[VAL_38]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_40:%.*]] = mhlo.multiply [[VAL_34]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_41:%.*]] = mhlo.add [[VAL_39]], [[VAL_40]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_42:%.*]] = mhlo.constant dense<-1.3331716628546209E-16> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_43:%.*]] = mhlo.constant dense<1.5027403968909828E-6> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_44:%.*]] = mhlo.select [[VAL_5]], [[VAL_42]], [[VAL_43]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_45:%.*]] = mhlo.constant dense<7.6157012080783394E-9> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_46:%.*]] = mhlo.select [[VAL_7]], [[VAL_44]], [[VAL_45]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_47:%.*]] = mhlo.multiply [[VAL_41]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_48:%.*]] = mhlo.add [[VAL_46]], [[VAL_47]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_49:%.*]] = mhlo.constant dense<2.0972767875968562E-17> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_50:%.*]] = mhlo.constant dense<-4.013867526981546E-6> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_51:%.*]] = mhlo.select [[VAL_5]], [[VAL_49]], [[VAL_50]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_52:%.*]] = mhlo.constant dense<-1.496002662714924E-8> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_53:%.*]] = mhlo.select [[VAL_7]], [[VAL_51]], [[VAL_52]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_54:%.*]] = mhlo.multiply [[VAL_48]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_55:%.*]] = mhlo.add [[VAL_53]], [[VAL_54]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_56:%.*]] = mhlo.constant dense<6.6376381343583238E-15> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_57:%.*]] = mhlo.constant dense<2.9234449089955446E-6> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_58:%.*]] = mhlo.select [[VAL_5]], [[VAL_56]], [[VAL_57]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_59:%.*]] = mhlo.constant dense<2.9147953450901081E-8> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_60:%.*]] = mhlo.select [[VAL_7]], [[VAL_58]], [[VAL_59]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_61:%.*]] = mhlo.multiply [[VAL_55]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_62:%.*]] = mhlo.add [[VAL_60]], [[VAL_61]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_63:%.*]] = mhlo.constant dense<-4.0545662729752069E-14> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_64:%.*]] = mhlo.constant dense<1.2475304481671779E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_65:%.*]] = mhlo.select [[VAL_5]], [[VAL_63]], [[VAL_64]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_66:%.*]] = mhlo.constant dense<-6.7711997758452339E-8> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_67:%.*]] = mhlo.select [[VAL_7]], [[VAL_65]], [[VAL_66]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_68:%.*]] = mhlo.multiply [[VAL_62]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_69:%.*]] = mhlo.add [[VAL_67]], [[VAL_68]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_70:%.*]] = mhlo.constant dense<-8.1519341976054721E-14> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_71:%.*]] = mhlo.constant dense<-4.7318229009055734E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_72:%.*]] = mhlo.select [[VAL_5]], [[VAL_70]], [[VAL_71]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_73:%.*]] = mhlo.constant dense<2.2900482228026655E-7> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_74:%.*]] = mhlo.select [[VAL_7]], [[VAL_72]], [[VAL_73]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_75:%.*]] = mhlo.multiply [[VAL_69]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_76:%.*]] = mhlo.add [[VAL_74]], [[VAL_75]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_77:%.*]] = mhlo.constant dense<2.6335093153082323E-12> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_78:%.*]] = mhlo.constant dense<6.8284851459573175E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_79:%.*]] = mhlo.select [[VAL_5]], [[VAL_77]], [[VAL_78]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_80:%.*]] = mhlo.constant dense<-9.9298272942317003E-7> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_81:%.*]] = mhlo.select [[VAL_7]], [[VAL_79]], [[VAL_80]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_82:%.*]] = mhlo.multiply [[VAL_76]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_83:%.*]] = mhlo.add [[VAL_81]], [[VAL_82]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_84:%.*]] = mhlo.constant dense<-1.2975133253453532E-11> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_85:%.*]] = mhlo.constant dense<2.4031110387097894E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_86:%.*]] = mhlo.select [[VAL_5]], [[VAL_84]], [[VAL_85]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_87:%.*]] = mhlo.constant dense<4.5260625972231537E-6> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_88:%.*]] = mhlo.select [[VAL_7]], [[VAL_86]], [[VAL_87]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_89:%.*]] = mhlo.multiply [[VAL_83]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_90:%.*]] = mhlo.add [[VAL_88]], [[VAL_89]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_91:%.*]] = mhlo.constant dense<-5.4154120542946279E-11> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_92:%.*]] = mhlo.constant dense<-3.5503752036284748E-4> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_93:%.*]] = mhlo.select [[VAL_5]], [[VAL_91]], [[VAL_92]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_94:%.*]] = mhlo.constant dense<-1.9681778105531671E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_95:%.*]] = mhlo.select [[VAL_7]], [[VAL_93]], [[VAL_94]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_96:%.*]] = mhlo.multiply [[VAL_90]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_97:%.*]] = mhlo.add [[VAL_95]], [[VAL_96]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_98:%.*]] = mhlo.constant dense<1.0512122733215323E-9> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_99:%.*]] = mhlo.constant dense<9.5328937973738049E-4> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_100:%.*]] = mhlo.select [[VAL_5]], [[VAL_98]], [[VAL_99]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_101:%.*]] = mhlo.constant dense<7.5995277030017761E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_102:%.*]] = mhlo.select [[VAL_7]], [[VAL_100]], [[VAL_101]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_103:%.*]] = mhlo.multiply [[VAL_97]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_104:%.*]] = mhlo.add [[VAL_102]], [[VAL_103]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_105:%.*]] = mhlo.constant dense<-4.1126339803469837E-9> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_106:%.*]] = mhlo.constant dense<-0.0016882755560235047> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_107:%.*]] = mhlo.select [[VAL_5]], [[VAL_105]], [[VAL_106]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_108:%.*]] = mhlo.constant dense<-2.1503011930044477E-4> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_109:%.*]] = mhlo.select [[VAL_7]], [[VAL_107]], [[VAL_108]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_110:%.*]] = mhlo.multiply [[VAL_104]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_111:%.*]] = mhlo.add [[VAL_109]], [[VAL_110]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_112:%.*]] = mhlo.constant dense<-2.9070369957882005E-8> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_113:%.*]] = mhlo.constant dense<0.0024914420961078508> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_114:%.*]] = mhlo.select [[VAL_5]], [[VAL_112]], [[VAL_113]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_115:%.*]] = mhlo.constant dense<-1.3871931833623122E-4> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_116:%.*]] = mhlo.select [[VAL_7]], [[VAL_114]], [[VAL_115]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_117:%.*]] = mhlo.multiply [[VAL_111]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_118:%.*]] = mhlo.add [[VAL_116]], [[VAL_117]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_119:%.*]] = mhlo.constant dense<4.2347877827932404E-7> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_120:%.*]] = mhlo.constant dense<-0.0037512085075692412> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_121:%.*]] = mhlo.select [[VAL_5]], [[VAL_119]], [[VAL_120]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_122:%.*]] = mhlo.constant dense<1.0103004648645344> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_123:%.*]] = mhlo.select [[VAL_7]], [[VAL_121]], [[VAL_122]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_124:%.*]] = mhlo.multiply [[VAL_118]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_125:%.*]] = mhlo.add [[VAL_123]], [[VAL_124]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_126:%.*]] = mhlo.constant dense<-1.3654692000834679E-6> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_127:%.*]] = mhlo.constant dense<0.0053709145535900636> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_128:%.*]] = mhlo.select [[VAL_5]], [[VAL_126]], [[VAL_127]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_129:%.*]] = mhlo.constant dense<4.8499064014085844> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_130:%.*]] = mhlo.select [[VAL_7]], [[VAL_128]], [[VAL_129]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_131:%.*]] = mhlo.multiply [[VAL_125]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_132:%.*]] = mhlo.add [[VAL_130]], [[VAL_131]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_133:%.*]] = mhlo.constant dense<-1.3882523362786469E-5> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_134:%.*]] = mhlo.constant dense<1.0052589676941592> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_135:%.*]] = mhlo.select [[VAL_5]], [[VAL_133]], [[VAL_134]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_136:%.*]] = mhlo.multiply [[VAL_132]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_137:%.*]] = mhlo.add [[VAL_135]], [[VAL_136]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_138:%.*]] = mhlo.select [[VAL_7]], [[VAL_137]], [[VAL_132]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_139:%.*]] = mhlo.constant dense<1.8673420803405714E-4> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_140:%.*]] = mhlo.constant dense<3.0838856104922208> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_141:%.*]] = mhlo.select [[VAL_5]], [[VAL_139]], [[VAL_140]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_142:%.*]] = mhlo.multiply [[VAL_138]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_143:%.*]] = mhlo.add [[VAL_141]], [[VAL_142]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_144:%.*]] = mhlo.select [[VAL_7]], [[VAL_143]], [[VAL_138]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_145:%.*]] = mhlo.constant dense<-7.4070253416626698E-4> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_146:%.*]] = mhlo.multiply [[VAL_144]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_147:%.*]] = mhlo.add [[VAL_145]], [[VAL_146]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_148:%.*]] = mhlo.select [[VAL_5]], [[VAL_147]], [[VAL_144]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_149:%.*]] = mhlo.constant dense<-0.0060336708714301491> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_150:%.*]] = mhlo.multiply [[VAL_148]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_151:%.*]] = mhlo.add [[VAL_149]], [[VAL_150]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_152:%.*]] = mhlo.select [[VAL_5]], [[VAL_151]], [[VAL_148]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_153:%.*]] = mhlo.constant dense<0.24015818242558962> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_154:%.*]] = mhlo.multiply [[VAL_152]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_155:%.*]] = mhlo.add [[VAL_153]], [[VAL_154]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_156:%.*]] = mhlo.select [[VAL_5]], [[VAL_155]], [[VAL_152]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_157:%.*]] = mhlo.constant dense<1.6536545626831027> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_158:%.*]] = mhlo.multiply [[VAL_156]], [[VAL_15]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_159:%.*]] = mhlo.add [[VAL_157]], [[VAL_158]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_160:%.*]] = mhlo.select [[VAL_5]], [[VAL_159]], [[VAL_156]] : tensor<16x16xi1>, tensor<16x16xf64> +// CHECK-DAG: [[VAL_161:%.*]] = mhlo.multiply [[VAL_160]], [[ARG_0]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_162:%.*]] = mhlo.abs [[ARG_0]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_163:%.*]] = mhlo.constant dense<1.000000e+00> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_164:%.*]] = mhlo.compare EQ, [[VAL_162]], [[VAL_163]], NOTYPE : (tensor<16x16xf64>, tensor<16x16xf64>) -> tensor<16x16xi1> +// CHECK-DAG: [[VAL_165:%.*]] = mhlo.constant dense<1.7976931348623157E+308> : tensor<16x16xf64> +// CHECK-DAG: [[VAL_166:%.*]] = mhlo.multiply [[ARG_0]], [[VAL_165]] : tensor<16x16xf64> +// CHECK-DAG: [[VAL_167:%.*]] = mhlo.select [[VAL_164]], [[VAL_166]], [[VAL_161]] : tensor<16x16xi1>, tensor<16x16xf64> +func.func @erf_inv_wide(%arg0 : tensor<16x16xf64>) { + %0 = chlo.erf_inv %arg0 : tensor<16x16xf64> -> tensor<16x16xf64> + return +}