/external/tensorflow/tensorflow/compiler/mlir/hlo/tests/ |
D | chlo_legalize_to_mhlo.mlir | 7 // CHECK: mhlo.constant dense<3.389{{.*}}e+38> 16 // CHECK: mhlo.constant dense<6.550{{.*}}e+04> 25 // CHECK: mhlo.constant dense<3.402{{.*}}E+38> 33 // CHECK: %[[TMP_0:.*]] = "mhlo.sign"(%[[ARG]]) 34 // CHECK: %[[TMP_1:.*]] = "mhlo.abs"(%[[ARG]]) 35 // CHECK: %[[TMP_2:.*]] = mhlo.constant dense<1.797{{.*}}E+308> 36 // CHECK: %[[TMP_3:.*]] = "mhlo.sqrt"(%[[TMP_2]]) 37 // CHECK: %[[TMP_4:.*]] = "mhlo.compare"(%[[TMP_1]], %[[TMP_3]]) {comparison_direction = "GE"} 38 // CHECK: %[[TMP_5:.*]] = "mhlo.abs"(%[[ARG]]) 39 // CHECK: %[[TMP_6:.*]] = "mhlo.log"(%[[TMP_5]]) [all …]
|
D | lower-complex.mlir | 1 // RUN: mlir-hlo-opt %s -chlo-legalize-to-hlo -mhlo-test-lower-complex | FileCheck %s 5 %2 = "mhlo.complex"(%arg0, %arg1) : (tensor<2xf32>, tensor<2xf32>) -> (tensor<2xcomplex<f32>>) 6 %3 = "mhlo.complex"(%arg2, %arg3) : (tensor<2xf32>, tensor<2xf32>) -> (tensor<2xcomplex<f32>>) 8 // CHECK-DAG: [[VAL0:%.+]] = mhlo.add %arg0, %arg2 9 // CHECK-DAG: [[VAL1:%.+]] = mhlo.add %arg1, %arg3 10 …%4 = "mhlo.add"(%2, %3) : (tensor<2xcomplex<f32>>, tensor<2xcomplex<f32>>) -> (tensor<2xcomplex<f3… 11 %5 = "mhlo.real"(%4) : (tensor<2xcomplex<f32>>) -> (tensor<2xf32>) 12 %6 = "mhlo.imag"(%4) : (tensor<2xcomplex<f32>>) -> (tensor<2xf32>) 20 %2 = "mhlo.complex"(%arg0, %arg1) : (tensor<*xf32>, tensor<*xf32>) -> (tensor<*xcomplex<f32>>) 21 %3 = "mhlo.complex"(%arg2, %arg3) : (tensor<*xf32>, tensor<*xf32>) -> (tensor<*xcomplex<f32>>) [all …]
|
D | canonicalize.mlir | 5 %0 = mhlo.constant dense<[1, 2, 3, 4]> : tensor<4xi64> 6 %1 = mhlo.constant dense<[5, 6, 7, 8]> : tensor<4xi64> 7 // CHECK: mhlo.constant dense<[6, 8, 10, 12]> 8 %2 = "mhlo.add"(%0, %1) : (tensor<4xi64>, tensor<4xi64>) -> (tensor<4xi64>) 14 %0 = mhlo.constant dense<1> : tensor<4xi64> 15 %1 = mhlo.constant dense<5> : tensor<4xi64> 16 // CHECK: mhlo.constant dense<6> 17 %2 = "mhlo.add"(%0, %1) : (tensor<4xi64>, tensor<4xi64>) -> (tensor<4xi64>) 23 %0 = mhlo.constant dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf64> 24 %1 = mhlo.constant dense<[5.0, 6.0, 7.0, 8.0]> : tensor<4xf64> [all …]
|
D | sink-constants-to-control-flow.mlir | 1 // RUN: mlir-hlo-opt %s -mhlo-sink-constants-to-control-flow | FileCheck %s 7 // CHECK-NEXT: mhlo.while 8 %c0 = mhlo.constant dense<1> : tensor<i64> 9 %c1 = mhlo.constant dense<2> : tensor<i64> 10 %0 = "mhlo.while"(%arg0) ( { 13 // CHECK: %[[C0:.+]] = mhlo.constant dense<1> : tensor<i64> 14 // CHECK: "mhlo.compare"(%[[C0]], %[[ARG1A]]) 15 …%1 = "mhlo.compare"(%c0, %arg1) {comparison_direction = "LT"} : (tensor<i64>, tensor<i64>) -> tens… 16 "mhlo.return"(%1) : (tensor<i1>) -> () 20 // CHECK-DAG: %[[C1:.+]] = mhlo.constant dense<2> : tensor<i64> [all …]
|
D | mhlo-fusion.mlir | 1 // RUN: mlir-hlo-opt %s -mhlo-fusion -split-input-file | FileCheck %s 5 %0 = "mhlo.add"(%arg0, %arg1) : (tensor<?x?xf32>, tensor<?x?xf32>) -> tensor<?x?xf32> 6 %1 = "mhlo.subtract"(%arg0, %0) : (tensor<?x?xf32>, tensor<?x?xf32>) -> tensor<?x?xf32> 7 %2 = "mhlo.add"(%1, %1) : (tensor<?x?xf32>, tensor<?x?xf32>) -> tensor<?x?xf32> 8 // CHECK: %[[RET:.*]]:2 = "mhlo.fusion" 9 // CHECK-NEXT: mhlo.add 10 // CHECK-NEXT: mhlo.subtract 11 // CHECK-NEXT: mhlo.add 12 // CHECK-NEXT: mhlo.return 20 %0 = "mhlo.abs"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32> [all …]
|
D | convert.mlir | 8 %0 = "mhlo.convert"(%arg) : (tensor<f32>) -> tensor<f32> 18 // CHECK-NEXT: [[RES:%.+]] = "mhlo.convert"([[ARG]]) : (tensor<i32>) -> tensor<i64> 19 %0 = "mhlo.convert"(%arg) : (tensor<i32>) -> tensor<i64> 29 // CHECK-NEXT: [[RES:%.+]] = "mhlo.convert"([[ARG]]) : (tensor<i32>) -> tensor<i16> 30 %0 = "mhlo.convert"(%arg) : (tensor<i32>) -> tensor<i16> 40 // CHECK-NEXT: [[RES:%.+]] = "mhlo.convert"([[ARG]]) : (tensor<f32>) -> tensor<i32> 41 %0 = "mhlo.convert"(%arg) : (tensor<f32>) -> tensor<i32> 51 // CHECK-NEXT: [[RES:%.+]] = "mhlo.convert"([[ARG]]) : (tensor<i32>) -> tensor<f32> 52 %0 = "mhlo.convert"(%arg) : (tensor<i32>) -> tensor<f32> 62 // CHECK-NEXT: [[RES:%.+]] = "mhlo.convert"([[ARG]]) : (tensor<2x3xi32>) -> tensor<2x3xf32> [all …]
|
D | reshape.mlir | 5 // CHECK-NEXT: [[CST:%.+]] = mhlo.constant dense<42> : tensor<i32> 6 %cst = mhlo.constant dense<42> : tensor<1x1xi32> 7 %0 = "mhlo.reshape"(%cst) : (tensor<1x1xi32>) -> tensor<i32> 16 // CHECK-NEXT: [[CST:%.+]] = mhlo.constant dense<42> : tensor<2xi32> 17 %cst = mhlo.constant dense<42> : tensor<1x2xi32> 18 %0 = "mhlo.reshape"(%cst) : (tensor<1x2xi32>) -> tensor<2xi32> 27 // CHECK-NEXT: [[CST:%.+]] = mhlo.constant dense<42> : tensor<1xi32> 28 %cst = mhlo.constant dense<42> : tensor<i32> 29 %0 = "mhlo.reshape"(%cst) : (tensor<i32>) -> tensor<1xi32> 38 // CHECK-NEXT: [[CST:%.+]] = mhlo.constant dense<42> : tensor<16xi64> [all …]
|
D | ops.mlir | 6 // CHECK-LABEL: func private @token_type() -> !mhlo.token 7 func private @token_type() -> !mhlo.token 11 // expected-error@+1 {{unknown mhlo type: foobar}} 12 func private @invalid_type() -> !mhlo.foobar 18 %0 = "mhlo.all_to_all"(%data) { 31 %0 = "mhlo.all_to_all"(%data) { 44 %0 = "mhlo.all_to_all"(%data) { 57 …%0 = "mhlo.broadcast"(%arg0) {broadcast_sizes = dense<[1, 2]> : tensor<2xi64>} : (tensor<3xi32>) -… 65 …%0 = "mhlo.broadcast"(%arg0) {broadcast_sizes = dense<[[1, 2]]> : tensor<1x2xi64>} : (tensor<3xi32… 73 …%0 = "mhlo.broadcast"(%arg0) {broadcast_sizes = dense<[2]> : tensor<1xi64>} : (tensor<3xi32>) -> t… [all …]
|
D | legalize-control-flow.mlir | 1 // RUN: mlir-hlo-opt -mhlo-legalize-control-flow %s -o - | FileCheck %s 7 //CHECK: [[VAL1:%.+]] = "mhlo.compare"([[VAL0]], [[VAL0]]) 11 //CHECK: [[VAL4:%.+]] = mhlo.add [[VAL3]], [[VAL3]] 14 %0 = "mhlo.while"(%arg0) ( { 16 …%1 = "mhlo.compare"(%arg1, %arg1) {comparison_direction = "LT", name = "compare.2"} : (tensor<i64>… 17 "mhlo.return"(%1) : (tensor<i1>) -> () 20 %1 = mhlo.add %arg1, %arg1 {name = "compare.0"} : tensor<i64> 21 "mhlo.return"(%1) : (tensor<i64>) -> () 33 …// CHECK: [[VAL0:%.+]] = "mhlo.compare"(%arg0, [[C0]]) {comparison_direction = "LT"} : (tensor<f… 34 …%0 = "mhlo.compare"(%arg0, %cst) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> ten… [all …]
|
D | optimize-hlo.mlir | 1 // RUN: mlir-hlo-opt %s -pass-pipeline='func(mhlo-test-optimize)' | FileCheck %s 5 // CHECK: [[CST:%.+]] = mhlo.constant dense<0> : tensor<i64> 6 …// CHECK: [[SLICE:%.+]] = "mhlo.dynamic-slice"(%arg0, %arg1, [[CST]], [[CST]]) {slice_sizes = dens… 7 // CHECK: [[RESHAPE:%.+]] = "mhlo.reshape"([[SLICE]]) 8 %res = "mhlo.gather"(%arg0, %arg1) { 24 // CHECK: [[CST:%.+]] = mhlo.constant dense<0> : tensor<i64> 25 // CHECK: [[RESHAPE:%.+]] = "mhlo.reshape"(%arg1) 26 …// CHECK: [[SLICE:%.+]] = "mhlo.dynamic-slice"(%arg0, [[RESHAPE]], [[CST]], [[CST]]) {slice_sizes … 27 // CHECK: [[RES:%.+]] = "mhlo.reshape"([[SLICE]]) 29 %res = "mhlo.gather"(%arg0, %arg1) { [all …]
|
D | unfuse_batch_norm.mlir | 1 // RUN: mlir-hlo-opt -split-input-file -mhlo-test-unfuse-batch-norm -verify-diagnostics %s | FILECH… 13 // CHECK-DAG: %[[EPS_BCAST:.+]] = mhlo.constant dense<1.001000e-05> : tensor<256xf32> 14 // CHECK-DAG: %[[VARIANCE_EPS:.+]] = mhlo.add %[[VARIANCE]], %[[EPS_BCAST]] : tensor<256xf32> 15 …// CHECK-DAG: %[[STDDEV:.+]] = "mhlo.sqrt"(%[[VARIANCE_EPS]]) : (tensor<256xf32>) -> tensor<256xf3… 16 …// CHECK-DAG: %[[STDDEV_BCAST:.+]] = "mhlo.broadcast_in_dim"(%[[STDDEV]]) {broadcast_dimensions = … 17 …// CHECK-DAG: %[[SCALE_BCAST:.+]] = "mhlo.broadcast_in_dim"(%[[SCALE]]) {broadcast_dimensions = de… 18 …// CHECK-DAG: %[[OFFSET_BCAST:.+]] = "mhlo.broadcast_in_dim"(%[[OFFSET]]) {broadcast_dimensions = … 19 …// CHECK-DAG: %[[MEAN_BCAST:.+]] = "mhlo.broadcast_in_dim"(%[[MEAN]]) {broadcast_dimensions = dens… 20 // CHECK-DAG: %[[X_CENTER:.+]] = mhlo.subtract %[[X]], %[[MEAN_BCAST]] : tensor<4x256xf32> 21 // CHECK-DAG: %[[X_SCALED:.+]] = mhlo.multiply %[[X_CENTER]], %[[SCALE_BCAST]] : tensor<4x256xf32> [all …]
|
D | legalize-to-std.mlir | 1 // RUN: mlir-hlo-opt -mhlo-legalize-to-std %s -o - | FileCheck %s 6 %0 = "mhlo.add"(%arg0, %arg1) {name = "add.3"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> 9 %1 = "mhlo.multiply"(%0, %arg1) {name = "mul.4"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> 12 %2 = "mhlo.subtract"(%1, %arg1) {name = "sub.5"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> 15 %3 = "mhlo.divide"(%2, %arg1) {name = "div.6"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> 18 %4 = "mhlo.remainder"(%3, %arg1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> 27 %0 = "mhlo.add"(%arg0, %arg1) {name = "add.3"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi32> 30 %1 = "mhlo.multiply"(%0, %arg1) {name = "mul.4"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi32> 33 %2 = "mhlo.subtract"(%1, %arg1) {name = "sub.5"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi32> 36 %3 = "mhlo.divide"(%2, %arg1) {name = "div.6"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi32> [all …]
|
/external/tensorflow/tensorflow/compiler/mlir/xla/tests/ |
D | legalize-tf-communication.mlir | 4 // `mhlo.send` per operand and `mhlo.recv` per result. Channel Id's are uniquely 5 // assigned per mhlo communcation op, and frontend attributes (modified keys) 20 // CHECK: [[INIT_TOKEN:%.*]] = "mhlo.create_token" 22 // CHECK: [[SEND_ARG0_TOKEN:%.*]] = "mhlo.send"([[ARG0]], [[INIT_TOKEN]]) 25 …// CHECK-SAME: mhlo.frontend_attributes = {_xla_host_transfer_original_type = "s32", _xla_host_tra… 26 // CHECK-SAME: mhlo.sharding = "\08\01\1A\01\01\22\01\00" 27 // CHECK-SAME: (tensor<i32>, !mhlo.token) -> !mhlo.token 29 // CHECK: [[SEND_ARG1_TOKEN:%.*]] = "mhlo.send"([[ARG1]], [[INIT_TOKEN]]) 32 …// CHECK-SAME: mhlo.frontend_attributes = {_xla_host_transfer_original_type = "s64", _xla_host_tra… 33 // CHECK-SAME: mhlo.sharding = "\08\01\1A\01\01\22\01\00" [all …]
|
D | legalize-tf-control-flow.mlir | 6 …// CHECK: [[VAL0:%.+]] = "mhlo.compare"([[ARG0]], [[ARG1]]) {comparison_direction = "GT"} : (tenso… 7 …%0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "GT"} : (tensor<f32>, tensor<f32>) -> te… 8 // CHECK: [[VAL1:%.+]] = "mhlo.tuple"([[ARG0]], [[ARG1]]) 9 // CHECK: [[VAL2:%.+]] = "mhlo.if"([[VAL0]], [[VAL1]], [[VAL1]]) ( { 11 // CHECK: [[VAL4:%.+]] = "mhlo.get_tuple_element"([[THEN_ARG]]) {index = 0 : i32} 12 // CHECK: [[VAL5:%.+]] = "mhlo.get_tuple_element"([[THEN_ARG]]) {index = 1 : i32} 14 // CHECK: [[VAL7:%.+]] = "mhlo.tuple"([[VAL6]]) 15 // CHECK: "mhlo.return"([[VAL7]]) : (tuple<tensor<f32>>) -> () 18 // CHECK: [[VAL4:%.+]] = "mhlo.get_tuple_element"([[ELSE_ARG]]) {index = 0 : i32} 19 // CHECK: [[VAL5:%.+]] = "mhlo.get_tuple_element"([[ELSE_ARG]]) {index = 1 : i32} [all …]
|
D | legalize-tf.mlir | 16 …// CHECK: "mhlo.batch_norm_inference"(%arg0, %arg1, %arg2, %arg3, %arg4) {epsilon = 1.000000e-03 :… 34 …// CHECK: "mhlo.batch_norm_inference"({{.*}}, %arg1, %arg2, %arg3, %arg4) {epsilon = 1.000000e-03 … 41 …// CHECK: %[[RESULT0:.*]] = "mhlo.batch_norm_training"({{.*}}, %arg1, %arg2) {epsilon = 1.000000e-… 43 …// CHECK: "mhlo.get_tuple_element"(%[[RESULT0]]) {index = 0 : i32} : (tuple<tensor<8x8x8x8xf32>, t… 44 …// CHECK: "mhlo.get_tuple_element"(%[[RESULT0]]) {index = 1 : i32} : (tuple<tensor<8x8x8x8xf32>, t… 45 …// CHECK: %[[VAR:.*]] = "mhlo.get_tuple_element"(%[[RESULT0]]) {index = 2 : i32} : (tuple<tensor<8… 46 // CHECK: mhlo.constant 53 …// CHECK: "mhlo.batch_norm_inference"({{.*}}, %arg1, %arg2, %arg3, %arg4) {epsilon = 1.000000e-03 … 61 …// CHECK: [[CONVERT_X:%.*]] = "mhlo.convert"([[X]]) : (tensor<8x8x8x8xbf16>) -> tensor<8x8x8x8xf32> 62 …// CHECK: [[Y:%.*]] = "mhlo.batch_norm_inference"([[CONVERT_X]], [[SCALE]], [[OFFSET]], [[MEAN]], … [all …]
|
/external/tensorflow/tensorflow/compiler/mlir/xla/tests/translate/ |
D | location_to_op_metadata.mlir | 4 func @main(%arg0: !mhlo.token) -> !mhlo.token { 5 %0 = "mhlo.after_all"(%arg0) : (!mhlo.token) -> !mhlo.token loc(unknown) 6 return %0 : !mhlo.token 15 func @main(%arg0: !mhlo.token) -> !mhlo.token { 16 %0 = "mhlo.after_all"(%arg0) : (!mhlo.token) -> !mhlo.token loc("AfterAll") 17 return %0 : !mhlo.token 26 func @main(%arg0: !mhlo.token) -> !mhlo.token { 27 %0 = "mhlo.after_all"(%arg0) : (!mhlo.token) -> !mhlo.token loc("name@function") 28 return %0 : !mhlo.token 37 func @main(%arg0: !mhlo.token) -> !mhlo.token { [all …]
|
D | dynamic_parameter_binding_invalid.mlir | 3 // Test bad `mhlo.padding_map` attribute type. 5 func @main(%arg0: tensor<i32>, %arg1: tensor<10xf32> {mhlo.padding_map = ""}) { 9 // CHECK: requires 'mhlo.padding_map' dict attribute at arg 1 13 // Test missing `shape_indices` attribute in `mhlo.padding_map`. 15 func @main(%arg0: tensor<i32>, %arg1: tensor<10xf32> {mhlo.padding_map = {}}) { 19 // CHECK: requires 'shape_indices' array attribute in 'mhlo.padding_map' dict at arg 1 23 // Test bad `shape_indices` attribute type in `mhlo.padding_map`. 25 func @main(%arg0: tensor<i32>, %arg1: tensor<10xf32> {mhlo.padding_map = {shape_indices = ""}}) { 29 // CHECK: requires 'shape_indices' array attribute in 'mhlo.padding_map' dict at arg 1 33 // Test missing `padding_arg_indices` attribute in `mhlo.padding_map`. [all …]
|
D | fully_connected_reference_model.hlotxt | 14 …// CHECK-NEXT: %[[VAL_2:.*]] = "mhlo.reshape"(%[[VAL_0]]) : (tensor<1x300xf32>) -> tensor<1x300xf3… 17 …// CHECK-NEXT: %[[VAL_3:.*]] = "mhlo.transpose"(%[[VAL_2]]) {permutation = dense<[1, 0]> : tensor<… 20 …// CHECK-NEXT: %[[VAL_4:.*]] = "mhlo.reshape"(%[[VAL_3]]) : (tensor<300x1xf32>) -> tensor<300x1x1x… 23 …// CHECK-NEXT: %[[VAL_5:.*]] = "mhlo.reshape"(%[[VAL_4]]) : (tensor<300x1x1xf32>) -> tensor<300x1x… 26 …// CHECK-NEXT: %[[VAL_6:.*]] = "mhlo.broadcast_in_dim"(%[[VAL_5]]) {broadcast_dimensions = dense<[… 29 // CHECK-NEXT: %[[VAL_7:.*]] = mhlo.constant dense<1.000000e+00> : tensor<f32> 32 …// CHECK-NEXT: %[[VAL_8:.*]] = "mhlo.broadcast_in_dim"(%[[VAL_7]]) {broadcast_dimensions = dense<>… 35 // CHECK-NEXT: %[[VAL_9:.*]] = mhlo.multiply %[[VAL_6]], %[[VAL_8]] : tensor<300x1x5xf32> 38 // CHECK-NEXT: %[[VAL_10:.*]] = mhlo.constant dense<0.000000e+00> : tensor<f32> 41 …// CHECK-NEXT: %[[VAL_11:.*]] = "mhlo.broadcast_in_dim"(%[[VAL_10]]) {broadcast_dimensions = dense… [all …]
|
D | if.mlir | 7 %0 = "mhlo.get_tuple_element"(%arg0) {index = 0 : i32} : (tuple<tensor<f32>>) -> tensor<f32> 10 %1 = "mhlo.log"(%0) : (tensor<f32>) -> tensor<f32> 13 %2 = "mhlo.tuple"(%1) : (tensor<f32>) -> tuple<tensor<f32>> 21 %0 = "mhlo.get_tuple_element"(%arg0) {index = 0 : i32} : (tuple<tensor<f32>>) -> tensor<f32> 24 %1 = "mhlo.exponential"(%0) : (tensor<f32>) -> tensor<f32> 27 %2 = "mhlo.tuple"(%1) : (tensor<f32>) -> tuple<tensor<f32>> 38 …%0 = "mhlo.compare"(%arg0, %cst) {comparison_direction = "LT"} : (tensor<f32>, tensor<f32>) -> ten… 41 %1 = "mhlo.tuple"(%arg0) : (tensor<f32>) -> tuple<tensor<f32>> 44 %2 = "mhlo.if"(%0, %1, %1) ( { 46 %6 = "mhlo.get_tuple_element"(%arg1) {index = 0 : i32} : (tuple<tensor<f32>>) -> tensor<f32> [all …]
|
D | export.mlir | 5 func @main(%arg0: !mhlo.token, %arg1: !mhlo.token) -> !mhlo.token { 6 %0 = "mhlo.after_all"(%arg0, %arg1) : (!mhlo.token, !mhlo.token) -> !mhlo.token 7 return %0 : !mhlo.token 19 %0 = "mhlo.all_reduce"(%arg0) ({ 22 %max = mhlo.maximum %lhs, %rhs : tensor<f32> 23 "mhlo.return"(%max) : (tensor<f32>) -> () 48 %0 = "mhlo.all_reduce"(%arg0) ({ 51 %max = mhlo.maximum %lhs, %rhs : tensor<f32> 52 "mhlo.return"(%max) : (tensor<f32>) -> () 76 …%0 = "mhlo.batch_norm_grad" (%input, %scale, %mean, %variance, %grad_output) {epsilon = 0.001 : f3… [all …]
|
/external/tensorflow/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/ |
D | chlo_legalize_to_hlo.cc | 59 rewriter.replaceOpWithNewOp<mhlo::ConstOp>( in matchAndRewrite() 68 Value constant = rewriter.create<mhlo::ConstOp>(loc, op.value()); in matchAndRewrite() 75 rewriter.replaceOpWithNewOp<mhlo::DynamicBroadcastInDimOp>( in matchAndRewrite() 87 poly = rewriter.create<mhlo::MulOp>(loc, x.getType(), poly, x); in MaterializePolynomialApproximation() 88 poly = rewriter.create<mhlo::AddOp>( in MaterializePolynomialApproximation() 128 Value x_sq = rewriter.create<mhlo::MulOp>(loc, x, x); in MaterializeErfcApproximationF64ForMagnituteGEOne() 129 Value z = rewriter.create<mhlo::NegOp>(loc, x_sq); in MaterializeErfcApproximationF64ForMagnituteGEOne() 133 Value exp_z = rewriter.create<mhlo::ExpOp>(loc, z); in MaterializeErfcApproximationF64ForMagnituteGEOne() 134 Value abs_x = rewriter.create<mhlo::AbsOp>(loc, x); in MaterializeErfcApproximationF64ForMagnituteGEOne() 137 Value exp_z_mul_poly_p = rewriter.create<mhlo::MulOp>(loc, exp_z, poly_p); in MaterializeErfcApproximationF64ForMagnituteGEOne() [all …]
|
D | hlo_legalize_to_lhlo.cc | 43 namespace mhlo { namespace 128 rewriter.create<mhlo::HloToLhloOp<HloOpTy>>(op->getLoc(), llvm::None, in matchAndRewrite() 142 class HloToLhloOpConverter<mhlo::DotOp> : public BaseOpConversion<mhlo::DotOp> { 144 using BaseOpConversion<mhlo::DotOp>::BaseOpConversion; 146 mhlo::DotOp hloOp, ArrayRef<Value> operands, in matchAndRewrite() 161 auto dimension_numbers = mhlo::DotDimensionNumbers::get( in matchAndRewrite() 171 : public BaseOpConversion<mhlo::CustomCallOp> { 173 using BaseOpConversion<mhlo::CustomCallOp>::BaseOpConversion; 176 mhlo::CustomCallOp hloOp, ArrayRef<Value> operands, in matchAndRewrite() 197 : public BaseOpConversion<mhlo::ReshapeOp> { [all …]
|
/external/tensorflow/tensorflow/compiler/mlir/xla/ |
D | hlo_utils.cc | 261 mlir::mhlo::GatherDimensionNumbers CreateGatherDimensionNumbers( in CreateGatherDimensionNumbers() 279 return mlir::mhlo::GatherDimensionNumbers::get( in CreateGatherDimensionNumbers() 287 if (isa<mlir::mhlo::ConstOp, mlir::lmhlo::ConstOp>(op)) { in MhloToHloOpcode() 289 } else if (isa<mlir::mhlo::IotaOp, mlir::lmhlo::IotaOp>(op)) { in MhloToHloOpcode() 291 } else if (isa<mlir::mhlo::ConvertOp, mlir::lmhlo::ConvertOp>(op)) { in MhloToHloOpcode() 293 } else if (isa<mlir::mhlo::AddOp, mlir::lmhlo::AddOp>(op)) { in MhloToHloOpcode() 295 } else if (isa<mlir::mhlo::Atan2Op, mlir::lmhlo::Atan2Op>(op)) { in MhloToHloOpcode() 297 } else if (isa<mlir::mhlo::DivOp, mlir::lmhlo::DivOp>(op)) { in MhloToHloOpcode() 299 } else if (isa<mlir::mhlo::MaxOp, mlir::lmhlo::MaxOp>(op)) { in MhloToHloOpcode() 301 } else if (isa<mlir::mhlo::MinOp, mlir::lmhlo::MinOp>(op)) { in MhloToHloOpcode() [all …]
|
/external/tensorflow/tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/transforms/ |
D | map_chlo_to_hlo_op.h | 29 static mhlo::ComplexOp CreateOp(BroadcastComplexOp from_op, Type result_type, in CreateOp() 32 return builder.create<mhlo::ComplexOp>(from_op.getLoc(), result_type, in CreateOp() 46 static mhlo::CompareOp CreateOp(BroadcastCompareOp from_op, Type result_type, in CreateOp() 49 return builder.create<mhlo::CompareOp>( in CreateOp() 67 POPULATE_BCAST(BroadcastAddOp, mhlo::AddOp); in PopulateForBroadcastingBinaryOp() 68 POPULATE_BCAST(BroadcastAndOp, mhlo::AndOp); in PopulateForBroadcastingBinaryOp() 69 POPULATE_BCAST(BroadcastAtan2Op, mhlo::Atan2Op); in PopulateForBroadcastingBinaryOp() 70 POPULATE_BCAST(BroadcastDivOp, mhlo::DivOp); in PopulateForBroadcastingBinaryOp() 71 POPULATE_BCAST(BroadcastMaxOp, mhlo::MaxOp); in PopulateForBroadcastingBinaryOp() 72 POPULATE_BCAST(BroadcastMinOp, mhlo::MinOp); in PopulateForBroadcastingBinaryOp() [all …]
|
/external/tensorflow/tensorflow/compiler/mlir/hlo/ |
D | BUILD | 38 exports_files(["include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.td"]) 40 exports_files(["include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td"]) 45 "include/mlir-hlo/Dialect/mhlo/IR/chlo_ops.td", 46 "include/mlir-hlo/Dialect/mhlo/IR/hlo_ops.td", 47 "include/mlir-hlo/Dialect/mhlo/IR/hlo_ops_base.td", 48 "include/mlir-hlo/Dialect/mhlo/IR/hlo_ops_base_enums.td", 49 "include/mlir-hlo/Dialect/mhlo/IR/hlo_ops_base_structs.td", 50 "include/mlir-hlo/Dialect/mhlo/IR/hlo_utils.td", 51 "include/mlir-hlo/Dialect/mhlo/IR/infer_fusibility_op_interface.td", 52 "include/mlir-hlo/Dialect/mhlo/IR/lhlo_ops.td", [all …]
|