Home
last modified time | relevance | path

Searched refs:lmhlo (Results 1 – 25 of 56) sorted by relevance

123

/external/tensorflow/tensorflow/compiler/mlir/xla/tests/hlo_to_lhlo_with_xla/
Dops.mlir4 // CHECK-SAME: %[[ARG0:.*]]: memref<2x2xf32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 0
8 // CHECK: lmhlo.abs
17 // CHECK-SAME: %[[ARG0:.*]]: memref<2x2xf32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 0
18 // CHECK-SAME: %[[ARG1:.*]]: memref<2x2xf32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 1
22 // CHECK: lmhlo.add
32 // CHECK-SAME: %[[ARG0:.*]]: memref<2x2xi32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 0
33 // CHECK-SAME: %[[ARG1:.*]]: memref<2x2xi32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 1
37 // CHECK: lmhlo.and
47 // CHECK-SAME: %[[ARG0:.*]]: memref<2x2xf32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 0
48 // CHECK-SAME: %[[ARG1:.*]]: memref<2x2xf32> {lmhlo.alloc = {{[0-9]+}} : index, lmhlo.params = 1
[all …]
Dpassthrough.mlir6 // CHECK-SAME: %[[ARG0:.*]]: memref<2x2xf32> {lmhlo.alloc = 1 : index, lmhlo.params = 0 : index},
7 // CHECK-SAME: %[[ARG1:.*]]: memref<16xi8> {lmhlo.alloc = 0 : index, lmhlo.liveout = true}
13 // CHECK: lmhlo.copy
Dhlo_text_to_lhlo_no_opt.hlotxt28 // CHECK: "lmhlo.scatter"
70 // CHECK: "lmhlo.select_and_scatter"(%{{.*}}, %{{.*}}, %[[GLOBAL_MEMREF]], %{{.*}})
93 // CHECK: "lmhlo.custom_call"
212 // CHECK: "lmhlo.all_reduce"([[INPUT]], {{.*}})
233 // CHECK: "lmhlo.all_reduce"
262 // CHECK: "lmhlo.all_reduce"([[INPUT]], {{.*}})
459 // CHECK: "lmhlo.infeed"
471 // CHECK: "lmhlo.outfeed"
485 // CHECK: "lmhlo.custom_call"
487 // CHECK: "lmhlo.outfeed"
[all …]
Dgpu_ops.mlir7 // CHECK-SAME: %[[ARG3:.*]]: memref<36xi8> {lmhlo.alloc = 0
9 // CHECK: "lmhlo.copy"(%[[ARG0]], %[[VIEW0]])
11 // CHECK: "lmhlo.scatter"(%[[VIEW0]], %[[ARG1]], %[[ARG2]], %[[VIEW1]])
/external/tensorflow/tensorflow/compiler/mlir/xla/
Dhlo_utils.cc287 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()
303 } else if (isa<mlir::mhlo::MulOp, mlir::lmhlo::MulOp>(op)) { in MhloToHloOpcode()
305 } else if (isa<mlir::mhlo::PowOp, mlir::lmhlo::PowOp>(op)) { in MhloToHloOpcode()
[all …]
/external/tensorflow/tensorflow/compiler/mlir/xla/transforms/
Dmhlo_to_lhlo_with_xla.h58 xla::StatusOr<lmhlo::SortOp> EmitSortOp(const xla::HloInstruction* instr);
59 xla::StatusOr<lmhlo::FusionOp> EmitFusionOp(const xla::HloInstruction* instr);
60 xla::StatusOr<lmhlo::ScatterOp> EmitScatterOp(
62 xla::StatusOr<lmhlo::SelectAndScatterOp> EmitSelectAndScatterOp(
75 xla::StatusOr<lmhlo::ReduceOp> EmitReduceOp(const xla::HloInstruction* instr);
79 xla::StatusOr<lmhlo::CompareOp> EmitCompareOp(
82 xla::StatusOr<lmhlo::InfeedOp> EmitInfeedOp(const xla::HloInstruction* instr);
83 xla::StatusOr<lmhlo::OutfeedOp> EmitOutfeedOp(
85 xla::StatusOr<lmhlo::MapOp> EmitMapOp(const xla::HloInstruction* instr);
87 xla::StatusOr<lmhlo::ReducePrecisionOp> EmitReducePrecisionOp(
[all …]
Dmhlo_to_lhlo_with_xla.cc160 mlir::lmhlo::LmhloDialect, mlir::lmhlo_gpu::LmhloGpuDialect>(); in getDependentDialects()
249 return CreateOpWithoutAttrs<lmhlo::AbsOp>(instr); in EmitOp()
251 return CreateOpWithoutAttrs<lmhlo::AddOp>(instr); in EmitOp()
259 return CreateOpWithoutAttrs<lmhlo::AndOp>(instr); in EmitOp()
261 return CreateOpWithoutAttrs<lmhlo::Atan2Op>(instr); in EmitOp()
265 return CreateOpWithoutAttrs<lmhlo::BitcastConvertOp>(instr); in EmitOp()
269 return CreateOpWithoutAttrs<lmhlo::CeilOp>(instr); in EmitOp()
271 return CreateOpWithoutAttrs<lmhlo::CbrtOp>(instr); in EmitOp()
273 return CreateOpWithoutAttrs<lmhlo::ClampOp>(instr); in EmitOp()
277 return CreateOpWithoutAttrs<lmhlo::ClzOp>(instr); in EmitOp()
[all …]
/external/tensorflow/tensorflow/compiler/mlir/hlo/tests/
Dlhlo_ops.mlir7 "lmhlo.all_reduce"(%input0, %input1, %input0, %input0) ({
22 "lmhlo.all_reduce"(%input0, %input1, %input0, %input1) ({
37 "lmhlo.all_gather"(%input0, %output)
48 "lmhlo.all_to_all"(%input0, %output)
59 "lmhlo.all_to_all"(%input0, %output)
70 "lmhlo.ceil"(%input, %result) : (memref<2x2xf32>, memref<2x2xf32>) -> ()
78 "lmhlo.ceil"(%input, %result) : (memref<2x2xi32>, memref<2x2xi32>) -> ()
86 "lmhlo.cosine"(%input, %result) : (memref<2x2xf32>, memref<2x2xf32>) -> ()
94 "lmhlo.cosine"(%input, %result) : (memref<2x2xcomplex<f32>>, memref<2x2xcomplex<f32>>) -> ()
102 "lmhlo.cosine"(%input, %result) : (memref<2x2xi32>, memref<2x2xi32>) -> ()
[all …]
Dhlo-legalize-to-lhlo.mlir10 …// CHECK: "lmhlo.exponential"(%{{.*}}, %{{.*}}) {some_attr_1 = "exp.1", some_attr_2 = dense<1> : t…
35 // CHECK-NEXT: "lmhlo.maximum"(%[[NEW_ARG0]], %[[NEW_ARG1]], %[[MAX_RESULT]])
37 // CHECK-NEXT: "lmhlo.add"(%[[NEW_ARG0]], %[[MAX_RESULT]], %[[ADD_RESULT]])
40 // CHECK-NEXT: "lmhlo.minimum"(%[[NEW_ARG0]], %[[NEW_ARG1]], %[[MIN_RESULT]])
42 //  CHECK-NEXT: "lmhlo.subtract"(%[[NEW_ARG1]], %[[MIN_RESULT]], %[[SUB_RESULT]])
45 // CHECK-NEXT: "lmhlo.multiply"(%[[ADD_RESULT]], %[[SUB_RESULT]], %[[MUL_RESULT]])
59 // CHECK-NEXT: "lmhlo.add"(%{{.*}}, %{{.*}}, %[[ADD_RESULT]])
63 // CHECK-NEXT: "lmhlo.multiply"(%[[ADD_RESULT]], %{{.*}}, %[[MUL_RESULT]])
75 // TODO-CHECK: "lmhlo.copy"(%{{.*}}, %{{.*}})
84 // CHECK: "lmhlo.exponential"(%{{.*}}, %{{.*}})
[all …]
Dlhlo-legalize-to-linalg.mlir7 "lmhlo.power"(%lhs, %rhs, %result)
22 "lmhlo.add"(%lhs, %rhs, %result)
37 "lmhlo.add"(%lhs, %rhs, %result)
51 "lmhlo.add"(%lhs, %rhs, %result)
66 "lmhlo.minimum"(%lhs, %rhs, %result)
84 "lmhlo.maximum"(%lhs, %rhs, %result)
99 "lmhlo.and"(%lhs, %rhs, %result)
112 "lmhlo.exponential"(%input, %result)
125 "lmhlo.log"(%input, %result) : (memref<2x2xf32>, memref<2x2xf32>) -> ()
137 "lmhlo.copy"(%in, %out) : (memref<2x4x8xf32>, memref<2x4x8xf32>) -> ()
[all …]
Dlhlo-legalize-to-affine.mlir20 "lmhlo.minimum"(%lhs, %rhs, %result) {name = "min.1"} :
30 "lmhlo.add"(%lhs, %rhs, %result) {name = "add.1"}
38 "lmhlo.add"(%lhs, %rhs, %result) {name = "add.1"}
48 "lmhlo.and"(%lhs, %rhs, %result) {name = "and.1"}
58 "lmhlo.divide"(%lhs, %rhs, %result) {name = "div.1"}
66 "lmhlo.divide"(%lhs, %rhs, %result) {name = "div.1"}
80 "lmhlo.maximum"(%lhs, %rhs, %result) {name = "max.1"}
90 "lmhlo.maximum"(%lhs, %rhs, %result) {name = "max.1"}
104 "lmhlo.minimum"(%lhs, %rhs, %result) {name = "min.1"}
114 "lmhlo.minimum"(%lhs, %rhs, %result) {name = "min.1"}
[all …]
Dlhlo-legalize-to-parallel-loops.mlir6 "lmhlo.reduce"(%arg, %init, %result) ( {
8 "lmhlo.add"(%lhs, %rhs, %res)
10 "lmhlo.terminator"() : () -> ()
38 // CHECK: "lmhlo.add"([[ELEM_BUF]], [[ACC_BUF]], [[ACC_OUT_BUF]])
52 "lmhlo.reduce"(%arg, %init, %result) ( {
54 "lmhlo.add"(%lhs, %rhs, %res)
56 "lmhlo.terminator"() : () -> ()
79 // CHECK: "lmhlo.add"([[ELEM_BUF]], [[ACC_BUF]], [[ACC_OUT_BUF]])
91 "lmhlo.reduce"(%arg, %init, %result) ( {
93 "lmhlo.add"(%lhs, %rhs, %res)
[all …]
Dlhlo-legalize-to-gpu.mlir6 "lmhlo.reduce"(%arg, %init, %result) ( {
8 "lmhlo.add"(%lhs, %rhs, %res)
10 "lmhlo.terminator"() : () -> ()
30 // CHECK: "lmhlo.add"(%[[LHS]], %[[RHS]], %[[LHS]]) : (memref<f32, {{.*}}>, memref<f32, {{.…
/external/tensorflow/tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/transforms/
Dmap_lmhlo_to_scalar_op.h35 namespace lmhlo {
44 struct LhloToScalarOp<lmhlo::AddOp> {
50 struct LhloToScalarOp<lmhlo::CompareOp> {
55 struct LhloToScalarOp<lmhlo::DivOp> {
60 struct LhloToScalarOp<lmhlo::MulOp> {
65 struct LhloToScalarOp<lmhlo::RemOp> {
70 struct LhloToScalarOp<lmhlo::SubOp> {
126 inline Value MapLhloOpToStdScalarOp<lmhlo::AbsOp>(Location loc,
147 auto neg_val = b->create<ScalarIOp<lmhlo::SubOp>>(loc, zero_intval, lhs);
153 inline Value MapLhloOpToStdScalarOp<lmhlo::AddOp>(Location loc,
[all …]
/external/tensorflow/tensorflow/compiler/mlir/hlo/lib/Dialect/mhlo/transforms/
Dlegalize_to_linalg.cc213 Value op_result = lmhlo::HloOpToStdScalarOp::map<OpTy>( in matchAndRewrite()
248 Value op_result = lmhlo::HloOpToStdScalarOp::map<LhloOp>( in matchAndRewrite()
262 struct ConvToLinalgConverter : public OpConversionPattern<lmhlo::ConvOp> {
264 using OpConversionPattern<lmhlo::ConvOp>::OpConversionPattern;
269 lmhlo::ConvOp op, ArrayRef<Value> args, in matchAndRewrite()
533 : public OpConversionPattern<lmhlo::BroadcastInDimOp> {
535 using OpConversionPattern<lmhlo::BroadcastInDimOp>::OpConversionPattern;
538 lmhlo::BroadcastInDimOp op, ArrayRef<Value> args, in matchAndRewrite()
540 lmhlo::BroadcastInDimOp::Adaptor operand_adaptor(args); in matchAndRewrite()
588 lmhlo::BroadcastInDimOp op, ArrayRef<Value> args, in InsertReshapeIfNecessary()
[all …]
Dlhlo_legalize_to_affine.cc28 namespace lmhlo { namespace
83 Value op_result = lmhlo::HloOpToStdScalarOp::map<DotOp>( in matchAndRewrite()
122 Value op_result = lmhlo::HloOpToStdScalarOp::map<LhloOpTy>( in matchAndRewrite()
141 BinaryOpConverter<lmhlo::AddOp>, in populateLHLOToAffineConversionPattern()
142 BinaryOpConverter<lmhlo::AndOp>, in populateLHLOToAffineConversionPattern()
143 BinaryOpConverter<lmhlo::DivOp>, in populateLHLOToAffineConversionPattern()
144 BinaryOpConverter<lmhlo::MaxOp>, in populateLHLOToAffineConversionPattern()
145 BinaryOpConverter<lmhlo::MinOp>, in populateLHLOToAffineConversionPattern()
146 BinaryOpConverter<lmhlo::MulOp>, in populateLHLOToAffineConversionPattern()
147 BinaryOpConverter<lmhlo::SubOp>, in populateLHLOToAffineConversionPattern()
Dlhlo_legalize_to_parallel_loops.cc28 namespace lmhlo { namespace
189 class ReduceOpConverter : public OpConversionPattern<lmhlo::ReduceOp> {
191 using OpConversionPattern<lmhlo::ReduceOp>::OpConversionPattern;
194 lmhlo::ReduceOp reduce_op, ArrayRef<Value> /*args*/, in matchAndRewrite()
228 lmhlo::ReduceOp reduce_op, ConversionPatternRewriter* rewriter) const { in CreateReduceOpInNestedParallelLoops()
360 : public OpConversionPattern<lmhlo::ReduceWindowOp> {
362 using OpConversionPattern<lmhlo::ReduceWindowOp>::OpConversionPattern;
365 lmhlo::ReduceWindowOp reduce_window_op, ArrayRef<Value> /*args*/, in matchAndRewrite()
384 lmhlo::ReduceWindowOp reduce_window_op, in CreateParallelLoopsToTraverseOutputAndWindow()
416 lmhlo::ReduceWindowOp reduce_window_op, scf::ParallelOp output_loop, in CreateReduceOpInNestedParallelLoops()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dir_emission_utils_test.cc36 "lmhlo.add" (%arg0, %arg1, %arg2) : (memref<f32>, memref<f32>, memref<f32>) -> () in TEST()
37 "lmhlo.terminator" () : () -> () in TEST()
53 "lmhlo.add" (%arg0, %arg1, %arg0) : (memref<f32>, memref<f32>, memref<f32>) -> () in TEST()
54 "lmhlo.terminator" () : () -> () in TEST()
70 "lmhlo.add" (%arg0, %arg1, %arg1) : (memref<f32>, memref<f32>, memref<f32>) -> () in TEST()
71 "lmhlo.terminator" () : () -> () in TEST()
Dnccl_all_reduce_thunk.cc40 mlir::lmhlo::AllReduceOp op) { in MatchReductionComputation()
89 static NcclAllReduceConfig GetNcclAllReduceConfig(mlir::lmhlo::AllReduceOp op, in GetNcclAllReduceConfig()
100 /*static*/ bool NcclAllReduceThunk::CanImplement(mlir::lmhlo::AllReduceOp op) { in CanImplement()
111 ThunkInfo thunk_info, mlir::lmhlo::AllReduceOp op, int64 replica_count, in NcclAllReduceThunk()
Dnccl_all_gather_thunk.cc33 static NcclAllGatherConfig GetNcclAllGatherConfig(mlir::lmhlo::AllGatherOp op, in GetNcclAllGatherConfig()
51 /*static*/ bool NcclAllGatherThunk::CanImplement(mlir::lmhlo::AllGatherOp op) { in CanImplement()
62 ThunkInfo thunk_info, mlir::lmhlo::AllGatherOp op, int64 replica_count, in NcclAllGatherThunk()
Dnccl_all_reduce_thunk.h38 NcclAllReduceThunk(ThunkInfo thunk_info, mlir::lmhlo::AllReduceOp op,
43 static bool CanImplement(mlir::lmhlo::AllReduceOp op);
Dnccl_all_to_all_thunk.h38 NcclAllToAllThunk(ThunkInfo thunk_info, mlir::lmhlo::AllToAllOp op,
43 static bool CanImplement(mlir::lmhlo::AllToAllOp op);
Dnccl_all_gather_thunk.h37 NcclAllGatherThunk(ThunkInfo thunk_info, mlir::lmhlo::AllGatherOp op,
45 static bool CanImplement(mlir::lmhlo::AllGatherOp op);
/external/tensorflow/tensorflow/compiler/xla/service/mlir_gpu/tests/passes/
Dfusion_op_remover.mlir6 // CHECK-NOT: lmhlo.fusion
7 "lmhlo.fusion"() ( {
16 // CHECK-NOT: lmhlo.terminator
17 "lmhlo.terminator"() : () -> ()
/external/tensorflow/tensorflow/compiler/mlir/hlo/tools/mlir-hlo-opt/
Dmlir-hlo-opt.cpp28 mlir::lmhlo::registerAllLmhloPasses(); in main()
34 registry.insert<mlir::lmhlo::LmhloDialect>(); in main()

123