static */
| 190 | |
| 191 | /* static */ |
| 192 | StatusOr<std::unique_ptr<GpuHloSchedule>> GpuHloSchedule::Build( |
| 193 | const HloModule& module, const StreamAssignment& stream_assignment, |
| 194 | int64 pointer_size) { |
| 195 | std::unique_ptr<GpuHloSchedule> schedule(new GpuHloSchedule); |
| 196 | |
| 197 | // Initialize thunk_launch_order_, the total order of thunk launches. |
| 198 | HloComputation* entry_computation = module.entry_computation(); |
| 199 | if (stream_assignment.StreamCount() == 1) { |
| 200 | // All kernels are launched on a single stream, so there's no loss of |
| 201 | // concurrency by optimizing for minimal memory usage. |
| 202 | TF_ASSIGN_OR_RETURN( |
| 203 | HloInstructionSequence sequence, |
| 204 | ScheduleComputation( |
| 205 | entry_computation, [pointer_size](const BufferValue& buffer) { |
| 206 | return ShapeUtil::ByteSizeOf(buffer.shape(), pointer_size); |
| 207 | })); |
| 208 | schedule->thunk_launch_order_ = sequence.instructions(); |
| 209 | } else { |
| 210 | // BFS tends to increase concurrency, but also increases memory usage. |
| 211 | BFSLaunchOrder(entry_computation, &schedule->thunk_launch_order_); |
| 212 | } |
| 213 | |
| 214 | schedule->hlo_ordering_ = absl::make_unique<GpuHloOrdering>( |
| 215 | &module, stream_assignment, schedule->thunk_launch_order_); |
| 216 | |
| 217 | return std::move(schedule); |
| 218 | } |
| 219 | |
| 220 | } // namespace gpu |
| 221 | } // namespace xla |