// Check that Conv2D computed in NCHW format, and all redundant transpose
  // operations removed from the function.

  // CHECK: %[[CONV:[0-9]*]] = "tf.Conv2D"(%arg0, %arg1)
  // CHECK-SAME: data_format = "NCHW"
  // CHECK-SAME: -> tensor<1x8x32x32xf32>

  // CHECK: return %[[CONV]]

  func.return %4 : tensor<1x8x32x32xf32>

  // CHECK-SAME: dilations = [1, 4, 2, 3]
  // CHECK-SAME: explicit_paddings = [1, 2, 7, 8, 3, 4, 5, 6]
  // CHECK-SAME: padding = "EXPLICIT"
  // CHECK-SAME: strides = [5, 8, 6, 7]
  // CHECK-SAME: (tensor<1x3x32x32xf32>, tensor<4xi32>, tensor<1x8x32x32xf32>)
  // CHECK-SAME: -> tensor<1x1x3x8xf32>

  // CHECK: return %[[CONV2D_BACKPROP]]

  %0 = "tf.Conv2DBackpropFilter"(%input, %filter_sizes, %out_backprop)
       {
         data_format = "NHWC",

  %27 = mhlo.subtract %26, %19 : tensor<1x64x32x32xf32>
  %28 = "mhlo.broadcast_in_dim"(%1) <{broadcast_dimensions = dense<3> : tensor<1xi64>}> : (tensor<32xf32>) -> tensor<1x64x32x32xf32>
  %29 = mhlo.multiply %27, %28 : tensor<1x64x32x32xf32>
  %30 = mhlo.add %19, %29 : tensor<1x64x32x32xf32>

// dilations, etc...). This test only verifies that changing convolution data
// layout will update all the attributes.

// CHECK-LABEL: func @transposeConv2D
func.func @transposeConv2D(%input: tensor<1x3x32x32xf32>, %filter: tensor<1x1x3x8xf32>) -> tensor<1x8x7x6xf32> {

  // CHECK: %[[ARG_PERM:.*]] = "tf.Const"() <{value = dense<[0, 2, 3, 1]> : tensor<4xi64>}>
  // CHECK: %[[ARG_TRANSPOSE:[0-9]*]] = "tf.Transpose"(%arg0, %[[ARG_PERM]])

  func.return %0 : tensor<?x32x32x16xf32>
}

func.func @testDepthwiseConv2D(tensor<1x112x112x3xf32>, tensor<1x3x3x32xf32>, tensor<32xf32>) -> tensor<1x112x112x32xf32> {
^bb0(%arg0: tensor<1x112x112x3xf32>, %arg1: tensor<1x3x3x32xf32>, %arg2: tensor<32xf32>):
  // CHECK: _arithmetic_count = 7626752 : i64

  func.return %0 : tensor<256x30x30x16xf32>
}

// -----

func.func @testDepthwiseConv(%arg0: tensor<1x112x112x32xf32>, %arg1: tensor<1x3x3x32xf32>, %arg2: tensor<32xf32>) -> tensor<1x112x112x32xf32> {
  // CHECK: tac.device = "GPU", tac.inference_type = "FLOAT"

    "tfl.yield"(%4, %5, %6, %7) : (tensor<1x32x32x32xf32>, tensor<1x32x32x32xf32>, tensor<1x32x32x32xf32>, tensor<1x32x32x32xf32>) -> ()
  }) : (tensor<1x64x64x32xf32>, tensor<1x64x64x32xf32>, tensor<1x64x64x32xf32>) -> (tensor<1x32x32x32xf32>, tensor<1x32x32x32xf32>, tensor<1x32x32x32xf32>, tensor<1x32x32x32xf32>)

  // CHECK: %[[CST_0:.*]] = "tfl.no_value"() <{value}> : () -> none
  // CHECK: %[[ARG1:.*]] = "tfl.transpose_conv"(%arg0, %[[ARG0]], %arg2, %[[CST_0]]) <{fused_activation_function = "NONE", padding = "SAME", stride_h = 2 : i32, stride_w = 2 : i32}> : (tensor<4xi32>, tensor<1x3x3x32xf32>, tensor<15x14x14x32xf32>, none) -> tensor<15x28x28x1xf32>