"func.func"() ({
  ^bb0(%arg0: tensor<i32>):
    %0 = "tf.Less"(%arg0, %arg0) {device = ""} : (tensor<i32>, tensor<i32>) -> tensor<*xi1>
    %1 = "tf.Less"(%arg0, %arg0) {device = ""} : (tensor<i32>, tensor<i32>) -> tensor<*xi1>
    %2 = "tf.LogicalAnd"(%0, %1) {device = ""} : (tensor<*xi1>, tensor<*xi1>) -> tensor<*xi1>
    "func.return"(%2) : (tensor<*xi1>) -> ()

  // op kernel.
  const PjRtBaseDevice::Metadata* pjrt_device_metadata_;

  // If the op associated with this XlaPlatformInfo is placed on an XLA device
  // then device_allocator_ is the xla::Backend's memory allocator.  If the op
  // is placed on a regular CPU or GPU device then device_allocator_ is null.
  // The allocator is of unknown provenance; keep it in a shared pointer to
  // set an artificial refcount of one.

    auto device = std::make_unique<XlaDevice>(session_options, options);

    Status status = device->UseAcceleratorDeviceInfo();
    if (!status.ok()) {
      LOG(INFO) << "Ignoring visible " << DEVICE_GPU_XLA_JIT
                << " device. Device number is " << i << ", reason: " << status;
      continue;
    }

    devices->push_back(std::move(device));
  }
  return absl::OkStatus();
}

  %0 = "tf.ReadVariableOp"(%arg1) {device = "/device:CPU:0"} : (tensor<*x!tf_type.resource>) -> tensor<1x3xf32>
  %1 = "tf.AddV2"(%arg0, %0) {device = "/device:CPU:0"} : (tensor<1x3xf32>, tensor<1x3xf32>) -> tensor<1x3xf32>
  %2 = "tf.Identity"(%1) {device = "/device:CPU:0"} : (tensor<1x3xf32>) -> tensor<1x3xf32>
  func.return %2 : tensor<1x3xf32>
}

// CHECK-LABEL: func @main

// CHECK:           return [[RESULT]] : tensor<3x4x4xf32>
// CHECK:         }

  %0 = "tf.BatchMatMulV2"(%arg0, %arg1) {T = f32, adj_x = false, adj_y = false, device = ""} : (tensor<1x4x2xf32>, tensor<3x2x4xf32>) -> tensor<3x4x4xf32>
  func.return %0 : tensor<3x4x4xf32>
}

func.func @batchmatmulv2_lhs_batch(%arg0: tensor<3x4x2xf32>, %arg1: tensor<2x4xf32>) -> tensor<3x4x4xf32> {

      // CHECK: %{{.*}} = stablehlo.custom_call @tf.ResizeBilinear(%arg0, %{{.*}}) {align_corners = false, device = "", half_pixel_centers = true} : (tensor<1x32x32x128xf32>, tensor<2xi32>) -> tensor<1x64x64x128xf32>
      // CHECK-OPT: %{{.*}} = stablehlo.custom_call @tf.ResizeBilinear(%arg0, %cst) {align_corners = false, device = "", half_pixel_centers = true} : (tensor<1x32x32x128xf32>, tensor<2xi32>) -> tensor<1x64x64x128xf32>

)

// HostUtils defines the set of methods for interacting with paths on a host.
type HostUtils interface {
	// DeviceOpened determines if the device (e.g. /dev/sdc) is in use elsewhere
	// on the system, i.e. still mounted.
	DeviceOpened(pathname string) (bool, error)
	// PathIsDevice determines if a path is a device.
	PathIsDevice(pathname string) (bool, error)

  %cst_1 = arith.constant dense<-1> : tensor<i32>
  %cst_2 = arith.constant dense<1> : tensor<i32>
  %0 = "tf.AddV2"(%arg2, %cst_2) {T = i32, device = ""} : (tensor<*xi32>, tensor<i32>) -> tensor<*xi32>
  %1 = "tf.Transpose"(%arg3, %cst) {T = f32, Tperm = i32, device = ""} : (tensor<*xf32>, tensor<2xi32>) -> tensor<?x?xf32>
  %2 = "tf.Rank"(%arg3) : (tensor<*xf32>) -> tensor<i32>

    const ResourceHandle& resource_handle =
        resource_tensor.scalar<ResourceHandle>()();

    Device* device;
    if (!(device_manager->LookupDevice(resource_handle.device(), &device)
              .ok())) {
      return module.emitOpError() << "failed to look up device";
    }

    tensorflow::Var* var_ptr;
    if (!(device->resource_manager()
              ->Lookup(resource_handle.container(), resource_handle.name(),

  func.func @main(%arg0: tensor<*x!tf_type.resource> {tf._composite_device = "/job:tpu_host_worker/replica:0/task:0/device:COMPOSITE:0", tf.device = "/job:tpu_host_worker/replica:0/task:0/device:COMPOSITE:0"}) {
    tf_executor.graph {
      %control = tf_executor.island {