| 96 | }; |
| 97 | |
| 98 | TEST_F(GpuConvRewriterTest, BackwardFilterConvolve) { |
| 99 | HloComputation::Builder builder(TestName()); |
| 100 | HloInstruction* activations = |
| 101 | builder.AddInstruction(HloInstruction::CreateParameter( |
| 102 | 0, ShapeUtil::MakeShape(F32, {1, 1, 3, 1}), "activations")); |
| 103 | HloInstruction* gradients = |
| 104 | builder.AddInstruction(HloInstruction::CreateParameter( |
| 105 | 1, ShapeUtil::MakeShape(F32, {1, 1, 2, 1}), "gradients")); |
| 106 | Window conv_window = default_conv_window_; |
| 107 | conv_window.mutable_dimensions(1)->set_size(2); |
| 108 | conv_window.mutable_dimensions(1)->set_window_dilation(2); |
| 109 | auto* conv = builder.AddInstruction(HloInstruction::CreateConvolve( |
| 110 | ShapeInference::InferConvolveShape( |
| 111 | activations->shape(), gradients->shape(), /*feature_group_count=*/1, |
| 112 | /*batch_group_count=*/1, conv_window, |
| 113 | tf_default_dnums_for_backward_filter_) |
| 114 | .ConsumeValueOrDie(), |
| 115 | activations, gradients, /*feature_group_count=*/1, |
| 116 | /*batch_group_count=*/1, conv_window, |
| 117 | tf_default_dnums_for_backward_filter_, DefaultPrecisionConfig(2))); |
| 118 | |
| 119 | OpMetadata metadata; |
| 120 | metadata.set_op_name("foo"); |
| 121 | conv->set_metadata(metadata); |
| 122 | |
| 123 | auto module = CreateNewVerifiedModule(); |
| 124 | HloComputation* entry_computation = |
| 125 | module->AddEntryComputation(builder.Build()); |
| 126 | EXPECT_TRUE(RunPass(module.get())); |
| 127 | ASSERT_THAT(entry_computation->root_instruction(), |
| 128 | op::GetTupleElement( |
| 129 | op::CustomCall(kCudnnConvBackwardFilterCallTarget), 0)); |
| 130 | |
| 131 | // Check that metadata was preserved. |
| 132 | const auto& md_after_opt = |
| 133 | entry_computation->root_instruction()->operand(0)->metadata(); |
| 134 | EXPECT_TRUE(protobuf_util::ProtobufEquals(md_after_opt, metadata)) |
| 135 | << md_after_opt.DebugString() << " vs " << metadata.DebugString(); |
| 136 | } |
| 137 | |
| 138 | TEST_F(GpuConvRewriterTest, |
| 139 | BackwardFilterConvolveEquivalentToForwardConvolution) { |
nothing calls this directly
no test coverage detected