shape_0,
          op.inputs[1],
          grad,
          strides=strides,
          padding=padding,
          dilations=dilations,
          data_format='NHWC'),
      tf.compat.v1.nn.conv2d_backprop_filter(
          op.inputs[0],
          shape_1,
          grad,
          strides=strides,
          padding=padding,
          dilations=dilations,
          data_format='NHWC'), bias_grad
  ]

  func.return %result : tensor<2x8x8x8x1xf32>
}

// -----

// CHECK-LABEL: @conv2d_backprop_filter
func.func @conv2d_backprop_filter(
    %input: tensor<100x28x28x1xf32>,
    %out_backprop: tensor<100x26x26x32xf32>
  ) -> tensor<3x3x1x32xf32> {
  // CHECK: %[[RESULT:.*]] = mhlo.convolution(%arg0, %arg1)

  %input:        tensor<1x28x28x64xf16>,
  %filter_size:  tensor<4xi32>,
  %out_backprop: tensor<1x28x28x64xf16>
) -> tensor<1x1x64x64xf16> {

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

  %input:        tensor<1x28x28x64xf32>,
  %filter_size:  tensor<4xi32>,
  %out_backprop: tensor<1x28x28x64xf32>
) -> tensor<1x1x64x64xf32> {

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

                                      .UseCudnnOnGpu(use_cudnn_on_gpu));
  grad_outputs->push_back(dx_1);
  auto dx_2 =
      Conv2DBackpropFilter(scope, op.input(0), Shape(scope, op.input(1)),
                           grad_inputs[0], strides, padding,
                           Conv2DBackpropFilter::DataFormat(data_format)
                               .UseCudnnOnGpu(use_cudnn_on_gpu));
  grad_outputs->push_back(dx_2);

 VarHandleOp 435
 _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

 ReadVariableOp 2
 Switch 1
 _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

  attr {
    key: "T"
    value {
      type: DT_FLOAT
    }
  }
}
node {
  name: "training/LossScaleOptimizer/gradients/res5c_branch2c_1/Conv2D_grad/Conv2DBackpropFilter"
  op: "Conv2DBackpropFilter"
  input: "activation_47_1/Relu"
  input: "ConstantFolding/training/LossScaleOptimizer/gradients/res5c_branch2c_1/Conv2D_grad/ShapeN-matshapes-1"

  // CHECK: %[[IN_TRANSPOSE:[0-9]*]] = "tf.Transpose"(%arg0, %[[ARG_PERM]])
  // CHECK: %[[OUT_BP_TRANSPOSE:[0-9]*]] = "tf.Transpose"(%arg2, %[[ARG_PERM]])

  // CHECK: %[[CONV2D_BACKPROP:[0-9]*]] = "tf.Conv2DBackpropFilter"
  // CHECK-SAME: (%[[IN_TRANSPOSE]], %[[FILTER_PERM]], %[[OUT_BP_TRANSPOSE]])
  // CHECK-SAME: data_format = "NCHW"
  // CHECK-SAME: dilations = [1, 4, 2, 3]
  // CHECK-SAME: explicit_paddings = [1, 2, 7, 8, 3, 4, 5, 6]

  }
  attr {
    key: "use_cudnn_on_gpu"
    value {
      b: true
    }
  }
}
node {
  name: "training/SGD/gradients/res5c_branch2c_1/Conv2D_grad/Conv2DBackpropFilter"
  op: "Conv2DBackpropFilter"
  input: "activation_47_1/Relu"
  input: "ConstantFolding/training/SGD/gradients/res5c_branch2c_1/Conv2D_grad/ShapeN-matshapes-1"