Home
last modified time | relevance | path

Searched refs:HloOpcode (Results 1 – 25 of 218) sorted by relevance

123456789

/external/tensorflow/tensorflow/compiler/xla/service/
Dbfloat16_support.cc26 case HloOpcode::kCall: in SupportsBF16Operand()
27 case HloOpcode::kConditional: in SupportsBF16Operand()
28 case HloOpcode::kCustomCall: in SupportsBF16Operand()
29 case HloOpcode::kDomain: in SupportsBF16Operand()
30 case HloOpcode::kGetTupleElement: in SupportsBF16Operand()
31 case HloOpcode::kTuple: in SupportsBF16Operand()
32 case HloOpcode::kWhile: in SupportsBF16Operand()
34 case HloOpcode::kConvert: in SupportsBF16Operand()
45 case HloOpcode::kCall: in SupportsBF16Output()
46 case HloOpcode::kConditional: in SupportsBF16Output()
[all …]
Dinstruction_fusion.cc50 return instruction.opcode() == HloOpcode::kConvert && in IsAlwaysDuplicable()
60 case HloOpcode::kAdd: in IsExpensive()
61 case HloOpcode::kAnd: in IsExpensive()
62 case HloOpcode::kBitcast: in IsExpensive()
63 case HloOpcode::kBitcastConvert: in IsExpensive()
64 case HloOpcode::kBroadcast: in IsExpensive()
65 case HloOpcode::kCeil: in IsExpensive()
66 case HloOpcode::kClamp: in IsExpensive()
67 case HloOpcode::kClz: in IsExpensive()
68 case HloOpcode::kCompare: in IsExpensive()
[all …]
Dhlo_instruction.cc67 HloOpcode opcode; in CreateFromProto()
90 opcode = HloOpcode::kCompare; in CreateFromProto()
145 case HloOpcode::kBatchNormTraining: in CreateFromProto()
150 case HloOpcode::kBatchNormInference: in CreateFromProto()
155 case HloOpcode::kBatchNormGrad: in CreateFromProto()
160 case HloOpcode::kFft: { in CreateFromProto()
167 case HloOpcode::kCompare: { in CreateFromProto()
178 case HloOpcode::kTriangularSolve: { in CreateFromProto()
183 case HloOpcode::kCholesky: { in CreateFromProto()
188 case HloOpcode::kSend: in CreateFromProto()
[all …]
Dhlo_opcode_test.cc26 ASSERT_EQ("multiply", HloOpcodeString(HloOpcode::kMultiply)); in TEST()
39 auto opcode = static_cast<HloOpcode>(i); in TEST()
45 case HloOpcode::kCompare: in TEST()
52 case HloOpcode::kAfterAll: in TEST()
53 case HloOpcode::kAllReduce: in TEST()
54 case HloOpcode::kAllToAll: in TEST()
55 case HloOpcode::kCall: in TEST()
56 case HloOpcode::kConcatenate: in TEST()
57 case HloOpcode::kConditional: in TEST()
58 case HloOpcode::kCustomCall: in TEST()
[all …]
Dhlo_graph_dumper.cc262 if (root->opcode() == HloOpcode()) { in MatchTrivialComputation()
277 case HloOpcode::kAdd: in MatchTrivialComputation()
279 case HloOpcode::kMultiply: in MatchTrivialComputation()
281 case HloOpcode::kMinimum: in MatchTrivialComputation()
283 case HloOpcode::kMaximum: in MatchTrivialComputation()
285 case HloOpcode::kCompare: { in MatchTrivialComputation()
538 if (to_node->IsFused() && to_node->opcode() == HloOpcode::kParameter) { in Header()
556 CHECK_EQ(instr->opcode(), HloOpcode::kFusion); in ShouldShowFusionSubcomputation()
586 if (parent_instr->opcode() != HloOpcode::kFusion) { in DumpSubcomputation()
609 if (parent_instr->opcode() == HloOpcode::kFusion) { in DumpSubcomputation()
[all …]
Dlayout_assignment.cc401 if (instruction->opcode() == HloOpcode::kSend || in BuildHostChannelConstraints()
402 instruction->opcode() == HloOpcode::kRecv) { in BuildHostChannelConstraints()
443 if (instruction->opcode() == HloOpcode::kInfeed) { in AddMandatoryConstraints()
450 } else if (instruction->opcode() == HloOpcode::kOutfeed) { in AddMandatoryConstraints()
455 } else if (instruction->opcode() == HloOpcode::kParameter) { in AddMandatoryConstraints()
476 } else if (instruction->opcode() == HloOpcode::kSend || in AddMandatoryConstraints()
477 instruction->opcode() == HloOpcode::kRecv) { in AddMandatoryConstraints()
485 if (instruction->opcode() == HloOpcode::kSend) { in AddMandatoryConstraints()
533 if (instruction->opcode() == HloOpcode::kCall) { in AddMandatoryConstraints()
547 } else if (instruction->opcode() == HloOpcode::kWhile) { in AddMandatoryConstraints()
[all …]
Ddynamic_padder.cc51 case HloOpcode::kReduce: in ChooseIdentityValue()
52 case HloOpcode::kReduceWindow: { in ChooseIdentityValue()
59 case HloOpcode::kConvolution: in ChooseIdentityValue()
60 case HloOpcode::kDot: { in ChooseIdentityValue()
67 case HloOpcode::kPad: { in ChooseIdentityValue()
71 case HloOpcode::kSelectAndScatter: { in ChooseIdentityValue()
74 case HloOpcode::kParameter: in ChooseIdentityValue()
75 case HloOpcode::kGetDimensionSize: in ChooseIdentityValue()
76 case HloOpcode::kReshape: in ChooseIdentityValue()
77 case HloOpcode::kTuple: in ChooseIdentityValue()
[all …]
Dhlo_element_type_converter.cc128 if (opcode == HloOpcode::kParameter || opcode == HloOpcode::kConstant || in Run()
129 opcode == HloOpcode::kTuple || opcode == HloOpcode::kConvert || in Run()
130 opcode == HloOpcode::kBitcastConvert || in Run()
131 opcode == HloOpcode::kGetTupleElement || in Run()
132 opcode == HloOpcode::kInfeed || opcode == HloOpcode::kOutfeed) { in Run()
138 if (opcode == HloOpcode::kCustomCall) { in Run()
144 if (opcode == HloOpcode::kWhile || opcode == HloOpcode::kCall || in Run()
145 opcode == HloOpcode::kAllReduce || opcode == HloOpcode::kFusion || in Run()
146 opcode == HloOpcode::kMap || opcode == HloOpcode::kReduce || in Run()
147 opcode == HloOpcode::kReduceWindow || opcode == HloOpcode::kScatter || in Run()
[all …]
Dhlo_opcode.cc23 string HloOpcodeString(HloOpcode opcode) { in HloOpcodeString()
26 case HloOpcode::enum_name: \ in HloOpcodeString()
33 StatusOr<HloOpcode> StringToHloOpcode(const string& opcode_name) { in StringToHloOpcode()
34 static auto* opcode_map = new absl::flat_hash_map<string, HloOpcode>({ in StringToHloOpcode()
36 {opcode_name, HloOpcode::enum_name}, in StringToHloOpcode()
47 bool HloOpcodeIsComparison(HloOpcode opcode) { in HloOpcodeIsComparison()
48 return opcode == HloOpcode::kCompare; in HloOpcodeIsComparison()
51 bool HloOpcodeIsVariadic(HloOpcode opcode) { in HloOpcodeIsVariadic()
54 case HloOpcode::enum_name: \ in HloOpcodeIsVariadic()
61 absl::optional<int> HloOpcodeArity(HloOpcode opcode) { in HloOpcodeArity()
[all …]
Dhlo_dataflow_analysis.cc66 user->opcode() != HloOpcode::kTuple) { in AreTransitiveUsesElementwiseOrTuple()
307 CHECK_EQ(bitcast->opcode(), HloOpcode::kBitcast); in UpdateBitcastValueSet()
319 CHECK_EQ(send->opcode(), HloOpcode::kSend); in UpdateSendValueSet()
341 CHECK_EQ(recv_done->opcode(), HloOpcode::kRecvDone); in UpdateRecvDoneValueSet()
363 CHECK_EQ(call->opcode(), HloOpcode::kCall); in UpdateCallValueSet()
376 CHECK_EQ(conditional->opcode(), HloOpcode::kConditional); in UpdateConditionalValueSet()
390 CHECK_EQ(copy->opcode(), HloOpcode::kCopy); in UpdateCopyValueSet()
414 CHECK_EQ(domain->opcode(), HloOpcode::kDomain); in UpdateDomainValueSet()
431 CHECK_EQ(add_dependency->opcode(), HloOpcode::kAddDependency); in UpdateAddDependencyValueSet()
444 CHECK_EQ(gte->opcode(), HloOpcode::kGetTupleElement); in UpdateGetTupleElementValueSet()
[all …]
Dbfloat16_normalization_test.cc39 if (hlo.opcode() == HloOpcode::kAdd || in SupportsBF16Operand()
40 hlo.opcode() == HloOpcode::kSubtract || in SupportsBF16Operand()
41 hlo.opcode() == HloOpcode::kReduce || in SupportsBF16Operand()
42 hlo.opcode() == HloOpcode::kTuple || in SupportsBF16Operand()
43 hlo.opcode() == HloOpcode::kGetTupleElement) { in SupportsBF16Operand()
46 if (hlo.opcode() == HloOpcode::kDot) { in SupportsBF16Operand()
54 if (hlo.opcode() == HloOpcode::kAdd || hlo.opcode() == HloOpcode::kReduce || in SupportsBF16Output()
55 hlo.opcode() == HloOpcode::kSubtract || in SupportsBF16Output()
56 hlo.opcode() == HloOpcode::kDot || hlo.opcode() == HloOpcode::kTuple || in SupportsBF16Output()
57 hlo.opcode() == HloOpcode::kGetTupleElement) { in SupportsBF16Output()
[all …]
Dreduce_precision_insertion_test.cc55 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a)); in TEST_F()
66 return instruction->opcode() == HloOpcode::kCos; in TEST_F()
82 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a)); in TEST_F()
93 return instruction->opcode() == HloOpcode::kCos; in TEST_F()
112 HloInstruction::CreateBinary(shape, HloOpcode::kAdd, a, b)); in TEST_F()
124 return instruction->opcode() == HloOpcode::kAdd; in TEST_F()
141 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a)); in TEST_F()
153 HloOpcode::kParameter; in TEST_F()
170 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a)); in TEST_F()
172 HloInstruction::CreateUnary(shape, HloOpcode::kSin, a)); in TEST_F()
[all …]
Dar_crs_combiner.cc48 case HloOpcode::kBitcast: in MatchesArCrsPattern()
49 case HloOpcode::kTranspose: in MatchesArCrsPattern()
50 case HloOpcode::kReshape: in MatchesArCrsPattern()
52 case HloOpcode::kConvert: in MatchesArCrsPattern()
57 case HloOpcode::kAdd: in MatchesArCrsPattern()
58 case HloOpcode::kSubtract: in MatchesArCrsPattern()
59 case HloOpcode::kMultiply: in MatchesArCrsPattern()
97 CHECK_EQ(HloOpcode::kParameter, instruction->opcode()); in WhileFromBodyParameter()
102 if (caller_instruction->opcode() == HloOpcode::kWhile) { in WhileFromBodyParameter()
111 if (instruction->opcode() == HloOpcode::kTuple) { in GetAllTuples()
[all …]
Dbfloat16_conversion_folding.cc72 CHECK_EQ(user->opcode(), HloOpcode::kConvert); in FoldOutputConversions()
83 CHECK_EQ(operand->opcode(), HloOpcode::kConvert); in FoldOperandConversion()
98 if (user->opcode() == HloOpcode::kConvert && in AllUsersAreF32ToBF16Converts()
116 if (operand->opcode() == HloOpcode::kConvert && in TryFoldBF16Conversions()
156 if (hlo->opcode() == HloOpcode::kTuple || // in DefaultAction()
157 hlo->opcode() == HloOpcode::kGetTupleElement || // in DefaultAction()
158 hlo->opcode() == HloOpcode::kConstant || // in DefaultAction()
159 hlo->opcode() == HloOpcode::kParameter || // in DefaultAction()
160 hlo->opcode() == HloOpcode::kFusion || // in DefaultAction()
161 hlo->opcode() == HloOpcode::kConvert || // in DefaultAction()
[all …]
Ddefuser_test.cc53 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1)); in TEST_F()
69 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1)); in TEST_F()
93 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1)); in TEST_F()
95 HloInstruction::CreateUnary(shape_, HloOpcode::kNegate, add)); in TEST_F()
121 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1)); in TEST_F()
123 HloInstruction::CreateUnary(shape_, HloOpcode::kNegate, add)); in TEST_F()
125 HloInstruction::CreateBinary(shape_, HloOpcode::kSubtract, add, negate)); in TEST_F()
127 HloInstruction::CreateBinary(shape_, HloOpcode::kMultiply, sub, param3)); in TEST_F()
129 HloInstruction::CreateBinary(shape_, HloOpcode::kDivide, mul, param3)); in TEST_F()
133 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, constant, div)); in TEST_F()
[all …]
Dwhile_loop_invariant_code_motion.cc85 old_instruction->opcode() != HloOpcode::kConstant); in CreateLoopInvariantCopy()
110 case HloOpcode::kConstant: in NotWorthHoistingIndividually()
113 case HloOpcode::kBitcast: in NotWorthHoistingIndividually()
114 case HloOpcode::kBroadcast: in NotWorthHoistingIndividually()
115 case HloOpcode::kIota: in NotWorthHoistingIndividually()
116 case HloOpcode::kReshape: in NotWorthHoistingIndividually()
117 case HloOpcode::kReverse: in NotWorthHoistingIndividually()
118 case HloOpcode::kSlice: in NotWorthHoistingIndividually()
119 case HloOpcode::kTranspose: in NotWorthHoistingIndividually()
120 case HloOpcode::kTuple: in NotWorthHoistingIndividually()
[all …]
Dbatchnorm_expander.cc90 shape, HloOpcode::kAdd, scalar_lhs, scalar_rhs)); in GetOrCreateScalarAddComputation()
98 return HloInstruction::CreateUnary(operand->shape(), HloOpcode::kRsqrt, in Rsqrt()
108 return HloInstruction::CreateBinary(operand->shape(), HloOpcode::kDivide, in Mean()
127 ShapeUtil::MakeShape(U32, {}), HloOpcode::kMultiply, in DynamicElementCountPerFeature()
198 auto add_binary = [&](const Shape& shape, const HloOpcode opcode, in HandleBatchNormTraining()
245 add_binary(operand_shape, HloOpcode::kMultiply, operand, operand); in HandleBatchNormTraining()
267 add_binary(feature_shape, HloOpcode::kMultiply, mean, mean); in HandleBatchNormTraining()
271 add_binary(feature_shape, HloOpcode::kSubtract, square_mean, mean_square); in HandleBatchNormTraining()
278 add_binary(operand_shape, HloOpcode::kAdd, var_broadcasted, epsilon); in HandleBatchNormTraining()
284 auto operand_minus_mean = add_binary(operand_shape, HloOpcode::kSubtract, in HandleBatchNormTraining()
[all …]
Dbfloat16_conversion_folding_test.cc37 if (hlo.opcode() == HloOpcode::kAdd || in SupportsBF16Operand()
38 hlo.opcode() == HloOpcode::kSubtract || in SupportsBF16Operand()
39 hlo.opcode() == HloOpcode::kTuple || in SupportsBF16Operand()
40 hlo.opcode() == HloOpcode::kGetTupleElement || in SupportsBF16Operand()
41 hlo.opcode() == HloOpcode::kAllReduce) { in SupportsBF16Operand()
48 if (hlo.opcode() == HloOpcode::kAdd || in SupportsBF16Output()
49 hlo.opcode() == HloOpcode::kSubtract || in SupportsBF16Output()
50 hlo.opcode() == HloOpcode::kTuple || in SupportsBF16Output()
51 hlo.opcode() == HloOpcode::kGetTupleElement || in SupportsBF16Output()
52 hlo.opcode() == HloOpcode::kAllReduce) { in SupportsBF16Output()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Dcpu_instruction_fusion_test.cc57 ShapeUtil::MakeShape(S32, {1024, 256}), HloOpcode::kExp, arg0)); in TEST_F()
76 ShapeUtil::MakeShape(S32, {256, 1024}), HloOpcode::kExp, arg1)); in TEST_F()
95 ShapeUtil::MakeShape(S32, {2, 512, 2, 128}), HloOpcode::kExp, arg0)); in TEST_F()
97 ShapeUtil::MakeShape(S32, {1024, 256}), HloOpcode::kBitcast, exp0)); in TEST_F()
115 ShapeUtil::MakeShape(S32, {2, 512, 2, 128}), HloOpcode::kExp, arg0)); in TEST_F()
137 ShapeUtil::MakeShape(S32, {256, 32 * 1024}), HloOpcode::kExp, arg1)); in TEST_F()
156 ShapeUtil::MakeShape(S32, {256, 1024}), HloOpcode::kExp, arg1)); in TEST_F()
263 HloModule* module, const std::multiset<HloOpcode>& expected_opcodes, in RunFusionAndCheckOpcodesWereFused()
275 std::vector<HloOpcode> fused_opcodes(root->fused_instruction_count()); in RunFusionAndCheckOpcodesWereFused()
281 std::multiset<HloOpcode>(fused_opcodes.begin(), fused_opcodes.end()), in RunFusionAndCheckOpcodesWereFused()
[all …]
Dcpu_instruction_fusion.cc34 hlo.opcode() == HloOpcode::kBroadcast || in CanBeLoopFused()
35 hlo.opcode() == HloOpcode::kConcatenate || in CanBeLoopFused()
36 hlo.opcode() == HloOpcode::kDynamicSlice || in CanBeLoopFused()
37 hlo.opcode() == HloOpcode::kDynamicUpdateSlice || in CanBeLoopFused()
38 hlo.opcode() == HloOpcode::kGather || in CanBeLoopFused()
39 hlo.opcode() == HloOpcode::kIota || hlo.opcode() == HloOpcode::kPad || in CanBeLoopFused()
40 hlo.opcode() == HloOpcode::kReshape || in CanBeLoopFused()
41 hlo.opcode() == HloOpcode::kReverse || in CanBeLoopFused()
42 hlo.opcode() == HloOpcode::kSlice || in CanBeLoopFused()
43 hlo.opcode() == HloOpcode::kTranspose; in CanBeLoopFused()
[all …]
Dparallel_task_assignment.cc139 if (opcode == HloOpcode::kParameter || opcode == HloOpcode::kConstant || in GetTargetParallelTaskCount()
140 opcode == HloOpcode::kCall || opcode == HloOpcode::kCustomCall || in GetTargetParallelTaskCount()
141 opcode == HloOpcode::kDot || opcode == HloOpcode::kSelectAndScatter || in GetTargetParallelTaskCount()
142 opcode == HloOpcode::kGetTupleElement || opcode == HloOpcode::kBitcast || in GetTargetParallelTaskCount()
143 opcode == HloOpcode::kFft || opcode == HloOpcode::kInfeed || in GetTargetParallelTaskCount()
144 opcode == HloOpcode::kOutfeed || opcode == HloOpcode::kRng || in GetTargetParallelTaskCount()
145 opcode == HloOpcode::kSort || in GetTargetParallelTaskCount()
146 (opcode == HloOpcode::kConvolution && in GetTargetParallelTaskCount()
149 (opcode == HloOpcode::kFusion && in GetTargetParallelTaskCount()
192 if (instruction->opcode() == HloOpcode::kWhile) { in AssignParallelTasksHelper()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/gpu/
Dinstruction_fusion.cc33 if (constant->opcode() != HloOpcode::kConstant || in IsIEEEFloatingPointScalarConstant()
47 case HloOpcode::kDivide: in IsExpensive()
122 if (producer->opcode() == HloOpcode::kDot || in ShouldFuseInexpensiveChecks()
123 (producer->opcode() == HloOpcode::kFusion && in ShouldFuseInexpensiveChecks()
124 producer->fused_expression_root()->opcode() == HloOpcode::kDot)) { in ShouldFuseInexpensiveChecks()
129 consumer->opcode() == HloOpcode::kFusion && in ShouldFuseInexpensiveChecks()
133 .WithOpcode(HloOpcode::kMultiply) in ShouldFuseInexpensiveChecks()
140 if (op1->opcode() != HloOpcode::kParameter) { in ShouldFuseInexpensiveChecks()
143 if (op1->opcode() != HloOpcode::kParameter || in ShouldFuseInexpensiveChecks()
144 op2->opcode() != HloOpcode::kBroadcast) { in ShouldFuseInexpensiveChecks()
[all …]
Dgpu_fusible.cc33 if (instr.opcode() == HloOpcode::kFusion) { in AppendParams()
75 } else if (instr.opcode() == HloOpcode::kFusion && in IsReduceInputFusion()
95 if (instr->opcode() == HloOpcode::kFusion) { in ShapesCompatibleForMultiOutputFusion()
142 if (instr.opcode() == HloOpcode::kScatter || in IsInputFusibleScatter()
143 (instr.opcode() == HloOpcode::kFusion && in IsInputFusibleScatter()
145 instr.fused_expression_root()->opcode() == HloOpcode::kScatter)) { in IsInputFusibleScatter()
164 instr.opcode() == HloOpcode::kBitcast || in IsLoopFusible()
165 instr.opcode() == HloOpcode::kBroadcast || in IsLoopFusible()
166 instr.opcode() == HloOpcode::kConcatenate || in IsLoopFusible()
167 instr.opcode() == HloOpcode::kDynamicSlice || in IsLoopFusible()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/cpu/tests/
Dcpu_fusion_test.cc63 HloInstruction::CreateBinary(vshape, HloOpcode::kAdd, input1, input2)); in TEST_F()
65 HloInstruction::CreateUnary(vshape, HloOpcode::kNegate, add1)); in TEST_F()
77 EXPECT_EQ(HloOpcode::kFusion, fusion_instruction->opcode()); in TEST_F()
78 EXPECT_EQ(HloOpcode::kNegate, in TEST_F()
99 HloInstruction::CreateUnary(vshape, HloOpcode::kNegate, input)); in TEST_F()
101 HloInstruction::CreateUnary(vshape, HloOpcode::kCeil, negate)); in TEST_F()
103 HloInstruction::CreateUnary(vshape, HloOpcode::kExp, ceil)); in TEST_F()
105 HloInstruction::CreateUnary(vshape, HloOpcode::kFloor, exp)); in TEST_F()
112 HloInstruction::CreateBinary(vshape, HloOpcode::kMultiply, two, floor)); in TEST_F()
124 EXPECT_EQ(HloOpcode::kFusion, fusion_instruction->opcode()); in TEST_F()
[all …]
/external/tensorflow/tensorflow/compiler/xla/tests/
Dbroadcast_simple_test.cc37 XlaOp BuildBinOp(HloOpcode op, const XlaOp& lhs, const XlaOp& rhs, in BuildBinOp()
40 case HloOpcode::kMinimum: { in BuildBinOp()
43 case HloOpcode::kMaximum: { in BuildBinOp()
46 case HloOpcode::kMultiply: { in BuildBinOp()
84 float ApplyOpToFloats(HloOpcode op, float lhs, float rhs) { in ApplyOpToFloats()
86 case HloOpcode::kMinimum: { in ApplyOpToFloats()
89 case HloOpcode::kMaximum: { in ApplyOpToFloats()
92 case HloOpcode::kMultiply: { in ApplyOpToFloats()
95 case HloOpcode::kAdd: { in ApplyOpToFloats()
308 HloOpcode op;
[all …]

123456789