}

# MLIR-LABEL: func @main
# MLIR-SAME:  (%[[ARG_0:[a-z0-9]+]]: tensor<1x1x1x256x!quant.uniform<i8:f32, 0.21632751372549019:27>>) -> tensor<1x6x31x!quant.uniform<i8:f32, 0.09363494573854933:22>>
# MLIR-SAME:  control_outputs = ""
# MLIR-SAME:  inputs = "input"
# MLIR-SAME:  outputs = "output"
# MLIR:         %[[shape:.*]] = arith.constant dense<[1, -1, 31]> : tensor<3xi32>

}

// CHECK-LABEL: func @testPow
// CHECK-SAME:(%[[ARG_0:.*]]: tensor<4xf32>, %[[ARG_1:.*]]: tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>, tensor<4xf32>)
func.func @testPow(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> (tensor<4xf32>, tensor<4xf32>, tensor<4xf32>) {

  %cst_zero = arith.constant dense<0.0> : tensor<f32>
  %cst_one = arith.constant dense<1.0> : tensor<f32>

    "mhlo.return"(%0, %1) : (tensor<f32>, tensor<f32>) -> ()
  },  {
    "mhlo.return"(%0, %1) : (tensor<f32>, tensor<f32>) -> ()
  // CHECK:      (tensor<i1>) -> (tensor<f32>, tensor<f32>, !mhlo.token)
  }) : (tensor<i1>) -> (tensor<f32>, tensor<f32>)

  // CHECK:      [[IF_SUBTUPLE_RESULT:%.*]] = mhlo.tuple [[IF]]#0, [[IF]]#1

  // copy-on-read mode is false.
  if (var->tensor()->RefCountIsOne()) {
    var->copy_on_read_mode.store(true);
    return absl::OkStatus();
  }
  Tensor tmp;
  if (variantType) {
    AllocatorAttributes attr;
    attr.set_on_host(true);
    TF_RETURN_IF_ERROR(context->allocate_temp(
        var->tensor()->dtype(), var->tensor()->shape(), &tmp, attr));

    const auto elements_in = var->tensor()->flat<Variant>();

// on-device shape. This is accurate for CPU and GPU devices that neither
// transpose nor pad tensors.
Status DefaultPaddedShapeFn(const Tensor& tensor, xla::Shape* shape) {
  const tensorflow::XlaTensor* xla_tensor =
      tensorflow::XlaTensor::FromTensor(&tensor);
  if (xla_tensor == nullptr) {
    return TensorShapeToXLAShape(tensor.dtype(), tensor.shape(), shape);
  }

  const xla::ShapedBuffer& shaped_buffer = xla_tensor->shaped_buffer();

// Legalize TensorList Ops iff all of them are supported.
std::unique_ptr<OperationPass<ModuleOp>> CreateLegalizeTensorListPass();

// Reduce the type precision of some tensor types if all values within that
// tensor are within the range of the reduced precision.
std::unique_ptr<OperationPass<ModuleOp>> CreateReduceTypePrecisionPass();

// Convervatively pushes transposes through elementwise ops to prepare

                                             TF_Status* status,
                                             void* device_info);

  // Method to copy a tensor from the custom device to a target device.
  TFE_TensorHandle* (*copy_tensor_from_device)(TFE_Context* context,
                                               TFE_TensorHandle* tensor,

  }

  return true;
}

// Calculates other_tensor_zp * tensor for zero point offset calculation.
// Things to do:
//  1. Reduce the tensor (which is an input of XlaDotV2) with contracting
//     dimensions of XlaDotV2.
//     - The resultant dimension must match with XlaDotV2 resultant dimension
//  2. Multiply it with zero point from the other tensor.
// We decided to use tf.Einsum for step 1, since it would require transposes/

  // criteria for determining quantizable ops differs by the inference type.
  bool getQuantizableOps(arith::ConstantOp op,
                         QuantizationUnits& quantizable_ops) const {
    // Non-float tensors do not need quantization.
    auto type = mlir::dyn_cast<ShapedType>(op.getType());
    if (!type || !type.getElementType().isF32()) return false;

    Value value = op.getResult();

        "TFE_TensorHandleDeviceMemorySize may not be called on a ",
        handle->TypeString(), " tensor handle.");
    return 0;
  }
  const tensorflow::Tensor* tensor;
  status->status = handle->Tensor(&tensor);
  if (!status->status.ok()) {
    return 0;
  }
  return tensor->TotalBytes();
}

TFE_Op* TFE_NewOp(TFE_Context* ctx, const char* op_or_function_name,