Home
last modified time | relevance | path

Searched refs:reduce_window (Results 1 – 25 of 36) sorted by relevance

12

/external/tensorflow/tensorflow/compiler/xla/mlir_hlo/stablehlo/tests/
Dverify_reduce_window.mlir3 // CHECK: func @reduce_window
4 func.func @reduce_window(%arg0: tensor<4x2xf32>, %arg1: tensor<4x2xi32>,
7 %0:2 = "stablehlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
25 %0:2 = "stablehlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
44 %0 = "stablehlo.reduce_window"(%arg0, %init0) ({
60 %0 = "stablehlo.reduce_window"(%arg0, %init0) ({
78 %0:2 = "stablehlo.reduce_window"() ({
98 %0:2 = "stablehlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
119 %0:2 = "stablehlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
140 %0:2 = "stablehlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
[all …]
/external/tensorflow/tensorflow/compiler/xla/mlir_hlo/tests/Dialect/mhlo/
Dverifier_reduce_window_op.mlir3 // CHECK: func @reduce_window
4 func.func @reduce_window(%arg0: tensor<4x2xf32>, %arg1: tensor<4x2xi32>,
7 %0:2 = "mhlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
25 %0:2 = "mhlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
44 %0 = "mhlo.reduce_window"(%arg0, %init0) ({
60 %0 = "mhlo.reduce_window"(%arg0, %init0) ({
78 %0:2 = "mhlo.reduce_window"() ({
98 %0:2 = "mhlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
119 %0:2 = "mhlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
140 %0:2 = "mhlo.reduce_window"(%arg0, %arg1, %init0, %init1) ({
[all …]
Dhlo-legalize-to-lhlo.mlir557 // CHECK-LABEL: func @reduce_window
558 func.func @reduce_window(%arg0: tensor<1x17x17x64xf32>, %arg1: tensor<f32>) -> tensor<1x8x8x64xf32>…
560 // CHECK: "lmhlo.reduce_window"(%{{.+}}, %{{.+}}, %[[OUT]]) ({
569 %0 = "mhlo.reduce_window"(%arg0, %arg1) ( {
585 // CHECK: "lmhlo.reduce_window"(%{{.+}}, %{{.+}}, %{{.+}}, %{{.+}}, %[[OUT_F]], %[[OUT_I]]) ({
597 %0:2 = "mhlo.reduce_window"(%arg0, %arg1, %arg2, %arg3) ( {
Dhlo-legalize-to-linalg.mlir3409 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3436 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3461 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3486 %1 = "mhlo.reduce_window"(%arg0, %0) ({
3511 %0:2 = "mhlo.reduce_window"(%arg0, %arg0, %arg1, %arg1) ({
3548 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3587 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3612 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3637 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
3665 %0 = "mhlo.reduce_window"(%arg0, %arg1) ({
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/
Dtree_reduction_rewriter.cc97 HloInstruction *reduce_window = in HandleReduce() local
103 HloInstruction::CreateReduce(reduce_shape, reduce_window, initial_value, in HandleReduce()
Dselect_and_scatter_expander.cc142 auto* reduce_window = in ExpandInstruction() local
160 HloInstruction::CreateGetTupleElement(reduce_window, i)); in ExpandInstruction()
Dalgebraic_simplifier.cc5604 auto* reduce_window = Cast<HloReduceWindowInstruction>(hlo); in HandleReduceWindow() local
5605 const bool multi_output_reduce_window = reduce_window->shape().IsTuple(); in HandleReduceWindow()
5606 auto inputs = reduce_window->inputs(); in HandleReduceWindow()
5607 auto init_values = reduce_window->init_values(); in HandleReduceWindow()
5608 auto input_count = reduce_window->input_count(); in HandleReduceWindow()
5609 auto input_shapes = reduce_window->input_shapes(); in HandleReduceWindow()
5610 auto output_shapes = reduce_window->output_shapes(); in HandleReduceWindow()
5614 return ReplaceWithNewInstruction(reduce_window, in HandleReduceWindow()
5617 return ReplaceInstruction(reduce_window, elements[0]); in HandleReduceWindow()
5634 reduce_window->to_apply()->root_instruction()->opcode() == in HandleReduceWindow()
[all …]
Dhlo_cost_analysis.cc410 const HloInstruction* reduce_window) { in HandleReduceWindow() argument
411 const Window& window = reduce_window->window(); in HandleReduceWindow()
412 auto function = reduce_window->to_apply(); in HandleReduceWindow()
425 ShapeUtil::ElementsIn(reduce_window->shape().IsArray() in HandleReduceWindow()
426 ? reduce_window->shape() in HandleReduceWindow()
427 : reduce_window->shape().tuple_shapes(0)); in HandleReduceWindow()
Dsharding_propagation.cc1419 auto* reduce_window = Cast<HloReduceWindowInstruction>(&user); in GetShardingFromUser() local
1420 if (!absl::c_linear_search(reduce_window->inputs(), &instruction)) { in GetShardingFromUser()
1423 if (reduce_window->shape().IsTuple()) { in GetShardingFromUser()
1424 auto sub_sharding = reduce_window->sharding().GetSubSharding( in GetShardingFromUser()
1425 reduce_window->shape(), in GetShardingFromUser()
1426 {reduce_window->operand_index(&instruction)}); in GetShardingFromUser()
1429 return reduce_window->sharding(); in GetShardingFromUser()
1997 auto* reduce_window = Cast<HloReduceWindowInstruction>(instruction); in InferShardingFromOperands() local
2005 << reduce_window->ToString(); in InferShardingFromOperands()
2009 for (HloInstruction* operand : reduce_window->inputs()) { in InferShardingFromOperands()
[all …]
Ddfs_hlo_visitor_with_default.h199 Status HandleReduceWindow(HloInstructionPtr reduce_window) override { in HandleReduceWindow() argument
200 return DefaultAction(reduce_window); in HandleReduceWindow()
Dconvolution_group_converter.cc402 auto reduce_window = add(HloInstruction::CreateReduceWindow( in HandleBatchGroupCount() local
406 Shape convert_back_shape = reduce_window->shape(); in HandleBatchGroupCount()
411 HloInstruction::CreateConvert(convert_back_shape, reduce_window); in HandleBatchGroupCount()
Dhlo_verifier.cc1262 Status ShapeVerifier::HandleReduceWindow(HloInstruction* reduce_window) { in HandleReduceWindow() argument
1263 VLOG(2) << "Verify reduce window:" << reduce_window->ToString() << "\n"; in HandleReduceWindow()
1264 auto reduce_window_instr = Cast<HloReduceWindowInstruction>(reduce_window); in HandleReduceWindow()
1268 VLOG(2) << "reduce instruction is :" << reduce_window->ToString() << "\n"; in HandleReduceWindow()
1270 reduce_window, ShapeInference::InferReduceWindowShape( in HandleReduceWindow()
1271 input_shapes, init_shapes, reduce_window->window(), in HandleReduceWindow()
1272 reduce_window->to_apply()->ComputeProgramShape()))); in HandleReduceWindow()
1277 *reduce_window, reduce_window->operand_count()); in HandleReduceWindow()
Ddynamic_dimension_inference.cc1236 auto* reduce_window = Cast<HloReduceWindowInstruction>(hlo); in HandleReduceWindow() local
1238 reduce_window->window().dimensions(dimension); in HandleReduceWindow()
1240 if (operand_index >= reduce_window->input_count()) { in HandleReduceWindow()
1257 reduce_window->shape(), in HandleReduceWindow()
1259 if (!ShapeUtil::IsLeafIndex(reduce_window->shape(), in HandleReduceWindow()
1263 parent_->SetDynamicSize(reduce_window, reduce_window_result_index, in HandleReduceWindow()
Dsharding_propagation_test.cc1526 auto* reduce_window = FindInstruction(module.get(), "reduce-window"); in TEST_P() local
1527 ASSERT_NE(reduce_window, nullptr); in TEST_P()
1528 EXPECT_THAT(reduce_window, op::Sharding("{devices=[2,1]0,1}")); in TEST_P()
1532 EXPECT_THAT(reduce_window->sharding(), in TEST_P()
1536 EXPECT_THAT(reduce_window->sharding(), ShardingMetadata({})); in TEST_P()
1579 auto* reduce_window = FindInstruction(module.get(), "reduce-window"); in TEST_P() local
1580 ASSERT_NE(reduce_window, nullptr); in TEST_P()
1581 EXPECT_THAT(reduce_window, in TEST_P()
1588 EXPECT_THAT(reduce_window->sharding().tuple_elements()[0], in TEST_P()
1590 EXPECT_THAT(reduce_window->sharding().tuple_elements()[1], in TEST_P()
[all …]
Dhlo_evaluator_typed_visitor.h1758 Status HandleReduceWindow(HloInstruction* reduce_window) override {
1759 auto* reduce_window_instr = Cast<HloReduceWindowInstruction>(reduce_window);
1760 const Window& window = reduce_window->window();
1761 HloComputation* function = reduce_window->to_apply();
1769 ShapeUtil::Compatible(reduce_window->shape(), inferred_return_shape))
1771 << ShapeUtil::HumanStringWithLayout(reduce_window->shape())
1900 parent_->evaluated_[reduce_window] = std::move(result);
Dhlo_cost_analysis.h154 Status HandleReduceWindow(const HloInstruction* reduce_window) override;
Delemental_ir_emitter.cc2784 const HloReduceWindowInstruction* reduce_window, in EmitElementalReduceWindow() argument
2797 int64_t input_count = reduce_window->input_count(); in EmitElementalReduceWindow()
2803 auto operand = reduce_window->inputs()[operand_index]; in EmitElementalReduceWindow()
2827 const Window& window = reduce_window->window(); in EmitElementalReduceWindow()
2828 llvm_ir::ForLoopNest loops(IrName(reduce_window), b_, index_type); in EmitElementalReduceWindow()
2874 reduce_window->inputs()[0]->shape().dimensions(i)))); in EmitElementalReduceWindow()
2882 std::vector<llvm::Value*> input_values(reduce_window->operand_count()); in EmitElementalReduceWindow()
2884 reduce_window->inputs()[0]->shape(), index_type); in EmitElementalReduceWindow()
2895 EmitThreadLocalCall(*reduce_window->to_apply(), in EmitElementalReduceWindow()
2906 reduce_window->shape().IsTuple()); in EmitElementalReduceWindow()
Dhlo_verifier.h195 Status HandleReduceWindow(HloInstruction* reduce_window) override;
Delemental_ir_emitter.h281 const HloReduceWindowInstruction* reduce_window,
/external/tensorflow/tensorflow/compiler/tests/
Dreduce_window_test.py34 output = xla.reduce_window(placeholder, init, reducer, **kwargs)
/external/tensorflow/tensorflow/compiler/xla/mlir_hlo/tests/Dialect/lhlo/
Dlhlo-legalize-to-parallel-loops.mlir135 func.func @reduce_window(%arg: memref<112x112xf32>,
138 "lmhlo.reduce_window"(%arg, %init, %result) ({
150 // CHECK-LABEL: func @reduce_window(
/external/tensorflow/tensorflow/compiler/mlir/xla/tests/
Dlegalize-tf-no-tf2xla-fallback.mlir1521 // CHECK: "mhlo.reduce_window"(%[[ARG]], %[[INIT]])
1547 // CHECK: "mhlo.reduce_window"(%[[ARG]], %[[INIT]])
5210 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5236 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5262 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5288 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5313 // CHECK: %[[DIVIDEND:.*]] = "mhlo.reduce_window"(%[[ARG0]], %[[ZERO]]) ({
5323 // CHECK: %[[DIVISOR:.*]] = "mhlo.reduce_window"(%[[ONES]], %[[ZERO]]) ({
5345 // CHECK: %[[DIVIDEND:.*]] = "mhlo.reduce_window"(%[[ARG0]], %[[ZERO]]) ({
5355 // CHECK: %[[DIVISOR:.*]] = "mhlo.reduce_window"(%[[ONES]], %[[ZERO]]) ({
[all …]
Dlegalize-tf.mlir1538 // CHECK: "mhlo.reduce_window"(%[[ARG]], %[[INIT]])
1564 // CHECK: "mhlo.reduce_window"(%[[ARG]], %[[INIT]])
5378 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5404 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5430 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5456 // CHECK: [[DIVIDEND:%.+]] = "mhlo.reduce_window"([[CONV32]], [[ZERO]]) ({
5481 // CHECK: %[[DIVIDEND:.*]] = "mhlo.reduce_window"(%[[ARG0]], %[[ZERO]]) ({
5491 // CHECK: %[[DIVISOR:.*]] = "mhlo.reduce_window"(%[[ONES]], %[[ZERO]]) ({
5513 // CHECK: %[[DIVIDEND:.*]] = "mhlo.reduce_window"(%[[ARG0]], %[[ZERO]]) ({
5523 // CHECK: %[[DIVISOR:.*]] = "mhlo.reduce_window"(%[[ONES]], %[[ZERO]]) ({
[all …]
/external/tensorflow/tensorflow/compiler/xla/service/cpu/
Dir_emitter.h157 Status HandleReduceWindow(HloInstruction* reduce_window) override;
/external/tensorflow/tensorflow/compiler/tf2xla/python/
Dxla.py416 def reduce_window(operand, function

12