xla::DeviceAssignment device_assignment(num_replicas, num_cores_per_replica);
  int pos = 0;
  for (int replica = 0; replica < num_replicas; ++replica) {
    for (int logical_core = 0; logical_core < num_cores_per_replica;
         ++logical_core) {
      int x = device_assignment_attr[pos++];
      int y = device_assignment_attr[pos++];
      int z = device_assignment_attr[pos++];
      int core = device_assignment_attr[pos++];

      %1:2 = tf_device.replicate {n = 2 : i32, devices = {TPU_REPLICATED_CORE_0 = ["/job:worker/replica:0/task:0/device:TPU:1", "/job:worker/replica:0/task:0/device:TPU:2"]}} {
        %2 = "tf._TPUDeviceOrdinalPlaceholder"() {logical_core = 0} : () -> tensor<i64>
        tf_device.return %2 : tensor<i64>
      }
      tf_executor.yield %1#0, %1#1 : tensor<i64>, tensor<i64>
    }
    tf_executor.fetch
  }
  func.return
}

  %105 = "tf.LogicalOr"(%104, %2) {device = ""} : (tensor<i1>, tensor<i1>) -> tensor<i1>
  %106 = "tf.Equal"(%103, %13) {device = "", incompatible_shape_error = true} : (tensor<i64>, tensor<i64>) -> tensor<i1>
  %107 = "tf.LogicalOr"(%105, %106) {device = ""} : (tensor<i1>, tensor<i1>) -> tensor<i1>

REGISTER_NO_GRADIENT_OP("Equal");
REGISTER_NO_GRADIENT_OP("ApproximateEqual");
REGISTER_NO_GRADIENT_OP("NotEqual");
REGISTER_NO_GRADIENT_OP("LogicalAnd");
REGISTER_NO_GRADIENT_OP("LogicalOr");
REGISTER_NO_GRADIENT_OP("LogicalNot");
REGISTER_NO_GRADIENT_OP("Floor");

// Conjugate helper function returns the conjugate of an Output if it
// is complex valued.
Output ConjugateHelper(const Scope& scope, const Output& out) {

  // CHECK: tf_device.replicate
  // CHECK:   tf._TPUDeviceOrdinalPlaceholder
  %0:2 = tf_device.replicate([%arg0, %arg1] as %ri: tensor<*xf32>) {n = 2: i32} {
    %1 = "tf._TPUDeviceOrdinalPlaceholder"() {logical_core = 0} : () -> tensor<i64>
    tf_device.return %1 : tensor<i64>
  }
  func.return
}


// CHECK-LABEL: func @replicate_resource_var_arg_shape

    (TF_ConstOp ConstantAttr<RankedF32ElementsAttr<[]>, "1.0f">))),
  (TF_Log1pOp:$dest $arg), [], [(CopyAttrs $src, $dest)]>;

//===----------------------------------------------------------------------===//
// LogicalNot op patterns.
//===----------------------------------------------------------------------===//

def LogicalNotOfEqual : Pat<
  (TF_LogicalNotOp:$src (TF_EqualOp $arg0, $arg1, $shape_error)),

// CHECK: return %arg0
}

// CHECK-LABEL: testDoubleLogicalNot
func.func @testDoubleLogicalNot(%arg0: tensor<8x16x32x64xi1>) -> tensor<8x16x32x64xi1> {
  %0 = "tf.LogicalNot"(%arg0) : (tensor<8x16x32x64xi1>) -> tensor<8x16x32x64xi1>
  %1 = "tf.LogicalNot"(%0) : (tensor<8x16x32x64xi1>) -> tensor<8x16x32x64xi1>
  func.return %1: tensor<8x16x32x64xi1>

// CHECK: return %arg0
}

// CHECK-LABEL: testDoubleNeg

  func.return %0: tensor<*xi1>
}

// CHECK-LABEL: func @or
func.func @or(%arg0: tensor<2xi1>, %arg1: tensor<2xi1>) -> tensor<2xi1> {
  // CHECK-NEXT:  mhlo.or
  %0 = "tf.LogicalOr"(%arg0, %arg1) : (tensor<2xi1>, tensor<2xi1>) -> tensor<2xi1>
  func.return %0: tensor<2xi1>
}

// CHECK-LABEL: func @bitwise_or
func.func @bitwise_or(%arg0: tensor<4xi32>, %arg1: tensor<4xi32>) -> tensor<4xi32> {

        `tf.atan2`, `tf.atanh`, `tf.cos`, `tf.cosh`, `tf.equal`, `tf.exp`,
        `tf.floor`, `tf.greater`, `tf.greater_equal`, `tf.less`,
        `tf.less_equal`, `tf.log`, `tf.logp1`, `tf.logical_and`,
        `tf.logical_not`, `tf.logical_or`, `tf.maximum`, `tf.minimum`,
        `tf.not_equal`, `tf.sin`, `tf.sinh`, `tf.tan`
    *   Deprecate `tf.data.Dataset.shard`.
    *   Deprecate `saved_model.loader.load` which is replaced by

}

// -----

// Nested IfRegions
// CHECK: func private @tf.IfRegion1_else
// CHECK-NEXT: "tf.Acos"
// CHECK-NEXT: "tf.Abs"

// CHECK: func private @tf.IfRegion1_then
// CHECK-NEXT: "tf.LogicalNot"
// CHECK-NEXT: "tf.Asin"
// CHECK-NEXT: "tf.If"({{.+}}) <{else_branch = @tf.IfRegion_else, {{.+}} then_branch = @tf.IfRegion_then}

// CHECK: func private @tf.IfRegion_else
// CHECK-NEXT: "tf.Neg"