Home
last modified time | relevance | path

Searched refs:operand_idx (Results 1 – 25 of 26) sorted by relevance

12

/aosp_15_r20/external/tensorflow/tensorflow/compiler/xla/service/cpu/
H A Dcpu_layout_assignment.cc47 optional<int64_t> operand_idx = in ShouldMakeAllUsersColMajor() local
49 if (!operand_idx || user->operand(*operand_idx) != instruction || in ShouldMakeAllUsersColMajor()
59 optional<int64_t> operand_idx = in ShouldMakeOperandColumnMajor() local
61 if (!operand_idx) { in ShouldMakeOperandColumnMajor()
65 const HloInstruction* operand = instruction.operand(*operand_idx); in ShouldMakeOperandColumnMajor()
78 return it->second ? operand_idx : nullopt; in ShouldMakeOperandColumnMajor()
/aosp_15_r20/external/tensorflow/tensorflow/compiler/xla/service/
H A Dmemory_space_propagation.cc38 for (int operand_idx = 0; operand_idx < instruction->operand_count(); in Run() local
39 ++operand_idx) { in Run()
42 instruction->operand(operand_idx)->shape())) { in Run()
45 instruction->fused_parameter(operand_idx), in Run()
H A Dtranspose_folding.cc254 const HloInstruction& dot, int64_t operand_idx) { in IsRowColumnTransposeDotOperand() argument
256 TF_RET_CHECK(dot.operand_count() > operand_idx); in IsRowColumnTransposeDotOperand()
258 const HloInstruction& transpose = *dot.operand(operand_idx); in IsRowColumnTransposeDotOperand()
263 auto batch_dims = (operand_idx == 0) ? dot_dims.lhs_batch_dimensions() in IsRowColumnTransposeDotOperand()
266 auto contracting_dims = (operand_idx == 0) in IsRowColumnTransposeDotOperand()
H A Dbfloat16_normalization.cc79 Status InsertConvertBeforeOperand(HloInstruction* hlo, int64_t operand_idx,
218 HloInstruction* hlo, int64_t operand_idx, PrimitiveType from, in InsertConvertBeforeOperand() argument
220 auto operand = hlo->mutable_operand(operand_idx); in InsertConvertBeforeOperand()
227 hlo->ReplaceOperandWithDifferentShape(operand_idx, new_operand)); in InsertConvertBeforeOperand()
H A Dhlo_instruction_test.cc1236 for (int64_t operand_idx = 0; operand_idx < fusion->operand_count(); in TEST_F() local
1237 ++operand_idx) { in TEST_F()
1238 const HloInstruction* operand = fusion->operand(operand_idx); in TEST_F()
1240 EXPECT_FALSE(fusion->IsElementwiseOnOperand(operand_idx)); in TEST_F()
1242 EXPECT_TRUE(fusion->IsElementwiseOnOperand(operand_idx)); in TEST_F()
1277 for (int64_t operand_idx = 0; operand_idx < fusion->operand_count(); in TEST_F() local
1278 ++operand_idx) { in TEST_F()
1279 if (fusion->operand(operand_idx) == y) { in TEST_F()
1280 EXPECT_FALSE(fusion->IsElementwiseOnOperand(operand_idx)); in TEST_F()
1282 EXPECT_TRUE(fusion->IsElementwiseOnOperand(operand_idx)); in TEST_F()
H A Delemental_ir_emitter.cc2885 for (int64_t operand_idx = 0; operand_idx < input_count; ++operand_idx) { in EmitElementalReduceWindow() local
2887 input_generators[operand_idx](input_index)); in EmitElementalReduceWindow()
2888 input_values[input_count + operand_idx] = input_value; in EmitElementalReduceWindow()
2889 input_values[operand_idx] = in EmitElementalReduceWindow()
2890 Load(llvm::cast<llvm::AllocaInst>(accum_ptrs[operand_idx]) in EmitElementalReduceWindow()
2892 accum_ptrs[operand_idx]); in EmitElementalReduceWindow()
2899 for (int64_t operand_idx = 0; operand_idx < accum_values.size(); in EmitElementalReduceWindow() local
2900 ++operand_idx) { in EmitElementalReduceWindow()
2901 Store(accum_values[operand_idx], accum_ptrs[operand_idx]); in EmitElementalReduceWindow()
H A Dtranspose_folding.h71 const HloInstruction& dot, int64_t operand_idx);
H A Dhlo_dataflow_analysis.cc196 int64_t operand_idx = concat.operand_index(&operand); in ConcatIsEffectivelyElementwise() local
204 if (!is_concat_identical || info.concat_opnd_idx != operand_idx) { in ConcatIsEffectivelyElementwise()
209 const HloInstruction* slice_to_recover_opnd = users.at(operand_idx); in ConcatIsEffectivelyElementwise()
211 ConcatUsageInfo{&concat, operand_idx, slice_to_recover_opnd}); in ConcatIsEffectivelyElementwise()
H A Dhlo_instructions.cc1352 const std::optional<int64_t>& operand_idx) const { in IsElementwiseImpl()
1468 const std::optional<int64_t>& operand_idx) const { in IsElementwiseImpl()
1879 const std::optional<int64_t>& operand_idx) const { in IsElementwiseImpl()
1880 if (!operand_idx.has_value()) { in IsElementwiseImpl()
1892 worklist.push_back(fused_parameter(operand_idx.value())); in IsElementwiseImpl()
1893 visited.insert(fused_parameter(operand_idx.value())); in IsElementwiseImpl()
2180 const std::optional<int64_t>& operand_idx) const { in IsElementwiseImpl()
H A Dhlo_instructions.h1030 const std::optional<int64_t>& operand_idx) const override;
1125 const std::optional<int64_t>& operand_idx) const override;
1320 const std::optional<int64_t>& operand_idx) const override;
1372 const std::optional<int64_t>& operand_idx) const override;
H A Dhlo_instruction.h1651 bool IsElementwiseOnOperand(int64_t operand_idx) const;
2274 const std::optional<int64_t>& operand_idx) const;
H A Dhlo_instruction.cc3036 const std::optional<int64_t>& operand_idx) const { in IsElementwiseImpl()
3038 return operand_idx.has_value() && operand_idx.value() == 0; in IsElementwiseImpl()
3881 bool HloInstruction::IsElementwiseOnOperand(int64_t operand_idx) const { in IsElementwiseOnOperand()
3882 return IsElementwiseImpl(operand_idx); in IsElementwiseOnOperand()
H A Dhlo_verifier.cc921 int64_t operand_idx = parameter_idx / 2; in HandleSort() local
923 sort->operand(operand_idx)->shape().element_type(), {}); in HandleSort()
/aosp_15_r20/external/tensorflow/tensorflow/compiler/xla/service/gpu/
H A Dkernel_mapping_scheme.h213 const HloInstruction* instruction, int operand_idx) const { in GetCalculationStateFor() argument
215 CHECK_LT(operand_idx, op_state.size()); in GetCalculationStateFor()
216 return op_state[operand_idx]; in GetCalculationStateFor()
221 const HloInstruction* instruction, int operand_idx) { in SetCalculationStateFor() argument
223 CHECK_EQ(operand_idx, op_state.size()); in SetCalculationStateFor()
H A Dreduction_layout_normalizer.cc42 int operand_idx = -1; in HandleReduce() local
51 operand_idx++; in HandleReduce()
53 if (operand_idx != 0 && in HandleReduce()
73 reduce->shape().IsTuple() ? reduce->shape().tuple_shapes(operand_idx) in HandleReduce()
H A Dir_emission_utils.cc787 for (int64_t operand_idx = 0; operand_idx < operand_shapes.size(); in FindTranspose021DimsAndParameters() local
788 ++operand_idx) { in FindTranspose021DimsAndParameters()
790 ShapeUtil::FindTranspose021(operand_shapes[operand_idx], output_shape); in FindTranspose021DimsAndParameters()
804 params_012.push_back(operand_idx); in FindTranspose021DimsAndParameters()
H A Dmatmul_utils.cc198 int64_t operand_idx) { in CanFoldTransposeOperandIntoDot() argument
200 TF_RET_CHECK(dot.operand_count() > operand_idx); in CanFoldTransposeOperandIntoDot()
202 const HloInstruction& transpose = *dot.operand(operand_idx); in CanFoldTransposeOperandIntoDot()
216 auto batch_dims = (operand_idx == 0) ? dot_dims.lhs_batch_dimensions() in CanFoldTransposeOperandIntoDot()
218 auto contracting_dims = (operand_idx == 0) in CanFoldTransposeOperandIntoDot()
H A Dmatmul_utils.h87 int64_t operand_idx);
/aosp_15_r20/external/tensorflow/tensorflow/compiler/mlir/tensorflow/ir/
H A Dtf_executor.cc143 for (int operand_idx : llvm::seq<int>(0, op->getNumOperands())) { in VerifyControlOperandsAfterAllData() local
144 if (op->getOperand(operand_idx).getType().isa<ControlType>()) { in VerifyControlOperandsAfterAllData()
149 return op->emitOpError() << "found non-control operand #" << operand_idx in VerifyControlOperandsAfterAllData()
294 for (int operand_idx : llvm::seq<int>(0, yield.getNumOperands())) { in verify() local
295 if (island.getResult(operand_idx).getType() != in verify()
296 yield.getOperand(operand_idx).getType()) in verify()
298 << "operand #" << operand_idx << " type mismatch island results"; in verify()
303 for (int operand_idx : llvm::seq<int>(0, island.getNumResults() - 1)) { in verify() local
304 if (island.getResult(operand_idx).getType() == control_type) in verify()
306 << "unexpected control type for operand #" << operand_idx; in verify()
H A Dtf_device.cc559 int32_t operand_idx = operand_type_and_idx.index(); in verify() local
560 for (int32_t i = n * operand_idx, e = i + n; i < e; ++i) in verify()
563 << " and terminator operand " << operand_idx; in verify()
H A Dtf_ops_a_m.cc1170 unsigned operand_idx = 0; in matchAndRewrite() local
1174 ++operand_idx; in matchAndRewrite()
1177 exceptions[arg] = operand_idx++; in matchAndRewrite()
1287 auto args = [&](int operand_idx) -> SmallVector<Value, 8> { in GetHoistParams() argument
1290 return arg.getDefiningOp()->getOperand(operand_idx); in GetHoistParams()
1297 auto is_all_tensors = [&](int operand_idx, int axis) -> bool { in GetHoistParams() argument
1306 operand = arg.getDefiningOp()->getOperand(operand_idx); in GetHoistParams()
1315 auto is_all_scalars = [&](int operand_idx) -> bool { in GetHoistParams() argument
1318 auto operand = arg.getDefiningOp()->getOperand(operand_idx); in GetHoistParams()
/aosp_15_r20/external/tensorflow/tensorflow/compiler/mlir/lite/ir/
H A Dtfl_ops.cc3872 for (const int operand_idx : llvm::seq<int>(0, yield.getNumOperands())) { in verify() local
3873 if (control_node.getResult(operand_idx).getType() != in verify()
3874 yield.getOperand(operand_idx).getType()) in verify()
3875 return yield.emitOpError() << "operand #" << operand_idx in verify()
/aosp_15_r20/external/cronet/third_party/rust/chromium_crates_io/vendor/syn-1.0.109/tests/common/
H A Deq.rs497 spanless_eq_enum!(InlineAsmTemplatePiece; String(0) Placeholder(operand_idx modifier span));
/aosp_15_r20/external/rust/android-crates-io/crates/syn/tests/common/
Deq.rs574 spanless_eq_enum!(InlineAsmTemplatePiece; String(0) Placeholder(operand_idx modifier span));
/aosp_15_r20/external/cronet/third_party/rust/chromium_crates_io/vendor/syn-2.0.55/tests/common/
H A Deq.rs574 spanless_eq_enum!(InlineAsmTemplatePiece; String(0) Placeholder(operand_idx modifier span));

12