Home
last modified time | relevance | path

Searched refs:HloInstruction (Results 1 – 25 of 472) sorted by relevance

12345678910>>...19

/external/tensorflow/tensorflow/compiler/xla/service/
Dhlo_creation_utils.h32 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 …]
Dbfloat16_propagation_test.cc39 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 …]
Dhlo_cost_analysis.h53 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 …]
Dhlo_instruction.h270 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 …]
Dtuple_simplifier_test.cc61 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 …]
Dhlo_computation.h71 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 …]
Dmemory_space_assignment_test.cc96 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 …]
Dhlo_verifier.h42 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 …]
Dbfloat16_conversion_folding_test.cc35 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 …]
Dinstruction_fusion.h42 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 …]
Dwhile_loop_invariant_code_motion_test.cc38 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 …]
Dmulti_output_fusion.h65 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 …]
Dhlo_dataflow_analysis.h59 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 …]
Dhlo_instructions.h28 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 …]
Dhlo_instruction_test.cc55 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 …]
Dcopy_insertion_test.cc90 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 …]
Dhlo_creation_utils.cc37 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 …]
Ddynamic_dimension_inference.cc52 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 …]
Dhlo_module_group_util.cc41 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 …]
Dbatchnorm_expander.cc51 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/
Dgpu_fusible.h29 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 …]
Dcudnn_batchnorm_rewriter.cc37 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/
Dcpu_instruction_fusion_test.cc38 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 …]
Dir_emitter.h86 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/
Dspmd_partitioner.h72 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 …]

12345678910>>...19