Home
last modified time | relevance | path

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

123456789

/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dfusion_merger.cc89 double CalculateBytesReadByFusionInstruction(HloInstruction* fusion) { in CalculateBytesReadByFusionInstruction() argument
91 for (auto* fused_instruction : fusion->fused_instructions()) { in CalculateBytesReadByFusionInstruction()
102 double GetCurrentBytesTransferred(HloInstruction* fusion) { in GetCurrentBytesTransferred() argument
103 CHECK_EQ(HloOpcode::kFusion, fusion->opcode()); in GetCurrentBytesTransferred()
104 const double bytes_read = CalculateBytesReadByFusionInstruction(fusion); in GetCurrentBytesTransferred()
106 if (fusion->IsMultiOutputFusion()) { in GetCurrentBytesTransferred()
107 for (auto& operand : fusion->fused_expression_root()->operands()) { in GetCurrentBytesTransferred()
112 ShapeUtil::ByteSizeOf(fusion->fused_expression_root()->shape()); in GetCurrentBytesTransferred()
117 return bytes_read + bytes_written * (fusion->user_count() + 1); in GetCurrentBytesTransferred()
121 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 …]
Dgpu_conv_runner.cc111 se::DeviceMemory<OutputType> side_input(params.fusion->side_input_buf); in RunGpuConvForwardActivation()
114 if (params.config.fusion->side_input_scale != 0) { in RunGpuConvForwardActivation()
132 params.config.fusion->side_input_scale, bias_desc, in RunGpuConvForwardActivation()
133 DeviceMemory<BiasType>(params.fusion->bias_buf), in RunGpuConvForwardActivation()
134 params.config.fusion->mode, params.config.output_descriptor, &output_buf, in RunGpuConvForwardActivation()
302 config.fusion.emplace(); in GetGpuConvConfig()
303 GpuConvConfig::FusionConfig& fusion = *config.fusion; in GetGpuConvConfig() local
308 fusion.mode = in GetGpuConvConfig()
310 fusion.side_input_scale = backend_config.side_input_scale(); in GetGpuConvConfig()
484 params.fusion.emplace(); in GetGpuConvParams()
[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/angle/third_party/vulkan-deps/spirv-tools/src/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/
Dfusion_node_indexing_evaluation_test.cc111 auto fusion = instruction_fusion.Fuse(slice1, sub); in TEST_F() local
112 EXPECT_EQ(instruction_fusion.EvaluateEmittedInstructions(slice2, fusion), 1); in TEST_F()
113 instruction_fusion.Fuse(slice2, fusion); in TEST_F()
151 auto fusion = instruction_fusion.Fuse(slice2_0, add2); in TEST_F() local
154 EXPECT_EQ(instruction_fusion.EvaluateEmittedInstructions(slice2_1, fusion), in TEST_F()
156 instruction_fusion.Fuse(slice2_1, fusion); in TEST_F()
157 HloInstruction* add1 = fusion->mutable_operand(0); in TEST_F()
160 EXPECT_EQ(instruction_fusion.EvaluateEmittedInstructions(add1, fusion), 2); in TEST_F()
161 instruction_fusion.Fuse(add1, fusion); in TEST_F()
162 HloInstruction* slice1_0 = fusion->mutable_operand(0); in TEST_F()
[all …]
Dmulti_output_fusion.cc201 HloInstruction* fusion = instr1; in GetNewFusibles() local
204 fusion = instr2; in GetNewFusibles()
208 FusionCandidate& fusion_node = candidates_[get_candidate_id(fusion)]; in GetNewFusibles()
217 if (is_fused(instr) || is_connected(fusion, instr)) { in GetNewFusibles()
231 if (instr == fusion || is_fused(instr) || is_connected(fusion, instr)) { in GetNewFusibles()
248 HloInstruction* fusion = instr1; in UpdateBeforeFuse() local
251 fusion = instr2; in UpdateBeforeFuse()
256 for (auto use : fusion->users()) { in UpdateBeforeFuse()
265 UpdateReachability(fusion, fused, all_fusion_candidates_, in UpdateBeforeFuse()
270 HloInstruction* fusion, in UpdateAfterFuse() argument
[all …]
Dhlo_cost_analysis_test.cc625 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
636 ASSERT_IS_OK(fusion->Accept(&fusion_analysis)); in TEST_F()
644 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 0), in TEST_F()
646 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 1), in TEST_F()
648 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 2), in TEST_F()
650 EXPECT_EQ(fusion_analysis.output_bytes_accessed(*fusion), in TEST_F()
687 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
691 ASSERT_IS_OK(fusion->Accept(&fusion_analysis)); in TEST_F()
695 EXPECT_EQ(fusion_analysis.bytes_accessed(*fusion), in TEST_F()
698 EXPECT_EQ(fusion_analysis.operand_bytes_accessed(*fusion, 0), in TEST_F()
[all …]
Dhlo_instruction_test.cc637 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
640 EXPECT_THAT(fusion->operands(), ElementsAre(constant)); in TEST_F()
641 EXPECT_THAT(constant->users(), ElementsAre(fusion)); in TEST_F()
655 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
658 EXPECT_THAT(fusion->operands(), ElementsAre(constant1, constant2)); in TEST_F()
659 EXPECT_THAT(constant1->users(), ElementsAre(fusion)); in TEST_F()
660 EXPECT_THAT(constant2->users(), ElementsAre(fusion)); in TEST_F()
677 auto* fusion = computation->CreateFusionInstruction( in TEST_F() local
680 EXPECT_THAT(fusion->operands(), ElementsAre(constant)); in TEST_F()
681 EXPECT_THAT(constant->users(), ElementsAre(fusion)); in TEST_F()
[all …]
/external/protobuf/
DREADME.md57 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fc…
58 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fj…
59 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fp…
60 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fmacos%2Fob…
61 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fc…
62 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fj…
63 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fr…
65 …fusion.corp.google.com/projectanalysis/current/KOKORO/prod:protobuf%2Fgithub%2Fmaster%2Fubuntu%2Fp…
/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/mlir/xla/tests/translate/
Dfusion.hlotxt6 // CHECK: %0 = "mhlo.fusion"(%[[ARG0:.*]], %[[ARG1:.*]]) ( {
9 // CHECK: %1 = "mhlo.fusion"(%[[ARG0:.*]], %[[ARG1:.*]]) ( {
31 %fusion.7 = f32[] fusion(f32[] %Arg_0.1, f32[] %Arg_1.2), kind=kLoop, calls=%region_0.3
32 %fusion.14 = (f32[], f32[]) fusion(f32[] %Arg_0.1, f32[] %Arg_1.2), kind=kLoop, calls=%region_1.8
33 %get-tuple-element.15 = f32[] get-tuple-element((f32[], f32[]) %fusion.14), index=0
34 ROOT %get-tuple-element.16 = f32[] get-tuple-element((f32[], f32[]) %fusion.14), index=1
/external/llvm-project/mlir/test/Transforms/
Dloop-fusion-dependence-check.mlir1 // RUN: mlir-opt -allow-unregistered-dialect %s -test-loop-fusion -test-loop-fusion-dependence-chec…
20 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 0 into loop nest 2 a…
29 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 2 into loop nest 0 a…
52 // Should fuse: no fusion preventing remarks should be emitted for this test.
85 // Should fuse: no fusion preventing remarks should be emitted for this test.
110 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 0 into loop nest 1 a…
119 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 1 into loop nest 0 a…
135 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 0 into loop nest 1 a…
144 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 1 into loop nest 0 a…
161 …// expected-remark@-1 {{block-level dependence preventing fusion of loop nest 0 into loop nest 1 a…
[all …]
/external/angle/third_party/vulkan-deps/spirv-tools/src/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/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/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/tensorflow/tensorflow/core/kernels/
Dfused_eigen_output_kernels.h137 static bool IsSupported(FusedComputationType fusion) { in IsSupported()
138 return fusion == FusedComputationType::kBiasAdd || in IsSupported()
139 fusion == FusedComputationType::kBiasAddWithRelu || in IsSupported()
140 fusion == FusedComputationType::kBiasAddWithRelu6 || in IsSupported()
141 fusion == FusedComputationType::kBiasAddWithElu || in IsSupported()
142 fusion == FusedComputationType::kBiasAddWithLeakyRelu; in IsSupported()
159 static bool IsSupported(FusedComputationType fusion) { in IsSupported()
160 return fusion == FusedComputationType::kFusedBatchNorm || in IsSupported()
161 fusion == FusedComputationType::kFusedBatchNormWithRelu || in IsSupported()
162 fusion == FusedComputationType::kFusedBatchNormWithRelu6 || in IsSupported()
[all …]

123456789