1 /* Copyright 2017 The TensorFlow Authors. All Rights Reserved. 2 3 Licensed under the Apache License, Version 2.0 (the "License"); 4 you may not use this file except in compliance with the License. 5 You may obtain a copy of the License at 6 7 http://www.apache.org/licenses/LICENSE-2.0 8 9 Unless required by applicable law or agreed to in writing, software 10 distributed under the License is distributed on an "AS IS" BASIS, 11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 See the License for the specific language governing permissions and 13 limitations under the License. 14 ==============================================================================*/ 15 16 #ifndef TENSORFLOW_COMPILER_XLA_SERVICE_CPU_PARALLEL_TASK_ASSIGNMENT_H_ 17 #define TENSORFLOW_COMPILER_XLA_SERVICE_CPU_PARALLEL_TASK_ASSIGNMENT_H_ 18 19 #include "tensorflow/compiler/xla/service/cpu/target_machine_features.h" 20 #include "tensorflow/compiler/xla/service/hlo_cost_analysis.h" 21 #include "tensorflow/compiler/xla/service/hlo_module.h" 22 #include "tensorflow/compiler/xla/service/hlo_pass_interface.h" 23 24 namespace xla { 25 namespace cpu { 26 27 // Simple interface for different parallel cost model implementations. 28 class ParallelCostModel { 29 public: 30 virtual ~ParallelCostModel() = default; 31 virtual int64 GetParallelTaskCount(HloInstruction* instruction) = 0; 32 }; 33 34 // ParallelTaskAssignment computes parallel task counts for HLOs in 'module'. 35 class ParallelTaskAssignment { 36 public: 37 // 'max_parallelism': the maximum parallel task count per instruction. 38 // 'shape_size': shape size function used by HloCostAnalysis during parallel 39 // task assignment. 40 // 'module': the containing HloModule. 41 ParallelTaskAssignment(const int64 max_parallelism, 42 const HloCostAnalysis::ShapeSizeFunction& shape_size, 43 HloModule* module, 44 const TargetMachineFeatures* target_machine_features); ~ParallelTaskAssignment()45 ~ParallelTaskAssignment() {} 46 47 // Computes and returns the target parallel task count for 'instruction'. 48 int64 GetTargetParallelTaskCount(HloInstruction* instruction); 49 50 private: 51 std::unique_ptr<ParallelCostModel> cost_model_; 52 const TargetMachineFeatures& target_machine_features_; 53 }; 54 55 // ParallelTaskAssigner computes target parallel task counts for all HLOs 56 // in the module, then assigns parallel task counts to HLOs in the entry 57 // computation, or to HLOs in embedded computations invoked by (potentially 58 // nested) kWhile or kCall instructions. 59 // Each HLO which is assigned parallel task counts is outlined into its 60 // own embedded computation, which is compiled as a parallel compute function, 61 // and which is invoked from a kCall instruction that is lowered in codegen to 62 // a runtime parallel fork/join call. 63 class ParallelTaskAssigner : public HloModulePass { 64 public: 65 // 'max_parallelism': the maximum parallel task count per instruction. 66 // 'shape_size': shape size function used by HloCostAnalysis during parallel 67 // task assignment. ParallelTaskAssigner(const int64 max_parallelism,const HloCostAnalysis::ShapeSizeFunction & shape_size,const TargetMachineFeatures * target_machine_features)68 ParallelTaskAssigner(const int64 max_parallelism, 69 const HloCostAnalysis::ShapeSizeFunction& shape_size, 70 const TargetMachineFeatures* target_machine_features) 71 : max_parallelism_(max_parallelism), 72 shape_size_function_(shape_size), 73 target_machine_features_(*target_machine_features) {} ~ParallelTaskAssigner()74 ~ParallelTaskAssigner() override {} 75 name()76 absl::string_view name() const override { 77 return "cpu-parallel-task-assigner"; 78 } 79 80 // Run parallel task assigner on 'module'. 81 // Returns true if the computation was changed, false otherwise. 82 StatusOr<bool> Run(HloModule* module) override; 83 84 private: 85 using HloToParallelTasks = std::unordered_map<const HloInstruction*, int64>; 86 87 // Assigns target parallel tasks from 'hlo_to_parallel_tasks' to HLOs in 88 // 'module'. 89 // Returns true if the computation was changed, false otherwise. 90 bool AssignParallelTasks(HloModule* module, 91 const HloToParallelTasks& hlo_to_parallel_tasks); 92 bool AssignParallelTasksHelper( 93 HloModule* module, HloComputation* computation, 94 const HloToParallelTasks& hlo_to_parallel_tasks); 95 96 // Computes target parallel task counts (returned in 'parallel_task_counts') 97 // for parallelizable instructions in 'module'. 98 void ComputeTargetParallelTasks(HloModule* module, 99 HloToParallelTasks* hlo_to_parallel_tasks); 100 101 int64 max_parallelism_; 102 HloCostAnalysis::ShapeSizeFunction shape_size_function_; 103 const TargetMachineFeatures& target_machine_features_; 104 }; 105 106 } // namespace cpu 107 } // namespace xla 108 109 #endif // TENSORFLOW_COMPILER_XLA_SERVICE_CPU_PARALLEL_TASK_ASSIGNMENT_H_ 110