MCPcopy Create free account
hub / github.com/DeepRec-AI/DeepRec / HandleFusion

Method HandleFusion

tensorflow/compiler/xla/service/gpu/fusion_merger.cc:181–293  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

179}
180
181Status FusionInstructionMerger::HandleFusion(HloInstruction* fusion) {
182 ++total_visited_;
183 // Skip 'fusion' instruction if there are no users into which we can merge.
184 if (fusion->users().empty()) {
185 VLOG(3) << "Not merging " << fusion->name() << ": Has no users.";
186 ++num_fail_no_users_;
187 return Status::OK();
188 }
189
190 // Skip 'fusion' instruction if it is not a loop fusion. Library fusion
191 // instructions match specific patterns, so they shouldn't be further fused.
192 // Input fusion instructions need to be rooted at a particular HLO (e.g.
193 // kReduce), so they shouldn't be further fused either.
194 if (!fusion->IsLoopFusion()) {
195 VLOG(3) << "Not merging " << fusion->name() << ": Is not loop fusion.";
196 ++num_fail_not_loop_fusion_;
197 return Status::OK();
198 }
199
200 // Skip 'fusion' instruction if we cannot merge into all of its users.
201 // Merging into all users enables the removal of 'fusion' from the
202 // computation.
203 if (!absl::c_all_of(fusion->users(), [&](const HloInstruction* user) {
204 return user->opcode() == HloOpcode::kFusion &&
205 IsProducerConsumerFusible(*fusion, *user);
206 })) {
207 VLOG(3) << "Not merging " << fusion->name()
208 << ": Some of its users are not loop/input fusion kernels.";
209 ++num_fail_merge_all_users_;
210 return Status::OK();
211 }
212
213 // Skip 'fusion' instruction if any of its fused instructions are expensive.
214 // This is done to avoid the duplication of expensive instructions, which
215 // would occur if 'fusion' were merged into multiple users.
216 //
217 // If 'fusion' has just one user, then an earlier fusion pass chose not to
218 // fuse this producer/consumer pair (likely because of expensive instruction
219 // re-use by the consumer), and so we honor that choice here as well.
220 if (absl::c_any_of(fusion->fused_instructions(),
221 [](const HloInstruction* instruction) {
222 return instruction->opcode() != HloOpcode::kParameter &&
223 GpuInstructionFusion::IsExpensive(*instruction);
224 })) {
225 VLOG(3) << "Not merging " << fusion->name()
226 << ": Contains one or more expensive instructions.";
227 ++num_fail_expensive_fused_instruction_;
228 return Status::OK();
229 }
230
231 // Skip 'fusion' instruction if merging it into all users would result in a
232 // net increase in bytes transferred (currently allowing the net bytes
233 // transferred to be exceeded up to ~10% in exchange for eliminating the
234 // overhead from a GPU kernel launch).
235 const double current_bytes_transferred = GetCurrentBytesTransferred(fusion);
236 const double merged_bytes_transferred = GetMergedBytesTransferred(fusion);
237 const double merged_to_current_bytes_ratio =
238 merged_bytes_transferred / std::max(1.0, current_bytes_transferred);

Callers

nothing calls this directly

Calls 14

FusionWouldBeTooLargeFunction · 0.85
IsLoopFusionMethod · 0.80
opcodeMethod · 0.80
user_countMethod · 0.80
nameMethod · 0.65
maxFunction · 0.50
StrAppendFunction · 0.50
emptyMethod · 0.45
fused_instructionsMethod · 0.45

Tested by

no test coverage detected