Home
last modified time | relevance | path

Searched refs:dynamic_update_slice (Results 1 – 18 of 18) sorted by relevance

/external/tensorflow/tensorflow/compiler/xla/service/llvm_ir/
Ddynamic_update_slice_util.cc57 bool CanUpdateDynamicSliceInPlace(HloInstruction* dynamic_update_slice, in CanUpdateDynamicSliceInPlace() argument
59 CHECK_EQ(HloOpcode::kDynamicUpdateSlice, dynamic_update_slice->opcode()); in CanUpdateDynamicSliceInPlace()
60 const HloInstruction* operand = dynamic_update_slice->operand(0); in CanUpdateDynamicSliceInPlace()
61 return assignment.HasTopLevelAllocation(dynamic_update_slice) && in CanUpdateDynamicSliceInPlace()
63 assignment.SharesTopLevelSlice(dynamic_update_slice, operand); in CanUpdateDynamicSliceInPlace()
197 auto* dynamic_update_slice = fusion->root_instruction(); in EmitFusedDynamicUpdateSliceInPlaceImpl() local
199 const auto* update = dynamic_update_slice->operand(1); in EmitFusedDynamicUpdateSliceInPlaceImpl()
200 const auto* start_indices = dynamic_update_slice->operand(2); in EmitFusedDynamicUpdateSliceInPlaceImpl()
217 dynamic_update_slice->shape(), &update_shape)); in EmitFusedDynamicUpdateSliceInPlaceImpl()
227 fused_emitter->GetGenerator(dynamic_update_slice->operand(2 + index))); in EmitFusedDynamicUpdateSliceInPlaceImpl()
[all …]
Ddynamic_update_slice_util.h55 bool CanUpdateDynamicSliceInPlace(HloInstruction* dynamic_update_slice,
/external/tensorflow/tensorflow/compiler/tests/
Ddynamic_slice_ops_test.py47 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/
Dhlo_cost_analysis.cc271 const HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
273 GetShapeSize(dynamic_update_slice->operand(1)->shape()) * 2 + in HandleDynamicUpdateSlice()
274 GetShapeSize(dynamic_update_slice->operand(2)->shape()); in HandleDynamicUpdateSlice()
277 GetShapeSize(dynamic_update_slice->operand(1)->shape())); in HandleDynamicUpdateSlice()
280 1, GetShapeSize(dynamic_update_slice->operand(1)->shape())); in HandleDynamicUpdateSlice()
282 2, GetShapeSize(dynamic_update_slice->operand(2)->shape())); in HandleDynamicUpdateSlice()
Ddfs_hlo_visitor_with_default.h174 HloInstructionPtr dynamic_update_slice) override { in HandleDynamicUpdateSlice() argument
175 return DefaultAction(dynamic_update_slice); in HandleDynamicUpdateSlice()
Dhlo_evaluator_typed_visitor.h1662 HloInstruction* dynamic_update_slice) override {
1663 auto operand = dynamic_update_slice->operand(0);
1664 auto update = dynamic_update_slice->operand(1);
1665 auto start_indices = dynamic_update_slice->operand(2);
1666 auto result_shape = dynamic_update_slice->shape();
1671 Cast<HloDynamicUpdateSliceInstruction>(dynamic_update_slice)
1687 parent_->evaluated_[dynamic_update_slice],
1690 absl::MakeConstSpan(dynamic_update_slice->operands())
1695 parent_->evaluated_[dynamic_update_slice],
1698 absl::MakeConstSpan(dynamic_update_slice->operands())
[all …]
Dhlo_dataflow_analysis_test.cc2131 auto dynamic_update_slice = in TEST_F() local
2136 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
2140 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
2174 auto dynamic_update_slice = in TEST_F() local
2179 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
2183 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
2412 auto dynamic_update_slice = in TEST_F() local
2417 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
2421 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
2453 auto dynamic_update_slice = in TEST_F() local
[all …]
Dhlo_cost_analysis.h107 const HloInstruction* dynamic_update_slice) override;
Dtuple_points_to_analysis_test.cc940 auto dynamic_update_slice = in TEST_F() local
944 HloInstruction::CreateTuple({gte0, dynamic_update_slice})); in TEST_F()
948 {dynamic_update_slice, starts, update, gte1}, in TEST_F()
Dhlo_verifier.cc855 HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
857 dynamic_update_slice, in HandleDynamicUpdateSlice()
859 dynamic_update_slice->operand(0)->shape(), in HandleDynamicUpdateSlice()
860 dynamic_update_slice->operand(1)->shape(), in HandleDynamicUpdateSlice()
861 Cast<HloDynamicUpdateSliceInstruction>(dynamic_update_slice) in HandleDynamicUpdateSlice()
Dhlo_verifier.h90 HloInstruction* dynamic_update_slice) override;
Dalgebraic_simplifier.cc358 HloInstruction* dynamic_update_slice) override;
4423 HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
4429 HloInstruction* updated = dynamic_update_slice->mutable_operand(0); in HandleDynamicUpdateSlice()
4430 HloInstruction* dus_update = dynamic_update_slice->mutable_operand(1); in HandleDynamicUpdateSlice()
4436 auto update_start_indx = dynamic_update_slice->operand(2); in HandleDynamicUpdateSlice()
4443 update_start_indx = dynamic_update_slice; in HandleDynamicUpdateSlice()
4489 VLOG(2) << dynamic_update_slice->ToString(); in HandleDynamicUpdateSlice()
4492 << dynamic_update_slice->parent()->ToString(); in HandleDynamicUpdateSlice()
4493 return ReplaceInstruction(dynamic_update_slice, pad); in HandleDynamicUpdateSlice()
4499 if (SameShape(dynamic_update_slice, dus_update)) { in HandleDynamicUpdateSlice()
[all …]
Dhlo_alias_analysis_test.cc1078 auto dynamic_update_slice = builder.AddInstruction( in TEST_F() local
1087 analysis.GetUniqueBufferAt(dynamic_update_slice)); in TEST_F()
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Dir_emitter.cc1978 HloInstruction* dynamic_update_slice) { in HandleDynamicUpdateSlice() argument
1979 auto update = dynamic_update_slice->operand(1); in HandleDynamicUpdateSlice()
1980 if (ShapeUtil::IsScalar(dynamic_update_slice->shape())) { in HandleDynamicUpdateSlice()
1981 TF_RETURN_IF_ERROR(EmitTargetAddressForOp(dynamic_update_slice)); in HandleDynamicUpdateSlice()
1982 return EmitMemcpy(*update, *dynamic_update_slice); in HandleDynamicUpdateSlice()
1983 } else if (llvm_ir::CanUpdateDynamicSliceInPlace(dynamic_update_slice, in HandleDynamicUpdateSlice()
1985 TF_RETURN_IF_ERROR(EmitTargetAddressForOp(dynamic_update_slice)); in HandleDynamicUpdateSlice()
1986 auto operands = GetIrArraysForOperandsOf(dynamic_update_slice); in HandleDynamicUpdateSlice()
1988 operands, GetIrArrayFor(dynamic_update_slice), in HandleDynamicUpdateSlice()
1989 IrName(dynamic_update_slice, "in_place"), &b_); in HandleDynamicUpdateSlice()
[all …]
Dir_emitter.h156 HloInstruction* dynamic_update_slice) override;
/external/tensorflow/tensorflow/compiler/mlir/xla/tests/
Dlegalize-tf-with-tf2xla.mlir169 // CHECK-LABEL: dynamic_update_slice
171 func @dynamic_update_slice(%arg0: tensor<3x4xi32>, %arg1: tensor<2x2xi32>, %arg2: tensor<2xi32>) ->…
/external/tensorflow/tensorflow/compiler/tf2xla/python/
Dxla.py321 dynamic_update_slice = gen_xla_ops.xla_dynamic_update_slice variable
/external/tensorflow/tensorflow/compiler/mlir/hlo/tests/
Dops.mlir805 // CHECK-LABEL: @dynamic_update_slice
806 func @dynamic_update_slice(%input: tensor<3x4xi64>, %update: tensor<2xi64>, %start1: tensor<i64>, %…