Home
last modified time | relevance | path

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

123456789

/external/angle/third_party/vulkan-deps/spirv-tools/src/test/opt/loop_optimizations/
Dfusion_legal.cpp168 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
170 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
171 EXPECT_TRUE(fusion.IsLegal()); in TEST_F()
173 fusion.Fuse(); in TEST_F()
304 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
306 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
307 EXPECT_TRUE(fusion.IsLegal()); in TEST_F()
309 fusion.Fuse(); in TEST_F()
437 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
439 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
[all …]
Dfusion_illegal.cpp141 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
143 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
144 EXPECT_FALSE(fusion.IsLegal()); in TEST_F()
260 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
262 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
263 EXPECT_FALSE(fusion.IsLegal()); in TEST_F()
427 LoopFusion fusion(context.get(), loop_0, loop_1); in TEST_F() local
428 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
432 LoopFusion fusion(context.get(), loop_0, loop_2); in TEST_F() local
433 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
[all …]
Dfusion_compatibility.cpp107 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
108 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
192 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
193 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
278 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
279 EXPECT_TRUE(fusion.AreCompatible()); in TEST_F()
365 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
366 EXPECT_FALSE(fusion.AreCompatible()); in TEST_F()
453 LoopFusion fusion(context.get(), loops[0], loops[1]); in TEST_F() local
454 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/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/tensorflow/tensorflow/compiler/xla/service/gpu/
Dfusion_merger.cc90 double CalculateBytesReadByFusionInstruction(HloInstruction* fusion) { in CalculateBytesReadByFusionInstruction() argument
92 for (auto* fused_instruction : fusion->fused_instructions()) { in CalculateBytesReadByFusionInstruction()
103 double GetCurrentBytesTransferred(HloInstruction* fusion) { in GetCurrentBytesTransferred() argument
104 CHECK_EQ(HloOpcode::kFusion, fusion->opcode()); in GetCurrentBytesTransferred()
105 const double bytes_read = CalculateBytesReadByFusionInstruction(fusion); in GetCurrentBytesTransferred()
107 if (fusion->IsMultiOutputFusion()) { in GetCurrentBytesTransferred()
108 for (auto& operand : fusion->fused_expression_root()->operands()) { in GetCurrentBytesTransferred()
113 ShapeUtil::ByteSizeOf(fusion->fused_expression_root()->shape()); in GetCurrentBytesTransferred()
118 return bytes_read + bytes_written * (fusion->user_count() + 1); in GetCurrentBytesTransferred()
122 double GetMergedBytesTransferred(HloInstruction* fusion) { in GetMergedBytesTransferred() argument
[all …]
Dmulti_output_fusion_test.cc82 const HloInstruction* fusion = in TEST_F() local
84 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
85 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
170 const HloInstruction* fusion = in TEST_F() local
172 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
173 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
203 const HloInstruction* fusion = in TEST_F() local
205 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
206 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
266 const HloInstruction* fusion = in TEST_F() local
[all …]
Dhorizontal_input_fusion_test.cc74 const HloInstruction* fusion = entry_root->operand(0)->operand(0); in TEST_F() local
75 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
76 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
242 const HloInstruction* fusion = entry_root->operand(0)->operand(0); in TEST_F() local
243 ASSERT_TRUE(fusion->IsMultiOutputFusion()); in TEST_F()
244 EXPECT_THAT(fusion->fused_expression_root(), in TEST_F()
Dfusion_bitcast_lift.cc82 if (auto* fusion = DynCast<HloFusionInstruction>(instr)) { in Run() local
84 if (fusion->fusion_kind() != HloInstruction::FusionKind::kInput || in Run()
86 fusion->fused_instructions_computation())) { in Run()
92 fusion->fused_instructions().begin(), in Run()
93 fusion->fused_instructions().end(), [](HloInstruction* inner) { in Run()
104 for (HloInstruction* fused_instr : fusion->fused_instructions()) { in Run()
129 fusion->Clone("bitcast"); in Run()
257 TF_RETURN_IF_ERROR(fusion->parent()->ReplaceWithNewInstruction( in Run()
258 fusion, std::move(cloned_fusion))); in Run()
/external/tensorflow/tensorflow/compiler/xla/service/
Dfusion_node_indexing_evaluation_test.cc108 auto fusion = in TEST_F() local
110 EXPECT_EQ(instruction_fusion.EvaluateEmittedInstructions(slice2, fusion), 1); in TEST_F()
111 instruction_fusion.Fuse(slice2, fusion, module->entry_computation()); in TEST_F()
149 auto fusion = in TEST_F() local
153 EXPECT_EQ(instruction_fusion.EvaluateEmittedInstructions(slice2_1, fusion), in TEST_F()
155 instruction_fusion.Fuse(slice2_1, fusion, module->entry_computation()); in TEST_F()
156 HloInstruction* add1 = fusion->mutable_operand(0); in TEST_F()
159 EXPECT_EQ(instruction_fusion.EvaluateEmittedInstructions(add1, fusion), 2); in TEST_F()
160 instruction_fusion.Fuse(add1, fusion, module->entry_computation()); in TEST_F()
161 HloInstruction* slice1_0 = fusion->mutable_operand(0); in TEST_F()
[all …]
Dmulti_output_fusion.cc199 HloInstruction* fusion = instr1; in GetNewFusibles() local
202 fusion = instr2; in GetNewFusibles()
206 FusionCandidate& fusion_node = candidates_[get_candidate_id(fusion)]; in GetNewFusibles()
215 if (is_fused(instr) || is_connected(fusion, instr)) { in GetNewFusibles()
229 if (instr == fusion || is_fused(instr) || is_connected(fusion, instr)) { in GetNewFusibles()
246 HloInstruction* fusion = instr1; in UpdateBeforeFuse() local
249 fusion = instr2; in UpdateBeforeFuse()
254 for (auto use : fusion->users()) { in UpdateBeforeFuse()
263 UpdateReachability(fusion, fused, all_fusion_candidates_, in UpdateBeforeFuse()
268 HloInstruction* fusion, in UpdateAfterFuse() argument
[all …]
Dhlo_cost_analysis_test.cc660 HloInstruction* fusion = module->entry_computation()->root_instruction(); in TEST_F() local
661 ASSERT_IS_OK(fusion->Accept(&fusion_analysis)); in TEST_F()
678 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 0), 0); in TEST_F()
680 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 0), in TEST_F()
682 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 1), in TEST_F()
684 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 2), in TEST_F()
686 EXPECT_EQ(fusion_analysis.output_bytes_accessed(*fusion), in TEST_F()
728 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
739 ASSERT_IS_OK(fusion->Accept(&fusion_analysis)); in TEST_F()
747 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 0), in TEST_F()
[all …]
Dhlo_instruction_test.cc638 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
641 EXPECT_THAT(fusion->operands(), ElementsAre(constant)); in TEST_F()
642 EXPECT_THAT(constant->users(), ElementsAre(fusion)); in TEST_F()
656 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
659 EXPECT_THAT(fusion->operands(), ElementsAre(constant1, constant2)); in TEST_F()
660 EXPECT_THAT(constant1->users(), ElementsAre(fusion)); in TEST_F()
661 EXPECT_THAT(constant2->users(), ElementsAre(fusion)); in TEST_F()
678 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
681 EXPECT_THAT(fusion->operands(), ElementsAre(constant)); in TEST_F()
682 EXPECT_THAT(constant->users(), ElementsAre(fusion)); in TEST_F()
[all …]
/external/tensorflow/tensorflow/compiler/mlir/xla/tests/translate/
Dfusion.hlotxt6 // CHECK: %[[F0:.+]] = "mhlo.fusion"(%[[ARG0:.*]], %[[ARG1:.*]]) ({
9 // CHECK: %[[F1:.+]]:2 = "mhlo.fusion"(%[[ARG0:.*]]) ({
12 // CHECK: %[[F2:.+]]:2 = "mhlo.fusion"(%[[ARG0:.*]], %[[ARG1:.*]]) ({
15 // CHECK: %[[F3:.+]]:2 = "mhlo.fusion"(%[[ARG0:.*]], %[[ARG1:.*]]) ({
58 %fusion.0 = f32[] fusion(f32[] %Arg_0.1, f32[] %Arg_1.2), kind=kLoop, calls=%region_0
60 %fusion.1 = (f32[], f32[]) fusion(f32[] %Arg_0.1, f32[] %Arg_1.2), kind=kLoop, calls=%region_1
63 %fusion.2 = (f32[], f32[]) fusion((f32[]) %tuple.0), kind=kLoop, calls=%region_2
66 %fusion.3 = (f32[], f32[]) fusion((f32[], (f32[])) %tuple.1), kind=kLoop, calls=%region_3
68 %get-tuple-element.15 = f32[] get-tuple-element((f32[], f32[]) %fusion.1), index=0
69 ROOT %get-tuple-element.16 = f32[] get-tuple-element((f32[], f32[]) %fusion.1), index=1
Dfusion.mlir10 // CHECK: %[[FUSION0:.*]] = f32[] fusion(f32[] %[[PARAM0]], f32[] %[[PARAM1]]), kind=kLoop, calls…
11 // CHECK: %[[FUSION1:.*]] = (f32[], f32[]) fusion(f32[] %[[PARAM0]], f32[] %[[PARAM1]]), kind=kLo…
14 // CHECK: %[[FUSION2:.*]] = (f32[], f32[]) fusion(f32[] %[[PARAM0]]), kind=kLoop, calls=%[[REGION…
19 %result = "mhlo.fusion"(%arg0, %arg1) ({
24 %result0, %result1 = "mhlo.fusion"(%arg0, %arg1) ({
30 %result2:2 = "mhlo.fusion"(%arg0) ( {
/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Ddynamic_update_slice_util.cc66 bool CanEmitFusedDynamicUpdateSliceInPlace(HloInstruction* fusion, in CanEmitFusedDynamicUpdateSliceInPlace() argument
68 CHECK_EQ(fusion->opcode(), HloOpcode::kFusion); in CanEmitFusedDynamicUpdateSliceInPlace()
69 if (!MayBeImplementedAsInPlaceDynamicUpdateSlice(fusion)) { in CanEmitFusedDynamicUpdateSliceInPlace()
75 HloInstruction* fused_root = fusion->fused_expression_root(); in CanEmitFusedDynamicUpdateSliceInPlace()
83 auto* operand = fusion->operand(fusion_operand->parameter_number()); in CanEmitFusedDynamicUpdateSliceInPlace()
85 assignment.HasAllocationAt(fusion, {}) && in CanEmitFusedDynamicUpdateSliceInPlace()
86 assignment.SharesSliceAtIndex(fusion, {}, operand, index); in CanEmitFusedDynamicUpdateSliceInPlace()
192 const HloComputation* fusion, const IrArray& fusion_output_array, in EmitFusedDynamicUpdateSliceInPlaceImpl() argument
195 VLOG(2) << "EmitFusedDynamicUpdateSliceInPlace for " << fusion->ToString(); in EmitFusedDynamicUpdateSliceInPlaceImpl()
197 auto* dynamic_update_slice = fusion->root_instruction(); in EmitFusedDynamicUpdateSliceInPlaceImpl()
[all …]
/external/tensorflow/tensorflow/compiler/xla/mlir_hlo/tests/
Dgpu_fusion_rewrite.mlir1 // RUN: mlir-hlo-opt --split-input-file --gpu-fusion-rewrite %s | FileCheck %s
13 "lmhlo.fusion"() ({
32 "lmhlo.fusion"() ({
50 "lmhlo.fusion"() ({
56 "lmhlo.fusion"() ({
73 "lmhlo.fusion"() ({
89 "lmhlo.fusion"() ({
107 "lmhlo.fusion"() ({
125 "lmhlo.fusion"() ({
143 "lmhlo.fusion"() ({
/external/tensorflow/tensorflow/compiler/mlir/xla/tests/hlo_to_lhlo_with_xla/
Dops.mlir8 // CHECK: lmhlo.fusion
22 // CHECK: lmhlo.fusion
37 // CHECK: lmhlo.fusion
52 // CHECK: lmhlo.fusion
66 // CHECK: lmhlo.fusion
79 // CHECK: lmhlo.fusion
92 // CHECK: lmhlo.fusion
107 // CHECK: lmhlo.fusion
121 // CHECK: lmhlo.fusion
135 // CHECK: lmhlo.fusion
[all …]
/external/angle/third_party/vulkan-deps/spirv-tools/src/source/opt/
Dloop_fusion_pass.cpp47 LoopFusion fusion(context(), &loop_0, &loop_1); in ProcessFunction() local
49 if (fusion.AreCompatible() && fusion.IsLegal()) { in ProcessFunction()
55 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/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()

123456789