Home
last modified time | relevance | path

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

12345678910>>...12

/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.cc55 return instruction.opcode() == HloOpcode::kConvert && in IsAlwaysDuplicable()
67 case HloOpcode::kAdd: in IsExpensive()
68 case HloOpcode::kAnd: in IsExpensive()
69 case HloOpcode::kBitcast: in IsExpensive()
70 case HloOpcode::kBitcastConvert: in IsExpensive()
71 case HloOpcode::kBroadcast: in IsExpensive()
72 case HloOpcode::kCeil: in IsExpensive()
73 case HloOpcode::kClamp: in IsExpensive()
74 case HloOpcode::kClz: in IsExpensive()
75 case HloOpcode::kCompare: in IsExpensive()
[all …]
Dhlo_instruction.cc73 HloOpcode opcode; in CreateFromProto()
96 opcode = HloOpcode::kCompare; in CreateFromProto()
151 case HloOpcode::kBatchNormTraining: in CreateFromProto()
156 case HloOpcode::kBatchNormInference: in CreateFromProto()
161 case HloOpcode::kBatchNormGrad: in CreateFromProto()
166 case HloOpcode::kFft: { in CreateFromProto()
173 case HloOpcode::kCopyStart: { in CreateFromProto()
178 case HloOpcode::kCompare: { in CreateFromProto()
200 case HloOpcode::kTriangularSolve: { in CreateFromProto()
205 case HloOpcode::kCholesky: { 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 …]
Dsharding_propagation.cc130 CHECK_EQ(instruction->opcode(), HloOpcode::kConvolution); in IsConvolutionKernelSmall()
150 case HloOpcode::kMap: in PickRepresentativeOperand()
151 case HloOpcode::kPad: in PickRepresentativeOperand()
152 case HloOpcode::kPower: in PickRepresentativeOperand()
153 case HloOpcode::kReverse: in PickRepresentativeOperand()
154 case HloOpcode::kSlice: in PickRepresentativeOperand()
155 case HloOpcode::kShiftLeft: in PickRepresentativeOperand()
156 case HloOpcode::kShiftRightArithmetic: in PickRepresentativeOperand()
157 case HloOpcode::kShiftRightLogical: in PickRepresentativeOperand()
165 case HloOpcode::kAbs: in PickRepresentativeOperand()
[all …]
Dhlo_graph_dumper.cc267 if (root->opcode() == HloOpcode()) { in MatchTrivialComputation()
282 case HloOpcode::kAdd: in MatchTrivialComputation()
284 case HloOpcode::kMultiply: in MatchTrivialComputation()
286 case HloOpcode::kMinimum: in MatchTrivialComputation()
288 case HloOpcode::kMaximum: in MatchTrivialComputation()
290 case HloOpcode::kCompare: { in MatchTrivialComputation()
544 if (to_node->IsFused() && to_node->opcode() == HloOpcode::kParameter) { in Header()
562 CHECK_EQ(instr->opcode(), HloOpcode::kFusion); in ShouldShowFusionSubcomputation()
593 if (parent_instr->opcode() != HloOpcode::kFusion) { in DumpSubcomputation()
616 if (parent_instr->opcode() == HloOpcode::kFusion) { in DumpSubcomputation()
[all …]
Dlayout_assignment.cc411 if (instruction->opcode() == HloOpcode::kSend || in BuildHostChannelConstraints()
412 instruction->opcode() == HloOpcode::kRecv) { in BuildHostChannelConstraints()
459 if (instruction->opcode() == HloOpcode::kInfeed) { in AddMandatoryConstraints()
466 } else if (instruction->opcode() == HloOpcode::kOutfeed) { in AddMandatoryConstraints()
471 } else if (instruction->opcode() == HloOpcode::kParameter) { in AddMandatoryConstraints()
490 } else if (instruction->opcode() == HloOpcode::kSend || in AddMandatoryConstraints()
491 instruction->opcode() == HloOpcode::kRecv) { in AddMandatoryConstraints()
499 if (instruction->opcode() == HloOpcode::kSend) { in AddMandatoryConstraints()
551 if (instruction->opcode() == HloOpcode::kCall) { in AddMandatoryConstraints()
565 } else if (instruction->opcode() == HloOpcode::kWhile) { in AddMandatoryConstraints()
[all …]
Dhlo_replication_analysis.cc58 if (hlo->opcode() == HloOpcode::kAllReduce || in DetermineHloInstructionIsReplicated()
59 hlo->opcode() == HloOpcode::kAllGather) { in DetermineHloInstructionIsReplicated()
75 if (hlo->opcode() == HloOpcode::kAllReduce) { in DetermineHloInstructionIsReplicated()
111 if (hlo->opcode() == HloOpcode::kReplicaId) { in DetermineHloInstructionIsReplicated()
115 if (hlo->opcode() == HloOpcode::kPartitionId) { in DetermineHloInstructionIsReplicated()
120 if (hlo->opcode() == HloOpcode::kParameter) { in DetermineHloInstructionIsReplicated()
128 if (hlo->opcode() == HloOpcode::kConstant) { in DetermineHloInstructionIsReplicated()
132 if (hlo->opcode() == HloOpcode::kCustomCall && in DetermineHloInstructionIsReplicated()
140 hlo->opcode() == HloOpcode::kConcatenate || // in DetermineHloInstructionIsReplicated()
141 hlo->opcode() == HloOpcode::kConvolution || // in DetermineHloInstructionIsReplicated()
[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.cc55 if (instruction->opcode() == HloOpcode::kCall || in CalculatePostOrderScheduleHelper()
56 instruction->opcode() == HloOpcode::kConditional) { in CalculatePostOrderScheduleHelper()
63 if (instruction->opcode() == HloOpcode::kWhile) { in CalculatePostOrderScheduleHelper()
111 user->opcode() != HloOpcode::kTuple) { in AreTransitiveUsesElementwiseOrTuple()
388 CHECK_EQ(bitcast->opcode(), HloOpcode::kBitcast); in UpdateBitcastValueSet()
401 CHECK_EQ(set_dimension_size->opcode(), HloOpcode::kSetDimensionSize); in UpdateSetDimensionSizeValueSet()
414 CHECK_EQ(send->opcode(), HloOpcode::kSend); in UpdateSendValueSet()
437 CHECK_EQ(custom_call->opcode(), HloOpcode::kCustomCall); in UpdateCustomCallValueSet()
453 CHECK_EQ(copy_start->opcode(), HloOpcode::kCopyStart); in UpdateCopyStartValueSet()
466 CHECK_EQ(copy_done->opcode(), HloOpcode::kCopyDone); in UpdateCopyDoneValueSet()
[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 …]
Dmemory_space_assignment_test.cc101 instruction->opcode() == HloOpcode::kParameter) { in AssignMemorySpace()
209 if (instruction->opcode() == HloOpcode::kCopyStart) { in CountMaximumOutstandingAsyncCopies()
218 } else if (instruction->opcode() == HloOpcode::kCopyDone) { in CountMaximumOutstandingAsyncCopies()
261 HloInstruction::CreateUnary(shape, HloOpcode::kTanh, p0)); in CreateEvictAndPrefetchModule()
268 HloInstruction::CreateBinary(shape, HloOpcode::kAdd, p0, tanh)); in CreateEvictAndPrefetchModule()
270 HloInstruction::CreateBinary(shape, HloOpcode::kSubtract, p0, p1)); in CreateEvictAndPrefetchModule()
272 HloInstruction::CreateBinary(shape, HloOpcode::kMultiply, p0, p1)); in CreateEvictAndPrefetchModule()
274 HloInstruction::CreateBinary(shape, HloOpcode::kSubtract, p0, p1)); in CreateEvictAndPrefetchModule()
276 HloInstruction::CreateBinary(shape, HloOpcode::kMultiply, a, b)); in CreateEvictAndPrefetchModule()
278 HloInstruction::CreateBinary(shape, HloOpcode::kMultiply, a, c)); in CreateEvictAndPrefetchModule()
[all …]
Dwhile_loop_all_reduce_code_motion.cc106 case HloOpcode::kBitcast: in IsAllReduceMovable()
107 case HloOpcode::kConvert: in IsAllReduceMovable()
108 case HloOpcode::kReshape: in IsAllReduceMovable()
109 case HloOpcode::kTranspose: in IsAllReduceMovable()
110 case HloOpcode::kDynamicReshape: in IsAllReduceMovable()
113 case HloOpcode::kGetTupleElement: { in IsAllReduceMovable()
125 case HloOpcode::kParameter: { in IsAllReduceMovable()
132 if (instruction->opcode() == HloOpcode::kParameter) { in IsAllReduceMovable()
153 case HloOpcode::kBitcast: in IsAllReduceMovable()
154 case HloOpcode::kConvert: in IsAllReduceMovable()
[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()
44 hlo.opcode() == HloOpcode::kAllToAll) { in SupportsBF16Operand()
47 if (hlo.opcode() == HloOpcode::kDot) { in SupportsBF16Operand()
55 if (hlo.opcode() == HloOpcode::kAdd || hlo.opcode() == HloOpcode::kReduce || in SupportsBF16Output()
56 hlo.opcode() == HloOpcode::kSubtract || in SupportsBF16Output()
57 hlo.opcode() == HloOpcode::kDot || hlo.opcode() == HloOpcode::kTuple || in SupportsBF16Output()
[all …]
Dhlo_cse.cc64 if (instruction->opcode() == HloOpcode::kConstant) { in CombineConstants()
108 hash, instruction->opcode() == HloOpcode::kGetTupleElement in CseHash()
120 case HloOpcode::kConstant: in CseHash()
122 case HloOpcode::kSlice: in CseHash()
126 case HloOpcode::kPad: in CseHash()
129 case HloOpcode::kDot: in CseHash()
132 case HloOpcode::kConvolution: in CseHash()
137 case HloOpcode::kReduceWindow: in CseHash()
139 case HloOpcode::kConcatenate: in CseHash()
140 case HloOpcode::kBroadcast: in CseHash()
[all …]
Dlogistic_expander.cc46 MakeUnaryHlo(HloOpcode::kTanh, in ExpandLogisticWithTanh()
47 MakeBinaryHlo(HloOpcode::kMultiply, half_constant, operand) in ExpandLogisticWithTanh()
51 HloOpcode::kAdd, half_constant, in ExpandLogisticWithTanh()
52 MakeBinaryHlo(HloOpcode::kMultiply, half_constant, tanh_instr) in ExpandLogisticWithTanh()
63 MakeUnaryHlo(HloOpcode::kExp, in ExpandLogisticWithExp()
64 MakeUnaryHlo(HloOpcode::kNegate, operand).ValueOrDie()) in ExpandLogisticWithExp()
67 MakeBinaryHlo(HloOpcode::kAdd, one_constant, exp_instr).ValueOrDie(); in ExpandLogisticWithExp()
68 return MakeBinaryHlo(HloOpcode::kDivide, one_constant, denominator) in ExpandLogisticWithExp()
75 return instruction->opcode() == HloOpcode::kLogistic; in InstructionMatchesPattern()
Dbatchnorm_expander.cc82 shape, HloOpcode::kAdd, scalar_lhs, scalar_rhs)); in GetOrCreateScalarAddComputation()
90 return HloInstruction::CreateUnary(operand->shape(), HloOpcode::kRsqrt, in Rsqrt()
100 return HloInstruction::CreateBinary(operand->shape(), HloOpcode::kDivide, in Mean()
119 ShapeUtil::MakeShape(S32, {}), HloOpcode::kMultiply, in DynamicElementCountPerFeature()
165 auto add_binary = [&](const Shape& shape, const HloOpcode opcode, in HandleBatchNormTraining()
212 add_binary(operand_shape, HloOpcode::kMultiply, operand, operand); in HandleBatchNormTraining()
234 add_binary(feature_shape, HloOpcode::kMultiply, mean, mean); in HandleBatchNormTraining()
238 add_binary(feature_shape, HloOpcode::kSubtract, square_mean, mean_square); in HandleBatchNormTraining()
245 add_binary(operand_shape, HloOpcode::kAdd, var_broadcasted, epsilon); in HandleBatchNormTraining()
251 auto operand_minus_mean = add_binary(operand_shape, HloOpcode::kSubtract, in HandleBatchNormTraining()
[all …]
Dbfloat16_conversion_folding.cc81 CHECK_EQ(user->opcode(), HloOpcode::kConvert); in FoldOutputConversions()
92 CHECK_EQ(operand->opcode(), HloOpcode::kConvert); in FoldOperandConversion()
107 if (user->opcode() == HloOpcode::kConvert && in AllUsersAreF32ToBF16Converts()
125 if (operand->opcode() == HloOpcode::kConvert && in TryFoldBF16Conversions()
165 if (hlo->opcode() == HloOpcode::kTuple || // in DefaultAction()
166 hlo->opcode() == HloOpcode::kGetTupleElement || // in DefaultAction()
167 hlo->opcode() == HloOpcode::kConstant || // in DefaultAction()
168 hlo->opcode() == HloOpcode::kParameter || // in DefaultAction()
169 hlo->opcode() == HloOpcode::kFusion || // in DefaultAction()
170 hlo->opcode() == HloOpcode::kBitcastConvert || // in DefaultAction()
[all …]
/external/tensorflow/tensorflow/compiler/mlir/xla/
Dhlo_utils.cc284 StatusOr<::xla::HloOpcode> MhloToHloOpcode(mlir::Operation* op) { in MhloToHloOpcode()
288 return xla::HloOpcode::kConstant; in MhloToHloOpcode()
290 return xla::HloOpcode::kIota; in MhloToHloOpcode()
292 return xla::HloOpcode::kConvert; in MhloToHloOpcode()
294 return xla::HloOpcode::kAdd; in MhloToHloOpcode()
296 return xla::HloOpcode::kAtan2; in MhloToHloOpcode()
298 return xla::HloOpcode::kDivide; in MhloToHloOpcode()
300 return xla::HloOpcode::kMaximum; in MhloToHloOpcode()
302 return xla::HloOpcode::kMinimum; in MhloToHloOpcode()
304 return xla::HloOpcode::kMultiply; in MhloToHloOpcode()
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Dcpu_instruction_fusion_test.cc58 ShapeUtil::MakeShape(F32, {1024, 256}), HloOpcode::kExp, arg0)); in TEST_F()
77 ShapeUtil::MakeShape(F32, {256, 1024}), HloOpcode::kExp, arg1)); in TEST_F()
96 ShapeUtil::MakeShape(F32, {2, 512, 2, 128}), HloOpcode::kExp, arg0)); in TEST_F()
98 ShapeUtil::MakeShape(F32, {1024, 256}), HloOpcode::kBitcast, exp0)); in TEST_F()
117 ShapeUtil::MakeShape(F32, {2, 512, 2, 128}), HloOpcode::kExp, arg0)); in TEST_F()
139 ShapeUtil::MakeShape(F32, {32 * 1024, 256}), HloOpcode::kExp, arg1)); in TEST_F()
158 ShapeUtil::MakeShape(F32, {256, 1024}), HloOpcode::kExp, arg1)); in TEST_F()
265 HloModule* module, const std::multiset<HloOpcode>& expected_opcodes, in RunFusionAndCheckOpcodesWereFused()
277 std::vector<HloOpcode> fused_opcodes(root->fused_instruction_count()); in RunFusionAndCheckOpcodesWereFused()
283 std::multiset<HloOpcode>(fused_opcodes.begin(), fused_opcodes.end()), in RunFusionAndCheckOpcodesWereFused()
[all …]
Dcpu_instruction_fusion.cc31 hlo.opcode() == HloOpcode::kBitcast || in CanBeLoopFused()
32 hlo.opcode() == HloOpcode::kBroadcast || in CanBeLoopFused()
33 hlo.opcode() == HloOpcode::kConcatenate || in CanBeLoopFused()
34 hlo.opcode() == HloOpcode::kDynamicSlice || in CanBeLoopFused()
35 hlo.opcode() == HloOpcode::kDynamicUpdateSlice || in CanBeLoopFused()
36 hlo.opcode() == HloOpcode::kGather || in CanBeLoopFused()
37 hlo.opcode() == HloOpcode::kIota || hlo.opcode() == HloOpcode::kPad || in CanBeLoopFused()
38 hlo.opcode() == HloOpcode::kReduce || in CanBeLoopFused()
39 hlo.opcode() == HloOpcode::kReshape || in CanBeLoopFused()
40 hlo.opcode() == HloOpcode::kReverse || in CanBeLoopFused()
[all …]
Dparallel_task_assignment.cc146 instruction->shape().IsTuple() || opcode == HloOpcode::kRng || in GetTargetParallelTaskCount()
147 opcode == HloOpcode::kConstant) { in GetTargetParallelTaskCount()
153 opcode == HloOpcode::kBroadcast || opcode == HloOpcode::kConcatenate || in GetTargetParallelTaskCount()
154 opcode == HloOpcode::kDynamicSlice || in GetTargetParallelTaskCount()
155 opcode == HloOpcode::kDynamicUpdateSlice || in GetTargetParallelTaskCount()
156 opcode == HloOpcode::kGather || opcode == HloOpcode::kIota || in GetTargetParallelTaskCount()
157 opcode == HloOpcode::kPad || opcode == HloOpcode::kReduce || in GetTargetParallelTaskCount()
158 opcode == HloOpcode::kReduceWindow || opcode == HloOpcode::kReshape || in GetTargetParallelTaskCount()
159 opcode == HloOpcode::kReverse || opcode == HloOpcode::kSlice || in GetTargetParallelTaskCount()
160 opcode == HloOpcode::kTranspose || in GetTargetParallelTaskCount()
[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/service/gpu/
Dgpu_fusible.cc42 if (instr.opcode() == HloOpcode::kFusion) { in AppendParams()
53 CHECK_NE(instr.opcode(), HloOpcode::kFusion) << "`instr` has to be unfused."; in IfFusedReadsElementsMultipleTimes()
54 if (instr.opcode() == HloOpcode::kReduce && in IfFusedReadsElementsMultipleTimes()
60 if (instr.opcode() == HloOpcode::kReduceWindow) { in IfFusedReadsElementsMultipleTimes()
125 } else if (instr.opcode() == HloOpcode::kFusion && in IsReduceInputFusion()
138 if (instr.opcode() == HloOpcode::kReduce && instr.shape().IsTuple()) { in IsInputFusibleReduction()
148 if (instr.opcode() != HloOpcode::kFusion) { in GetRealHeroForMultiOutputFusion()
195 if (instr.opcode() == HloOpcode::kScatter || in IsInputFusibleScatter()
196 (instr.opcode() == HloOpcode::kFusion && in IsInputFusibleScatter()
198 instr.fused_expression_root()->opcode() == HloOpcode::kScatter)) { in IsInputFusibleScatter()
[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 …]

12345678910>>...12