| 179 | } |
| 180 | |
| 181 | Status 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); |
nothing calls this directly
no test coverage detected