// CHECK: %[[BUFFER:.*]] = "tf.BroadcastTo"(%[[ZERO_F32]], %[[MAX_SIZE]])
  // CHECK: %[[BUFFER_SIZE:.*]] = "tf.Const"() <{value = dense<0> : tensor<1xi32>}>
  // CHECK-NOT: tf.EmptyTensorList
  %if_op = "tf.IfRegion"(%arg0) ({
      %elem = "tf._SomeOp"() : () -> tensor<f32>
      // CHECK: %[[UPDATE:.*]] = "tf.XlaDynamicUpdateSlice"
      // CHECK: %[[ONE:.*]] = "tf.Const"() <{value = dense<1> : tensor<1xi32>}> : () -> tensor<1xi32>

  std::string retvals_communication_key =
      llvm::formatv("host_compute_channel_{0}_retvals",
                    (communication_key_index))
          .str();

  // Use a unique name when sending just the IfRegion predicate.  This is
  // for readable and to match the key in the TF2XLA bridge.
  if (clustered_ops.size() == 1 && llvm::isa<mlir::TF::IfRegionOp>(op) &&
      external_operands.size() == 1) {
    args_communication_key =

  %18 = "tf.Equal"(%17, %4) {device = "", incompatible_shape_error = true} : (tensor<i64>, tensor<i64>) -> tensor<i1>
  %19 = "tf.All"(%18, %5) {device = "", keep_dims = false} : (tensor<i1>, tensor<0xi32>) -> tensor<i1>
  %20 = "tf.IfRegion"(%19) ({
    %72 = "func.call"(%19, %17, %4) {callee = @RaggedFromNestedRowSplits_RaggedFromRowSplits_RowPartitionFromRowSplits_assert_equal_1_Assert_AssertGuard_true_27770} : (tensor<i1>, tensor<i64>, tensor<i64>) -> tensor<i1>

  bool InferShapeForRestore(Operation* op);

  // Infers the shape IfOp outputs based on the shapes of the then and else
  // function result types.
  bool InferShapeForIf(IfOp op);

  // Infers the shape IfRegion outputs based on the shapes of the then and else
  // yields.
  bool InferShapeForIfRegion(IfRegionOp op);

  // Infers the shape CaseOp outputs based on the shapes of branch function
  // result types.

    types of inputs must match the signature of the operation that contains the
    region.
  }];

  let arguments = (ins Variadic<AnyType>);
}

def TF_IfRegionOp : TF_Op<"IfRegion",
      [SingleBlockImplicitTerminator<"YieldOp">, NoRegionArguments,
       DeclareOpInterfaceMethods<RegionBranchOpInterface, [
           "areTypesCompatible",
           "getEntrySuccessorOperands",

  func.return %0, %1 : tensor<4x4xf32>, tensor<3x4xi32>
}

// CHECK-LABEL: @ifRegion
// CHECK-SAME:  ([[ARG0:%.+]]: tensor<f32>, [[ARG1:%.+]]: tensor<f32>)
func.func @ifRegion(%arg0: tensor<f32>, %arg1: tensor<f32>) -> (tensor<f32>) {
  // CHECK: [[VAL0:%.+]] = mhlo.compare GT, [[ARG0]], [[ARG1]]

  func.func @cluster(%arg0 : tensor<i1>) {
    "tf_device.parallel_execute"() ({
      "tf.IfRegion"(%arg0) ({
        // expected-error@+1 {{The ParallelExecute ancestor of a ClusterFunc must be its direct parent.}}

// CHECK-LABEL:  func @if
// CHECK-DAG:      %[[CST_0:.*]] = arith.constant dense<0> : tensor<i32>
// CHECK-DAG:      %[[CST_1:.*]] = arith.constant dense<1000> : tensor<i32>
// CHECK:          %[[RES:.*]]  = "tf.IfRegion"(%arg0) <{is_stateless = false}> ({
// CHECK:            "tf.Yield"(%[[CST_0]]) : (tensor<i32>) -> ()
// CHECK:          }, {
// CHECK:            "tf.Yield"(%[[CST_1]]) : (tensor<i32>) -> ()