v->set_tag(tag.data(), tag.size());
  histo.EncodeToProto(v->mutable_histo(), false /* Drop zero buckets */);

  Safe_TF_TensorPtr summary_tensor(TF_AllocateOutput(
      /*context=*/ctx, /*index=*/0, /*dtype=*/TF_ExpectedOutputDataType(ctx, 0),
      /*dims=*/nullptr, /*num_dims=*/0,
      /*len=*/sizeof(tensorflow::tstring), status.get()));

  if (TF_GetCode(status.get()) != TF_OK) {

//     positive for any real input.
//
//     ```python
//     x = tf.constant(2.0)
//     tf.math.exp(x) ==> 7.389056
//
//     x = tf.constant([2.0, 8.0])
//     tf.math.exp(x) ==> array([7.389056, 2980.958], dtype=float32)
//     ```
//
//     For complex numbers, the exponential value is calculated as follows:
//
//     ```
//     e^(x+iy) = e^x * e^iy = e^x * (cos y + i sin y)
//     ```
//

  for (auto* n : arg_nodes) {
    int index;
    TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "index", &index));
    DataType dtype;
    TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "T", &dtype));
    (*recv_at_host_dtypes)[index] = dtype;
  }
  for (int i = 0, end = recv_at_host_dtypes->size(); i < end; i++) {
    if ((*recv_at_host_dtypes)[i] == DT_INVALID) {

  }

  tensorflow::AbstractTensorInterface* t =
      tensorflow::unwrap(ctx)->CreateTensor(
          static_cast<tensorflow::DataType>(dtype), dimvec);

  if (t == nullptr) {
    status->status =
        tensorflow::errors::InvalidArgument("Unsupported dtype: ", dtype);
    return nullptr;
  }

  return new TF_Tensor{t};
}

TFE_TensorHandle* TFE_NewTensorHandleFromTensor(TFE_Context* ctx, TF_Tensor* t,

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();
  *shape = shaped_buffer.on_device_shape();
  return absl::OkStatus();
}

      "XTimesY",
      // Args
      {"x: float", "y: resource"},
      // Return values
      {"z: float"},
      // Attr def
      {},
      // Nodes
      {
          {{"y0"}, "ReadVariableOp", {"y"}, {{"dtype", DT_FLOAT}}},
          {{"z"}, "Mul", {"x", "y0"}, {{"T", DT_FLOAT}}},
      });
}

class XlaKernelCreatorTest : public ::testing::Test {
 protected:
  void Init(const std::vector<FunctionDef>& flib) {

  auto build_result = BuildHloFromGraph(*graph, /*use_output_shapes=*/true);
  ASSERT_FALSE(build_result.ok());
  EXPECT_THAT(build_result.message(),
              HasSubstr("op operand type 'tensor<2x3x4x5xi32>' and result type "
                        "'tensor<1xi32>' are cast incompatible"));
}

}  // namespace

                         kSavedModelVariablesFilename);
        ImmediateTensorHandlePtr restored_output;
        TF_RETURN_IF_ERROR(internal::SingleRestore(
            context, variables_path_prefix, checkpoint_key, variable->dtype(),
            &restored_output));

        // Assign the restored tensor's value to the variable
        return variable->Assign(restored_output.get());
      }));

  return Status();
}

      // If the arg type already matches the global_tensor type, we don't need
      // to do anything.
      if (!underlying_type.empty() &&
          underlying_type[0] == global_tensor.getType()) {
        assert(underlying_type.size() == 1);
        continue;
      }

      // Otherwise, set this argument's type to the global_tensor's type.
      auto new_arg_type = mlir::RankedTensorType::get(

  CreateContext();

  std::vector<XlaCompiler::Argument> args(2);
  args[0].kind = XlaCompiler::Argument::kParameter;
  args[0].type = DT_INT32;
  args[0].shape = TensorShape({1, 3});
  args[1].kind = XlaCompiler::Argument::kParameter;
  args[1].type = DT_INT32;
  args[1].shape = TensorShape({1, 3});

  const XlaCompiler::CompilationResult* result;
  xla::PjRtLoadedExecutable* executable;