| /external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/ |
| D | dynamic_update_slice_util.cc | 26 bool CanUpdateDynamicSliceInPlace(HloInstruction* dynamic_update_slice, in CanUpdateDynamicSliceInPlace() argument
28 CHECK_EQ(HloOpcode::kDynamicUpdateSlice, dynamic_update_slice->opcode()); in CanUpdateDynamicSliceInPlace()
29 const HloInstruction* operand = dynamic_update_slice->operand(0); in CanUpdateDynamicSliceInPlace()
30 return assignment.HasTopLevelAllocation(dynamic_update_slice) && in CanUpdateDynamicSliceInPlace()
32 assignment.SharesTopLevelSlice(dynamic_update_slice, operand); in CanUpdateDynamicSliceInPlace()
146 auto* dynamic_update_slice = fusion->fused_expression_root(); in EmitFusedDynamicUpdateSliceInPlaceImpl() local
148 const auto* update = dynamic_update_slice->operand(1); in EmitFusedDynamicUpdateSliceInPlaceImpl()
149 const auto* start_indices = dynamic_update_slice->operand(2); in EmitFusedDynamicUpdateSliceInPlaceImpl()
171 TF_RETURN_IF_ERROR(dynamic_update_slice->Accept(&fused_emitter)); in EmitFusedDynamicUpdateSliceInPlaceImpl()
176 fused_emitter.GetGenerator(dynamic_update_slice->operand(2 + index)); in EmitFusedDynamicUpdateSliceInPlaceImpl()
|
| D | dynamic_update_slice_util.h | 38 bool CanUpdateDynamicSliceInPlace(HloInstruction* dynamic_update_slice,
|
| /external/tensorflow/tensorflow/compiler/tests/ |
| D | dynamic_slice_ops_test.py | 47 xla.dynamic_update_slice, [
55 xla.dynamic_update_slice, [
63 xla.dynamic_update_slice, [
73 xla.dynamic_update_slice, [
83 xla.dynamic_update_slice, [
|
| /external/tensorflow/tensorflow/compiler/xla/service/ |
| D | buffer_liveness_test.cc | 642 auto dynamic_update_slice = in BuildModule() local
647 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in BuildModule()
655 {dynamic_update_slice, starts, update, CHECK_NOTNULL(slice), gte1}, in BuildModule()
659 {dynamic_update_slice, starts, update, gte1}, in BuildModule()
798 auto dynamic_update_slice = in Run() local
803 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in Run()
|
| D | dfs_hlo_visitor_with_default.h | 155 HloInstructionPtr dynamic_update_slice) override { in HandleDynamicUpdateSlice() argument
156 return DefaultAction(dynamic_update_slice); in HandleDynamicUpdateSlice()
|
| D | hlo_evaluator_typed_visitor.h | 1518 HloInstruction* dynamic_update_slice) override { in HandleDynamicUpdateSlice() argument
1519 auto operand = dynamic_update_slice->operand(0); in HandleDynamicUpdateSlice()
1520 auto update = dynamic_update_slice->operand(1); in HandleDynamicUpdateSlice()
1521 auto start_indices = dynamic_update_slice->operand(2); in HandleDynamicUpdateSlice()
1522 auto result_shape = dynamic_update_slice->shape(); in HandleDynamicUpdateSlice()
1527 Cast<HloDynamicUpdateSliceInstruction>(dynamic_update_slice) in HandleDynamicUpdateSlice()
1543 parent_->evaluated_[dynamic_update_slice], in HandleDynamicUpdateSlice()
1546 absl::MakeConstSpan(dynamic_update_slice->operands()) in HandleDynamicUpdateSlice()
1551 parent_->evaluated_[dynamic_update_slice], in HandleDynamicUpdateSlice()
1554 absl::MakeConstSpan(dynamic_update_slice->operands()) in HandleDynamicUpdateSlice()
[all …]
|
| D | tuple_points_to_analysis_test.cc | 889 auto dynamic_update_slice = in TEST_F() local
893 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
897 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
983 auto dynamic_update_slice = in TEST_F() local
987 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
991 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
|
| D | hlo_dataflow_analysis_test.cc | 2005 auto dynamic_update_slice = in TEST_F() local
2010 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
2014 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
2048 auto dynamic_update_slice = in TEST_F() local
2053 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
2057 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
2316 auto dynamic_update_slice = in TEST_F() local
2321 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
2325 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
2357 auto dynamic_update_slice = in TEST_F() local
[all …]
|
| D | hlo_verifier.cc | 585 HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
587 dynamic_update_slice, in HandleDynamicUpdateSlice()
589 dynamic_update_slice->operand(0)->shape(), in HandleDynamicUpdateSlice()
590 dynamic_update_slice->operand(1)->shape(), in HandleDynamicUpdateSlice()
591 Cast<HloDynamicUpdateSliceInstruction>(dynamic_update_slice) in HandleDynamicUpdateSlice()
|
| D | hlo_cost_analysis.h | 98 const HloInstruction* dynamic_update_slice) override;
|
| D | hlo_cost_analysis.cc | 204 const HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
206 GetShapeSize(dynamic_update_slice->operand(1)->shape()) * 2; in HandleDynamicUpdateSlice()
|
| D | hlo_verifier.h | 81 HloInstruction* dynamic_update_slice) override;
|
| D | algebraic_simplifier.cc | 236 HloInstruction* dynamic_update_slice) override;
2889 HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
2890 auto update = dynamic_update_slice->mutable_operand(1); in HandleDynamicUpdateSlice()
2894 if (SameShape(dynamic_update_slice, update)) { in HandleDynamicUpdateSlice()
2895 return ReplaceInstruction(dynamic_update_slice, update); in HandleDynamicUpdateSlice()
2902 return ReplaceInstruction(dynamic_update_slice, in HandleDynamicUpdateSlice()
2903 dynamic_update_slice->mutable_operand(0)); in HandleDynamicUpdateSlice()
|
| /external/tensorflow/tensorflow/compiler/xla/service/cpu/ |
| D | ir_emitter.cc | 2064 HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
2065 auto update = dynamic_update_slice->operand(1); in HandleDynamicUpdateSlice()
2066 if (ShapeUtil::IsScalar(dynamic_update_slice->shape())) { in HandleDynamicUpdateSlice()
2067 TF_RETURN_IF_ERROR(EmitTargetAddressForOp(dynamic_update_slice)); in HandleDynamicUpdateSlice()
2068 return EmitMemcpy(*update, *dynamic_update_slice); in HandleDynamicUpdateSlice()
2069 } else if (llvm_ir::CanUpdateDynamicSliceInPlace(dynamic_update_slice, in HandleDynamicUpdateSlice()
2071 TF_RETURN_IF_ERROR(EmitTargetAddressForOp(dynamic_update_slice)); in HandleDynamicUpdateSlice()
2072 auto operands = GetIrArraysForOperandsOf(dynamic_update_slice); in HandleDynamicUpdateSlice()
2074 operands, GetIrArrayFor(dynamic_update_slice), in HandleDynamicUpdateSlice()
2075 IrName(dynamic_update_slice, "in_place"), &b_); in HandleDynamicUpdateSlice()
[all …]
|
| D | ir_emitter.h | 170 HloInstruction* dynamic_update_slice) override;
|
| /external/tensorflow/tensorflow/compiler/tf2xla/python/ |
| D | xla.py | 306 dynamic_update_slice = gen_xla_ops.xla_dynamic_update_slice variable
|