| /external/tensorflow/tensorflow/compiler/xla/service/ |
| bfloat16_support.cc | 25 case HloOpcode::kCall:
26 case HloOpcode::kConditional:
27 case HloOpcode::kCustomCall:
28 case HloOpcode::kGetTupleElement:
29 case HloOpcode::kTuple:
30 case HloOpcode::kWhile:
32 case HloOpcode::kConvert:
43 case HloOpcode::kCall:
44 case HloOpcode::kConditional:
45 case HloOpcode::kCustomCall
[all...] |
| hlo_opcode_test.cc | 24 // This test verifies that an example HloOpcode stringifies as expected.
26 ASSERT_EQ("multiply", HloOpcodeString(HloOpcode::kMultiply));
39 auto opcode = static_cast<HloOpcode>(i);
45 case HloOpcode::kEq:
46 case HloOpcode::kNe:
47 case HloOpcode::kGt:
48 case HloOpcode::kLt:
49 case HloOpcode::kGe:
50 case HloOpcode::kLe:
57 case HloOpcode::kCall
[all...] |
| instruction_fusion.cc | 35 case HloOpcode::kAdd:
36 case HloOpcode::kAnd:
37 case HloOpcode::kBitcast:
38 case HloOpcode::kBitcastConvert:
39 case HloOpcode::kBroadcast:
40 case HloOpcode::kCeil:
41 case HloOpcode::kClamp:
42 case HloOpcode::kComplex:
43 case HloOpcode::kConcatenate:
44 case HloOpcode::kConstant
[all...] |
| hlo_opcode.cc | 23 string HloOpcodeString(HloOpcode opcode) {
26 case HloOpcode::enum_name: \
33 StatusOr<HloOpcode> StringToHloOpcode(const string& opcode_name) {
34 static auto* opcode_map = new tensorflow::gtl::FlatMap<string, HloOpcode>({
36 {opcode_name, HloOpcode::enum_name},
54 bool HloOpcodeIsComparison(HloOpcode opcode) {
57 case HloOpcode::enum_name: \
64 bool HloOpcodeIsVariadic(HloOpcode opcode) {
67 case HloOpcode::enum_name: \
|
| hlo_instruction.cc | 59 TF_ASSIGN_OR_RETURN(HloOpcode opcode, StringToHloOpcode(proto.opcode()));
78 if (instruction->opcode() == HloOpcode::kFusion) {
159 WrapUnique(new HloInstruction(HloOpcode::kParameter, shape));
168 WrapUnique(new HloInstruction(HloOpcode::kTrace, ShapeUtil::MakeNil()));
177 WrapUnique(new HloInstruction(HloOpcode::kConstant, literal->shape()));
186 WrapUnique(new HloInstruction(HloOpcode::kGetTupleElement, shape));
195 auto instruction = WrapUnique(new HloInstruction(HloOpcode::kRng, shape));
205 const Shape& shape, HloOpcode opcode,
207 if (opcode == HloOpcode::kCopy) {
219 const Shape& shape, HloOpcode opcode, HloInstruction* operand)
[all...] |
| hlo_constant_folding.cc | 55 if (instruction->opcode() == HloOpcode::kParameter ||
56 instruction->opcode() == HloOpcode::kConstant ||
57 instruction->opcode() == HloOpcode::kTuple ||
58 instruction->opcode() == HloOpcode::kReduce) {
69 if (instruction->opcode() == HloOpcode::kBroadcast) {
|
| bfloat16_conversion_folding.cc | 69 CHECK_EQ(user->opcode(), HloOpcode::kConvert);
80 CHECK_EQ(operand->opcode(), HloOpcode::kConvert);
94 if (operand->opcode() == HloOpcode::kConvert &&
113 if (user->opcode() == HloOpcode::kConvert &&
146 if (hlo->opcode() == HloOpcode::kTuple || //
147 hlo->opcode() == HloOpcode::kGetTupleElement || //
148 hlo->opcode() == HloOpcode::kInfeed || //
149 hlo->opcode() == HloOpcode::kOutfeed || //
150 hlo->opcode() == HloOpcode::kConstant || //
151 hlo->opcode() == HloOpcode::kParameter || /
[all...] |
| tuple_simplifier.cc | 38 if (instruction->opcode() == HloOpcode::kTuple ||
39 instruction->opcode() == HloOpcode::kGetTupleElement) {
58 if (instruction->opcode() == HloOpcode::kTuple) {
76 if (operand->opcode() != HloOpcode::kGetTupleElement ||
100 CHECK_EQ(instruction->opcode(), HloOpcode::kGetTupleElement);
110 if (instruction->operand(0)->opcode() == HloOpcode::kTuple) {
117 if (user->opcode() == HloOpcode::kTuple ||
118 user->opcode() == HloOpcode::kGetTupleElement) {
|
| hlo_graph_dumper.cc | 132 CHECK_EQ(instr->opcode(), HloOpcode::kFusion);
251 if (operand0->opcode() != HloOpcode::kParameter ||
252 operand1->opcode() != HloOpcode::kParameter) {
268 case HloOpcode::kLe:
269 case HloOpcode::kGe:
270 case HloOpcode::kGt:
271 case HloOpcode::kLt:
287 case HloOpcode::kAdd:
289 case HloOpcode::kMultiply:
291 case HloOpcode::kMinimum
[all...] |
| hlo_query.cc | 26 if (instruction->opcode() == HloOpcode::kConstant &&
37 if (operand->opcode() != HloOpcode::kParameter &&
38 operand->opcode() != HloOpcode::kConstant) {
47 if (operand->opcode() != HloOpcode::kParameter) {
56 if (operand->opcode() != HloOpcode::kConstant) {
92 bool MatchBinaryInstructionOperandOpcode(HloOpcode opcode,
|
| bfloat16_normalization_test.cc | 37 if (hlo.opcode() == HloOpcode::kAdd ||
38 hlo.opcode() == HloOpcode::kSubtract ||
39 hlo.opcode() == HloOpcode::kReduce ||
40 hlo.opcode() == HloOpcode::kTuple ||
41 hlo.opcode() == HloOpcode::kGetTupleElement) {
48 if (hlo.opcode() == HloOpcode::kAdd || hlo.opcode() == HloOpcode::kReduce ||
49 hlo.opcode() == HloOpcode::kSubtract ||
50 hlo.opcode() == HloOpcode::kTuple ||
51 hlo.opcode() == HloOpcode::kGetTupleElement)
[all...] |
| bfloat16_conversion_folding_test.cc | 37 if (hlo.opcode() == HloOpcode::kAdd ||
38 hlo.opcode() == HloOpcode::kSubtract ||
39 hlo.opcode() == HloOpcode::kTuple ||
40 hlo.opcode() == HloOpcode::kGetTupleElement) {
47 if (hlo.opcode() == HloOpcode::kAdd ||
48 hlo.opcode() == HloOpcode::kSubtract ||
49 hlo.opcode() == HloOpcode::kTuple ||
50 hlo.opcode() == HloOpcode::kGetTupleElement) {
57 if (hlo.opcode() == HloOpcode::kAdd || hlo.opcode() == HloOpcode::kTuple |
[all...] |
| liveness_util.cc | 37 if (user->opcode() == HloOpcode::kGetTupleElement && !index.empty()) {
41 } else if (user->opcode() == HloOpcode::kFusion &&
78 if (user->opcode() == HloOpcode::kFusion &&
141 CHECK_EQ(HloOpcode::kFusion, fusion->opcode());
197 if (user->opcode() == HloOpcode::kFusion) {
200 HloOpcode::kDynamicUpdateSlice) {
209 user->fused_expression_root()->opcode() == HloOpcode::kAdd) {
218 return operand->opcode() == HloOpcode::kConvolution ||
219 operand->opcode() == HloOpcode::kDot ||
220 (operand->opcode() == HloOpcode::kFusion &
[all...] |
| reduce_precision_insertion_test.cc | 55 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a));
66 return instruction->opcode() == HloOpcode::kCos;
85 HloInstruction::CreateBinary(shape, HloOpcode::kAdd, a, b));
97 return instruction->opcode() == HloOpcode::kAdd;
114 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a));
126 HloOpcode::kParameter;
143 HloInstruction::CreateUnary(shape, HloOpcode::kCos, a));
145 HloInstruction::CreateUnary(shape, HloOpcode::kSin, a));
147 HloInstruction::CreateBinary(shape, HloOpcode::kAdd, b, c));
159 return instruction->opcode() == HloOpcode::kCos |
[all...] |
| hlo_element_type_converter.cc | 127 if (opcode == HloOpcode::kParameter || opcode == HloOpcode::kConstant ||
128 opcode == HloOpcode::kTuple || opcode == HloOpcode::kConvert ||
129 opcode == HloOpcode::kGetTupleElement ||
130 opcode == HloOpcode::kInfeed || opcode == HloOpcode::kOutfeed) {
136 if (opcode == HloOpcode::kCustomCall) {
142 if (opcode == HloOpcode::kWhile || opcode == HloOpcode::kCall |
[all...] |
| reshape_mover.cc | 52 return instruction->opcode() == HloOpcode::kReshape ||
53 instruction->opcode() == HloOpcode::kTranspose;
80 if (instruction->opcode() == HloOpcode::kConstant) {
87 if (instruction->opcode() == HloOpcode::kRng &&
120 case HloOpcode::kTranspose:
122 case HloOpcode::kReshape:
200 case HloOpcode::kConstant: {
201 if (first_reshape_operand->opcode() == HloOpcode::kReshape) {
206 CHECK(first_reshape_operand->opcode() == HloOpcode::kTranspose);
214 case HloOpcode::kRng:
[all...] |
| defuser_test.cc | 52 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1));
67 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1));
90 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1));
92 HloInstruction::CreateUnary(shape_, HloOpcode::kNegate, add));
117 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, param0, param1));
119 HloInstruction::CreateUnary(shape_, HloOpcode::kNegate, add));
121 HloInstruction::CreateBinary(shape_, HloOpcode::kSubtract, add, negate));
123 HloInstruction::CreateBinary(shape_, HloOpcode::kMultiply, sub, param3));
125 HloInstruction::CreateBinary(shape_, HloOpcode::kDivide, mul, param3));
129 HloInstruction::CreateBinary(shape_, HloOpcode::kAdd, constant, div))
[all...] |
| batchnorm_expander.cc | 85 HloOpcode opcode) {
202 GetScalarBinaryComputation(ptype, HloOpcode::kAdd);
206 operand_shape, HloOpcode::kMultiply, operand, operand));
233 feature_shape, HloOpcode::kDivide, sum, elements_per_feature));
240 feature_shape, HloOpcode::kDivide, squared_sum, elements_per_feature));
244 feature_shape, HloOpcode::kMultiply, mean, mean));
248 feature_shape, HloOpcode::kSubtract, square_mean, mean_square));
255 operand_shape, HloOpcode::kAdd, var_broadcasted, epsilon));
264 operand_shape, HloOpcode::kPower, var_add_epsilon, neg_half));
268 operand_shape, HloOpcode::kSubtract, operand, mean_broadcasted))
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/gpu/ |
| instruction_fusion.cc | 29 hlo.opcode() == HloOpcode::kBroadcast ||
30 hlo.opcode() == HloOpcode::kConcatenate ||
31 hlo.opcode() == HloOpcode::kDynamicSlice ||
32 hlo.opcode() == HloOpcode::kDynamicUpdateSlice ||
33 hlo.opcode() == HloOpcode::kFusion ||
34 hlo.opcode() == HloOpcode::kGetTupleElement ||
35 hlo.opcode() == HloOpcode::kPad ||
36 hlo.opcode() == HloOpcode::kReduce ||
37 hlo.opcode() == HloOpcode::kReduceWindow ||
38 hlo.opcode() == HloOpcode::kReshape |
[all...] |
| stream_assignment_test.cc | 44 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, x, y));
46 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, dot1, z));
63 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, x, y));
65 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, y, x));
67 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kAdd, dot1, dot2));
95 f32_2x2_, HloOpcode::kDot, params[2], params[3]));
97 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, params[1], d00));
99 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, d00, params[4]));
101 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, params[0], d10));
103 HloInstruction::CreateBinary(f32_2x2_, HloOpcode::kDot, d10, d11))
[all...] |
| while_transformer.cc | 39 // Each ExprTree node is comprised of an HloOpcode, and a set of operands (each
41 // HloOpcode of the operand.
53 // ExprTree add(HloOpcode::kAdd,
54 // ExprTree(HloOpcode::kConstant),
55 // ExprTree(HloOpcode::kGetTupleElement,
56 // tuple_index, ExprTree(HloOpcode::kParameter)));
70 explicit ExprTree(HloOpcode opcode) : opcode_(opcode) {}
71 ExprTree(HloOpcode opcode, const string& tag) : opcode_(opcode), tag_(tag) {}
72 ExprTree(HloOpcode opcode, const ExprTree& operand0) : opcode_(opcode) {
75 ExprTree(HloOpcode opcode, int64 index0, const ExprTree& operand0
[all...] |
| fusion_merger_test.cc | 66 data_shape_, HloOpcode::kAdd, one_vec, gte0));
71 HloInstruction::CreateBinary(data_shape_, HloOpcode::kAdd, gte1, gte2));
77 data_shape_, HloOpcode::kAdd, add1, one_vec));
79 data_shape_, HloOpcode::kMultiply, add2, one_vec));
83 data_shape_, HloOpcode::kMultiply, add1, one_vec));
85 data_shape_, HloOpcode::kAdd, mul0, one_vec));
133 data_shape_, HloOpcode::kMultiply, gte0, gte1));
136 data_shape_, HloOpcode::kMultiply, gte0, mul0));
145 data_shape_, HloOpcode::kAdd, mul0, one_vec));
147 data_shape_, HloOpcode::kMultiply, add0, one_vec))
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/service/cpu/ |
| cpu_instruction_fusion.cc | 33 hlo.opcode() == HloOpcode::kBitcast ||
34 hlo.opcode() == HloOpcode::kBroadcast ||
35 hlo.opcode() == HloOpcode::kConcatenate ||
36 hlo.opcode() == HloOpcode::kDynamicSlice ||
37 hlo.opcode() == HloOpcode::kDynamicUpdateSlice ||
38 hlo.opcode() == HloOpcode::kPad ||
39 hlo.opcode() == HloOpcode::kReshape ||
40 hlo.opcode() == HloOpcode::kReverse ||
41 hlo.opcode() == HloOpcode::kSlice ||
42 hlo.opcode() == HloOpcode::kTranspose
[all...] |
| cpu_instruction_fusion_test.cc | 50 ShapeUtil::MakeShape(S32, {1024, 256}), HloOpcode::kExp, arg0));
69 ShapeUtil::MakeShape(S32, {256, 1024}), HloOpcode::kExp, arg1));
88 ShapeUtil::MakeShape(S32, {2, 512, 2, 128}), HloOpcode::kExp, arg0));
90 ShapeUtil::MakeShape(S32, {1024, 256}), HloOpcode::kBitcast, exp0));
109 ShapeUtil::MakeShape(S32, {2, 512, 2, 128}), HloOpcode::kExp, arg0));
131 ShapeUtil::MakeShape(S32, {256, 32 * 1024}), HloOpcode::kExp, arg1));
150 ShapeUtil::MakeShape(S32, {256, 1024}), HloOpcode::kExp, arg1));
169 ShapeUtil::MakeShape(S32, {1024, 256}), HloOpcode::kExp, arg1));
185 EXPECT_EQ(computation->root_instruction()->opcode(), HloOpcode::kFusion);
189 EXPECT_EQ(computation->root_instruction()->opcode(), HloOpcode::kFusion)
[all...] |
| /external/tensorflow/tensorflow/compiler/xla/legacy_flags/ |
| debug_options_parsers_test.cc | 50 EXPECT_EQ(static_cast<HloOpcode>(proto.opcodes_to_suffix(0)),
51 HloOpcode::kAdd);
52 EXPECT_EQ(static_cast<HloOpcode>(proto.opcodes_to_suffix(1)),
53 HloOpcode::kDot);
65 EXPECT_EQ(static_cast<HloOpcode>(proto.opcodes_to_suffix(0)),
66 HloOpcode::kAdd);
67 EXPECT_EQ(static_cast<HloOpcode>(proto.opcodes_to_suffix(1)),
68 HloOpcode::kDot);
93 EXPECT_EQ(static_cast<HloOpcode>(proto.opcodes_to_suffix(0)),
94 HloOpcode::kSubtract)
[all...] |
