Home
last modified time | relevance | path

Searched refs:fusion_kind (Results 1 – 25 of 32) sorted by relevance

12

/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dmulti_output_fusion.cc55 instr->fusion_kind() == HloInstruction::FusionKind::kLoop) || in IsFusible()
94 instr1->fusion_kind() != instr2->fusion_kind()) || in LegalToFuse()
96 instr1->fusion_kind() == HloInstruction::FusionKind::kLoop)) { in LegalToFuse()
160 producer->fusion_kind() == HloInstruction::FusionKind::kLoop; in DoProducerConsumerMultiOutputFusion()
Dgpu_fusible.cc68 CHECK(instr.fusion_kind() == HloInstruction::FusionKind::kInput) in IsReduceInputFusion()
77 CHECK(instr.fusion_kind() == HloInstruction::FusionKind::kInput) in IsReduceInputFusion()
144 instr.fusion_kind() == HloInstruction::FusionKind::kInput && in IsInputFusibleScatter()
170 instr.fusion_kind() == HloInstruction::FusionKind::kLoop) || in IsLoopFusible()
Dinstruction_fusion.cc130 consumer->fusion_kind() == HloInstruction::FusionKind::kLoop && in ShouldFuseInexpensiveChecks()
170 consumer->fusion_kind() == HloInstruction::FusionKind::kOutput && in ShouldFuseInexpensiveChecks()
291 return consumer->fusion_kind(); in ChooseKind()
Dfusion_merger.cc192 if (fusion->fusion_kind() != HloInstruction::FusionKind::kLoop) { in HandleFusion()
209 (user->fusion_kind() == HloInstruction::FusionKind::kLoop || in HandleFusion()
Dinstruction_fusion_test.cc327 EXPECT_EQ(root->fusion_kind(), HloInstruction::FusionKind::kOutput); in TEST_F()
355 EXPECT_EQ(root->fusion_kind(), HloInstruction::FusionKind::kOutput); in TEST_F()
462 EXPECT_EQ(root->fusion_kind(), HloInstruction::FusionKind::kLoop); in TEST_F()
604 EXPECT_EQ(root->operand(0)->fusion_kind(), in TEST_F()
Dir_emission_utils.cc93 hlo.fusion_kind() == HloInstruction::FusionKind::kOutput && in ImplementedAsGemm()
Dgemm_thunk.cc276 CHECK_EQ(hlo_instruction.fusion_kind(), HloInstruction::FusionKind::kOutput); in GetDimensionNumbers()
Dir_emitter.cc725 CHECK_EQ(HloInstruction::FusionKind::kLoop, fusion->fusion_kind()); in HandleFusion()
Dir_emitter_unnested.cc577 if (HloInstruction::FusionKind::kInput == fusion->fusion_kind()) { in HandleFusion()
680 CHECK_EQ(fusion->fusion_kind(), HloInstruction::FusionKind::kLoop); in HandleFusion()
1792 CHECK_EQ(inst->fusion_kind(), HloInstruction::FusionKind::kOutput); in BuildGemmThunk()
3369 hlo->fusion_kind() == HloInstruction::FusionKind::kLoop) in CheckAndEmitHloWithTile021()
/external/tensorflow/tensorflow/compiler/xla/service/
Dreduce_precision_insertion.cc125 (instruction->fusion_kind() == HloInstruction::FusionKind::kLoop || in insert_on_inputs()
126 instruction->fusion_kind() == HloInstruction::FusionKind::kInput)) { in insert_on_inputs()
176 (instruction->fusion_kind() == HloInstruction::FusionKind::kLoop || in insert_on_outputs()
177 instruction->fusion_kind() == HloInstruction::FusionKind::kOutput)) { in insert_on_outputs()
Dinstruction_fusion.cc541 if (kind != fusion_instruction->fusion_kind()) { in AddFusionInstruction()
622 consumer->fusion_kind() != HloInstruction::FusionKind::kLoop && in ShouldFuse()
623 consumer->fusion_kind() != HloInstruction::FusionKind::kInput && in ShouldFuse()
624 consumer->fusion_kind() != HloInstruction::FusionKind::kOutput) { in ShouldFuse()
Dtuple_points_to_analysis.cc602 user->fusion_kind() == HloInstruction::FusionKind::kLoop) { in DoesNotUseOperandBuffer()
720 if (user->fusion_kind() == HloInstruction::FusionKind::kLoop || in CanShareOperandBufferWithUser()
721 user->fusion_kind() == HloInstruction::FusionKind::kInput) { in CanShareOperandBufferWithUser()
736 } else if (user->fusion_kind() == HloInstruction::FusionKind::kOutput && in CanShareOperandBufferWithUser()
Dhlo_dataflow_analysis.cc933 user->fusion_kind() == HloInstruction::FusionKind::kLoop) { in DoesNotUseOperandBuffer()
982 auto kind = fusion->fusion_kind(); in CanDoInPlaceDynamicUpdateSlice()
1059 if (user->fusion_kind() == HloInstruction::FusionKind::kLoop || in CanShareOperandBufferWithUser()
1060 user->fusion_kind() == HloInstruction::FusionKind::kInput) { in CanShareOperandBufferWithUser()
1064 if (user->fusion_kind() == HloInstruction::FusionKind::kOutput && in CanShareOperandBufferWithUser()
Dhlo_instructions.cc1134 FusionKind fusion_kind, in HloFusionInstruction() argument
1136 : HloInstruction(HloOpcode::kFusion, shape), fusion_kind_(fusion_kind) { in HloFusionInstruction()
1145 const Shape& shape, FusionKind fusion_kind, in HloFusionInstruction() argument
1148 : HloInstruction(HloOpcode::kFusion, shape), fusion_kind_(fusion_kind) { in HloFusionInstruction()
1158 switch (fusion_kind()) { in ToCategory()
1172 proto.set_fusion_kind(xla::ToString(fusion_kind())); in ToProto()
1540 return {StrCat("kind=", xla::ToString(fusion_kind()))}; in ExtraAttributesToStringImpl()
1547 return fusion_kind() == other.fusion_kind() && in IdenticalSlowPath()
1576 shape, fusion_kind(), new_operands, new_fused_computation); in CloneWithNewOperandsImpl()
Dhlo_instruction.cc322 TF_RET_CHECK(!proto.fusion_kind().empty()); in CreateFromProto()
323 TF_ASSIGN_OR_RETURN(FusionKind fusion_kind, in CreateFromProto()
324 StringToFusionKind(proto.fusion_kind())); in CreateFromProto()
336 CreateFusion(shape, fusion_kind, all_operands(), fused_computation); in CreateFromProto()
1201 const Shape& shape, FusionKind fusion_kind, HloInstruction* fused_root) { in CreateFusion() argument
1202 return absl::make_unique<HloFusionInstruction>(shape, fusion_kind, in CreateFusion()
1207 const Shape& shape, FusionKind fusion_kind, in CreateFusion() argument
1210 return absl::make_unique<HloFusionInstruction>(shape, fusion_kind, operands, in CreateFusion()
3355 HloInstruction::FusionKind HloInstruction::fusion_kind() const { in fusion_kind() function in xla::HloInstruction
3356 return Cast<HloFusionInstruction>(this)->fusion_kind(); in fusion_kind()
Dhlo_instruction.h749 const Shape& shape, FusionKind fusion_kind, HloInstruction* fused_root);
752 const Shape& shape, FusionKind fusion_kind,
1480 FusionKind fusion_kind() const;
Dhlo_computation.h243 HloInstruction::FusionKind fusion_kind);
Dhlo_instructions.h707 explicit HloFusionInstruction(const Shape& shape, FusionKind fusion_kind,
710 explicit HloFusionInstruction(const Shape& shape, FusionKind fusion_kind,
789 FusionKind fusion_kind() const { return fusion_kind_; } in fusion_kind() function
Dhlo_computation.cc616 HloInstruction::FusionKind fusion_kind) { in CreateFusionInstruction() argument
619 HloInstruction::CreateFusion(root->shape(), fusion_kind, root)); in CreateFusionInstruction()
Dhlo.proto66 string fusion_kind = 11; field
/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Ddynamic_update_slice_util.h48 fusion->fusion_kind() != HloInstruction::FusionKind::kLoop) { in CanEmitFusedDynamicUpdateSliceInPlace()
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Dcpu_instruction_fusion.cc157 consumer->fusion_kind() == HloInstruction::FusionKind::kLoop) { in ShouldFuse()
Dparallel_task_assignment.cc150 instruction->fusion_kind() != HloInstruction::FusionKind::kLoop) || in GetTargetParallelTaskCount()
Dcpu_instruction_fusion_test.cc264 HloInstruction::FusionKind fusion_kind = in RunFusionAndCheckOpcodesWereFused() argument
273 EXPECT_EQ(root->fusion_kind(), fusion_kind); in RunFusionAndCheckOpcodesWereFused()
/external/tensorflow/tensorflow/compiler/xla/tools/
Dinteractive_graphviz.cc316 std::cout << " Fusion kind: " << xla::ToString(instr->fusion_kind()) in DoInfoCommand()

12