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)

 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

  // 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]

  return builder->create<TF::SliceOp>(input.getLoc(), slice_result_type, input,
                                      start_position, slice_size_op);
}

// Transforms Conv2DBackPropFilter for space to depth.
void HandleConv2DBackPropFilter(TF::Conv2DBackpropFilterOp backprop,
                                ArrayRef<int32_t> old_filter_shape,

    // CHECK: %[[BACKPROP:.*]] = "tf.Conv2DBackpropFilter"
    // CHECK-SAME: strides = [1, 1, 1, 1]
    // CHECK-SAME: (tensor<2x115x115x12xf32>, tensor<4xi32>, tensor<2x112x112x64xf32>) -> tensor<4x4x12x64xf32>

    ([CVE-2021-29536](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2021-29536))
*   Fixes a division by zero in `Conv2DBackpropFilter`
    ([CVE-2021-29538](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2021-29538))
*   Fixes a heap buffer overflow in `Conv2DBackpropFilter`
    ([CVE-2021-29540](https://cve.mitre.org/cgi-bin/cvename.cgi?name=CVE-2021-29540))
*   Fixes a heap buffer overflow in `StringNGrams`

      "BiasAddV1",
      "Bincount",
      "Bucketize",
      "Case",
      "CheckNumerics",
      "Cholesky",
      "ControlTrigger",
      "Conv",
      "Conv2D",
      "Conv2DBackpropFilter",
      "Conv2DBackpropInput",
      "Conv3D",
      "Conv3DBackpropFilterV2",
      "Conv3DBackpropInputV2",
      "Cross",
      "Cumprod",
      "Cumsum",
      "CumulativeLogsumexp",

    std::vector<int64_t> explicit_paddings;
    if (padding == tensorflow::Padding::EXPLICIT) {
      // EXPLICIT padding mode and the associated attribute is limited to
      // Conv2DBackpropFilter. So, fetch attribute by identifier instead of the
      // op.explicit_paddings() attribute getter.
      ArrayRef<Attribute> explicit_paddings_attr =