message: [XLA] Remove some unnecessary host/device synchronization in GenericTransferManager.
  Use a host callback to release memory for tuple index tables rather than blocking the host on the stream.
  Call the async version of WriteTupleIndexTables inside TransferLiteralToDeviceAsync ? the caller has asynchronous semantics anyway.
diff: @@ -45,14 +45,18 @@ Status GenericTransferManager::WriteSingleTupleIndexTable(
  const Shape& shape, se::DeviceMemoryBase* region) {
  TF_RET_CHECK(elements.size() == ShapeUtil::TupleElementCount(shape));
  - std::vector<const void*> element_pointers;
  + auto element_pointers = std::make_shared<std::vector<const void*>>();
  + element_pointers->reserve(elements.size());
  for (const se::DeviceMemoryBase& element : elements) {
  - element_pointers.push_back(element.opaque());
  + element_pointers->push_back(element.opaque());
  }
  TF_RETURN_IF_ERROR(TransferBufferToDevice(
  - stream, GetByteSizeRequirement(shape), element_pointers.data(), region));
  + stream, GetByteSizeRequirement(shape), element_pointers->data(), region));
  // Ensure the buffer is transferred before we destroy element_pointers.
  - return stream->BlockHostUntilDone();
  + stream->ThenDoHostCallback([element_pointers]() {
  + /* holds reference to element_pointers in closure */
  + });
  + return Status::OK();
  }
  void GenericTransferManager::TransferLiteralFromDevice(
  @@ -115,7 +119,7 @@ Status GenericTransferManager::TransferLiteralToDeviceAsync(
  TF_RET_CHECK(stream->parent()->device_ordinal() ==
  device_buffer.device_ordinal());
  - TF_RETURN_IF_ERROR(WriteTupleIndexTables(stream, device_buffer));
  + TF_RETURN_IF_ERROR(WriteTupleIndexTablesAsync(stream, device_buffer));
  return ShapeUtil::ForEachSubshapeWithStatus(
  device_buffer.on_host_shape(),
  