Searched refs:all_reduce_start (Results 1 – 5 of 5) sorted by relevance
| /external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
| D | gpu_hlo_schedule_test.cc | 447 HloInstruction* all_reduce_start = in TEST_F() local
453 TF_CHECK_OK(add1->AddControlDependencyTo(all_reduce_start)); in TEST_F()
457 f32_2x2_, HloOpcode::kAllReduceDone, all_reduce_start)); in TEST_F()
472 EXPECT_TRUE(order->ExecutesBefore(add0, all_reduce_start)); in TEST_F()
474 EXPECT_TRUE(order->ExecutesBefore(add1, all_reduce_start)); in TEST_F()
476 EXPECT_TRUE(order->ExecutesBefore(all_reduce_start, add2)); in TEST_F()
477 EXPECT_TRUE(order->ExecutesBefore(all_reduce_start, add3)); in TEST_F()
478 EXPECT_TRUE(order->ExecutesBefore(all_reduce_start, add4)); in TEST_F()
|
| /external/tensorflow/tensorflow/compiler/xla/service/ |
| D | hlo_dataflow_analysis.cc | 953 HloInstruction* all_reduce_start) { in UpdateAllReduceStartValueSet() argument
954 CHECK_EQ(all_reduce_start->opcode(), HloOpcode::kAllReduceStart); in UpdateAllReduceStartValueSet()
957 for (int64_t i = 0; i < all_reduce_start->operand_count(); ++i) { in UpdateAllReduceStartValueSet()
959 GetValueSet(all_reduce_start->operand(i)); in UpdateAllReduceStartValueSet()
962 if (all_reduce_start->operand_count() > 1) { in UpdateAllReduceStartValueSet()
966 HloValueSet& value_set = GetValueSet(all_reduce_start, output_index); in UpdateAllReduceStartValueSet()
|
| D | hlo_dataflow_analysis.h | 240 bool UpdateAllReduceStartValueSet(HloInstruction* all_reduce_start);
|
| /external/tensorflow/tensorflow/compiler/mlir/hlo/include/mlir-hlo/Dialect/mhlo/IR/ |
| D | lhlo_gpu_ops.td | 239 LHLOGPU_Op<"all_reduce_start", [SameOperandsElementType, SameVariadicOperandSize]> {
|
| /external/tensorflow/tensorflow/compiler/mlir/xla/tests/hlo_to_lhlo_with_xla/ |
| D | hlo_text_to_lhlo_no_opt.hlotxt | 290 // CHECK: [[TOKEN:%.*]] = "lmhlo_gpu.all_reduce_start"([[INPUT]], [[OUTPUT]])
324 …// CHECK: [[TOKEN:%.*]] = "lmhlo_gpu.all_reduce_start"([[INPUT0]], [[INPUT1]], [[OUTPUT0]], [[OUT…
|