%0 = "tf.Bitcast"(%arg0) : (tensor<8x16x32x64xf32>) -> tensor<8x16x32x64xf32>
  // CHECK: return %arg0
  func.return %0: tensor<8x16x32x64xf32>
}

// CHECK-LABEL: testDifferentBitcastType
func.func @testDifferentBitcastType(%arg0: tensor<8x16x32x64xf32>) -> tensor<8x16x32x64xi32> {

  // Build the XLA computation graph.
  // "arg0" (T=DT_RESOURCE), "arg1" (T=DT_INT32)
  // "arg0", "arg1" -> "if" (If) -> "ret0", "ret1"
  // "arg0", "arg1" -> "while" (While) -> "ret2", "ret3"
  tensorflow::Scope s = tensorflow::Scope::NewRootScope();
  Output arg0 = ops::_Arg(s.WithOpName("arg0"), DT_RESOURCE, 0);
  Output arg1 = ops::_Arg(s.WithOpName("arg1"), DT_BOOL, 1);
  NameAttrList f;
  f.set_name("f1");

  %1:2 = "tf.WhileRegion"(%cst_0, %0) <{is_stateless = false}> ({
      ^bb0(%carg0: tensor<i32>, %carg1: tensor<!tf_type.variant>):
       %cst_2 = arith.constant dense<2> : tensor<i32>
       %1 = "tf.Less"(%carg0, %cst_2) : (tensor<i32>, tensor<i32>) -> tensor<i1>
       "tf.Yield"(%1) : (tensor<i1>) -> ()

      // verify condition types
      // CHECK: ^bb0(%[[CARG0:.*]]: tensor<i32>, %[[CARG1:.*]]: tensor<*xf32>):

    %2:2 = "tf.WhileRegion"(%1, %arg0) ({
      ^bb0(%carg0: tensor<f32>, %carg1: tensor<i32>):
         %limit = arith.constant dense<5> : tensor<i32>
         %cond = "tf.NotEqual"(%carg1, %limit) : (tensor<i32>, tensor<i32>) -> tensor<i1>
         "tf.Yield"(%cond) : (tensor<i1>) -> ()
    },  {
      ^bb0(%barg0: tensor<f32>, %barg1: tensor<i32>):

#ifdef GOOS_windows
#define RARG0 CX
#define RARG1 DX
#define RARG2 R8
#define RARG3 R9
#else
#define RARG0 DI
#define RARG1 SI
#define RARG2 DX
#define RARG3 CX
#endif

// func runtime·domsanread(addr unsafe.Pointer, sz uintptr)
// Called from msanread.
TEXT	runtime·domsanread(SB), NOSPLIT, $0-16
	MOVQ	addr+0(FP), RARG0
	MOVQ	sz+8(FP), RARG1

		// atomic compare and swap.
		// arg0 = pointer, arg1 = old value, arg2 = new value, arg3 = memory.
		// if *arg0 == arg1 {
		//   *arg0 = arg2
		//   return (true, memory)
		// } else {
		//   return (false, memory)
		// }
		// DBAR
		// MOVV $0, Rout
		// LL	(Rarg0), Rtmp
		// BNE	Rtmp, Rarg1, 4(PC)
		// MOVV Rarg2, Rout
		// SC	Rout, (Rarg0)
		// BEQ	Rout, -4(PC)
		// DBAR

// CHECK-LABEL: @squeeze_folder
func.func @squeeze_folder(%arg0 : tensor<?x?xf32>) -> tensor<?x?xf32> {
  %0 = "tfl.squeeze"(%arg0) : (tensor<?x?xf32>) -> tensor<?x?xf32>
  // CHECK: return %arg0
  func.return %0 : tensor<?x?xf32>
}

// -----

// CHECK-LABEL: @squeeze_folder
func.func @squeeze_folder(%arg0 : tensor<?x?xf32>) -> tensor<*xf32> {
  %0 = "tfl.squeeze"(%arg0) : (tensor<?x?xf32>) -> tensor<*xf32>
  // CHECK: "tfl.squeeze"

		{name: "MODSD", argLength: 2, reg: gp21, asm: "MODSD", typ: "Int64"},  // arg0 % arg1 (signed 64-bit)
		{name: "MODUW", argLength: 2, reg: gp21, asm: "MODUW", typ: "UInt32"}, // arg0 % arg1 (unsigned 32-bit)
		{name: "MODSW", argLength: 2, reg: gp21, asm: "MODSW", typ: "Int32"},  // arg0 % arg1 (signed 32-bit)
		// MOD is implemented as rem := arg0 - (arg0/arg1) * arg1

  // expected-remark@below {{Region #1, Arg #2, ID 7 : 1, 7, 8}}
  %w:3 = "tf.WhileRegion"(%arg0, %arg1, %arg2) ({
          ^bb0(%carg0: !tf_res, %carg1: !tf_res, %carg2: !tf_res):
          %0 = arith.constant dense<false> : tensor<i1>
          "tf.Yield"(%0) : (tensor<i1>) -> ()
         },{
          ^bb0(%barg0: !tf_res, %barg1: !tf_res, %barg2: !tf_res):
          // expected-remark@below {{Result #0, ID 1 : Unknown}}

  // CHECK: cond

  // CHECK: do
  // CHECK: mhlo.return %[[ARG0]], %[[ARG1]] : tensor<2x?xi8>, tensor<?x2xi8>

  %2:2 = "tf.WhileRegion"(%0, %1) ({
  ^bb0(%carg0: tensor<2x2x!tf_type.qint8>, %carg1: tensor<2x2x!tf_type.qint8>):
    %check = "tf.Equal"(%arg1, %one) : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "tf.Yield"(%check) : (tensor<i1>) -> ()
    }, {
  ^bb0(%barg0: tensor<2x?x!tf_type.qint8>, %barg1: tensor<?x2x!tf_type.qint8>):