/external/tensorflow/tensorflow/compiler/xla/service/ |
D | hlo_creation_utils.h | 32 StatusOr<HloInstruction*> MakeUnaryHlo(HloOpcode opcode, 33 HloInstruction* operand); 37 StatusOr<HloInstruction*> MakeBinaryHlo(HloOpcode opcode, HloInstruction* lhs, 38 HloInstruction* rhs); 42 StatusOr<HloInstruction*> MakeCompareHlo(Comparison::Direction direction, 43 HloInstruction* lhs, 44 HloInstruction* rhs); 49 StatusOr<HloInstruction*> MakePadHlo(HloInstruction* operand, 50 HloInstruction* padding_value, 55 StatusOr<HloInstruction*> MakeSliceHlo(HloInstruction* operand, [all …]
|
D | bfloat16_propagation_test.cc | 39 bool SupportsBF16Operand(const HloInstruction& hlo, in SupportsBF16Operand() 44 bool SupportsBF16Output(const HloInstruction& hlo) const override { in SupportsBF16Output() 48 bool SupportsMixedPrecisions(const HloInstruction& hlo) const override { in SupportsMixedPrecisions() 52 bool EffectiveOperandPrecisionIsBF16(const HloInstruction& hlo, in EffectiveOperandPrecisionIsBF16() 76 bool OutputsBF16(const HloInstruction* inst) { in OutputsBF16() 85 std::unique_ptr<HloInstruction> CreateDot(const Shape& shape, in CreateDot() 86 HloInstruction* lhs, in CreateDot() 87 HloInstruction* rhs) { in CreateDot() 91 return HloInstruction::CreateDot(shape, lhs, rhs, dot_dnums, in CreateDot() 102 HloInstruction* a = in TEST_F() [all …]
|
D | hlo_cost_analysis.h | 53 Status HandleElementwiseUnary(const HloInstruction* hlo) override; 54 Status HandleElementwiseBinary(const HloInstruction* hlo) override; 55 Status HandleConstant(const HloInstruction* constant) override; 56 Status HandleIota(const HloInstruction* iota) override; 58 const HloInstruction* get_tuple_element) override; 59 Status HandleSelect(const HloInstruction* hlo) override; 60 Status HandleTupleSelect(const HloInstruction* hlo) override; 61 Status HandleCompare(const HloInstruction* compare) override; 62 Status HandleClamp(const HloInstruction* clamp) override; 63 Status HandleReducePrecision(const HloInstruction* hlo) override; [all …]
|
D | hlo_instruction.h | 270 std::function<string(const HloInstruction*, const string&, int, bool)>; 277 using HloInstructionPredicate = std::function<bool(const HloInstruction*)>; 322 string format_instruction(const HloInstruction* instr, in format_instruction() 327 bool print_instruction(const HloInstruction* instr) const { in print_instruction() 354 FormatInstructionFunc format_instruction_ = [](const HloInstruction* instr, 360 HloInstructionPredicate print_instruction_ = [](const HloInstruction* instr) { 410 class HloInstruction { 483 virtual ~HloInstruction() { DetachFromOperandsAndUsers(); } in ~HloInstruction() 497 static StatusOr<std::unique_ptr<HloInstruction>> CreateFromProto( 499 const absl::flat_hash_map<int64, HloInstruction*>& instruction_map, [all …]
|
D | tuple_simplifier_test.cc | 61 HloInstruction* param0 = builder.AddInstruction( in TEST_F() 62 HloInstruction::CreateParameter(0, scalar_shape_, "param0")); in TEST_F() 63 HloInstruction* param1 = builder.AddInstruction( in TEST_F() 64 HloInstruction::CreateParameter(1, scalar_shape_, "param1")); in TEST_F() 65 HloInstruction* param2 = builder.AddInstruction( in TEST_F() 66 HloInstruction::CreateParameter(2, scalar_shape_, "param2")); in TEST_F() 67 builder.AddInstruction(HloInstruction::CreateTuple({param0, param1, param2})); in TEST_F() 77 HloInstruction* param = builder.AddInstruction( in TEST_F() 78 HloInstruction::CreateParameter(0, tuple_shape_, "param")); in TEST_F() 80 HloInstruction::CreateGetTupleElement(scalar_shape_, param, 1)); in TEST_F() [all …]
|
D | hlo_computation.h | 71 HloInstruction* fusion_instruction = nullptr) 81 HloInstruction* root_instruction = nullptr); 83 HloInstruction* AddInstruction( in AddInstruction() 84 std::unique_ptr<HloInstruction> instruction) { in AddInstruction() 91 const std::function<Status(const HloInstruction*)>& func) const { in ForEachInstruction() 100 HloInstruction* last_added_instruction_; 101 HloInstruction* fusion_instruction_; 102 std::vector<std::unique_ptr<HloInstruction>> instructions_; 112 HloInstruction* AddInstruction( in AddInstruction() 113 std::unique_ptr<HloInstruction> instruction) { in AddInstruction() [all …]
|
D | memory_space_assignment_test.cc | 96 HloInstruction* instruction = value.instruction(); in AssignMemorySpace() 110 for (const HloInstruction* parameter : in AssignMemorySpace() 180 for (const HloInstruction* parameter : in CheckParametersInDefaultMemory() 206 for (HloInstruction* instruction : module.schedule() in CountMaximumOutstandingAsyncCopies() 236 const HloInstruction* instruction, in GetAlternateMemoryOffset() 256 HloInstruction* p0 = in CreateEvictAndPrefetchModule() 257 builder.AddInstruction(HloInstruction::CreateParameter(0, shape, "p0")); in CreateEvictAndPrefetchModule() 258 HloInstruction* p1 = in CreateEvictAndPrefetchModule() 259 builder.AddInstruction(HloInstruction::CreateParameter(1, shape, "p1")); in CreateEvictAndPrefetchModule() 260 HloInstruction* tanh = builder.AddInstruction( in CreateEvictAndPrefetchModule() [all …]
|
D | hlo_verifier.h | 42 Status Preprocess(HloInstruction* hlo) override; 44 Status HandleElementwiseUnary(HloInstruction* hlo) override; 45 Status HandleElementwiseBinary(HloInstruction* hlo) override; 46 Status HandleClamp(HloInstruction* clamp) override; 47 Status HandleSelect(HloInstruction* select) override; 48 Status HandleTupleSelect(HloInstruction* tuple_select) override; 49 Status HandleConcatenate(HloInstruction* concatenate) override; 50 Status HandleIota(HloInstruction* hlo) override; 51 Status HandleConvert(HloInstruction* convert) override; 52 Status HandleBitcastConvert(HloInstruction* convert) override; [all …]
|
D | bfloat16_conversion_folding_test.cc | 35 bool SupportsBF16Operand(const HloInstruction& hlo, in SupportsBF16Operand() 47 bool SupportsBF16Output(const HloInstruction& hlo) const override { in SupportsBF16Output() 58 bool SupportsMixedPrecisions(const HloInstruction& hlo) const override { in SupportsMixedPrecisions() 88 HloInstruction* a = builder.AddInstruction( in TEST_F() 89 HloInstruction::CreateParameter(0, f32_shape, "a")); in TEST_F() 90 HloInstruction* b = builder.AddInstruction( in TEST_F() 91 HloInstruction::CreateParameter(1, f32_shape, "b")); in TEST_F() 92 HloInstruction* c = builder.AddInstruction( in TEST_F() 93 HloInstruction::CreateParameter(2, f32_shape, "c")); in TEST_F() 95 HloInstruction* add0 = builder.AddInstruction( in TEST_F() [all …]
|
D | instruction_fusion.h | 42 std::function<bool(const HloInstruction& instruction)> is_expensive, 59 static bool IsExpensive(const HloInstruction& instruction); 79 virtual bool ShouldFuse(HloInstruction* consumer, int64 operand_index); 84 virtual bool ShouldFuseIntoMultiOutput(HloInstruction* consumer, in ShouldFuseIntoMultiOutput() 91 virtual HloInstruction::FusionKind ChooseKind(const HloInstruction* producer, 92 const HloInstruction* consumer); 97 virtual HloInstruction* FuseInstruction(HloInstruction* fusion_instruction, 98 HloInstruction* producer); 101 virtual HloInstruction* Fuse(HloInstruction* producer, 102 HloInstruction* consumer); [all …]
|
D | while_loop_invariant_code_motion_test.cc | 38 HloInstruction** while_instruction) { in FindOnlyWhileInstruction() 54 HloInstruction::CreateParameter(0, param_shape, "param")); in MakeAlwaysTrueComputation() 56 HloInstruction::CreateConstant(LiteralUtil::CreateR0<bool>(true))); in MakeAlwaysTrueComputation() 68 HloInstruction* param = builder.AddInstruction( in TEST_F() 69 HloInstruction::CreateParameter(0, while_shape, "param")); in TEST_F() 70 HloInstruction* gte_0 = builder.AddInstruction( in TEST_F() 71 HloInstruction::CreateGetTupleElement(scalar_s32, param, 0)); in TEST_F() 72 HloInstruction* gte_1 = builder.AddInstruction( in TEST_F() 73 HloInstruction::CreateGetTupleElement(scalar_s32, param, 1)); in TEST_F() 74 HloInstruction* add_result = in TEST_F() [all …]
|
D | multi_output_fusion.h | 65 virtual bool ShapesCompatibleForFusion(HloInstruction* instr1, 66 HloInstruction* instr2) = 0; 69 virtual bool IsFusible(HloInstruction* instr) = 0; 74 virtual int64 GetProfit(HloInstruction* instr1, HloInstruction* instr2) = 0; 77 virtual bool IsProfitableOperand(HloInstruction* instr); 80 virtual bool LegalToFuse(HloInstruction* instr1, HloInstruction* instr2); 84 bool LegalToFuseMainConstraints(HloInstruction* instr1, 85 HloInstruction* instr2); 89 virtual HloInstruction* Fuse(HloInstruction* instr1, HloInstruction* instr2); 102 HloInstruction* instr1, HloInstruction* instr2, [all …]
|
D | hlo_dataflow_analysis.h | 59 const HloInstruction* instr, const HloInstruction* operand, 85 static bool AreTransitiveUsesElementwiseOrTuple(const HloInstruction* inst); 89 bool ValueIsDefinedAt(const HloInstruction* instruction, 96 const HloValue& GetValueDefinedAt(const HloInstruction* instruction, 98 HloValue& GetValueDefinedAt(const HloInstruction* instruction, 103 const HloInstruction* instruction) const; 105 const HloInstruction* instruction); 109 HloValueSet GetFlattenedValueSet(const HloInstruction* instruction) const; 113 const HloValueSet& GetValueSet(const HloInstruction* instruction, 117 HloValueSet& GetValueSet(const HloInstruction* instruction, [all …]
|
D | hlo_instructions.h | 28 class HloBatchNormInstruction : public HloInstruction { 43 HloInstruction* operand, 44 HloInstruction* scale, float epsilon, 51 const HloInstruction& other, 64 HloInstruction* operand, 65 HloInstruction* scale, 66 HloInstruction* offset, 71 std::unique_ptr<HloInstruction> CloneWithNewOperandsImpl( 72 const Shape& shape, absl::Span<HloInstruction* const> new_operands, 79 const Shape& shape, HloInstruction* operand, HloInstruction* scale, [all …]
|
D | hlo_instruction_test.cc | 55 Status DefaultAction(HloInstruction* hlo_instruction) override { in DefaultAction() 60 Status HandleParameter(HloInstruction* parameter) override { in HandleParameter() 66 Status HandleConstant(HloInstruction* constant) override { in HandleConstant() 72 Status HandleAdd(HloInstruction* add) override { in HandleAdd() 82 Status HandleNegate(HloInstruction* negate) override { in HandleNegate() 90 Status HandleMap(HloInstruction* map) override { in HandleMap() 92 for (HloInstruction* arg : map->operands()) { in HandleMap() 99 Status HandleReduce(HloInstruction* reduce) override { in HandleReduce() 109 int64 NumOperands(const HloInstruction* node) { in NumOperands() 115 int64 NumUsers(const HloInstruction* node) { in NumUsers() [all …]
|
D | copy_insertion_test.cc | 90 HloInstruction* x = builder.AddInstruction( in TEST_F() 91 HloInstruction::CreateParameter(0, ShapeUtil::MakeShape(F32, {}), "x")); in TEST_F() 92 HloInstruction* tuple = in TEST_F() 93 builder.AddInstruction(HloInstruction::CreateTuple({x})); in TEST_F() 110 HloInstruction* constant = builder.AddInstruction( in TEST_F() 111 HloInstruction::CreateConstant(LiteralUtil::CreateR0<float>(1.0))); in TEST_F() 112 HloInstruction* tuple = in TEST_F() 113 builder.AddInstruction(HloInstruction::CreateTuple({constant})); in TEST_F() 133 HloInstruction* constant = in TEST_F() 134 builder.AddInstruction(HloInstruction::CreateConstant( in TEST_F() [all …]
|
D | hlo_creation_utils.cc | 37 StatusOr<HloInstruction*> MakeUnaryHlo(HloOpcode opcode, in MakeUnaryHlo() 38 HloInstruction* operand) { in MakeUnaryHlo() 43 HloInstruction::CreateUnary(unary_op_shape, opcode, operand)); in MakeUnaryHlo() 46 StatusOr<HloInstruction*> MakeBinaryHlo(HloOpcode opcode, HloInstruction* lhs, in MakeBinaryHlo() 47 HloInstruction* rhs) { in MakeBinaryHlo() 53 HloInstruction::CreateBinary(binary_op_shape, opcode, lhs, rhs)); in MakeBinaryHlo() 56 StatusOr<HloInstruction*> MakeCompareHlo(ComparisonDirection direction, in MakeCompareHlo() 57 HloInstruction* lhs, in MakeCompareHlo() 58 HloInstruction* rhs) { in MakeCompareHlo() 65 HloInstruction::CreateCompare(binary_op_shape, lhs, rhs, direction)); in MakeCompareHlo() [all …]
|
D | dynamic_dimension_inference.cc | 52 HloInstruction::CreateParameter(0, wide_shape, "wide_param")); in WidenComputation() 56 HloInstruction* wide_parameter = wide_comp->parameter_instruction(0); in WidenComputation() 57 HloInstruction* truncated_parameter = TupleUtil::ExtractPrefix( in WidenComputation() 59 HloInstruction* call_narrow_comp = wide_comp->AddInstruction( in WidenComputation() 60 HloInstruction::CreateCall(narrow_comp->root_instruction()->shape(), in WidenComputation() 79 Status DefaultAction(HloInstruction* hlo) override; 91 Status HandleParameter(HloInstruction* hlo) override; 93 Status HandleReduce(HloInstruction* hlo) override; 95 Status HandleDot(HloInstruction* hlo) override; 97 Status HandleTuple(HloInstruction* hlo) override; [all …]
|
D | hlo_module_group_util.cc | 41 std::vector<HloInstruction*> HloModuleGroupUtil::GlobalPredecessors( in GlobalPredecessors() 42 HloInstruction* instruction) { in GlobalPredecessors() 43 std::vector<HloInstruction*> in GlobalPredecessors() 45 absl::flat_hash_set<HloInstruction*> unique; in GlobalPredecessors() 51 auto add_unique_predecessor = [&](HloInstruction* predecessor) { in GlobalPredecessors() 56 for (HloInstruction* instr : metadata_.Companions(predecessor)) { in GlobalPredecessors() 64 for (HloInstruction* instr : in GlobalPredecessors() 79 std::vector<HloInstruction*> instruction_group; in GlobalPredecessors() 81 for (HloInstruction* companion : metadata_.Companions(instruction)) { in GlobalPredecessors() 90 for (HloInstruction* hlo : instruction_group) { in GlobalPredecessors() [all …]
|
D | batchnorm_expander.cc | 51 Status HandleBatchNormTraining(HloInstruction* batch_norm) override; 53 Status HandleBatchNormInference(HloInstruction* batch_norm) override; 55 Status HandleBatchNormGrad(HloInstruction* batch_norm) override; 78 HloInstruction::CreateParameter(0, shape, "scalar_lhs")); in GetOrCreateScalarAddComputation() 80 HloInstruction::CreateParameter(1, shape, "scalar_rhs")); in GetOrCreateScalarAddComputation() 81 auto scalar_op = b.AddInstruction(HloInstruction::CreateBinary( in GetOrCreateScalarAddComputation() 86 std::unique_ptr<HloInstruction> Rsqrt( in Rsqrt() 87 HloInstruction* operand, in Rsqrt() 88 const std::function<HloInstruction*(std::unique_ptr<HloInstruction>)>& in Rsqrt() 90 return HloInstruction::CreateUnary(operand->shape(), HloOpcode::kRsqrt, in Rsqrt() [all …]
|
/external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
D | gpu_fusible.h | 29 bool IsInputFusible(const HloInstruction& instr); 31 bool IsLoopFusible(const HloInstruction& instr); 39 bool LayoutsAreReduceInputFusionFriendly(const HloInstruction& producer, 40 const HloInstruction& reduce); 49 bool IsReduceInputFusion(const HloInstruction& instr); 53 bool IsInputFusibleReduction(const HloInstruction& instr); 57 bool IsInputFusibleScatter(const HloInstruction& instr); 65 bool FusionWouldBeTooLarge(const HloInstruction& instr1, 66 const HloInstruction& instr2, 71 bool CreatesNestedLoop(const HloInstruction& producer, [all …]
|
D | cudnn_batchnorm_rewriter.cc | 37 Status DefaultAction(HloInstruction* /*hlo_instruction*/) override { in DefaultAction() argument 41 Status HandleBatchNormInference(HloInstruction* batch_norm) override; 42 Status HandleBatchNormTraining(HloInstruction* batch_norm) override; 43 Status HandleBatchNormGrad(HloInstruction* batch_norm) override; 48 HloInstruction* AddConvert(HloInstruction* hlo, PrimitiveType elem_type) { in AddConvert() 51 HloInstruction::CreateConvert(shape, hlo)); in AddConvert() 57 bool EpsilonInRange(HloInstruction* batch_norm) { in EpsilonInRange() 61 bool IsF32BatchNormWithFP16Inputs(HloInstruction* batch_norm) { in IsF32BatchNormWithFP16Inputs() 69 Status Visitor::HandleBatchNormInference(HloInstruction* batch_norm) { in HandleBatchNormInference() 85 HloInstruction* epsilon = in HandleBatchNormInference() [all …]
|
/external/tensorflow/tensorflow/compiler/xla/service/cpu/ |
D | cpu_instruction_fusion_test.cc | 38 std::unique_ptr<HloInstruction> MakeDot(const Shape& shape, HloInstruction* lhs, in MakeDot() 39 HloInstruction* rhs) { in MakeDot() 46 return HloInstruction::CreateDot(shape, lhs, rhs, dot_dnums, in MakeDot() 52 HloInstruction* arg0 = builder.AddInstruction(HloInstruction::CreateParameter( in TEST_F() 54 HloInstruction* arg1 = builder.AddInstruction(HloInstruction::CreateParameter( in TEST_F() 57 HloInstruction* exp0 = builder.AddInstruction(HloInstruction::CreateUnary( in TEST_F() 59 HloInstruction* dot = builder.AddInstruction( in TEST_F() 71 HloInstruction* arg0 = builder.AddInstruction(HloInstruction::CreateParameter( in TEST_F() 73 HloInstruction* arg1 = builder.AddInstruction(HloInstruction::CreateParameter( in TEST_F() 76 HloInstruction* exp1 = builder.AddInstruction(HloInstruction::CreateUnary( in TEST_F() [all …]
|
D | ir_emitter.h | 86 std::unordered_map<const HloInstruction*, int64> 114 absl::Span<HloInstruction* const> instruction_order); 130 Status DefaultAction(HloInstruction* hlo) override; 132 Status HandleAllToAll(HloInstruction* instruction) override; 133 Status HandleBitcast(HloInstruction* bitcast) override; 134 Status HandleConstant(HloInstruction* constant) override; 135 Status HandleCopy(HloInstruction* copy) override; 136 Status HandleGetTupleElement(HloInstruction* get_tuple_element) override; 137 Status HandleSelect(HloInstruction* select) override; 138 Status HandleTupleSelect(HloInstruction* tuple_select) override; [all …]
|
/external/tensorflow/tensorflow/compiler/xla/service/spmd/ |
D | spmd_partitioner.h | 72 SpmdBuilder(const std::string& name, HloInstruction* hlo) in SpmdBuilder() 76 HloInstruction* AddInstruction(std::unique_ptr<HloInstruction> instruction); 78 const std::vector<HloInstruction*>& derived_instructions( in derived_instructions() 79 HloInstruction* hlo) { in derived_instructions() 83 void set_visiting_hlo(HloInstruction* hlo) { visiting_hlo_ = hlo; } in set_visiting_hlo() 85 HloInstruction* visiting_hlo() const { return visiting_hlo_; } in visiting_hlo() 89 const HloInstruction* hlo) { in BroadcastDimsForCreatedHlo() 99 HloInstruction* visiting_hlo_; 103 HloInstructionMap<std::vector<HloInstruction*>> instructions_; 108 absl::flat_hash_map<const HloInstruction*, absl::flat_hash_set<int64>> [all …]
|