}

func.func @real(%arg0: tensor<1 x complex<f64>>) -> tensor<1xf64> {
  %0 = "tf.Real"(%arg0) : (tensor<1 x complex<f64>>) -> tensor<1xf64>
  func.return %0: tensor<1xf64>

// CHECK-LABEL: real
// CHECK:  "tfl.real"(%arg0) : (tensor<1xcomplex<f64>>) -> tensor<1xf64>
// CHECK:  return
}

func.func @imag(%arg0: tensor<1 x complex<f64>>) -> tensor<1xf64> {

func.return %0 : tensor<2x?x2x4xf32>

// CHECK-LABEL:     convert_dot_general_dynamic_batch_dim
// CHECK:           %0 = "tfl.pseudo_const"() <{value = dense<[0, 1, 3, 2]> : tensor<4xi64>}> : () -> tensor<4xi64>
// CHECK-NEXT:      %1 = "tfl.cast"(%0) : (tensor<4xi64>) -> tensor<4xi32>
// CHECK-NEXT:      %2 = "tfl.transpose"(%arg1, %1) : (tensor<2x?x4x3xf32>, tensor<4xi32>) -> tensor<2x?x3x4xf32>

  // CHECK-NEXT:   mhlo.return %[[CMP]]
  // CHECK-NEXT: }) : (tensor<8xf64>) -> tensor<8xf64>
  // CHECK-NEXT: return %[[SORT]]
  %output = "tf.XlaSort"(%input) : (tensor<8xf64>) -> (tensor<8xf64>)
  func.return %output : tensor<8xf64>
}

// -----

// CHECK-LABEL: @xlasort_const
func.func @xlasort_const() -> (tensor<2x3xi64>) {

//CHECK-NEXT: return %0 : tensor<1x161x1xf32>
//CHECK-NEXT:}

func.func @convert(%arg0: tensor<2xf64>) -> tensor<2xf32> {
  %0 = "vhlo.convert_v1" (%arg0) : (tensor<2xf64>) -> tensor<2xf32>
  return %0 : tensor<2xf32>
}

//CHECK:func.func private @convert(%arg0: tensor<2xf64>) -> tensor<2xf32> {
//CHECK-NEXT: %0 = "vhlo.convert_v1"(%arg0) : (tensor<2xf64>) -> tensor<2xf32>
//CHECK-NEXT: return %0 : tensor<2xf32>
//CHECK-NEXT:}

    // CHECK-NEXT: %15 = mhlo.constant dense<1.000000e+00> : tensor<f32>
    // CHECK-NEXT: %16 = mhlo.constant dense<[32, 12, 12, 64]> : tensor<4xi64>
    // CHECK-NEXT: %17 = "mhlo.rng"(%14, %15, %16) <{rng_distribution = #mhlo.rng_distribution<UNIFORM>}> : (tensor<f32>, tensor<f32>, tensor<4xi64>) -> tensor<32x12x12x64xf32>
  %cst = "tf.Const"() {value = dense<[32, 12, 12, 64]> : tensor<4xi32>} : () -> tensor<4xi32>

func.func @testMaximumOfZeroToReluInt64OnGpu(%arg0: tensor<4xi64>) -> tensor<4xi64> {
  // CHECK: %[[RESULT:.*]] = "tf.Relu"(%arg0) {device = "/job:localhost/replica:0/task:0/device:GPU:0"} : (tensor<4xi64>) -> tensor<4xi64>
  // CHECK: return %[[RESULT]]
  %cst_0 = arith.constant dense<0> : tensor<i64>
  %0 = "tf.Maximum"(%arg0, %cst_0) {device = "/job:localhost/replica:0/task:0/device:GPU:0"} : (tensor<4xi64>, tensor<i64>) -> tensor<4xi64>

  %begin0 = arith.constant dense<0> : tensor<4xi64>
  %shape0 = arith.constant dense<[2,3,4,4]> : tensor<4xi64>
  %begin1 = arith.constant dense<1> : tensor<4xi64>
  %shape1 = arith.constant dense<[1,2,3,4]> : tensor<4xi64>
  %0 = "tfl.slice"(%arg0, %begin0, %shape0) : (tensor<2x3x4x5xf32>, tensor<4xi64>, tensor<4xi64>) -> tensor<2x3x4x4xf32>

  // CHECK-DAG: [[VAL1:%.+]] = "tf.Const"() <{value = dense<[1, 5, 3, 2]> : tensor<4xi64>}>
  // CHECK-DAG: [[VAL2:%.+]] = "tf.Const"() <{value = dense<{{\[\[}}0, 0], [3, 4], [0, 0{{\]\]}}> : tensor<3x2xi64>}>
  // CHECK-DAG: [[VAL3:%.+]] = "tf.Const"() <{value = dense<0.000000e+00> : tensor<f32>}>
  // CHECK-DAG: [[VAL4:%.+]] = "tf.Const"() <{value = dense<[2, 0, 1, 3]> : tensor<4xi64>}>
  // CHECK-DAG: [[VAL5:%.+]] = "tf.PadV2"(%arg0, [[VAL2]], [[VAL3]])

  func.return %0 : tensor<3xf32>
}

// -----

// CHECK-LABEL: testReal
func.func @testReal(%arg0: tensor<? x complex<f64>>) -> tensor<?xf64> {
  // CHECK: "tfl.real"(%arg0)
  %0 = "tfl.real"(%arg0): (tensor<? x complex<f64>>) -> tensor<?xf64>
  func.return %0 : tensor<?xf64>
}

// -----

func.func @testRealWrongShape(%arg0: tensor<3 x complex<f64>>) -> tensor<4xf32> {

        %bdy_cast = "tf.Cast"(%barg0) : (tensor<*xi64>) -> tensor<4xf32>
        %bdy:2 = func.call @while_body(%bdy_cast, %barg1) : (tensor<4xf32>, tensor<i32>) -> (tensor<4xi64>, tensor<i32>)
        "tf.Yield"(%bdy#0, %bdy#1) : (tensor<4xi64>, tensor<i32>) -> ()
    }
  ) { is_stateless = false } : (tensor<*xi64>, tensor<i32>) -> (tensor<*xi64>, tensor<i32>)
  // CHECK: return [[Result]]#0
  func.return %0#0 : tensor<*xi64>