return %2 : tensor<1x8x4x4xf32>
}
// CHECK-DAG: %[[CONST:.+]] = stablehlo.constant {{.*}} : tensor<3x3x8x8xf32>
// CHECK-DAG: %[[TRANSPOSE_0:.+]] = stablehlo.transpose %[[ARG]], dims = [0, 2, 3, 1] : (tensor<1x8x4x4xf32>) -> tensor<1x4x4x8xf32>

}

// CHECK-LABEL:   func.func @avg_pool2d_1(
// CHECK-SAME:                            %[[VAL_0:.*]]: tensor<1x3x6x6xf32>) -> tensor<*xf32> {
// CHECK:           %[[VAL_1:.*]] = arith.constant dense<[0, 2, 3, 1]> : tensor<4xi32>
// CHECK:           %[[VAL_2:.*]] = "tfl.transpose"(%[[VAL_0]], %[[VAL_1]]) : (tensor<1x3x6x6xf32>, tensor<4xi32>) -> tensor<1x6x6x3xf32>
// CHECK:           %[[VAL_3:.*]] = arith.constant dense<0> : tensor<4x2xi32>

  %input_scale = "tf.Const"() { value = dense<0.588> : tensor<f32> } : () -> tensor<f32>
  %input_zp = "tf.Const"() { value = dense<42> : tensor<i32> } : () -> tensor<i32>
  %filter_scale = "tf.Const"() { value = dense<0.0235> : tensor<f32> } : () -> tensor<f32>
  %filter_zp = "tf.Const"() { value = dense<0> : tensor<i32> } : () -> tensor<i32>
  %accum_scale = "tf.Const"() { value = dense<0.013818> : tensor<f32> } : () -> tensor<f32>

}
// CHECK-SAME: (%[[ARG_0:.+]]: tensor<1x3x3x4xf32>) -> tensor<1x4x3x3xf32>
// CHECK-DAG: %[[CONST_0:.+]] = stablehlo.constant
// CHECK-DAG: %[[TRANSPOSE_0:.+]] = stablehlo.transpose %[[CONST_0]], dims = [0, 2, 3, 1] : (tensor<1x4x3x3xf32>) -> tensor<1x3x3x4xf32>

// Check that the shape of the add is changed to reflect the deferred transpose.
// CHECK: %[[ADD_0:.+]] = stablehlo.add %[[ARG_0]], %[[TRANSPOSE_0]] : tensor<1x3x3x4xf32>

00000040  00 00 00 00 00 00 00 00  00 00 00 00 00 32 cc a9  |.............2..|
00000050  cc a8 c0 2b c0 2f c0 2c  c0 30 c0 09 c0 13 c0 0a  |...+./.,.0......|
00000060  c0 14 00 9c 00 9d 00 2f  00 35 c0 12 00 0a c0 23  |......./.5.....#|
00000070  c0 27 00 3c c0 07 c0 11  00 05 13 03 13 01 13 02  |.'.<............|
00000080  01 00 00 7f 00 0b 00 02  01 00 ff 01 00 01 00 00  |................|

  func.return %0 : tensor<2x7x5x4xf32>
  // CHECK-LABEL: einsum_4D
  // CHECK-DAG: %[[cst:.*]] = arith.constant dense<[0, 2, 1, 3]> : tensor<4xi32>
  // CHECK-DAG: %[[cst_1:.*]] = arith.constant dense<[0, 2, 3, 1]> : tensor<4xi32>
  // CHECK: %[[v0:.*]] = "tf.Transpose"(%arg0, %[[cst]]) : (tensor<2x5x7x3xf32>, tensor<4xi32>) -> tensor<2x7x5x3xf32>

              /*stride_h=*/ExtractI32At<2>:$strides,
              /*stride_w=*/ExtractI32At<3>:$strides,
              /*fused_activation_function=*/TFL_AF_None,
              (returnType (GetTransposedType<"{0, 2, 3, 1}"> $src_op))),
              (Arith_ConstantOp
                  ConstantAttr<I32VectorElementsAttr<4>, "{0, 3, 1, 2}">))>;

def LegalizeArgMax : Pat<(TF_ArgMaxOp $input, $dim),

// Permutation from the NCHW tensor format to NHWC. This is an inverse
// permutation of `kNchwToNhwcPermutation`.
inline constexpr std::array<int64_t, 4> kNchwToNhwcPermutation = {0, 2, 3, 1};

// Permutation from the OIHW (== (output features, input features, height,
// width)) tensor format to HWIO. This is commonly used to transpose convolution

    // Read a delimiter.
    if (b != addressOffset) {
      if (input[i] != '.') return false // Wrong delimiter.
      i++
    }

    // Read 1 or more decimal digits for a value in 0..255.
    var value = 0
    val groupOffset = i
    while (i < limit) {
      val c = input[i]
      if (c < '0' || c > '9') break

	// YCbCr.RGBA method (16 bit color output).
	//
	// The solution is to make the rounding adjustment non-constant, and equal
	// to 257*Y', which ranges over [0, 1<<16-1] as Y' ranges over [0, 255].
	// YY1 is then defined as:
	//	YY1 = 65536*Y' + 257*Y'
	// or equivalently:
	//	YY1 = Y' * 0x10101
	yy1 := int32(y) * 0x10101
	cb1 := int32(cb) - 128
	cr1 := int32(cr) - 128