Home
last modified time | relevance | path

Searched refs:fusion (Results 1 – 25 of 103) sorted by relevance

12345

/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dfusion_merger.cc88 double CalculateBytesReadByFusionInstruction(HloInstruction* fusion) { in CalculateBytesReadByFusionInstruction() argument
90 for (auto* fused_instruction : fusion->fused_instructions()) { in CalculateBytesReadByFusionInstruction()
101 double GetCurrentBytesTransferred(HloInstruction* fusion) { in GetCurrentBytesTransferred() argument
102 CHECK_EQ(HloOpcode::kFusion, fusion->opcode()); in GetCurrentBytesTransferred()
103 const double bytes_read = CalculateBytesReadByFusionInstruction(fusion); in GetCurrentBytesTransferred()
105 if (fusion->IsMultiOutputFusion()) { in GetCurrentBytesTransferred()
106 for (auto& operand : fusion->fused_expression_root()->operands()) { in GetCurrentBytesTransferred()
111 ShapeUtil::ByteSizeOf(fusion->fused_expression_root()->shape()); in GetCurrentBytesTransferred()
116 return bytes_read + bytes_written * (fusion->user_count() + 1); in GetCurrentBytesTransferred()
120 double GetMergedBytesTransferred(HloInstruction* fusion) { in GetMergedBytesTransferred() argument
[all …]
Dmulti_output_fusion_test.cc69 const HloInstruction* fusion = in TEST_F() local
71 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
72 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
157 const HloInstruction* fusion = in TEST_F() local
159 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
160 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
190 const HloInstruction* fusion = in TEST_F() local
192 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
193 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
252 const HloInstruction* fusion = in TEST_F() local
[all …]
Dcudnn_conv_runner.cc75 absl::optional<FusionParams> fusion; member
281 se::DeviceMemory<T> side_input(params.fusion->side_input_buf); in RunCudnnConvImpl()
284 if (params.fusion->side_input_scale != 0) { in RunCudnnConvImpl()
301 params.fusion->side_input_scale, bias_desc, in RunCudnnConvImpl()
302 DeviceMemory<T>(params.fusion->bias_buf), params.fusion->mode, in RunCudnnConvImpl()
371 params.fusion.emplace(); in GetCudnnConvParams()
372 auto& fusion = *params.fusion; in GetCudnnConvParams() local
377 fusion.mode = static_cast<se::dnn::ActivationMode>( in GetCudnnConvParams()
379 fusion.side_input_scale = backend_config.side_input_scale(); in GetCudnnConvParams()
383 params.fusion->bias_buf = operand_buffers[2]; in GetCudnnConvParams()
[all …]
Dir_emitter.h97 Status HandleFusion(HloInstruction* fusion) override;
188 HloInstruction* fusion) { in GetGeneratorForOperandIrArrays() argument
191 ir_arrays.reserve(fusion->operand_count()); in GetGeneratorForOperandIrArrays()
192 absl::c_transform(fusion->operands(), std::back_inserter(ir_arrays), in GetGeneratorForOperandIrArrays()
194 return GetIrArray(*operand, *fusion); in GetGeneratorForOperandIrArrays()
Dnvptx_compiler.cc355 HloPassFix<HloPassPipeline> fusion("fusion"); in OptimizeHloModule() local
358 fusion.AddPass<VariadicOpSplitter>(); in OptimizeHloModule()
361 fusion.AddInvariantChecker<HloVerifier>( in OptimizeHloModule()
365 fusion.AddPass<GpuInstructionFusion>(/*may_duplicate=*/false); in OptimizeHloModule()
366 fusion.AddPass<GpuInstructionFusion>(/*may_duplicate=*/true); in OptimizeHloModule()
367 fusion.AddPass<FusionMerger>(); in OptimizeHloModule()
368 fusion.AddPass<GpuMultiOutputFusion>(); in OptimizeHloModule()
369 fusion.AddPass<HloCSE>(/*is_layout_sensitive=*/true, in OptimizeHloModule()
371 fusion.AddPass<HloDCE>(); in OptimizeHloModule()
372 TF_RETURN_IF_ERROR(fusion.Run(hlo_module).status()); in OptimizeHloModule()
[all …]
/external/deqp-deps/SPIRV-Tools/test/opt/loop_optimizations/
Dfusion_legal.cpp171 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
173 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
174 EXPECT_TRUE(fusion.IsLegal()); in TEST_F()
176 fusion.Fuse(); in TEST_F()
307 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
309 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
310 EXPECT_TRUE(fusion.IsLegal()); in TEST_F()
312 fusion.Fuse(); in TEST_F()
440 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
442 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
[all …]
Dfusion_illegal.cpp144 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
146 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
147 EXPECT_FALSE(fusion.IsLegal()); in TEST_F()
263 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
265 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
266 EXPECT_FALSE(fusion.IsLegal()); in TEST_F()
430 LoopFusion fusion(context.get(), loop_0, loop_1); in TEST_F() local
431 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
435 LoopFusion fusion(context.get(), loop_0, loop_2); in TEST_F() local
436 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
[all …]
Dfusion_compatibility.cpp110 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
111 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
195 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
196 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
281 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
282 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
368 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
369 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
456 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
457 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
[all …]
/external/swiftshader/third_party/SPIRV-Tools/test/opt/loop_optimizations/
Dfusion_legal.cpp171 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
173 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
174 EXPECT_TRUE(fusion.IsLegal()); in TEST_F()
176 fusion.Fuse(); in TEST_F()
307 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
309 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
310 EXPECT_TRUE(fusion.IsLegal()); in TEST_F()
312 fusion.Fuse(); in TEST_F()
440 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
442 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
[all …]
Dfusion_illegal.cpp144 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
146 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
147 EXPECT_FALSE(fusion.IsLegal()); in TEST_F()
263 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
265 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
266 EXPECT_FALSE(fusion.IsLegal()); in TEST_F()
430 LoopFusion fusion(context.get(), loop_0, loop_1); in TEST_F() local
431 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
435 LoopFusion fusion(context.get(), loop_0, loop_2); in TEST_F() local
436 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
[all …]
Dfusion_compatibility.cpp110 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
111 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
195 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
196 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
281 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
282 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
368 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
369 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
456 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
457 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Ddynamic_update_slice_util.h44 HloInstruction* fusion, const BufferAssignment& assignment) { in CanEmitFusedDynamicUpdateSliceInPlace() argument
45 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion); in CanEmitFusedDynamicUpdateSliceInPlace()
46 HloInstruction* fused_root = fusion->fused_expression_root(); in CanEmitFusedDynamicUpdateSliceInPlace()
48 fusion->fusion_kind() != HloInstruction::FusionKind::kLoop) { in CanEmitFusedDynamicUpdateSliceInPlace()
60 auto* operand = fusion->operand(fusion_operand->parameter_number()); in CanEmitFusedDynamicUpdateSliceInPlace()
62 assignment.HasAllocationAt(fusion, {}) && in CanEmitFusedDynamicUpdateSliceInPlace()
63 assignment.SharesSliceAtIndex(fusion, {}, operand, index); in CanEmitFusedDynamicUpdateSliceInPlace()
79 HloInstruction* fusion,
87 HloInstruction* fusion,
Ddynamic_update_slice_util.cc138 HloInstruction* fusion, in EmitFusedDynamicUpdateSliceInPlaceImpl() argument
142 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion); in EmitFusedDynamicUpdateSliceInPlaceImpl()
144 << fusion->ToShortString(); in EmitFusedDynamicUpdateSliceInPlaceImpl()
146 auto* dynamic_update_slice = fusion->fused_expression_root(); in EmitFusedDynamicUpdateSliceInPlaceImpl()
166 LayoutUtil::CopyLayoutBetweenShapes(fusion->shape(), &update_shape)); in EmitFusedDynamicUpdateSliceInPlaceImpl()
182 fusion_output_array, launch_dimensions, IrName(fusion), b); in EmitFusedDynamicUpdateSliceInPlaceImpl()
186 HloInstruction* fusion, in EmitFusedDynamicUpdateSliceInPlace() argument
191 fusion, std::move(operand_arrays_generator), fusion_output_array, in EmitFusedDynamicUpdateSliceInPlace()
197 HloInstruction* fusion, in EmitParallelFusedDynamicUpdateSliceInPlace() argument
202 fusion, std::move(operand_arrays_generator), fusion_output_array, in EmitParallelFusedDynamicUpdateSliceInPlace()
/external/tensorflow/tensorflow/compiler/xla/service/
Dmulti_output_fusion.cc151 HloInstruction* fusion = instr1; in Update() local
154 fusion = instr2; in Update()
159 for (auto use : fusion->users()) { in Update()
166 FusionCandidate& fusion_node = candidates_[get_candidate_id(fusion)]; in Update()
170 UpdateReachability(fusion, fused, all_fusion_candidates_, in Update()
180 if (is_fused(instr) || is_connected(fusion, instr)) { in Update()
185 int64 profit = GetProfit(instr, fusion); in Update()
198 if (instr == fusion || is_fused(instr) || is_connected(fusion, instr)) { in Update()
204 int64 profit = GetProfit(instr, fusion); in Update()
212 worklist_.emplace(fusion, it.first, it.second); in Update()
Dhlo_instruction_test.cc635 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
638 EXPECT_THAT(fusion->operands(), ElementsAre(constant)); in TEST_F()
639 EXPECT_THAT(constant->users(), ElementsAre(fusion)); in TEST_F()
653 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
656 EXPECT_THAT(fusion->operands(), ElementsAre(constant1, constant2)); in TEST_F()
657 EXPECT_THAT(constant1->users(), ElementsAre(fusion)); in TEST_F()
658 EXPECT_THAT(constant2->users(), ElementsAre(fusion)); in TEST_F()
675 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
678 EXPECT_THAT(fusion->operands(), ElementsAre(constant)); in TEST_F()
679 EXPECT_THAT(constant->users(), ElementsAre(fusion)); in TEST_F()
[all …]
Dhlo_verifier.cc513 Status ShapeVerifier::HandleFusion(HloInstruction* fusion) { in HandleFusion() argument
514 auto& fused_parameters = fusion->fused_parameters(); in HandleFusion()
515 if (fused_parameters.size() != fusion->operand_count()) { in HandleFusion()
519 fused_parameters.size(), fusion->operand_count(), in HandleFusion()
520 fusion->ToString().c_str()); in HandleFusion()
524 if (!ShapesSame(fused_param->shape(), fusion->operand(param_no)->shape())) { in HandleFusion()
528 param_no, fusion->ToString().c_str()); in HandleFusion()
1179 Status CheckFusionInstruction(HloInstruction* fusion) { in CheckFusionInstruction() argument
1181 HloComputation* fused_computation = fusion->fused_instructions_computation(); in CheckFusionInstruction()
1182 if (fusion != fused_computation->FusionInstruction()) { in CheckFusionInstruction()
[all …]
Dinstruction_fusion_test.cc60 HloInstruction* fusion = in TEST_F() local
63 ASSERT_THAT(fusion, op::Fusion()) << module->ToString(); in TEST_F()
64 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
84 HloInstruction* fusion = in TEST_F() local
87 ASSERT_THAT(fusion, op::Fusion()) << module->ToString(); in TEST_F()
88 EXPECT_THAT(fusion->fused_expression_root(), op::Add(op::Abs(), op::Abs())) in TEST_F()
105 HloInstruction* fusion = in TEST_F() local
108 ASSERT_THAT(fusion, op::Fusion()) << module->ToString(); in TEST_F()
109 EXPECT_THAT(fusion->fused_expression_root(), op::Tuple(op::Tanh(), op::Abs())) in TEST_F()
Dbfloat16_propagation.cc38 HloInstruction* fusion) { in DetermineFusionComputationPrecision() argument
39 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion); in DetermineFusionComputationPrecision()
40 if (!bfloat16_support_->SupportsMixedPrecisions(*fusion)) { in DetermineFusionComputationPrecision()
49 auto root = fusion->fused_instructions_computation()->root_instruction(); in DetermineFusionComputationPrecision()
57 if (OutputTypeAfterChange(fusion, index) == BF16) { in DetermineFusionComputationPrecision()
61 << fusion->ToString(); in DetermineFusionComputationPrecision()
67 fusion->fused_instructions_computation()->MakeInstructionPostOrder(); in DetermineFusionComputationPrecision()
72 fusion->fused_instructions_computation()); in DetermineFusionComputationPrecision()
74 RevertIfFusionInternalBF16Changes(fusion); in DetermineFusionComputationPrecision()
78 HloInstruction* fusion) { in RevertIfFusionInternalBF16Changes() argument
[all …]
Dtuple_points_to_analysis_test.cc670 auto* fusion = module_->entry_computation()->root_instruction(); in Run() local
671 EXPECT_THAT(fusion, op::Fusion(tuple_param0)); in Run()
676 auto* fusion_param = GetFusionParameterForOperand(fusion, tuple_param0); in Run()
723 HloInstruction* GetFusionParameterForOperand(HloInstruction* fusion, in GetFusionParameterForOperand() argument
726 fusion->fused_instructions(), [&](const HloInstruction* fused) { in GetFusionParameterForOperand()
728 fusion->operand(fused->parameter_number()) == operand; in GetFusionParameterForOperand()
730 CHECK(it != fusion->fused_instructions().end()); in GetFusionParameterForOperand()
896 auto fusion = computation_->CreateFusionInstruction( in TEST_F() local
902 EXPECT_TRUE(points_to_analysis_->DoesNotUseOperandBuffer(tuple, {0}, fusion)); in TEST_F()
904 points_to_analysis_->DoesNotUseOperandBuffer(tuple, {1}, fusion)); in TEST_F()
[all …]
/external/swiftshader/third_party/SPIRV-Tools/source/opt/
Dloop_fusion_pass.cpp48 LoopFusion fusion(context(), &loop_0, &loop_1); in ProcessFunction() local
50 if (fusion.AreCompatible() && fusion.IsLegal()) { in ProcessFunction()
56 fusion.Fuse(); in ProcessFunction()
/external/deqp-deps/SPIRV-Tools/source/opt/
Dloop_fusion_pass.cpp48 LoopFusion fusion(context(), &loop_0, &loop_1); in ProcessFunction() local
50 if (fusion.AreCompatible() && fusion.IsLegal()) { in ProcessFunction()
56 fusion.Fuse(); in ProcessFunction()
/external/tensorflow/tensorflow/compiler/xla/service/cpu/tests/
Dcpu_fusion_test.cc70 CpuInstructionFusion fusion; in TEST_F() local
71 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie()); in TEST_F()
117 CpuInstructionFusion fusion; in TEST_F() local
118 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie()); in TEST_F()
192 CpuInstructionFusion fusion; in TEST_F() local
193 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie()); in TEST_F()
264 CpuInstructionFusion fusion; in TEST_F() local
265 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie()); in TEST_F()
320 CpuInstructionFusion fusion; in TEST_F() local
321 EXPECT_TRUE(fusion.Run(module.get()).ValueOrDie()); in TEST_F()
/external/boringssl/src/util/fipstools/delocate/testdata/ppc64le-Sample2/
Din.s49 addis 31,2,.LC0@toc@ha # gpr load fusion, type long
51 addis 19,2,.LC3@toc@ha # gpr load fusion, type long
55 addis 20,2,.LC6@toc@ha # gpr load fusion, type long
118 addis 31,2,.LC11@toc@ha # gpr load fusion, type long
132 addis 6,2,.LC12@toc@ha # gpr load fusion, type long
148 addis 6,2,.LC13@toc@ha # gpr load fusion, type long
/external/swiftshader/third_party/llvm-7.0/llvm/test/CodeGen/AArch64/
Dmacro-fusion-last.mir1 # RUN: llc -o - %s -mtriple=aarch64-- -mattr=+arith-bcc-fusion -run-pass postmisched | FileCheck %s…
2 # RUN: llc -o - %s -mtriple=aarch64-- -mattr=-arith-bcc-fusion -run-pass postmisched | FileCheck %s…
20 ; move the SUBSXri in between them. However doing so breaks macro fusion.
/external/llvm/test/CodeGen/NVPTX/
Dfp-contract.ll7 ;; If fusion is allowed, we try to form fma.rn at the PTX level, and emit
9 ;; is free to fuse with a multiply if it is able. If fusion is not allowed,
11 ;; for all adds to prevent ptxas from fusion the ops.

12345