strides = [1, op.get_attr('stride_w'), op.get_attr('stride_h'), 1]
  padding = op.get_attr('padding')
  shape_0, shape_1 = tf.shape_n([op.inputs[0], op.inputs[1]])
  return [
      tf.compat.v1.nn.conv2d_backprop_input(
          shape_0,
          op.inputs[1],
          grad,
          strides=strides,
          padding=padding,
          dilations=dilations,
          data_format='NHWC'),

}

func.func @conv2d_backprop_input(%arg0: tensor<4xi32>, %arg1: tensor<3x3x1x32xf32>, %arg2: tensor<15x14x14x32xf32>) -> tensor<15x28x28x1xf32> {
  %0 = "tf.Conv2DBackpropInput"(%arg0, %arg1, %arg2) {strides = [1, 2, 2, 1], padding="SAME", dilations=[1, 1, 1, 1]}: (tensor<4xi32>, tensor<3x3x1x32xf32>, tensor<15x14x14x32xf32>) -> tensor<15x28x28x1xf32>

  return %result : tensor<?x512x512x1xf32>
}

// -----

// CHECK-LABEL: @conv2d_backprop_input
func.func @conv2d_backprop_input(
    %filter: tensor<3x3x1x32xf32>,
    %out_backprop: tensor<100x26x26x32xf32>
  ) -> tensor<100x28x28x1xf32> {
    // CHECK: %[[REV_FILTER:.*]] = "mhlo.reverse"(%arg0) <{dimensions = dense<[0, 1]> : tensor<2xi64>}>

}

// CHECK-LABEL: conv2d_backprop_input_with_add
func.func @conv2d_backprop_input_with_add(%arg0: tensor<4xi32>, %arg1: tensor<3x3x1x32xf32>, %arg2: tensor<15x14x14x32xf32>) -> tensor<15x28x28x1xf32> {
  %0 = "tf.Conv2DBackpropInput"(%arg0, %arg1, %arg2) {strides = [1, 2, 2, 1], padding="SAME", dilations=[1, 1, 1, 1]}: (tensor<4xi32>, tensor<3x3x1x32xf32>, tensor<15x14x14x32xf32>) -> tensor<15x28x28x1xf32>

}

// CHECK-LABEL: conv2d_backprop_input_with_add
func.func @conv2d_backprop_input_with_add(%arg0: tensor<4xi32>, %arg1: tensor<3x3x1x32xf32>, %arg2: tensor<15x14x14x32xf32>) -> tensor<15x28x28x1xf32> {
  %0 = "tf.Conv2DBackpropInput"(%arg0, %arg1, %arg2) {strides = [1, 2, 2, 1], padding="SAME", dilations=[1, 1, 1, 1]}: (tensor<4xi32>, tensor<3x3x1x32xf32>, tensor<15x14x14x32xf32>) -> tensor<15x28x28x1xf32>

  TF_RETURN_IF_ERROR(GetNodeAttr(attrs, "use_cudnn_on_gpu", &use_cudnn_on_gpu));
  auto dx_1 = Conv2DBackpropInput(scope, Shape(scope, op.input(0)), op.input(1),
                                  grad_inputs[0], strides, padding,
                                  Conv2DBackpropInput::DataFormat(data_format)
                                      .UseCudnnOnGpu(use_cudnn_on_gpu));
  grad_outputs->push_back(dx_1);

  %input_size:   tensor<4xi32>,
  %filter:       tensor<1x28x28x64xf32>,
  %out_backprop: tensor<1x28x28x64xf32>
) -> tensor<1x28x28x64xf32> {

  // CHECK: "tf.Conv2DBackpropInput"
  // CHECK-SAME: data_format = "NCHW"
  %0 = "tf.Conv2DBackpropInput"(%input_size, %filter, %out_backprop)
       {
         data_format = "NHWC",
         padding = "VALID",
         strides = [1, 1, 1, 1]

  %input_size:   tensor<4xi32>,
  %filter:       tensor<1x28x28x64xf16>,
  %out_backprop: tensor<1x28x28x64xf16>
) -> tensor<1x28x28x64xf16> {

  // CHECK: "tf.Conv2DBackpropInput"
  // CHECK-SAME: data_format = "NCHW"
  %0 = "tf.Conv2DBackpropInput"(%input_size, %filter, %out_backprop)
       {
         data_format = "NHWC",
         padding = "VALID",
         strides = [1, 1, 1, 1]

 _Retval 2
cluster 0 size 2178
 Add 17
 AddN 72
 ArgMax 1
 AssignAddVariableOp 1
 AssignSubVariableOp 106
 BiasAdd 1
 BiasAddGrad 1
 Cast 3
 Const 357
 Conv2D 53
 Conv2DBackpropFilter 53
 Conv2DBackpropInput 52
 DivNoNan 1
 Equal 1
 FusedBatchNorm 53
 FusedBatchNormGrad 53
 Identity 2
 MatMul 3
 MaxPool 1
 MaxPoolGrad 1
 Mean 1
 Mul 164
 Pad 1
 ReadVariableOp 646
 Relu 49
 ReluGrad 49

 _Arg 435
 _Retval 2
cluster 0 size 1910
 Add 16
 AddN 71
 ArgMax 1
 AssignAddVariableOp 1
 BiasAdd 1
 BiasAddGrad 1
 Cast 115
 Const 407
 Conv2D 53
 Conv2DBackpropFilter 53
 Conv2DBackpropInput 52
 Equal 1
 FusedBatchNormGradV2 53
 FusedBatchNormV2 53
 MatMul 3
 MaxPool 1
 MaxPoolGrad 1
 Mean 1
 Mul 218
 Pad 2
 ReadVariableOp 538
 Relu 49
 ReluGrad 49
 Reshape 2