diff --git a/.github/ISSUE_TEMPLATE/40-tflite-op-request.md b/.github/ISSUE_TEMPLATE/40-tflite-op-request.md
new file mode 100644
index 0000000000000000000000000000000000000000..7b391279e479ade4ed5327728f19be8752e11507
--- /dev/null
+++ b/.github/ISSUE_TEMPLATE/40-tflite-op-request.md
@@ -0,0 +1,24 @@
+---
+name: TensorFlow Lite Op Request
+about: Use this template for reporting ops you are using or missing.
+
+---
+
+
+**System information**
+- OS Platform and Distribution (e.g., Linux Ubuntu 16.04):
+- TensorFlow installed from (source or binary):
+- TensorFlow version (or github SHA if from source):
+
+
+**Provide the text output from tflite_convert**
+
+```
+# Copy and paste here
+```
+
+Also, please include a link to a GraphDef or the model if possible.
+
+**Any other info / logs**
+
+Include any logs or source code that would be helpful to diagnose the problem. If including tracebacks, please include the full traceback. Large logs and files should be attached.
diff --git a/CODEOWNERS b/CODEOWNERS
index 54a61a4d72c40d297d90d53e223f64f813d9167d..cb3fa2312405ce44d5dfc30ea4164740f436e07e 100644
--- a/CODEOWNERS
+++ b/CODEOWNERS
@@ -1,7 +1,7 @@
# Where component owners are known, add them here.
/tenosrflow/core/debug @caisq
-/tensorflow/core/nccl/ @azaks @csigg
+/tensorflow/core/nccl/ @azaks2 @chsigg
/tensorflow/core/platform/windows/ @mrry
/tensorflow/core/platform/s3 @yongtang
/tensorflow/go @asimshankar
@@ -51,13 +51,13 @@
/tensorflow/contrib/pi_examples/ @maciekcc
/tensorflow/contrib/quantization/ @petewarden
/tensorflow/contrib/rnn/ @ebrevdo @scottzhu
-/tensorflow/contrib/saved_model/ @nfiedel @sukritiramesh @allenl
+/tensorflow/contrib/saved_model/ @nfiedel @sukritiramesh @allenlavoie
/tensorflow/contrib/seq2seq/ @ebrevdo @lmthang
/tensorflow/contrib/session_bundle/ @nfiedel @sukritiramesh
/tensorflow/contrib/slim/ @sguada @thenbasilmanran
/tensorflow/contrib/stateless/ @girving @alextp
/tensorflow/contrib/tensor_forest/ @gilberthendry @thomascolthurst @yupbank
-/tensorflow/contrib/tensorrt/ @aaroey
+/tensorflow/contrib/tensorrt/ @aaroey @smit-hinsu @azaks2
# NEED OWNER: /tensorflow/contrib/testing/
/tensorflow/contrib/timeseries/ @allenlavoie
/tensorflow/contrib/tpu/ @frankchn @saeta @jhseu @sourabhbajaj
diff --git a/README.md b/README.md
index 8af5370befbb090966a8b3af54d80c84a969aaa5..68d7e180d1e757693dfd8d4d73003983eafd4ade 100644
--- a/README.md
+++ b/README.md
@@ -9,12 +9,14 @@
|-----------------|
| [](https://www.tensorflow.org/api_docs/) |
-**TensorFlow** is an open source software library for numerical computation using
-data flow graphs. The graph nodes represent mathematical operations, while
+**TensorFlow** is an open source software library for numerical computation
+using data flow graphs. The graph nodes represent mathematical operations, while
the graph edges represent the multidimensional data arrays (tensors) that flow
-between them. This flexible architecture enables you to deploy computation to one
-or more CPUs or GPUs in a desktop, server, or mobile device without rewriting
-code. TensorFlow also includes [TensorBoard](https://www.tensorflow.org/guide/summaries_and_tensorboard), a data visualization toolkit.
+between them. This flexible architecture enables you to deploy computation to
+one or more CPUs or GPUs in a desktop, server, or mobile device without
+rewriting code. TensorFlow also includes
+[TensorBoard](https://github.com/tensorflow/tensorboard), a data visualization
+toolkit.
TensorFlow was originally developed by researchers and engineers
working on the Google Brain team within Google's Machine Intelligence Research
@@ -111,22 +113,25 @@ The TensorFlow project strives to abide by generally accepted best practices in
Build Type | Status | Artifacts
----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | ---------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | ---------
**IBM s390x** | [](http://ibmz-ci.osuosl.org/job/TensorFlow_IBMZ_CI/) | TBA
-**IBM ppc64le CPU** | [](http://powerci.osuosl.org/job/TensorFlow_Ubuntu_16.04_CPU/) | TBA
-**IBM ppc64le GPU** Nightly | [](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Nightly_Artifact/) | [Nightly](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Nightly_Artifact/)
-**IBM ppc64le GPU** Stable Release | [](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Release_Build/) | [Release](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Release_Build/)
+**Linux ppc64le CPU** Nightly | [](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_CPU_Build/) | [Nightly](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_CPU_Nightly_Artifact/)
+**Linux ppc64le CPU** Stable Release | [](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_CPU_Release_Build/) | [Release](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_CPU_Release_Build/)
+**Linux ppc64le GPU** Nightly | [](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Build/) | [Nightly](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Nightly_Artifact/)
+**Linux ppc64le GPU** Stable Release | [](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Release_Build/) | [Release](https://powerci.osuosl.org/job/TensorFlow_PPC64LE_GPU_Release_Build/)
**Linux CPU with Intel® MKL-DNN** Nightly | [](https://tensorflow-ci.intel.com/job/tensorflow-mkl-linux-cpu/) | [Nightly](https://tensorflow-ci.intel.com/job/tensorflow-mkl-build-whl-nightly/)
**Linux CPU with Intel® MKL-DNN** Python 2.7
**Linux CPU with Intel® MKL-DNN** Python 3.4
**Linux CPU with Intel® MKL-DNN** Python 3.5
**Linux CPU with Intel® MKL-DNN** Python 3.6 | [](https://tensorflow-ci.intel.com/job/tensorflow-mkl-build-release-whl/lastStableBuild) | [1.11.0 py2.7](https://storage.googleapis.com/intel-optimized-tensorflow/tensorflow-1.11.0-cp27-cp27mu-linux_x86_64.whl)
[1.11.0 py3.4](https://storage.googleapis.com/intel-optimized-tensorflow/tensorflow-1.11.0-cp34-cp34m-linux_x86_64.whl)
[1.11.0 py3.5](https://storage.googleapis.com/intel-optimized-tensorflow/tensorflow-1.11.0-cp35-cp35m-linux_x86_64.whl)
[1.11.0 py3.6](https://storage.googleapis.com/intel-optimized-tensorflow/tensorflow-1.11.0-cp36-cp36m-linux_x86_64.whl)
## For more information
-* [TensorFlow Website](https://www.tensorflow.org)
-* [TensorFlow Tutorials](https://www.tensorflow.org/tutorials/)
-* [TensorFlow Model Zoo](https://github.com/tensorflow/models)
-* [TensorFlow Twitter](https://twitter.com/tensorflow)
-* [TensorFlow Blog](https://medium.com/tensorflow)
-* [TensorFlow Course at Stanford](https://web.stanford.edu/class/cs20si)
-* [TensorFlow Roadmap](https://www.tensorflow.org/community/roadmap)
-* [TensorFlow White Papers](https://www.tensorflow.org/about/bib)
-* [TensorFlow YouTube Channel](https://www.youtube.com/channel/UC0rqucBdTuFTjJiefW5t-IQ)
+
+* [TensorFlow Website](https://www.tensorflow.org)
+* [TensorFlow Tutorials](https://www.tensorflow.org/tutorials/)
+* [TensorFlow Model Zoo](https://github.com/tensorflow/models)
+* [TensorFlow Twitter](https://twitter.com/tensorflow)
+* [TensorFlow Blog](https://medium.com/tensorflow)
+* [TensorFlow Course at Stanford](https://web.stanford.edu/class/cs20si)
+* [TensorFlow Roadmap](https://www.tensorflow.org/community/roadmap)
+* [TensorFlow White Papers](https://www.tensorflow.org/about/bib)
+* [TensorFlow YouTube Channel](https://www.youtube.com/channel/UC0rqucBdTuFTjJiefW5t-IQ)
+* [TensorFlow Visualization Toolkit](https://github.com/tensorflow/tensorboard)
Learn more about the TensorFlow community at the [community page of tensorflow.org](https://www.tensorflow.org/community) for a few ways to participate.
diff --git a/WORKSPACE b/WORKSPACE
index 0c7bc085b512b084b9470abe17326d7c119aa327..7cc08e0164a202581ad7ebbe107a9e19410e70e4 100644
--- a/WORKSPACE
+++ b/WORKSPACE
@@ -1,5 +1,7 @@
workspace(name = "org_tensorflow")
+load("@bazel_tools//tools/build_defs/repo:http.bzl", "http_archive")
+
http_archive(
name = "io_bazel_rules_closure",
sha256 = "a38539c5b5c358548e75b44141b4ab637bba7c4dc02b46b1f62a96d6433f56ae",
@@ -57,9 +59,9 @@ android_workspace()
# Please add all new TensorFlow dependencies in workspace.bzl.
tf_workspace()
-new_http_archive(
+http_archive(
name = "inception_v1",
- build_file = "models.BUILD",
+ build_file = "//:models.BUILD",
sha256 = "7efe12a8363f09bc24d7b7a450304a15655a57a7751929b2c1593a71183bb105",
urls = [
"http://storage.googleapis.com/download.tensorflow.org/models/inception_v1.zip",
@@ -67,9 +69,9 @@ new_http_archive(
],
)
-new_http_archive(
+http_archive(
name = "mobile_ssd",
- build_file = "models.BUILD",
+ build_file = "//:models.BUILD",
sha256 = "bddd81ea5c80a97adfac1c9f770e6f55cbafd7cce4d3bbe15fbeb041e6b8f3e8",
urls = [
"http://storage.googleapis.com/download.tensorflow.org/models/object_detection/ssd_mobilenet_v1_android_export.zip",
@@ -77,9 +79,9 @@ new_http_archive(
],
)
-new_http_archive(
+http_archive(
name = "mobile_multibox",
- build_file = "models.BUILD",
+ build_file = "//:models.BUILD",
sha256 = "859edcddf84dddb974c36c36cfc1f74555148e9c9213dedacf1d6b613ad52b96",
urls = [
"http://storage.googleapis.com/download.tensorflow.org/models/mobile_multibox_v1a.zip",
@@ -87,9 +89,9 @@ new_http_archive(
],
)
-new_http_archive(
+http_archive(
name = "stylize",
- build_file = "models.BUILD",
+ build_file = "//:models.BUILD",
sha256 = "3d374a730aef330424a356a8d4f04d8a54277c425e274ecb7d9c83aa912c6bfa",
urls = [
"http://storage.googleapis.com/download.tensorflow.org/models/stylize_v1.zip",
@@ -97,9 +99,9 @@ new_http_archive(
],
)
-new_http_archive(
+http_archive(
name = "speech_commands",
- build_file = "models.BUILD",
+ build_file = "//:models.BUILD",
sha256 = "c3ec4fea3158eb111f1d932336351edfe8bd515bb6e87aad4f25dbad0a600d0c",
urls = [
"http://storage.googleapis.com/download.tensorflow.org/models/speech_commands_v0.01.zip",
diff --git a/configure.py b/configure.py
index 234561d94a46f57c4de5ca487360e2d5a3dfdb2f..6c905a0be3d685b5921dfbc5bddfbe6471a82625 100644
--- a/configure.py
+++ b/configure.py
@@ -238,6 +238,13 @@ def setup_python(environ_cp):
write_to_bazelrc('build --python_path=\"%s"' % python_bin_path)
environ_cp['PYTHON_BIN_PATH'] = python_bin_path
+ # If choosen python_lib_path is from a path specified in the PYTHONPATH
+ # variable, need to tell bazel to include PYTHONPATH
+ if environ_cp.get('PYTHONPATH'):
+ python_paths = environ_cp.get('PYTHONPATH').split(':')
+ if python_lib_path in python_paths:
+ write_action_env_to_bazelrc('PYTHONPATH', environ_cp.get('PYTHONPATH'))
+
# Write tools/python_bin_path.sh
with open(
os.path.join(_TF_WORKSPACE_ROOT, 'tools', 'python_bin_path.sh'),
@@ -445,11 +452,12 @@ def convert_version_to_int(version):
return int(version_str)
-def check_bazel_version(min_version):
- """Check installed bazel version is at least min_version.
+def check_bazel_version(min_version, max_version):
+ """Check installed bazel version is between min_version and max_version.
Args:
min_version: string for minimum bazel version.
+ max_version: string for maximum bazel version.
Returns:
The bazel version detected.
@@ -467,6 +475,7 @@ def check_bazel_version(min_version):
min_version_int = convert_version_to_int(min_version)
curr_version_int = convert_version_to_int(curr_version)
+ max_version_int = convert_version_to_int(max_version)
# Check if current bazel version can be detected properly.
if not curr_version_int:
@@ -480,6 +489,10 @@ def check_bazel_version(min_version):
print('Please upgrade your bazel installation to version %s or higher to '
'build TensorFlow!' % min_version)
sys.exit(0)
+ if curr_version_int > max_version_int:
+ print('Please downgrade your bazel installation to version %s or lower to '
+ 'build TensorFlow!' % max_version)
+ sys.exit(0)
return curr_version
@@ -859,7 +872,7 @@ def set_tf_cuda_version(environ_cp):
cuda_toolkit_paths_full = [
os.path.join(cuda_toolkit_path, x) for x in cuda_rt_lib_paths
]
- if any([os.path.exists(x) for x in cuda_toolkit_paths_full]):
+ if any(os.path.exists(x) for x in cuda_toolkit_paths_full):
break
# Reset and retry
@@ -1552,7 +1565,7 @@ def main():
# environment variables.
environ_cp = dict(os.environ)
- check_bazel_version('0.15.0')
+ check_bazel_version('0.15.0', '0.20.0')
reset_tf_configure_bazelrc()
# Explicitly import tools/bazel.rc, this is needed for Bazel 0.19.0 or later
@@ -1694,6 +1707,7 @@ def main():
config_info_line('nohdfs', 'Disable HDFS support.')
config_info_line('noignite', 'Disable Apacha Ignite support.')
config_info_line('nokafka', 'Disable Apache Kafka support.')
+ config_info_line('nonccl', 'Disable NVIDIA NCCL support.')
if __name__ == '__main__':
diff --git a/tensorflow/BUILD b/tensorflow/BUILD
index 17577afecb74b7008db5a282255278b35ed138a6..fd4b94202aad24a82abef8abd16431f61a8326f0 100644
--- a/tensorflow/BUILD
+++ b/tensorflow/BUILD
@@ -246,6 +246,12 @@ config_setting(
visibility = ["//visibility:public"],
)
+config_setting(
+ name = "no_nccl_support",
+ define_values = {"no_nccl_support": "true"},
+ visibility = ["//visibility:public"],
+)
+
# Crosses between platforms and file system libraries not supported on those
# platforms due to limitations in nested select() statements.
config_setting(
diff --git a/tensorflow/api_template.__init__.py b/tensorflow/api_template.__init__.py
index 2efb8846c6837a3935e0a8439a18838cb2bea804..4eba763129a6aef40e3c130d56bf8ab19638b7ca 100644
--- a/tensorflow/api_template.__init__.py
+++ b/tensorflow/api_template.__init__.py
@@ -20,18 +20,14 @@ from __future__ import print_function as _print_function
import os as _os
-# pylint: disable=g-bad-import-order
-from tensorflow.python import pywrap_tensorflow # pylint: disable=unused-import
+# API IMPORTS PLACEHOLDER
+# pylint: disable=g-bad-import-order
from tensorflow.python.tools import component_api_helper as _component_api_helper
_component_api_helper.package_hook(
parent_package_str=__name__,
child_package_str=('tensorflow_estimator.python.estimator.api.estimator'))
-# API IMPORTS PLACEHOLDER
-
-from tensorflow.python.platform import flags # pylint: disable=g-import-not-at-top
-
# Make sure directory containing top level submodules is in
# the __path__ so that "from tensorflow.foo import bar" works.
# We're using bitwise, but there's nothing special about that.
@@ -39,8 +35,9 @@ _tf_api_dir = _os.path.dirname(_os.path.dirname(bitwise.__file__)) # pylint: di
if _tf_api_dir not in __path__:
__path__.append(_tf_api_dir)
-# Calls to enable and disable features.
-enable_eager_execution() # pylint: disable=undefined-variable
+# Enable TF2 behaviors
+from tensorflow.python.compat import compat as _compat # pylint: disable=g-import-not-at-top
+_compat.enable_v2_behavior()
# These symbols appear because we import the python package which
# in turn imports from tensorflow.core and tensorflow.python. They
diff --git a/tensorflow/api_template_v1.__init__.py b/tensorflow/api_template_v1.__init__.py
index 65bdb6cb1b5e6fb0656a12b932d767aeacfccd29..21b5277614667bdbd7271ac3e57f5b69d5a19264 100644
--- a/tensorflow/api_template_v1.__init__.py
+++ b/tensorflow/api_template_v1.__init__.py
@@ -23,13 +23,13 @@ import os as _os
# pylint: disable=g-bad-import-order
from tensorflow.python import pywrap_tensorflow # pylint: disable=unused-import
+# API IMPORTS PLACEHOLDER
+
from tensorflow.python.tools import component_api_helper as _component_api_helper
_component_api_helper.package_hook(
parent_package_str=__name__,
child_package_str=('tensorflow_estimator.python.estimator.api.estimator'))
-# API IMPORTS PLACEHOLDER
-
from tensorflow.python.util.lazy_loader import LazyLoader # pylint: disable=g-import-not-at-top
contrib = LazyLoader('contrib', globals(), 'tensorflow.contrib')
del LazyLoader
diff --git a/tensorflow/c/BUILD b/tensorflow/c/BUILD
index 84238ffc1f2b73c59055461fbeba33687d756208..25df970ecab0757f23465ab19e7f45de0c759458 100644
--- a/tensorflow/c/BUILD
+++ b/tensorflow/c/BUILD
@@ -121,6 +121,7 @@ tf_cuda_library(
":c_api",
":c_api_internal",
"//tensorflow/c/eager:c_api",
+ "//tensorflow/c/eager:c_api_internal",
"//tensorflow/compiler/jit:flags",
"//tensorflow/contrib/tpu:all_ops",
"//tensorflow/core:core_cpu",
@@ -174,6 +175,34 @@ tf_cuda_library(
],
)
+tf_cuda_library(
+ name = "env",
+ srcs = [
+ "env.cc",
+ ],
+ hdrs = [
+ "env.h",
+ ],
+ copts = tf_copts(),
+ visibility = ["//visibility:public"],
+ deps = select({
+ "//tensorflow:android": [
+ ":c_api",
+ ":tf_status_helper",
+ "//tensorflow/core:android_tensorflow_lib_lite",
+ "//tensorflow/core:platform_env",
+ "//tensorflow/core:lib",
+ ],
+ "//conditions:default": [
+ ":c_api",
+ ":tf_status_helper",
+ "//tensorflow/core:framework",
+ "//tensorflow/core:platform_env",
+ "//tensorflow/core:lib",
+ ],
+ }) + [":c_api_internal"],
+)
+
tf_cuda_library(
name = "kernels",
srcs = [
@@ -187,10 +216,14 @@ tf_cuda_library(
deps = select({
"//tensorflow:android": [
":c_api",
+ ":c_api_internal",
+ ":tf_status_helper",
"//tensorflow/core:android_tensorflow_lib_lite",
],
"//conditions:default": [
":c_api",
+ ":c_api_internal",
+ ":tf_status_helper",
"//tensorflow/core:framework",
],
}),
@@ -263,7 +296,7 @@ tf_cuda_cc_test(
tf_cc_test(
name = "c_api_experimental_test",
- size = "small",
+ size = "medium",
srcs = ["c_api_experimental_test.cc"],
data = ["testdata/tf_record"],
linkopts = select({
@@ -274,8 +307,11 @@ tf_cc_test(
# the shared library must be able to use core:framework.
# linkstatic = tf_kernel_tests_linkstatic(),
deps = [
+ ":c_api",
":c_api_experimental",
":c_test_util",
+ "//tensorflow/c/eager:c_api",
+ "//tensorflow/c/eager:c_api_test_util",
"//tensorflow/core:lib",
"//tensorflow/core:protos_all_cc",
"//tensorflow/core:test",
@@ -326,6 +362,27 @@ tf_kernel_library(
alwayslink = 1,
)
+tf_cuda_cc_test(
+ name = "env_test",
+ size = "small",
+ srcs = ["env_test.cc"],
+ linkopts = select({
+ "//tensorflow:darwin": ["-headerpad_max_install_names"],
+ "//conditions:default": [],
+ }),
+ tags = ["noasan"],
+ # We must ensure that the dependencies can be dynamically linked since
+ # the shared library must be able to use core:framework.
+ # linkstatic = tf_kernel_tests_linkstatic(),
+ deps = [
+ ":c_api",
+ ":env",
+ "//tensorflow/core:lib",
+ "//tensorflow/core:test",
+ "//tensorflow/core:test_main",
+ ],
+)
+
tf_cuda_cc_test(
name = "kernels_test",
size = "small",
diff --git a/tensorflow/c/c_api.cc b/tensorflow/c/c_api.cc
index f13e8777dff164bcd8eedf46310ae846abd0c804..9580215a317b1a6b1cdacbd430a1764af61be990 100644
--- a/tensorflow/c/c_api.cc
+++ b/tensorflow/c/c_api.cc
@@ -136,16 +136,22 @@ const char* TF_Message(const TF_Status* s) {
namespace {
class TF_ManagedBuffer : public TensorBuffer {
public:
- void* data_;
- size_t len_;
- void (*deallocator_)(void* data, size_t len, void* arg);
- void* deallocator_arg_;
+ TF_ManagedBuffer(void* data, size_t len,
+ void (*deallocator)(void* data, size_t len, void* arg),
+ void* deallocator_arg)
+ : TensorBuffer(data),
+ len_(len),
+ deallocator_(deallocator),
+ deallocator_arg_(deallocator_arg) {}
+
+ const size_t len_;
+ void (*const deallocator_)(void* data, size_t len, void* arg);
+ void* const deallocator_arg_;
~TF_ManagedBuffer() override {
- (*deallocator_)(data_, len_, deallocator_arg_);
+ (*deallocator_)(data(), len_, deallocator_arg_);
}
- void* data() const override { return data_; }
size_t size() const override { return len_; }
TensorBuffer* root_buffer() override { return this; }
void FillAllocationDescription(AllocationDescription* proto) const override {
@@ -199,8 +205,7 @@ TF_Tensor* TF_NewTensor(TF_DataType dtype, const int64_t* dims, int num_dims,
dimvec[i] = static_cast(dims[i]);
}
- TF_ManagedBuffer* buf = new TF_ManagedBuffer;
- buf->len_ = len;
+ TF_ManagedBuffer* buf = nullptr;
if (dtype != TF_STRING && dtype != TF_RESOURCE &&
tensorflow::DataTypeCanUseMemcpy(static_cast(dtype)) &&
reinterpret_cast(data) % std::max(1, EIGEN_MAX_ALIGN_BYTES) !=
@@ -212,17 +217,15 @@ TF_Tensor* TF_NewTensor(TF_DataType dtype, const int64_t* dims, int num_dims,
//
// Other types have the same representation, so copy only if it is safe to
// do so.
- buf->data_ = allocate_tensor("TF_NewTensor", len);
- std::memcpy(buf->data_, data, len);
- buf->deallocator_ = deallocate_buffer;
- buf->deallocator_arg_ = nullptr;
+ buf = new TF_ManagedBuffer(allocate_tensor("TF_NewTensor", len), len,
+ deallocate_buffer, nullptr);
+ std::memcpy(buf->data(), data, len);
// Free the original buffer.
deallocator(data, len, deallocator_arg);
} else {
- buf->data_ = data;
- buf->deallocator_ = deallocator;
- buf->deallocator_arg_ = deallocator_arg;
+ buf = new TF_ManagedBuffer(data, len, deallocator, deallocator_arg);
}
+
TF_Tensor* ret = new TF_Tensor{dtype, TensorShape(dimvec), buf};
size_t elem_size = TF_DataTypeSize(dtype);
if (elem_size > 0 && len < (elem_size * ret->shape.num_elements())) {
@@ -477,14 +480,15 @@ static TF_Tensor* EmptyTensor(TF_DataType dtype, const TensorShape& shape) {
CHECK_EQ(nelems, 0);
static_assert(sizeof(int64_t) == sizeof(tensorflow::int64),
"64-bit int types should match in size");
- return TF_NewTensor(dtype, reinterpret_cast(dims.data()),
- shape.dims(), reinterpret_cast(&empty), 0,
- [](void*, size_t, void*) {}, nullptr);
+ return TF_NewTensor(
+ dtype, reinterpret_cast(dims.data()), shape.dims(),
+ reinterpret_cast(&empty), 0, [](void*, size_t, void*) {}, nullptr);
}
// Non-static for testing.
TF_Tensor* TF_TensorFromTensor(const tensorflow::Tensor& src,
TF_Status* status) {
+ TF_SetStatus(status, TF_OK, "");
if (!src.IsInitialized()) {
status->status = FailedPrecondition(
"attempt to use a tensor with an uninitialized value");
@@ -1592,18 +1596,20 @@ TF_AttrMetadata TF_OperationGetAttrMetadata(TF_Operation* oper,
break; \
}
- LIST_CASE(s, TF_ATTR_STRING, metadata.total_size = 0;
- for (int i = 0; i < attr->list().s_size();
- ++i) { metadata.total_size += attr->list().s(i).size(); });
+ LIST_CASE(
+ s, TF_ATTR_STRING, metadata.total_size = 0;
+ for (int i = 0; i < attr->list().s_size();
+ ++i) { metadata.total_size += attr->list().s(i).size(); });
LIST_CASE(i, TF_ATTR_INT);
LIST_CASE(f, TF_ATTR_FLOAT);
LIST_CASE(b, TF_ATTR_BOOL);
LIST_CASE(type, TF_ATTR_TYPE);
- LIST_CASE(shape, TF_ATTR_SHAPE, metadata.total_size = 0;
- for (int i = 0; i < attr->list().shape_size(); ++i) {
- const auto& s = attr->list().shape(i);
- metadata.total_size += s.unknown_rank() ? 0 : s.dim_size();
- });
+ LIST_CASE(
+ shape, TF_ATTR_SHAPE, metadata.total_size = 0;
+ for (int i = 0; i < attr->list().shape_size(); ++i) {
+ const auto& s = attr->list().shape(i);
+ metadata.total_size += s.unknown_rank() ? 0 : s.dim_size();
+ });
LIST_CASE(tensor, TF_ATTR_TENSOR);
LIST_CASE(tensor, TF_ATTR_FUNC);
#undef LIST_CASE
diff --git a/tensorflow/c/c_api.h b/tensorflow/c/c_api.h
index 3d56268110edbe96616201d15a69cc8c84d3115a..c7abba85521fccec07983cd5ab4f94a8368d6181 100644
--- a/tensorflow/c/c_api.h
+++ b/tensorflow/c/c_api.h
@@ -91,7 +91,7 @@ extern "C" {
// --------------------------------------------------------------------------
// TF_Version returns a string describing version information of the
// TensorFlow library. TensorFlow using semantic versioning.
-TF_CAPI_EXPORT extern const char* TF_Version();
+TF_CAPI_EXPORT extern const char* TF_Version(void);
// --------------------------------------------------------------------------
// TF_DataType holds the type for a scalar value. E.g., one slot in a tensor.
@@ -157,7 +157,7 @@ typedef enum TF_Code {
typedef struct TF_Status TF_Status;
// Return a new status object.
-TF_CAPI_EXPORT extern TF_Status* TF_NewStatus();
+TF_CAPI_EXPORT extern TF_Status* TF_NewStatus(void);
// Delete a previously created status object.
TF_CAPI_EXPORT extern void TF_DeleteStatus(TF_Status*);
@@ -196,7 +196,7 @@ TF_CAPI_EXPORT extern TF_Buffer* TF_NewBufferFromString(const void* proto,
size_t proto_len);
// Useful for passing *out* a protobuf.
-TF_CAPI_EXPORT extern TF_Buffer* TF_NewBuffer();
+TF_CAPI_EXPORT extern TF_Buffer* TF_NewBuffer(void);
TF_CAPI_EXPORT extern void TF_DeleteBuffer(TF_Buffer*);
@@ -305,7 +305,7 @@ TF_CAPI_EXPORT extern size_t TF_StringEncodedSize(size_t len);
typedef struct TF_SessionOptions TF_SessionOptions;
// Return a new options object.
-TF_CAPI_EXPORT extern TF_SessionOptions* TF_NewSessionOptions();
+TF_CAPI_EXPORT extern TF_SessionOptions* TF_NewSessionOptions(void);
// Set the target in TF_SessionOptions.options.
// target can be empty, a single entry, or a comma separated list of entries.
@@ -338,7 +338,7 @@ TF_CAPI_EXPORT extern void TF_DeleteSessionOptions(TF_SessionOptions*);
typedef struct TF_Graph TF_Graph;
// Return a new graph object.
-TF_CAPI_EXPORT extern TF_Graph* TF_NewGraph();
+TF_CAPI_EXPORT extern TF_Graph* TF_NewGraph(void);
// Destroy an options object. Graph will be deleted once no more
// TFSession's are referencing it.
@@ -890,7 +890,8 @@ TF_CAPI_EXPORT extern void TF_GraphVersions(TF_Graph* graph,
// TF_GraphImportGraphDef.
typedef struct TF_ImportGraphDefOptions TF_ImportGraphDefOptions;
-TF_CAPI_EXPORT extern TF_ImportGraphDefOptions* TF_NewImportGraphDefOptions();
+TF_CAPI_EXPORT extern TF_ImportGraphDefOptions* TF_NewImportGraphDefOptions(
+ void);
TF_CAPI_EXPORT extern void TF_DeleteImportGraphDefOptions(
TF_ImportGraphDefOptions* opts);
@@ -1611,7 +1612,7 @@ TF_CAPI_EXPORT extern void TF_DeleteLibraryHandle(TF_Library* lib_handle);
//
// The data in the buffer will be the serialized OpList proto for ops registered
// in this address space.
-TF_CAPI_EXPORT extern TF_Buffer* TF_GetAllOpList();
+TF_CAPI_EXPORT extern TF_Buffer* TF_GetAllOpList(void);
// TF_ApiDefMap encapsulates a collection of API definitions for an operation.
//
diff --git a/tensorflow/c/c_api_experimental.cc b/tensorflow/c/c_api_experimental.cc
index f160f204dec50b6495ed11c12c48918611206b01..38e29aa74a90f4e85d1369b6928a5a58c531b2da 100644
--- a/tensorflow/c/c_api_experimental.cc
+++ b/tensorflow/c/c_api_experimental.cc
@@ -15,7 +15,10 @@ limitations under the License.
#include "tensorflow/c/c_api_experimental.h"
+#include "tensorflow/c/c_api.h"
#include "tensorflow/c/c_api_internal.h"
+#include "tensorflow/c/eager/c_api.h"
+#include "tensorflow/c/eager/c_api_internal.h"
#include "tensorflow/compiler/jit/flags.h"
#include "tensorflow/core/common_runtime/eager/attr_builder.h"
#include "tensorflow/core/framework/tensor.pb.h"
@@ -23,6 +26,7 @@ limitations under the License.
#include "tensorflow/core/graph/node_builder.h"
#include "tensorflow/core/lib/strings/strcat.h"
#include "tensorflow/core/platform/init_main.h"
+#include "tensorflow/core/platform/net.h"
#include "tensorflow/core/platform/platform.h"
#include "tensorflow/core/protobuf/config.pb.h"
#include "tensorflow/core/protobuf/tensorflow_server.pb.h"
@@ -6526,7 +6530,7 @@ library {
}
}
node_def {
- name: "ParallelInterleaveDataset/cycle_length"
+ name: "ExperimentalParallelInterleaveDataset/cycle_length"
op: "Const"
attr {
key: "dtype"
@@ -6547,7 +6551,7 @@ library {
}
}
node_def {
- name: "ParallelInterleaveDataset/block_length"
+ name: "ExperimentalParallelInterleaveDataset/block_length"
op: "Const"
attr {
key: "dtype"
@@ -6568,7 +6572,7 @@ library {
}
}
node_def {
- name: "ParallelInterleaveDataset/sloppy"
+ name: "ExperimentalParallelInterleaveDataset/sloppy"
op: "Const"
attr {
key: "dtype"
@@ -6589,7 +6593,7 @@ library {
}
}
node_def {
- name: "ParallelInterleaveDataset/buffer_output_elements"
+ name: "ExperimentalParallelInterleaveDataset/buffer_output_elements"
op: "Const"
attr {
key: "dtype"
@@ -6610,7 +6614,7 @@ library {
}
}
node_def {
- name: "ParallelInterleaveDataset/prefetch_input_elements"
+ name: "ExperimentalParallelInterleaveDataset/prefetch_input_elements"
op: "Const"
attr {
key: "dtype"
@@ -6631,14 +6635,14 @@ library {
}
}
node_def {
- name: "ParallelInterleaveDataset"
- op: "ParallelInterleaveDataset"
+ name: "ExperimentalParallelInterleaveDataset"
+ op: "ExperimentalParallelInterleaveDataset"
input: "RepeatDataset:handle:0"
- input: "ParallelInterleaveDataset/cycle_length:output:0"
- input: "ParallelInterleaveDataset/block_length:output:0"
- input: "ParallelInterleaveDataset/sloppy:output:0"
- input: "ParallelInterleaveDataset/buffer_output_elements:output:0"
- input: "ParallelInterleaveDataset/prefetch_input_elements:output:0"
+ input: "ExperimentalParallelInterleaveDataset/cycle_length:output:0"
+ input: "ExperimentalParallelInterleaveDataset/block_length:output:0"
+ input: "ExperimentalParallelInterleaveDataset/sloppy:output:0"
+ input: "ExperimentalParallelInterleaveDataset/buffer_output_elements:output:0"
+ input: "ExperimentalParallelInterleaveDataset/prefetch_input_elements:output:0"
attr {
key: "Targuments"
value {
@@ -6738,7 +6742,7 @@ library {
node_def {
name: "ShuffleDataset_2"
op: "ShuffleDataset"
- input: "ParallelInterleaveDataset:handle:0"
+ input: "ExperimentalParallelInterleaveDataset:handle:0"
input: "ShuffleDataset_2/buffer_size_1:output:0"
input: "ShuffleDataset_2/seed_2:output:0"
input: "ShuffleDataset_2/seed2_2:output:0"
@@ -8740,14 +8744,65 @@ void TFE_TensorHandlePrintDebugString(TFE_TensorHandle* handle) {
TF_DeleteStatus(status);
}
-TF_CAPI_EXPORT extern void TF_MakeInternalErrorStatus(TF_Status* status,
- const char* errMsg) {
+struct TFE_ExecuteOpNotification {
+ TFE_ExecuteOpNotification() : status(TF_NewStatus(), TF_DeleteStatus) {}
+ tensorflow::Notification n;
+ std::unique_ptr thread;
+ std::unique_ptr status;
+};
+
+TFE_ExecuteOpNotification* TFE_ExecuteOpInNewThread(TFE_Op* op,
+ TFE_TensorHandle** retvals,
+ int* num_retvals,
+ TF_Status* status) {
+ TFE_ExecuteOpNotification* n = new TFE_ExecuteOpNotification;
+
+ n->thread.reset(op->operation.EagerContext()->TFEnv()->StartThread(
+ tensorflow::ThreadOptions(), "ExecuteOpThread",
+ [op, retvals, num_retvals, n]() {
+ TFE_Execute(op, retvals, num_retvals, n->status.get());
+ n->n.Notify();
+ }));
+
+ return n;
+}
+
+void TFE_ExecuteOpNotificationWaitAndDelete(
+ TFE_ExecuteOpNotification* notification, TF_Status* status) {
+ if (notification == nullptr) {
+ status->status = tensorflow::errors::InvalidArgument(
+ "Passed in notification is a nullptr.");
+
+ return;
+ }
+ if (notification->thread == nullptr) {
+ status->status = tensorflow::errors::InvalidArgument(
+ "Passed in notification didn't start a thread correctly. Cleaning up "
+ "this notification. Please re-execute the operation to get a new "
+ "notification.");
+
+ delete notification;
+ return;
+ }
+
+ notification->n.WaitForNotification();
+
+ status->status = notification->status->status;
+
+ delete notification;
+}
+
+void TF_MakeInternalErrorStatus(TF_Status* status, const char* errMsg) {
status->status = tensorflow::errors::Internal(errMsg);
}
// This builder is used in the eager API to build a NodeDef.
struct TF_AttrBuilder : public tensorflow::AttrBuilder {
using tensorflow::AttrBuilder::AttrBuilder;
+ // The string buffers to make sure that any `attr_name` we pass into
+ // `builder->Set()` will outlive the subsequent
+ // `TF_AttrBuilderCheckCanRunOnDevice()` call(s) on the same `builder`.
+ std::set attr_names;
};
TF_AttrBuilder* TF_NewAttrBuilder(const char* op_name) {
@@ -8758,13 +8813,15 @@ void TF_DeleteAttrBuilder(TF_AttrBuilder* builder) { delete builder; }
void TF_AttrBuilderSetType(TF_AttrBuilder* builder, const char* attr_name,
TF_DataType value) {
- builder->Set(attr_name, static_cast(value));
+ auto iter = builder->attr_names.insert(attr_name).first;
+ builder->Set((*iter).c_str(), static_cast(value));
}
void TF_AttrBuilderSetTypeList(TF_AttrBuilder* builder, const char* attr_name,
const TF_DataType* values, int num_values) {
+ auto iter = builder->attr_names.insert(attr_name).first;
builder->Set(
- attr_name,
+ (*iter).c_str(),
tensorflow::gtl::ArraySlice(
reinterpret_cast(values), num_values));
}
@@ -8815,3 +8872,17 @@ int TF_OpIsStateful(const char* op_type, TF_Status* status) {
void TF_InitMain(const char* usage, int* argc, char*** argv) {
tensorflow::port::InitMain(usage, argc, argv);
}
+
+int TF_PickUnusedPortOrDie() {
+ return tensorflow::internal::PickUnusedPortOrDie();
+}
+
+TFE_TensorHandle* TFE_NewTensorHandleFromScalar(TF_DataType dtype_arg,
+ void* data, size_t len) {
+ auto dtype = static_cast(dtype_arg);
+ DCHECK(tensorflow::DataTypeCanUseMemcpy(dtype));
+
+ tensorflow::Tensor tensor(dtype, tensorflow::TensorShape({}));
+ std::memcpy(tensorflow::TensorCApi::Buffer(tensor)->data(), data, len);
+ return new TFE_TensorHandle(tensor, nullptr, nullptr);
+}
diff --git a/tensorflow/c/c_api_experimental.h b/tensorflow/c/c_api_experimental.h
index 25c03df51890a6a599528645aad6ed9ff5b49ff0..3e3a485eb763b871b0551414c4ef04746b2ed9a3 100644
--- a/tensorflow/c/c_api_experimental.h
+++ b/tensorflow/c/c_api_experimental.h
@@ -180,6 +180,25 @@ TF_CAPI_EXPORT extern TFE_TensorHandle* TFE_DequeueVariantTensor(
TF_CAPI_EXPORT extern void TFE_TensorHandlePrintDebugString(
TFE_TensorHandle* handle);
+typedef struct TFE_ExecuteOpNotification TFE_ExecuteOpNotification;
+
+// Allows invoking a kernel asynchronously, and explicitly returns a
+// notification that can be waited upon. This always executes the kernel in a
+// new thread.
+// 1. `retvals` and `num_retvals` can only be consumed after
+// `TFE_ExecuteOp` returns successfully. They shouldn't be used
+// if the return is unsuccessful
+// 2. These new APIs cannot be used together with the TFE context level async
+// support.
+TF_CAPI_EXPORT extern TFE_ExecuteOpNotification* TFE_ExecuteOpInNewThread(
+ TFE_Op* op, TFE_TensorHandle** retvals, int* num_retvals,
+ TF_Status* status);
+
+// Waits to complete the op execution, and cleans up the notification.
+// Errors reported by op execution are set in `status`.
+TF_CAPI_EXPORT extern void TFE_ExecuteOpNotificationWaitAndDelete(
+ TFE_ExecuteOpNotification* notification, TF_Status* status);
+
TF_CAPI_EXPORT extern void TF_MakeInternalErrorStatus(TF_Status* status,
const char* errMsg);
@@ -218,6 +237,15 @@ TF_CAPI_EXPORT extern int TF_OpIsStateful(const char* op_type,
// this to be called.
TF_CAPI_EXPORT void TF_InitMain(const char* usage, int* argc, char*** argv);
+// Platform-specific implementation to return an unused port. (This should used
+// in tests only.)
+TF_CAPI_EXPORT int TF_PickUnusedPortOrDie(void);
+
+// Fast path method that makes constructing a single scalar tensor require less
+// overhead and copies.
+TF_CAPI_EXPORT extern TFE_TensorHandle* TFE_NewTensorHandleFromScalar(
+ TF_DataType dtype, void* scalar, size_t len);
+
#ifdef __cplusplus
} /* end extern "C" */
#endif
diff --git a/tensorflow/c/c_api_experimental_test.cc b/tensorflow/c/c_api_experimental_test.cc
index 881dbaf35a5ec470a7e359312e33c4a27752a727..daa7701b7fe7e8ce757b6504329cf6434ad39778 100644
--- a/tensorflow/c/c_api_experimental_test.cc
+++ b/tensorflow/c/c_api_experimental_test.cc
@@ -15,6 +15,8 @@ limitations under the License.
#include "tensorflow/c/c_api_experimental.h"
#include "tensorflow/c/c_test_util.h"
+#include "tensorflow/c/eager/c_api.h"
+#include "tensorflow/c/eager/c_api_test_util.h"
#include "tensorflow/core/lib/io/path.h"
#include "tensorflow/core/platform/env.h"
#include "tensorflow/core/platform/logging.h"
@@ -173,5 +175,126 @@ TEST(CAPI_EXPERIMENTAL, IsStateful) {
EXPECT_EQ(id, 0);
}
+TEST(CAPI_EXPERIMENTAL, TFE_ExecuteOpInNewThreadTest_Simple) {
+ TF_Status* status = TF_NewStatus();
+ TFE_ContextOptions* opts = TFE_NewContextOptions();
+ TFE_Context* ctx = TFE_NewContext(opts, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TFE_DeleteContextOptions(opts);
+
+ TFE_TensorHandle* m = TestMatrixTensorHandle();
+
+ TFE_Op* matmul_op = MatMulOp(ctx, m, m);
+
+ TFE_TensorHandle* retvals[1] = {nullptr};
+ int num_retvals = 1;
+
+ auto* r =
+ TFE_ExecuteOpInNewThread(matmul_op, &retvals[0], &num_retvals, status);
+
+ TFE_ExecuteOpNotificationWaitAndDelete(r, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+
+ TF_Tensor* t = TFE_TensorHandleResolve(retvals[0], status);
+ ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ float product[4] = {0};
+ EXPECT_EQ(sizeof(product), TF_TensorByteSize(t));
+ memcpy(&product[0], TF_TensorData(t), TF_TensorByteSize(t));
+ TF_DeleteTensor(t);
+ EXPECT_EQ(7, product[0]);
+ EXPECT_EQ(10, product[1]);
+ EXPECT_EQ(15, product[2]);
+ EXPECT_EQ(22, product[3]);
+
+ TFE_DeleteOp(matmul_op);
+ TFE_DeleteTensorHandle(m);
+
+ TFE_DeleteTensorHandle(retvals[0]);
+ TFE_DeleteContext(ctx);
+ TF_DeleteStatus(status);
+}
+
+// Perform a send/recv test. Recv blocks, so they need to be executed
+// asynchronously.
+TEST(CAPI_EXPERIMENTAL, TFE_ExecuteOpInNewThreadTest_Blocking) {
+ TF_Status* status = TF_NewStatus();
+ TFE_ContextOptions* opts = TFE_NewContextOptions();
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TFE_Context* ctx = TFE_NewContext(opts, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TFE_DeleteContextOptions(opts);
+
+ // Returns a 2x2 float32 Tensor on the CPU, with data 1., 2., 3., 4.
+ TFE_TensorHandle* m = TestMatrixTensorHandle();
+
+ // Build a send op.
+ TFE_Op* send_op = TFE_NewOp(ctx, "_Send", status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TFE_OpAddInput(send_op, m, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+
+ string tensor_name = "Tensor";
+ TFE_OpSetAttrType(send_op, "T", TF_FLOAT);
+ TFE_OpSetAttrString(send_op, "tensor_name", tensor_name.c_str(),
+ tensor_name.size());
+ string send_device = "/job:localhost/replica:0/task:0/device:CPU:0";
+ TFE_OpSetAttrString(send_op, "send_device", send_device.c_str(),
+ send_device.size());
+ TFE_OpSetAttrInt(send_op, "send_device_incarnation", 1234);
+ string recv_device = "/job:localhost/replica:0/task:0/device:CPU:0";
+ TFE_OpSetAttrString(send_op, "recv_device", recv_device.c_str(),
+ recv_device.size());
+ TFE_OpSetAttrBool(send_op, "client_terminated", true);
+
+ // Build a recv op.
+ TFE_Op* recv_op = TFE_NewOp(ctx, "_Recv", status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+
+ TFE_OpSetAttrType(recv_op, "tensor_type", TF_FLOAT);
+ TFE_OpSetAttrString(recv_op, "tensor_name", tensor_name.c_str(),
+ tensor_name.size());
+ TFE_OpSetAttrString(recv_op, "send_device", send_device.c_str(),
+ send_device.size());
+ TFE_OpSetAttrInt(recv_op, "send_device_incarnation", 1234);
+ TFE_OpSetAttrString(recv_op, "recv_device", recv_device.c_str(),
+ recv_device.size());
+ TFE_OpSetAttrBool(recv_op, "client_terminated", true);
+
+ TFE_TensorHandle* send_retvals;
+ int send_num_retvals = 0;
+ auto* send_result = TFE_ExecuteOpInNewThread(send_op, &send_retvals,
+ &send_num_retvals, status);
+
+ TFE_TensorHandle* recv_retvals[1] = {nullptr};
+ int recv_num_retvals = 1;
+ auto* recv_result = TFE_ExecuteOpInNewThread(recv_op, &recv_retvals[0],
+ &recv_num_retvals, status);
+
+ TFE_ExecuteOpNotificationWaitAndDelete(send_result, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TFE_ExecuteOpNotificationWaitAndDelete(recv_result, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+
+ TF_Tensor* t = TFE_TensorHandleResolve(recv_retvals[0], status);
+ ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+
+ float product[4] = {0};
+ EXPECT_EQ(sizeof(product), TF_TensorByteSize(t));
+ memcpy(&product[0], TF_TensorData(t), TF_TensorByteSize(t));
+ TF_DeleteTensor(t);
+ EXPECT_EQ(1, product[0]);
+ EXPECT_EQ(2, product[1]);
+ EXPECT_EQ(3, product[2]);
+ EXPECT_EQ(4, product[3]);
+
+ TFE_DeleteOp(send_op);
+ TFE_DeleteOp(recv_op);
+ TFE_DeleteTensorHandle(m);
+
+ TFE_DeleteTensorHandle(recv_retvals[0]);
+ TFE_DeleteContext(ctx);
+ TF_DeleteStatus(status);
+}
+
} // namespace
} // namespace tensorflow
diff --git a/tensorflow/c/eager/BUILD b/tensorflow/c/eager/BUILD
index ba3d8533db7623b8fa7fdf35093abcd1450776b1..c34a84fcfee9b6ba9a7be86ae16e2856a2d343c7 100644
--- a/tensorflow/c/eager/BUILD
+++ b/tensorflow/c/eager/BUILD
@@ -50,6 +50,7 @@ tf_cuda_library(
],
"//conditions:default": [],
}) + [
+ "@com_google_absl//absl/memory",
"//tensorflow/core/common_runtime/eager:eager_operation",
"//tensorflow/core/distributed_runtime/eager:eager_client",
"//tensorflow/core/distributed_runtime/rpc/eager:grpc_eager_client",
@@ -143,6 +144,7 @@ tf_cuda_cc_test(
"//tensorflow/core:test",
"//tensorflow/core:test_main",
"//tensorflow/core/distributed_runtime/rpc:grpc_server_lib",
+ "@com_google_absl//absl/strings",
],
)
diff --git a/tensorflow/c/eager/c_api.cc b/tensorflow/c/eager/c_api.cc
index 192044915f06e3644aebb200a229cce5f220752b..027d752f420238da867cb9d8c116640e1730caaa 100755
--- a/tensorflow/c/eager/c_api.cc
+++ b/tensorflow/c/eager/c_api.cc
@@ -21,6 +21,7 @@ limitations under the License.
#include
#include
+#include "absl/memory/memory.h"
#include "tensorflow/c/c_api.h"
#include "tensorflow/c/c_api_internal.h"
#include "tensorflow/c/eager/c_api_internal.h"
@@ -80,7 +81,7 @@ tensorflow::Status GetAllRemoteDevices(
const std::vector& remote_workers,
tensorflow::WorkerCacheInterface* worker_cache,
std::unique_ptr* device_mgr) {
- std::vector remote_devices;
+ std::vector> remote_devices;
tensorflow::Status status;
// TODO(nareshmodi) do this in parallel instead of serially.
for (const string& remote_worker : remote_workers) {
@@ -93,7 +94,7 @@ tensorflow::Status GetAllRemoteDevices(
status = s;
if (s.ok()) {
for (tensorflow::Device* d : *devices) {
- remote_devices.push_back(d);
+ remote_devices.emplace_back(d);
}
}
n.Notify();
@@ -101,7 +102,7 @@ tensorflow::Status GetAllRemoteDevices(
n.WaitForNotification();
}
std::unique_ptr remote_device_mgr(
- new tensorflow::DeviceMgr(remote_devices));
+ new tensorflow::DeviceMgr(std::move(remote_devices)));
TF_RETURN_IF_ERROR(status);
@@ -262,13 +263,13 @@ TF_CAPI_EXPORT extern void TFE_ContextSetAsyncForThread(TFE_Context* ctx,
void TFE_DeleteContextOptions(TFE_ContextOptions* options) { delete options; }
TFE_Context* TFE_NewContext(const TFE_ContextOptions* opts, TF_Status* status) {
- std::vector devices;
+ std::vector> devices;
status->status = tensorflow::DeviceFactory::AddDevices(
opts->session_options.options, "/job:localhost/replica:0/task:0",
&devices);
if (!status->status.ok()) return nullptr;
std::unique_ptr device_mgr(
- new tensorflow::DeviceMgr(devices));
+ new tensorflow::DeviceMgr(std::move(devices)));
tensorflow::Rendezvous* r =
new tensorflow::IntraProcessRendezvous(device_mgr.get());
@@ -410,6 +411,18 @@ const char* TFE_TensorHandleDeviceName(TFE_TensorHandle* h, TF_Status* status) {
: d->name().c_str();
}
+const char* TFE_TensorHandleBackingDeviceName(TFE_TensorHandle* h,
+ TF_Status* status) {
+ if (h == nullptr || h->handle == nullptr) {
+ status->status = tensorflow::errors::InvalidArgument(
+ "The passed in handle is a nullptr");
+ return nullptr;
+ }
+ tensorflow::Device* d = h->handle->device();
+ return (d == nullptr) ? "/job:localhost/replica:0/task:0/device:CPU:0"
+ : d->name().c_str();
+}
+
TF_CAPI_EXPORT extern TFE_TensorHandle* TFE_TensorHandleCopySharingTensor(
TFE_TensorHandle* h, TF_Status* status) {
if (h == nullptr || h->handle == nullptr) {
diff --git a/tensorflow/c/eager/c_api.h b/tensorflow/c/eager/c_api.h
index b2454d872207e26feb3764671474a5d87c01f84d..120748ab763a3358b6e38e64bb3b6fd2ea32f7c3 100755
--- a/tensorflow/c/eager/c_api.h
+++ b/tensorflow/c/eager/c_api.h
@@ -48,7 +48,7 @@ extern "C" {
typedef struct TFE_ContextOptions TFE_ContextOptions;
// Return a new options object.
-TF_CAPI_EXPORT extern TFE_ContextOptions* TFE_NewContextOptions();
+TF_CAPI_EXPORT extern TFE_ContextOptions* TFE_NewContextOptions(void);
// Set the config in TF_ContextOptions.options.
// config should be a serialized tensorflow.ConfigProto proto.
@@ -169,10 +169,21 @@ TF_CAPI_EXPORT extern int64_t TFE_TensorHandleNumElements(TFE_TensorHandle* h,
TF_CAPI_EXPORT extern int64_t TFE_TensorHandleDim(TFE_TensorHandle* h,
int dim_index,
TF_Status* status);
-// This function will block till the operation that produces `h` has completed.
+
+// Returns the device of the operation that produced `h`. If `h` was produced by
+// a copy, returns the destination device of the copy. Note that the returned
+// device name is not always the device holding the tensor handle's memory. If
+// you want the latter, use TFE_TensorHandleBackingDeviceName. This function
+// will block till the operation that produces `h` has completed.
TF_CAPI_EXPORT extern const char* TFE_TensorHandleDeviceName(
TFE_TensorHandle* h, TF_Status* status);
+// Returns the name of the device in whose memory `h` resides.
+//
+// This function will block till the operation that produces `h` has completed.
+TF_CAPI_EXPORT extern const char* TFE_TensorHandleBackingDeviceName(
+ TFE_TensorHandle* h, TF_Status* status);
+
// Return a pointer to a new TFE_TensorHandle that shares the underlying tensor
// with `h`. On success, `status` is set to OK. On failure, `status` reflects
// the error and a nullptr is returned.
diff --git a/tensorflow/c/eager/c_api_test.cc b/tensorflow/c/eager/c_api_test.cc
index 0045bb5622647974a3c9f2cdf35bc21e126b4f52..6b39b79ee82f9c7baaf856e573a42b7da65691e5 100644
--- a/tensorflow/c/eager/c_api_test.cc
+++ b/tensorflow/c/eager/c_api_test.cc
@@ -16,6 +16,7 @@ limitations under the License.
#include "tensorflow/c/eager/c_api.h"
#include
+#include "absl/strings/match.h"
#include "tensorflow/c/eager/c_api_test_util.h"
#include "tensorflow/core/distributed_runtime/rpc/grpc_server_lib.h"
#include "tensorflow/core/framework/function.pb.h"
@@ -794,6 +795,14 @@ TEST(CAPI, TensorHandleNullptr) {
TF_SetStatus(status.get(), TF_OK, "");
+ device_name = TFE_TensorHandleBackingDeviceName(h, status.get());
+ ASSERT_EQ(TF_INVALID_ARGUMENT, TF_GetCode(status.get()));
+ ASSERT_EQ(device_name, nullptr);
+ ASSERT_EQ("The passed in handle is a nullptr",
+ string(TF_Message(status.get())));
+
+ TF_SetStatus(status.get(), TF_OK, "");
+
int num_dims = TFE_TensorHandleNumDims(h, status.get());
ASSERT_EQ(TF_INVALID_ARGUMENT, TF_GetCode(status.get()));
ASSERT_EQ(num_dims, -1);
@@ -809,6 +818,62 @@ TEST(CAPI, TensorHandleNullptr) {
string(TF_Message(status.get())));
}
+TEST(CAPI, TensorHandleDevices) {
+ std::unique_ptr status(
+ TF_NewStatus(), TF_DeleteStatus);
+ TFE_ContextOptions* opts = TFE_NewContextOptions();
+ TFE_Context* ctx = TFE_NewContext(opts, status.get());
+ TFE_DeleteContextOptions(opts);
+ ASSERT_EQ(TF_OK, TF_GetCode(status.get())) << TF_Message(status.get());
+
+ TFE_TensorHandle* hcpu = TestMatrixTensorHandle();
+ const char* device_name = TFE_TensorHandleDeviceName(hcpu, status.get());
+ ASSERT_EQ(TF_OK, TF_GetCode(status.get())) << TF_Message(status.get());
+ ASSERT_TRUE(absl::StrContains(device_name, "CPU:0")) << device_name;
+ const char* backing_device_name =
+ TFE_TensorHandleBackingDeviceName(hcpu, status.get());
+ ASSERT_EQ(TF_OK, TF_GetCode(status.get())) << TF_Message(status.get());
+ ASSERT_TRUE(absl::StrContains(backing_device_name, "CPU:0"))
+ << backing_device_name;
+
+ // Disable the test if no GPU is present.
+ string gpu_device_name;
+ if (GetDeviceName(ctx, &gpu_device_name, "GPU")) {
+ TFE_TensorHandle* hgpu = TFE_TensorHandleCopyToDevice(
+ hcpu, ctx, gpu_device_name.c_str(), status.get());
+ ASSERT_TRUE(TF_GetCode(status.get()) == TF_OK) << TF_Message(status.get());
+
+ TFE_Op* shape_op = ShapeOp(ctx, hgpu);
+ TFE_OpSetDevice(shape_op, gpu_device_name.c_str(), status.get());
+ ASSERT_TRUE(TF_GetCode(status.get()) == TF_OK) << TF_Message(status.get());
+ TFE_TensorHandle* retvals[1];
+ int num_retvals = 1;
+ TFE_Execute(shape_op, &retvals[0], &num_retvals, status.get());
+ ASSERT_TRUE(TF_GetCode(status.get()) == TF_OK) << TF_Message(status.get());
+
+ // .device of shape is GPU since the op is executed on GPU
+ device_name = TFE_TensorHandleDeviceName(retvals[0], status.get());
+ ASSERT_EQ(TF_OK, TF_GetCode(status.get())) << TF_Message(status.get());
+ ASSERT_TRUE(absl::StrContains(device_name, "GPU:0")) << device_name;
+
+ // .backing_device of shape is CPU since the tensor is backed by CPU
+ backing_device_name =
+ TFE_TensorHandleBackingDeviceName(retvals[0], status.get());
+ ASSERT_EQ(TF_OK, TF_GetCode(status.get())) << TF_Message(status.get());
+ ASSERT_TRUE(absl::StrContains(backing_device_name, "CPU:0"))
+ << backing_device_name;
+
+ TFE_DeleteOp(shape_op);
+ TFE_DeleteTensorHandle(retvals[0]);
+ TFE_DeleteTensorHandle(hgpu);
+ }
+
+ TFE_DeleteTensorHandle(hcpu);
+ TFE_ContextAsyncWait(ctx, status.get());
+ EXPECT_EQ(TF_OK, TF_GetCode(status.get())) << TF_Message(status.get());
+ TFE_DeleteContext(ctx);
+}
+
void Execute_MatMul_CPU(bool async) {
TF_Status* status = TF_NewStatus();
TFE_ContextOptions* opts = TFE_NewContextOptions();
diff --git a/tensorflow/c/eager/c_api_test_util.cc b/tensorflow/c/eager/c_api_test_util.cc
index 008f088c2dcdd7d9114103516a4702e47a55c6de..bd38127d50c171af801dd1b937acefdba491b4a6 100644
--- a/tensorflow/c/eager/c_api_test_util.cc
+++ b/tensorflow/c/eager/c_api_test_util.cc
@@ -104,6 +104,19 @@ TFE_Op* MatMulOp(TFE_Context* ctx, TFE_TensorHandle* a, TFE_TensorHandle* b) {
return op;
}
+TFE_Op* ShapeOp(TFE_Context* ctx, TFE_TensorHandle* a) {
+ TF_Status* status = TF_NewStatus();
+
+ TFE_Op* op = TFE_NewOp(ctx, "Shape", status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TFE_OpAddInput(op, a, status);
+ CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
+ TF_DeleteStatus(status);
+ TFE_OpSetAttrType(op, "T", TFE_TensorHandleDataType(a));
+
+ return op;
+}
+
TFE_TensorHandle* TestAxisTensorHandle() {
int64_t dims[] = {1};
int data[] = {1};
diff --git a/tensorflow/c/eager/c_api_test_util.h b/tensorflow/c/eager/c_api_test_util.h
index 474cae67c89249af3a62707f0db00ba458ca8f31..75ef9459e93b4f2ed471c423a34565594efc1714 100644
--- a/tensorflow/c/eager/c_api_test_util.h
+++ b/tensorflow/c/eager/c_api_test_util.h
@@ -37,6 +37,9 @@ TFE_TensorHandle* TestMatrixTensorHandle3X2();
// Return a matmul op multiplying `a` by `b`.
TFE_Op* MatMulOp(TFE_Context* ctx, TFE_TensorHandle* a, TFE_TensorHandle* b);
+// Return a shape op fetching the shape of `a`.
+TFE_Op* ShapeOp(TFE_Context* ctx, TFE_TensorHandle* a);
+
// Return an 1-D INT32 tensor containing a single value 1.
TFE_TensorHandle* TestAxisTensorHandle();
diff --git a/tensorflow/c/env.cc b/tensorflow/c/env.cc
new file mode 100644
index 0000000000000000000000000000000000000000..1c35ff9001d0ee1ab0fbae9e1bcc07116fab1065
--- /dev/null
+++ b/tensorflow/c/env.cc
@@ -0,0 +1,183 @@
+/* Copyright 2018 The TensorFlow Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+#include "tensorflow/c/env.h"
+
+#include "tensorflow/c/c_api_internal.h"
+#include "tensorflow/c/tf_status_helper.h"
+#include "tensorflow/core/platform/env.h"
+#include "tensorflow/core/platform/types.h"
+
+struct TF_StringStream {
+ std::vector<::tensorflow::string>* list;
+ size_t position;
+};
+
+void TF_CreateDir(const char* dirname, TF_Status* status) {
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(
+ status, ::tensorflow::Env::Default()->CreateDir(dirname));
+}
+
+void TF_DeleteDir(const char* dirname, TF_Status* status) {
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(
+ status, ::tensorflow::Env::Default()->DeleteDir(dirname));
+}
+
+void TF_DeleteRecursively(const char* dirname, int64_t* undeleted_file_count,
+ int64_t* undeleted_dir_count, TF_Status* status) {
+ ::tensorflow::int64 f, d;
+
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(
+ status, ::tensorflow::Env::Default()->DeleteRecursively(dirname, &f, &d));
+ *undeleted_file_count = f;
+ *undeleted_dir_count = d;
+}
+
+void TF_FileStat(const char* filename, TF_FileStatistics* stats,
+ TF_Status* status) {
+ ::tensorflow::FileStatistics cc_stats;
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Status s =
+ ::tensorflow::Env::Default()->Stat(filename, &cc_stats);
+ ::tensorflow::Set_TF_Status_from_Status(status, s);
+ if (s.ok()) {
+ stats->length = cc_stats.length;
+ stats->mtime_nsec = cc_stats.mtime_nsec;
+ stats->is_directory = cc_stats.is_directory;
+ }
+}
+
+void TF_NewWritableFile(const char* filename, TF_WritableFileHandle** handle,
+ TF_Status* status) {
+ std::unique_ptr<::tensorflow::WritableFile> f;
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Status s =
+ ::tensorflow::Env::Default()->NewWritableFile(filename, &f);
+ ::tensorflow::Set_TF_Status_from_Status(status, s);
+
+ if (s.ok()) {
+ *handle = reinterpret_cast(f.release());
+ }
+}
+
+void TF_CloseWritableFile(TF_WritableFileHandle* handle, TF_Status* status) {
+ auto* cc_file = reinterpret_cast<::tensorflow::WritableFile*>(handle);
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(status, cc_file->Close());
+ delete cc_file;
+}
+
+void TF_SyncWritableFile(TF_WritableFileHandle* handle, TF_Status* status) {
+ auto* cc_file = reinterpret_cast<::tensorflow::WritableFile*>(handle);
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(status, cc_file->Sync());
+}
+
+void TF_FlushWritableFile(TF_WritableFileHandle* handle, TF_Status* status) {
+ auto* cc_file = reinterpret_cast<::tensorflow::WritableFile*>(handle);
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(status, cc_file->Flush());
+}
+
+void TF_AppendWritableFile(TF_WritableFileHandle* handle, const char* data,
+ size_t length, TF_Status* status) {
+ auto* cc_file = reinterpret_cast<::tensorflow::WritableFile*>(handle);
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(
+ status, cc_file->Append(::tensorflow::StringPiece{data, length}));
+}
+
+void TF_DeleteFile(const char* filename, TF_Status* status) {
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(
+ status, ::tensorflow::Env::Default()->DeleteFile(filename));
+}
+
+bool TF_StringStreamNext(TF_StringStream* list, const char** result) {
+ if (list->position >= list->list->size()) {
+ *result = nullptr;
+ return false;
+ }
+
+ *result = list->list->at(list->position++).c_str();
+ return true;
+}
+
+void TF_StringStreamDone(TF_StringStream* list) {
+ delete list->list;
+ delete list;
+}
+TF_StringStream* TF_GetChildren(const char* dirname, TF_Status* status) {
+ auto* children = new std::vector<::tensorflow::string>;
+
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(
+ status, ::tensorflow::Env::Default()->GetChildren(dirname, children));
+
+ auto* list = new TF_StringStream;
+ list->list = children;
+ list->position = 0;
+ return list;
+}
+
+TF_StringStream* TF_GetLocalTempDirectories() {
+ auto* tmpdirs = new std::vector<::tensorflow::string>;
+
+ ::tensorflow::Env::Default()->GetLocalTempDirectories(tmpdirs);
+
+ auto* list = new TF_StringStream;
+ list->list = tmpdirs;
+ list->position = 0;
+ return list;
+}
+
+TF_CAPI_EXPORT extern uint64_t TF_NowNanos(void) {
+ return ::tensorflow::Env::Default()->NowNanos();
+}
+
+// Returns the number of microseconds since the Unix epoch.
+TF_CAPI_EXPORT extern uint64_t TF_NowMicros(void) {
+ return ::tensorflow::Env::Default()->NowMicros();
+}
+
+// Returns the number of seconds since the Unix epoch.
+TF_CAPI_EXPORT extern uint64_t TF_NowSeconds(void) {
+ return ::tensorflow::Env::Default()->NowSeconds();
+}
+
+void TF_DefaultThreadOptions(TF_ThreadOptions* options) {
+ options->stack_size = 0;
+ options->guard_size = 0;
+ options->numa_node = -1;
+}
+
+TF_Thread* TF_StartThread(const TF_ThreadOptions* options,
+ const char* thread_name, void (*work_func)(void*),
+ void* param) {
+ ::tensorflow::ThreadOptions cc_options;
+ cc_options.stack_size = options->stack_size;
+ cc_options.guard_size = options->guard_size;
+ cc_options.numa_node = options->numa_node;
+ return reinterpret_cast(::tensorflow::Env::Default()->StartThread(
+ cc_options, thread_name, [=]() { (*work_func)(param); }));
+}
+
+void TF_JoinThread(TF_Thread* thread) {
+ // ::tensorflow::Thread joins on destruction
+ delete reinterpret_cast<::tensorflow::Thread*>(thread);
+}
diff --git a/tensorflow/c/env.h b/tensorflow/c/env.h
new file mode 100644
index 0000000000000000000000000000000000000000..15652353cd7e1f1e7d7a4c665703c0166682d790
--- /dev/null
+++ b/tensorflow/c/env.h
@@ -0,0 +1,194 @@
+/* Copyright 2018 The TensorFlow Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+#include
+#include
+
+#ifndef TENSORFLOW_C_ENV_H_
+#define TENSORFLOW_C_ENV_H_
+
+#include "tensorflow/c/c_api.h"
+
+// --------------------------------------------------------------------------
+// C API for tensorflow::Env.
+
+struct TF_WritableFileHandle;
+struct TF_StringStream;
+struct TF_Thread;
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef struct TF_FileStatistics {
+ // The length of the file in bytes.
+ int64_t length;
+ // The last modified time in nanoseconds.
+ int64_t mtime_nsec;
+ // Whether the name refers to a directory.
+ bool is_directory;
+} TF_FileStatistics;
+
+typedef struct TF_ThreadOptions {
+ // Thread stack size to use (in bytes), zero implies that the system default
+ // will be used.
+ size_t stack_size;
+
+ // Guard area size to use near thread stacks to use (in bytes), zero implies
+ // that the system default will be used.
+ size_t guard_size;
+
+ // The NUMA node to use, -1 implies that there should be no NUMA affinity for
+ // this thread.
+ int numa_node;
+} TF_ThreadOptions;
+
+// Creates the specified directory. Typical status code are:
+// * TF_OK - successfully created the directory
+// * TF_ALREADY_EXISTS - directory already exists
+// * TF_PERMISSION_DENIED - dirname is not writable
+TF_CAPI_EXPORT extern void TF_CreateDir(const char* dirname, TF_Status* status);
+
+// Deletes the specified directory. Typical status codes are:
+// * TF_OK - successfully deleted the directory
+// * TF_FAILED_PRECONDITION - the directory is not empty
+TF_CAPI_EXPORT extern void TF_DeleteDir(const char* dirname, TF_Status* status);
+
+// Deletes the specified directory and all subdirectories and files underneath
+// it. This is accomplished by traversing the directory tree rooted at dirname
+// and deleting entries as they are encountered.
+//
+// If dirname itself is not readable or does not exist, *undeleted_dir_count is
+// set to 1, *undeleted_file_count is set to 0 and an appropriate status (e.g.
+// TF_NOT_FOUND) is returned.
+//
+// If dirname and all its descendants were successfully deleted, TF_OK is
+// returned and both error counters are set to zero.
+//
+// Otherwise, while traversing the tree, undeleted_file_count and
+// undeleted_dir_count are updated if an entry of the corresponding type could
+// not be deleted. The returned error status represents the reason that any one
+// of these entries could not be deleted.
+//
+// Typical status codes:
+// * TF_OK - dirname exists and we were able to delete everything underneath
+// * TF_NOT_FOUND - dirname doesn't exist
+// * TF_PERMISSION_DENIED - dirname or some descendant is not writable
+// * TF_UNIMPLEMENTED - some underlying functions (like Delete) are not
+// implemented
+TF_CAPI_EXPORT extern void TF_DeleteRecursively(const char* dirname,
+ int64_t* undeleted_file_count,
+ int64_t* undeleted_dir_count,
+ TF_Status* status);
+
+// Obtains statistics for the given path. If status is TF_OK, *stats is
+// updated, otherwise it is not touched.
+TF_CAPI_EXPORT extern void TF_FileStat(const char* filename,
+ TF_FileStatistics* stats,
+ TF_Status* status);
+
+// Creates or truncates the given filename and returns a handle to be used for
+// appending data to the file. If status is TF_OK, *handle is updated and the
+// caller is responsible for freeing it (see TF_CloseWritableFile).
+TF_CAPI_EXPORT extern void TF_NewWritableFile(const char* filename,
+ TF_WritableFileHandle** handle,
+ TF_Status* status);
+
+// Closes the given handle and frees its memory. If there was a problem closing
+// the file, it is indicated by status. Memory is freed in any case.
+TF_CAPI_EXPORT extern void TF_CloseWritableFile(TF_WritableFileHandle* handle,
+ TF_Status* status);
+
+// Syncs content of the handle to the filesystem. Blocks waiting for the
+// filesystem to indicate that the content has been persisted.
+TF_CAPI_EXPORT extern void TF_SyncWritableFile(TF_WritableFileHandle* handle,
+ TF_Status* status);
+
+// Flush local buffers to the filesystem. If the process terminates after a
+// successful flush, the contents may still be persisted, since the underlying
+// filesystem may eventually flush the contents. If the OS or machine crashes
+// after a successful flush, the contents may or may not be persisted, depending
+// on the implementation.
+TF_CAPI_EXPORT extern void TF_FlushWritableFile(TF_WritableFileHandle* handle,
+ TF_Status* status);
+
+// Appends the given bytes to the file. Any failure to do so is indicated in
+// status.
+TF_CAPI_EXPORT extern void TF_AppendWritableFile(TF_WritableFileHandle* handle,
+ const char* data,
+ size_t length,
+ TF_Status* status);
+
+// Deletes the named file and indicates whether successful in *status.
+TF_CAPI_EXPORT extern void TF_DeleteFile(const char* filename,
+ TF_Status* status);
+
+// Retrieves the next item from the given TF_StringStream and places a pointer
+// to it in *result. If no more items are in the list, *result is set to NULL
+// and false is returned.
+//
+// Ownership of the items retrieved with this function remains with the library.
+// Item points are invalidated after a call to TF_StringStreamDone.
+TF_CAPI_EXPORT extern bool TF_StringStreamNext(TF_StringStream* list,
+ const char** result);
+
+// Frees the resources associated with given string list. All pointers returned
+// by TF_StringStreamNext are invalid after this call.
+TF_CAPI_EXPORT extern void TF_StringStreamDone(TF_StringStream* list);
+
+// Retrieves the list of children of the given directory. You can iterate
+// through the list with TF_StringStreamNext. The caller is responsible for
+// freeing the list (see TF_StringStreamDone).
+TF_CAPI_EXPORT extern TF_StringStream* TF_GetChildren(const char* filename,
+ TF_Status* status);
+
+// Retrieves a list of directory names on the local machine that may be used for
+// temporary storage. You can iterate through the list with TF_StringStreamNext.
+// The caller is responsible for freeing the list (see TF_StringStreamDone).
+TF_CAPI_EXPORT extern TF_StringStream* TF_GetLocalTempDirectories(void);
+
+// Returns the number of nanoseconds since the Unix epoch.
+TF_CAPI_EXPORT extern uint64_t TF_NowNanos(void);
+
+// Returns the number of microseconds since the Unix epoch.
+TF_CAPI_EXPORT extern uint64_t TF_NowMicros(void);
+
+// Returns the number of seconds since the Unix epoch.
+TF_CAPI_EXPORT extern uint64_t TF_NowSeconds(void);
+
+// Populates a TF_ThreadOptions struct with system-default values.
+TF_CAPI_EXPORT extern void TF_DefaultThreadOptions(TF_ThreadOptions* options);
+
+// Returns a new thread that is running work_func and is identified
+// (for debugging/performance-analysis) by thread_name.
+//
+// The given param (which may be null) is passed to work_func when the thread
+// starts. In this way, data may be passed from the thread back to the caller.
+//
+// Caller takes ownership of the result and must call TF_JoinThread on it
+// eventually.
+TF_CAPI_EXPORT extern TF_Thread* TF_StartThread(const TF_ThreadOptions* options,
+ const char* thread_name,
+ void (*work_func)(void*),
+ void* param);
+
+// Waits for the given thread to finish execution, then deletes it.
+TF_CAPI_EXPORT extern void TF_JoinThread(TF_Thread* thread);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif // TENSORFLOW_C_ENV_H_
diff --git a/tensorflow/c/env_test.cc b/tensorflow/c/env_test.cc
new file mode 100644
index 0000000000000000000000000000000000000000..687ad024137352662759ec1f43df87e89faca353
--- /dev/null
+++ b/tensorflow/c/env_test.cc
@@ -0,0 +1,127 @@
+/* Copyright 2018 The TensorFlow Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+#include "tensorflow/c/env.h"
+
+#include "tensorflow/c/c_api.h"
+#include "tensorflow/core/lib/io/path.h"
+#include "tensorflow/core/platform/mutex.h"
+#include "tensorflow/core/platform/test.h"
+#include "tensorflow/core/platform/types.h"
+
+#define ASSERT_TF_OK(x) ASSERT_EQ(TF_OK, TF_GetCode(x))
+
+TEST(TestEnv, TestDirHandling) {
+ TF_StringStream* tempdirs = TF_GetLocalTempDirectories();
+ const char* tempdir;
+ bool found = false;
+ while (TF_StringStreamNext(tempdirs, &tempdir)) {
+ found = true;
+
+ TF_Status* s = TF_NewStatus();
+
+ ::tensorflow::string dirpath =
+ ::tensorflow::io::JoinPath(tempdir, "somedir");
+ TF_CreateDir(dirpath.c_str(), s);
+ ASSERT_TF_OK(s) << "TF_CreateDir failed for " << dirpath << ": "
+ << TF_Message(s);
+
+ ::tensorflow::string filepath =
+ ::tensorflow::io::JoinPath(dirpath, "somefile.txt");
+ TF_WritableFileHandle* handle;
+ TF_NewWritableFile(filepath.c_str(), &handle, s);
+ ASSERT_TF_OK(s) << "NewWritableFile failed for " << filepath << ": "
+ << TF_Message(s);
+
+ const char* data = "Hello, world!\n";
+ TF_AppendWritableFile(handle, data, strlen(data), s);
+ ASSERT_TF_OK(s) << "TF_AppendWritableFile failed to append data to file at "
+ << filepath << ": " << TF_Message(s);
+
+ TF_CloseWritableFile(handle, s);
+ ASSERT_TF_OK(s) << "TF_CloseWritableFile failed to close handle to "
+ << filepath << ": " << TF_Message(s);
+
+ TF_StringStream* children = TF_GetChildren(dirpath.c_str(), s);
+ ASSERT_TF_OK(s) << "TF_GetChildren failed for " << dirpath;
+ const char* childpath;
+ ASSERT_TRUE(TF_StringStreamNext(children, &childpath));
+ ASSERT_EQ(::tensorflow::string(childpath), "somefile.txt");
+ // There should only be one file in this directory.
+ ASSERT_FALSE(TF_StringStreamNext(children, &childpath));
+ ASSERT_EQ(childpath, nullptr);
+ TF_StringStreamDone(children);
+
+ TF_FileStatistics stats;
+ TF_FileStat(filepath.c_str(), &stats, s);
+ ASSERT_EQ(stats.length, strlen(data));
+ ASSERT_FALSE(stats.is_directory);
+ ASSERT_GT(stats.mtime_nsec, 0);
+
+ // Trying to delete a non-empty directory should fail.
+ TF_DeleteDir(dirpath.c_str(), s);
+ ASSERT_NE(TF_OK, TF_GetCode(s))
+ << "TF_DeleteDir unexpectedly succeeded with a non-empty directory "
+ << dirpath;
+
+ TF_DeleteFile(filepath.c_str(), s);
+ ASSERT_TF_OK(s) << "TF_DeleteFile failed for " << filepath << ": "
+ << TF_Message(s);
+
+ // Now deleting the directory should work.
+ TF_DeleteDir(dirpath.c_str(), s);
+ ASSERT_TF_OK(s) << "TF_DeleteDir failed for " << dirpath << ": "
+ << TF_Message(s);
+
+ TF_DeleteStatus(s);
+ break;
+ }
+
+ ASSERT_TRUE(found) << "expected at least one temp dir";
+
+ TF_StringStreamDone(tempdirs);
+}
+
+TEST(TestEnv, TestTimeFunctions) {
+ ASSERT_GE(TF_NowSeconds(), 946684800); // Midnight Jan 1, 2000
+ ASSERT_GE(TF_NowMicros(), 946684800 * 1e6);
+ ASSERT_GE(TF_NowNanos(), 946684800 * 1e9);
+}
+
+namespace {
+
+struct SomeThreadData {
+ ::tensorflow::mutex mu;
+ bool did_work = false;
+};
+
+void SomeThreadFunc(void* data) {
+ auto* real_data = static_cast(data);
+ ::tensorflow::mutex_lock l(real_data->mu);
+ real_data->did_work = true;
+}
+
+} // namespace
+
+TEST(TestEnv, TestThreads) {
+ TF_ThreadOptions options;
+ TF_DefaultThreadOptions(&options);
+ SomeThreadData data;
+ TF_Thread* thread =
+ TF_StartThread(&options, "SomeThreadName", &SomeThreadFunc, &data);
+ TF_JoinThread(thread);
+ ::tensorflow::mutex_lock l(data.mu);
+ ASSERT_TRUE(data.did_work);
+}
diff --git a/tensorflow/c/kernels.cc b/tensorflow/c/kernels.cc
index ca69345264607ac689fb556b4f5c9bc08ea5eb88..2a4eaecb6cf2740a522b1e849d1306ebde6c4577 100644
--- a/tensorflow/c/kernels.cc
+++ b/tensorflow/c/kernels.cc
@@ -15,7 +15,9 @@ limitations under the License.
#include
+#include "tensorflow/c/c_api_internal.h"
#include "tensorflow/c/kernels.h"
+#include "tensorflow/c/tf_status_helper.h"
#include "tensorflow/core/framework/kernel_def_builder.h"
#include "tensorflow/core/framework/op_kernel.h"
@@ -116,3 +118,43 @@ void TF_RegisterKernelBuilder(const char* name, TF_KernelBuilder* builder,
TF_SetStatus(status, TF_OK, "");
}
+
+int TF_NumInputs(TF_OpKernelContext* ctx) {
+ auto* cc_ctx = reinterpret_cast<::tensorflow::OpKernelContext*>(ctx);
+ return cc_ctx->num_inputs();
+}
+
+int TF_NumOutputs(TF_OpKernelContext* ctx) {
+ auto* cc_ctx = reinterpret_cast<::tensorflow::OpKernelContext*>(ctx);
+ return cc_ctx->num_outputs();
+}
+
+void TF_GetInput(TF_OpKernelContext* ctx, int i, TF_Tensor** tensor,
+ TF_Status* status) {
+ auto* cc_ctx = reinterpret_cast<::tensorflow::OpKernelContext*>(ctx);
+ if (i < 0 || i >= cc_ctx->num_inputs()) {
+ TF_SetStatus(status, TF_OUT_OF_RANGE, "input index out of range");
+ return;
+ }
+ const ::tensorflow::Tensor& cc_tensor(cc_ctx->input(i));
+ TF_Tensor* result = ::tensorflow::TF_TensorFromTensor(cc_tensor, status);
+ if (TF_GetCode(status) == TF_OK) {
+ *tensor = result;
+ }
+}
+
+void TF_SetOutput(TF_OpKernelContext* ctx, int i, const TF_Tensor* tensor,
+ TF_Status* status) {
+ auto* cc_ctx = reinterpret_cast<::tensorflow::OpKernelContext*>(ctx);
+ if (i < 0 || i >= cc_ctx->num_inputs()) {
+ TF_SetStatus(status, TF_OUT_OF_RANGE, "input index out of range");
+ return;
+ }
+ ::tensorflow::Tensor cc_tensor;
+ ::tensorflow::Status s = ::tensorflow::TF_TensorToTensor(tensor, &cc_tensor);
+ TF_SetStatus(status, TF_OK, "");
+ ::tensorflow::Set_TF_Status_from_Status(status, s);
+ if (s.ok()) {
+ cc_ctx->set_output(i, cc_tensor);
+ }
+}
diff --git a/tensorflow/c/kernels.h b/tensorflow/c/kernels.h
index 2518789a3c141755d0b3373d53642c487331f68b..1a91aa184f11ac8e45b38a1d106c7b445747a7c1 100644
--- a/tensorflow/c/kernels.h
+++ b/tensorflow/c/kernels.h
@@ -85,6 +85,32 @@ TF_CAPI_EXPORT extern void TF_RegisterKernelBuilder(const char* kernel_name,
// builder is not registered with TensorFlow via TF_RegisterKernelBuilder.
TF_CAPI_EXPORT extern void TF_DeleteKernelBuilder(TF_KernelBuilder* builder);
+// --------------------------------------------------------------------------
+// OpKernelContext routines
+
+// TF_NumInputs returns the number of inputs available in ctx.
+TF_CAPI_EXPORT extern int TF_NumInputs(TF_OpKernelContext* ctx);
+
+// TF_NumOutputs returns the number of outputs to be placed in *ctx by the
+// kernel.
+TF_CAPI_EXPORT extern int TF_NumOutputs(TF_OpKernelContext* ctx);
+
+// Retrieves the ith input from ctx. If TF_GetCode(status) is TF_OK, *tensor is
+// populated and its ownership is passed to the caller. In any other case,
+// *tensor is not modified.
+//
+// If i < 0 or i >= TF_NumInputs(ctx), *status is set to TF_OUT_OF_RANGE.
+TF_CAPI_EXPORT extern void TF_GetInput(TF_OpKernelContext* ctx, int i,
+ TF_Tensor** tensor, TF_Status* status);
+
+// Sets the ith output of ctx to tensor. If TF_GetCode(status) is anything but
+// TF_OK, ctx is left unmodified.
+//
+// If i < 0 or i >= TF_NumOutputs(ctx), *status is set to TF_OUT_OF_RANGE.
+TF_CAPI_EXPORT extern void TF_SetOutput(TF_OpKernelContext* ctx, int i,
+ const TF_Tensor* tensor,
+ TF_Status* status);
+
#ifdef __cplusplus
} /* end extern "C" */
#endif
diff --git a/tensorflow/c/kernels_test.cc b/tensorflow/c/kernels_test.cc
index e706c7c1d96ee1781d8efc0f28c5e0cbcbc80861..e659ee3c3d258a626ccf03a782ec031b5a703a48 100644
--- a/tensorflow/c/kernels_test.cc
+++ b/tensorflow/c/kernels_test.cc
@@ -15,6 +15,7 @@ limitations under the License.
#include "tensorflow/c/kernels.h"
+#include "tensorflow/c/c_api.h"
#include "tensorflow/core/framework/kernel_def.pb.h"
#include "tensorflow/core/framework/node_def.pb_text.h"
#include "tensorflow/core/framework/op.h"
@@ -31,7 +32,6 @@ struct MyCustomKernel {
static bool delete_called = false;
static void* MyCreateFunc(TF_OpKernelConstruction* ctx) {
- LOG(INFO) << "Wow, actually got into creation";
struct MyCustomKernel* s = new struct MyCustomKernel;
s->created = true;
s->compute_called = false;
@@ -51,12 +51,31 @@ static void MyDeleteFunc(void* kernel) {
delete s;
}
+namespace tensorflow {
+
+static std::unique_ptr GetFakeKernel(const char* device_name,
+ const char* op_name,
+ Status* status) {
+ NodeDef def;
+ def.set_op(op_name);
+ def.set_device(device_name);
+ def.add_input("input1");
+ def.add_input("input2");
+ return CreateOpKernel(DeviceType(device_name), nullptr, nullptr, def, 1,
+ status);
+}
+
// Tests registration of a single C kernel and checks that calls through the
// C/C++ boundary are being made.
TEST(TestKernel, TestRegisterKernelBuilder) {
const char* kernel_name = "SomeKernelName";
const char* op_name = "FooOp";
- const char* device_name = "barDev";
+ const char* device_name = "FakeDeviceName1";
+
+ REGISTER_OP(op_name)
+ .Input("input1: double")
+ .Input("input2: uint8")
+ .Output("output1: uint8");
TF_KernelBuilder* builder = TF_NewKernelBuilder(
op_name, device_name, &MyCreateFunc, &MyComputeFunc, &MyDeleteFunc);
@@ -65,35 +84,120 @@ TEST(TestKernel, TestRegisterKernelBuilder) {
TF_Status* status = TF_NewStatus();
TF_RegisterKernelBuilder(kernel_name, builder, status);
EXPECT_EQ(TF_OK, TF_GetCode(status));
- TF_Buffer* buf = TF_GetRegisteredKernelsForOp("FooOp", status);
+ TF_Buffer* buf = TF_GetRegisteredKernelsForOp(op_name, status);
EXPECT_EQ(TF_OK, TF_GetCode(status));
- ::tensorflow::KernelList list;
+ KernelList list;
list.ParseFromArray(buf->data, buf->length);
ASSERT_EQ(1, list.kernel_size());
- ASSERT_EQ("barDev", list.kernel(0).device_type());
+ ASSERT_EQ(device_name, list.kernel(0).device_type());
TF_DeleteBuffer(buf);
TF_DeleteStatus(status);
}
- REGISTER_OP("FooOp")
+ {
+ Status status;
+ std::unique_ptr kernel =
+ GetFakeKernel(device_name, op_name, &status);
+ TF_EXPECT_OK(status);
+ ASSERT_NE(nullptr, kernel.get());
+ kernel->Compute(nullptr);
+ }
+
+ ASSERT_TRUE(delete_called);
+}
+
+class DummyDevice : public DeviceBase {
+ public:
+ DummyDevice(Env* env, bool save) : DeviceBase(env), save_(save) {}
+ bool RequiresRecordingAccessedTensors() const override { return save_; }
+ Allocator* GetAllocator(AllocatorAttributes /*attr*/) override {
+ return cpu_allocator();
+ }
+
+ private:
+ bool save_;
+};
+
+TEST(TestKernel, TestInputAndOutputCount) {
+ const char* kernel_name = "InputOutputCounterKernel";
+ const char* op_name = "BarOp";
+ const char* device_name = "FakeDeviceName2";
+
+ REGISTER_OP(op_name)
.Input("input1: double")
.Input("input2: uint8")
.Output("output1: uint8");
+ static int num_inputs = 0;
+ static int num_outputs = 0;
+
+ // A kernel whose Compute function has a side-effect of updating num_inputs
+ // and num_outputs. Various functions on TF_OpKernelContext are also
+ // exercised.
+ auto my_compute_func = [](void* kernel, TF_OpKernelContext* ctx) {
+ num_inputs = TF_NumInputs(ctx);
+ num_outputs = TF_NumOutputs(ctx);
+
+ TF_Tensor* input = nullptr;
+ TF_Status* s = TF_NewStatus();
+ TF_GetInput(ctx, 0, &input, s);
+ EXPECT_EQ(TF_OK, TF_GetCode(s)) << "Failed to get input: " << TF_Message(s);
+ EXPECT_EQ(123, *static_cast(TF_TensorData(input)));
+ TF_GetInput(ctx, -1, &input, s);
+ EXPECT_EQ(TF_OUT_OF_RANGE, TF_GetCode(s));
+ TF_GetInput(ctx, 3, &input, s);
+ EXPECT_EQ(TF_OUT_OF_RANGE, TF_GetCode(s));
+
+ // Copy the input tensor to output.
+ TF_SetOutput(ctx, 0, input, s);
+ EXPECT_EQ(TF_OK, TF_GetCode(s));
+
+ TF_SetOutput(ctx, 24, input, s);
+ EXPECT_EQ(TF_OUT_OF_RANGE, TF_GetCode(s));
+
+ TF_DeleteStatus(s);
+ if (input != nullptr) {
+ TF_DeleteTensor(input);
+ }
+ };
+
+ TF_KernelBuilder* builder = TF_NewKernelBuilder(op_name, device_name, nullptr,
+ my_compute_func, nullptr);
+
+ {
+ TF_Status* status = TF_NewStatus();
+ TF_RegisterKernelBuilder(kernel_name, builder, status);
+ EXPECT_EQ(TF_OK, TF_GetCode(status));
+ TF_DeleteStatus(status);
+ }
+
{
- ::tensorflow::NodeDef def;
- def.set_op("FooOp");
- def.set_device("bar");
- def.add_input("input1");
- def.add_input("input2");
- ::tensorflow::Status status;
- std::unique_ptr<::tensorflow::OpKernel> kernel =
- ::tensorflow::CreateOpKernel(::tensorflow::DeviceType("barDev"),
- nullptr, nullptr, def, 1, &status);
+ OpKernelContext::Params p;
+ DummyDevice dummy_device(nullptr, false);
+ p.device = &dummy_device;
+
+ Tensor t(tensorflow::uint8(123));
+
+ gtl::InlinedVector inputs;
+ // Simulate 2 inputs
+ inputs.emplace_back(&t);
+ inputs.emplace_back();
+ p.inputs = &inputs;
+
+ Status status;
+ std::unique_ptr kernel =
+ GetFakeKernel(device_name, op_name, &status);
TF_EXPECT_OK(status);
ASSERT_NE(nullptr, kernel.get());
- kernel->Compute(nullptr);
- }
- ASSERT_TRUE(delete_called);
+ p.op_kernel = kernel.get();
+ OpKernelContext ctx(&p);
+ kernel->Compute(&ctx);
+
+ ASSERT_EQ(2, num_inputs);
+ ASSERT_EQ(1, num_outputs);
+ ASSERT_EQ(123, ctx.mutable_output(0)->scalar()());
+ }
}
+
+} // namespace tensorflow
diff --git a/tensorflow/c/python_api.cc b/tensorflow/c/python_api.cc
index 247236b760dd8c07bbb08426100b6a4d34296d2e..98d8393332269ae349cf8aa5c0b612c6f17172e6 100644
--- a/tensorflow/c/python_api.cc
+++ b/tensorflow/c/python_api.cc
@@ -160,4 +160,17 @@ void SetHandleShapeAndType(TF_Graph* graph, TF_Output output, const void* proto,
ic->set_output_handle_shapes_and_types(output.index, shapes_and_types);
}
+void AddWhileInputHack(TF_Graph* graph, TF_Output new_src, TF_Operation* dst,
+ TF_Status* status) {
+ mutex_lock l(graph->mu);
+ status->status = graph->graph.AddWhileInputHack(&new_src.oper->node,
+ new_src.index, &dst->node);
+ if (status->status.ok()) {
+ // This modification only updates the destination node for
+ // the purposes of running this graph in a session. Thus, we don't
+ // record the source node as being modified.
+ RecordMutation(graph, *dst, "adding input tensor");
+ }
+}
+
} // namespace tensorflow
diff --git a/tensorflow/c/python_api.h b/tensorflow/c/python_api.h
index 5cce84020bc68d912d259f51512341eb5f464a2c..44779ca656165dd65590cb5e9ea3ccf71165ed63 100644
--- a/tensorflow/c/python_api.h
+++ b/tensorflow/c/python_api.h
@@ -34,6 +34,7 @@ void SetAttr(TF_Graph* graph, TF_Operation* op, const char* attr_name,
void SetRequestedDevice(TF_Graph* graph, TF_Operation* op, const char* device);
+// Updates 'dst' to consume 'new_src'.
void UpdateEdge(TF_Graph* graph, TF_Output new_src, TF_Input dst,
TF_Status* status);
@@ -65,6 +66,13 @@ std::string GetHandleShapeAndType(TF_Graph* graph, TF_Output output);
// because I couldn't get SWIG to work otherwise.
void SetHandleShapeAndType(TF_Graph* graph, TF_Output output, const void* proto,
size_t proto_len, TF_Status* status);
+
+// This method is used to add a new input edge to 'dst', which must be a While
+// op. The While op's "T" attribute must have already been updated to include
+// the new edge. This is used to construct tf.while_loop gradients.
+void AddWhileInputHack(TF_Graph* graph, TF_Output new_src, TF_Operation* dst,
+ TF_Status* status);
+
} // namespace tensorflow
#endif // TENSORFLOW_C_PYTHON_API_H_
diff --git a/tensorflow/cc/BUILD b/tensorflow/cc/BUILD
index 83353b79f722f0a95f508b32d4a49b14b35624fb..a09becc49b10d2c58f98fbcc11df5190f794c1d4 100644
--- a/tensorflow/cc/BUILD
+++ b/tensorflow/cc/BUILD
@@ -489,6 +489,7 @@ tf_gen_op_wrappers_cc(
"image_ops",
"io_ops",
"linalg_ops",
+ "list_ops",
"logging_ops",
"lookup_ops",
"manip_ops",
diff --git a/tensorflow/cc/saved_model/BUILD b/tensorflow/cc/saved_model/BUILD
index 3d3895c8fa82c3c0e2974228e9cad767d0e00df4..52345a376cc29ee47ccb9888c9bb26292468b5a9 100644
--- a/tensorflow/cc/saved_model/BUILD
+++ b/tensorflow/cc/saved_model/BUILD
@@ -133,5 +133,6 @@ filegroup(
"testdata/half_plus_two_pbtxt/**",
"testdata/half_plus_two_main_op/**",
"testdata/half_plus_two/**",
+ "testdata/half_plus_two_v2/**",
]),
)
diff --git a/tensorflow/cc/saved_model/constants.h b/tensorflow/cc/saved_model/constants.h
index 645a3f101d1ae7dda88ec4ca622c694dc5a7a919..6f00dc324bd7054b28de2c35023581e1666bfa01 100644
--- a/tensorflow/cc/saved_model/constants.h
+++ b/tensorflow/cc/saved_model/constants.h
@@ -33,10 +33,10 @@ constexpr char kSavedModelFilenamePb[] = "saved_model.pb";
/// SavedModel text format proto filename.
constexpr char kSavedModelFilenamePbTxt[] = "saved_model.pbtxt";
-/// SavedModel legacy init op key.
+/// SavedModel legacy init op collection key. Used in v1 SavedModels.
constexpr char kSavedModelLegacyInitOpKey[] = "legacy_init_op";
-/// SavedModel main op key.
+/// SavedModel main op collection key. Used in v1 SavedModels.
constexpr char kSavedModelMainOpKey[] = "saved_model_main_op";
/// Directory in which to save the SavedModel variables.
@@ -45,6 +45,11 @@ constexpr char kSavedModelVariablesDirectory[] = "variables";
/// SavedModel variables filename.
constexpr char kSavedModelVariablesFilename[] = "variables";
+/// SavedModel SignatureDef keys for the initialization and train ops. Used in
+/// V2 SavedModels.
+constexpr char kSavedModelInitOpSignatureKey[] = "__saved_model_init_op";
+constexpr char kSavedModelTrainOpSignatureKey[] = "__saved_model_train_op";
+
} // namespace tensorflow
#endif // TENSORFLOW_CC_SAVED_MODEL_CONSTANTS_H_
diff --git a/tensorflow/cc/saved_model/loader.cc b/tensorflow/cc/saved_model/loader.cc
index ec116f68cf4b61c9b2d15065916ad9169017b659..85d3dd01fa51b3c3ba6fcbf5faac03f1ff5630e2 100644
--- a/tensorflow/cc/saved_model/loader.cc
+++ b/tensorflow/cc/saved_model/loader.cc
@@ -122,34 +122,54 @@ Status RunOnce(const RunOptions& run_options,
return run_status;
}
-bool HasMainOp(const MetaGraphDef& meta_graph_def) {
+// RunInitOp will return OK if the initialization op was run successfully.
+// An empty init_op_name indicates that there are no init ops to run.
+Status RunInitOp(const RunOptions& run_options, const string& export_dir,
+ const MetaGraphDef& meta_graph_def,
+ const std::vector& asset_file_defs,
+ Session* session, const string& init_op_name) {
+ if (!init_op_name.empty()) {
+ LOG(INFO) << "Running initialization op on SavedModel bundle.";
+ std::vector> inputs;
+ AddAssetsTensorsToInputs(export_dir, asset_file_defs, &inputs);
+ RunMetadata run_metadata;
+ return RunOnce(run_options, inputs, {}, {init_op_name},
+ nullptr /* outputs */, &run_metadata, session);
+ }
+ return Status::OK();
+}
+
+// A SavedModel may store the name of the initialization op to run in the
+// in the SignatureDef (v2) or a collection (v1). If an init_op collection
+// exists, then the collection must contain exactly one op.
+Status GetInitOp(const string& export_dir, const MetaGraphDef& meta_graph_def,
+ string* init_op_name) {
+ const auto& sig_def_map = meta_graph_def.signature_def();
+ const auto& init_op_sig_it =
+ meta_graph_def.signature_def().find(kSavedModelInitOpSignatureKey);
+ if (init_op_sig_it != sig_def_map.end()) {
+ *init_op_name = init_op_sig_it->second.outputs()
+ .find(kSavedModelInitOpSignatureKey)
+ ->second.name();
+ return Status::OK();
+ }
+
const auto& collection_def_map = meta_graph_def.collection_def();
+ string init_op_collection_key;
if (collection_def_map.find(kSavedModelMainOpKey) !=
collection_def_map.end()) {
- return true;
+ init_op_collection_key = kSavedModelMainOpKey;
+ } else {
+ init_op_collection_key = kSavedModelLegacyInitOpKey;
}
- return false;
-}
-Status RunMainOp(const RunOptions& run_options, const string& export_dir,
- const MetaGraphDef& meta_graph_def,
- const std::vector& asset_file_defs,
- Session* session, const string& main_op_key) {
- LOG(INFO) << "Running MainOp with key " << main_op_key
- << " on SavedModel bundle.";
- const auto& collection_def_map = meta_graph_def.collection_def();
- const auto main_op_it = collection_def_map.find(main_op_key);
- if (main_op_it != collection_def_map.end()) {
- if (main_op_it->second.node_list().value_size() != 1) {
+ const auto init_op_it = collection_def_map.find(init_op_collection_key);
+ if (init_op_it != collection_def_map.end()) {
+ if (init_op_it->second.node_list().value_size() != 1) {
return errors::FailedPrecondition(
strings::StrCat("Expected exactly one main op in : ", export_dir));
}
- std::vector> inputs;
- AddAssetsTensorsToInputs(export_dir, asset_file_defs, &inputs);
- RunMetadata run_metadata;
- const StringPiece main_op_name = main_op_it->second.node_list().value(0);
- return RunOnce(run_options, inputs, {}, {string(main_op_name)},
- nullptr /* outputs */, &run_metadata, session);
+ *init_op_name = init_op_it->second.node_list().value(0);
}
return Status::OK();
}
@@ -236,15 +256,12 @@ Status LoadSavedModelInternal(const SessionOptions& session_options,
bundle->meta_graph_def.saver_def().restore_op_name(),
bundle->meta_graph_def.saver_def().filename_tensor_name(),
asset_file_defs, bundle->session.get()));
- if (HasMainOp(bundle->meta_graph_def)) {
- TF_RETURN_IF_ERROR(RunMainOp(run_options, export_dir,
- bundle->meta_graph_def, asset_file_defs,
- bundle->session.get(), kSavedModelMainOpKey));
- } else {
- TF_RETURN_IF_ERROR(RunMainOp(
- run_options, export_dir, bundle->meta_graph_def, asset_file_defs,
- bundle->session.get(), kSavedModelLegacyInitOpKey));
- }
+ string init_op_name;
+ TF_RETURN_IF_ERROR(
+ GetInitOp(export_dir, bundle->meta_graph_def, &init_op_name));
+ TF_RETURN_IF_ERROR(RunInitOp(run_options, export_dir, bundle->meta_graph_def,
+ asset_file_defs, bundle->session.get(),
+ init_op_name));
return Status::OK();
}
diff --git a/tensorflow/cc/saved_model/loader_test.cc b/tensorflow/cc/saved_model/loader_test.cc
index 72b8bc18710b0ee77cb01ed3ad0c2abb5183efb2..597e42bb65ab5536664089f7e65ec52d77fc8f23 100644
--- a/tensorflow/cc/saved_model/loader_test.cc
+++ b/tensorflow/cc/saved_model/loader_test.cc
@@ -36,6 +36,8 @@ constexpr char kTestDataMainOp[] =
"cc/saved_model/testdata/half_plus_two_main_op/00000123";
constexpr char kTestDataSharded[] =
"cc/saved_model/testdata/half_plus_two/00000123";
+constexpr char kTestDataInitOpV2[] =
+ "cc/saved_model/testdata/half_plus_two_v2/00000123";
class LoaderTest : public ::testing::Test {
protected:
@@ -227,5 +229,17 @@ TEST_F(LoaderTest, MaybeSavedModelDirectory) {
EXPECT_FALSE(MaybeSavedModelDirectory(invalid_export_dir));
}
+TEST_F(LoaderTest, SavedModelInitOpV2Format) {
+ SavedModelBundle bundle;
+ SessionOptions session_options;
+ RunOptions run_options;
+
+ const string export_dir =
+ io::JoinPath(testing::TensorFlowSrcRoot(), kTestDataInitOpV2);
+ TF_ASSERT_OK(LoadSavedModel(session_options, run_options, export_dir,
+ {kSavedModelTagServe}, &bundle));
+ CheckSavedModelBundle(export_dir, bundle);
+}
+
} // namespace
} // namespace tensorflow
diff --git a/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/assets/foo.txt b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/assets/foo.txt
new file mode 100644
index 0000000000000000000000000000000000000000..f9ff036688007836524129e23f5cf82edd1e8910
--- /dev/null
+++ b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/assets/foo.txt
@@ -0,0 +1 @@
+asset-file-contents
\ No newline at end of file
diff --git a/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/saved_model.pb b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/saved_model.pb
new file mode 100644
index 0000000000000000000000000000000000000000..a10bbf8fb6bca0fcee6414b2927d2f706de85ebc
Binary files /dev/null and b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/saved_model.pb differ
diff --git a/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/variables/variables.data-00000-of-00001 b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/variables/variables.data-00000-of-00001
new file mode 100644
index 0000000000000000000000000000000000000000..15b75d6ef6bffc336d138d923badb3928b8c4c13
Binary files /dev/null and b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/variables/variables.data-00000-of-00001 differ
diff --git a/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/variables/variables.index b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/variables/variables.index
new file mode 100644
index 0000000000000000000000000000000000000000..7ec9fb4fe2dd21d0a6c324aecd7658fc37cf2326
Binary files /dev/null and b/tensorflow/cc/saved_model/testdata/half_plus_two_v2/00000123/variables/variables.index differ
diff --git a/tensorflow/compat_template_v1.__init__.py b/tensorflow/compat_template_v1.__init__.py
index 7df80ec01245a7fe820c79d5879458c4cd0a93cb..d58acde09f007bc9df40b08b0ef79c6031ca7941 100644
--- a/tensorflow/compat_template_v1.__init__.py
+++ b/tensorflow/compat_template_v1.__init__.py
@@ -23,12 +23,12 @@ import os as _os
# pylint: disable=g-bad-import-order
from tensorflow.python import pywrap_tensorflow # pylint: disable=unused-import
+# API IMPORTS PLACEHOLDER
+
from tensorflow.python.tools import component_api_helper as _component_api_helper
_component_api_helper.package_hook(
parent_package_str=__name__,
child_package_str=('tensorflow_estimator.python.estimator.api.estimator'))
-# API IMPORTS PLACEHOLDER
-
from tensorflow.python.platform import flags # pylint: disable=g-import-not-at-top
app.flags = flags # pylint: disable=undefined-variable
diff --git a/tensorflow/compiler/aot/codegen.cc b/tensorflow/compiler/aot/codegen.cc
index b17bc658fa06b9feb7edb292bd89ef31e6309169..ab1c1be344e2257721507543bc7647d4ff4becb2 100644
--- a/tensorflow/compiler/aot/codegen.cc
+++ b/tensorflow/compiler/aot/codegen.cc
@@ -164,7 +164,8 @@ string RewriteWithName(const string& name, string code,
}
// Generate methods for args (inputs).
-Status GenArgMethods(const tf2xla::Config& config, const xla::ProgramShape& ps,
+Status GenArgMethods(const tf2xla::Config& config,
+ const xla::ProgramShapeProto& ps,
const CompileResult& compile_result, string* methods) {
size_t num_args = ps.parameters_size();
if (config.feed_size() != num_args) {
@@ -174,9 +175,10 @@ Status GenArgMethods(const tf2xla::Config& config, const xla::ProgramShape& ps,
}
for (int i = 0; i < num_args; ++i) {
std::vector> rewrites;
- TF_RETURN_IF_ERROR(AddRewritesForShape(i, ps.parameters(i), &rewrites));
+ TF_RETURN_IF_ERROR(
+ AddRewritesForShape(i, xla::Shape(ps.parameters(i)), &rewrites));
const string code = R"(
- void set_arg{{NAME}}_data(void* data) {
+ void set_arg{{NAME}}_data(const void* data) {
set_arg_data({{I}}, data);
}
{{TYPE}}* arg{{NAME}}_data() {
@@ -204,7 +206,7 @@ Status GenArgMethods(const tf2xla::Config& config, const xla::ProgramShape& ps,
// Generate methods for results (outputs).
Status GenResultMethods(const tf2xla::Config& config,
- const xla::ProgramShape& ps, string* methods) {
+ const xla::ProgramShapeProto& ps, string* methods) {
if (ps.result().element_type() != xla::TUPLE) {
// The XlaCompiler we use to build the xla computation always generates a
// tuple result, and we rely on this to simplify code generation.
@@ -217,8 +219,8 @@ Status GenResultMethods(const tf2xla::Config& config,
}
for (int i = 0; i < ps.result().tuple_shapes_size(); ++i) {
std::vector> rewrites;
- TF_RETURN_IF_ERROR(
- AddRewritesForShape(i, ps.result().tuple_shapes(i), &rewrites));
+ TF_RETURN_IF_ERROR(AddRewritesForShape(
+ i, xla::Shape(ps.result().tuple_shapes(i)), &rewrites));
string code = R"(
{{TYPE}}* result{{NAME}}_data() {
return static_cast<{{TYPE}}*>(result_data({{I}}));
@@ -336,7 +338,7 @@ Status GenerateHeader(const CodegenOpts& opts, const tf2xla::Config& config,
ExtractEntryParamBufferInfos(buffer_infos);
std::vector buffer_infos_for_temps =
ExtractTempBufferInfos(buffer_infos);
- const xla::ProgramShape& ps = compile_result.program_shape;
+ const xla::ProgramShapeProto& ps = compile_result.program_shape;
string methods_arg, methods_result;
TF_RETURN_IF_ERROR(GenArgMethods(config, ps, compile_result, &methods_arg));
TF_RETURN_IF_ERROR(GenResultMethods(config, ps, &methods_result));
@@ -548,8 +550,8 @@ class {{CLASS}} : public tensorflow::XlaCompiledCpuFunction {
static const char** StaticResultNames() {{RESULT_NAMES_CODE}}
// Shape of the args and results.
- static const xla::ProgramShape* StaticProgramShape() {
- static const xla::ProgramShape* kShape = {{PROGRAM_SHAPE_SHIM_EXPRESSION}};
+ static const xla::ProgramShapeProto* StaticProgramShape() {
+ static const xla::ProgramShapeProto* kShape = {{PROGRAM_SHAPE_SHIM_EXPRESSION}};
return kShape;
}
@@ -587,7 +589,7 @@ class {{CLASS}} : public tensorflow::XlaCompiledCpuFunction {
{"{{METHODS_RESULT}}\n", methods_result},
{"{{NS_END}}\n", ns_end},
{"{{NS_START}}\n", ns_start},
- {"{{PROGRAM_SHAPE}}", xla::ShapeUtil::HumanString(ps)},
+ {"{{PROGRAM_SHAPE}}", xla::ShapeUtil::HumanString(xla::ProgramShape(ps))},
{"{{PROGRAM_SHAPE_SHIM_EXPRESSION}}",
metadata_result.program_shape_access_shim},
{"{{RESULT_INDEX}}", absl::StrCat(result_index)},
@@ -615,11 +617,11 @@ static string CreateUniqueIdentifier(const CodegenOpts& opts,
Status GenerateMetadata(const CodegenOpts& opts,
const CompileResult& compile_result,
MetadataResult* metadata_result) {
- std::unique_ptr program_shape;
+ std::unique_ptr program_shape;
if (opts.gen_program_shape) {
program_shape =
- absl::make_unique(compile_result.program_shape);
+ absl::make_unique(compile_result.program_shape);
// The parameter names are currently meaningless, and redundant with the
// rest of our metadata, so clear them out to avoid confusion and save
@@ -631,8 +633,8 @@ Status GenerateMetadata(const CodegenOpts& opts,
// a shim that evaluates to nullptr, which is what we want.
ProtobufToEmbed program_shape_protobuf{
- CreateUniqueIdentifier(opts, "ProgramShape"), "xla::ProgramShape",
- program_shape.get()};
+ CreateUniqueIdentifier(opts, "ProgramShapeProto"),
+ "xla::ProgramShapeProto", program_shape.get()};
ProtobufToEmbed hlo_profile_printer_data_protobuf{
CreateUniqueIdentifier(opts, "HloProfilePrinterData"),
diff --git a/tensorflow/compiler/aot/codegen.h b/tensorflow/compiler/aot/codegen.h
index 90410c46a8e36e44454f1219ad76d0fb0937070d..9485e86b10e225a3c9c12eafd9905bdf7c15c9fa 100644
--- a/tensorflow/compiler/aot/codegen.h
+++ b/tensorflow/compiler/aot/codegen.h
@@ -57,7 +57,7 @@ struct MetadataResult {
std::vector header_variable_decls;
// program_shape_access_shim is a C++ expression that constructs the
- // xla::ProgramShape instance for the CompileResult passed to
+ // xla::ProgramShapeProto instance for the CompileResult passed to
// GenerateMetadata.
string program_shape_access_shim;
diff --git a/tensorflow/compiler/aot/codegen_test.cc b/tensorflow/compiler/aot/codegen_test.cc
index bb288d23000527be74f01630d20bbf82e50007ce..c1788ca32a1d099284eeb870f9513891051fd29e 100644
--- a/tensorflow/compiler/aot/codegen_test.cc
+++ b/tensorflow/compiler/aot/codegen_test.cc
@@ -181,13 +181,15 @@ TEST(CodegenTest, Golden) {
BufferInfo::MakeEntryParameter(/*size=*/96, /*param_number=*/1),
BufferInfo::MakeTempBuffer(3), BufferInfo::MakeTempBuffer(120)},
5, {}));
- compile_result.program_shape = xla::ShapeUtil::MakeProgramShape(
- {
- xla::ShapeUtil::MakeShape(xla::F32, {1, 2}),
- xla::ShapeUtil::MakeShape(xla::S64, {3, 4}),
- },
- xla::ShapeUtil::MakeTupleShape(
- {xla::ShapeUtil::MakeShape(xla::U32, {5, 6})}));
+ compile_result.program_shape =
+ xla::ShapeUtil::MakeProgramShape(
+ {
+ xla::ShapeUtil::MakeShape(xla::F32, {1, 2}),
+ xla::ShapeUtil::MakeShape(xla::S64, {3, 4}),
+ },
+ xla::ShapeUtil::MakeTupleShape(
+ {xla::ShapeUtil::MakeShape(xla::U32, {5, 6})}))
+ .ToProto();
compile_result.entry_point = "entry_point";
compile_result.pointer_size = 8;
diff --git a/tensorflow/compiler/aot/codegen_test_h.golden b/tensorflow/compiler/aot/codegen_test_h.golden
index e4d8a02877c75fa72c5747650ab9c7ac229955b3..968afad65ed6d4b5510687df484b7ce6743f6a85 100644
--- a/tensorflow/compiler/aot/codegen_test_h.golden
+++ b/tensorflow/compiler/aot/codegen_test_h.golden
@@ -22,7 +22,7 @@ extern "C" void entry_point(
void* result, const xla::ExecutableRunOptions* run_options,
const void** args, void** temps, tensorflow::int64* profile_counters);
-extern "C" char __tfcompile_foo_bar_MyClass_ProgramShape_protobuf_array_contents[];
+extern "C" char __tfcompile_foo_bar_MyClass_ProgramShapeProto_protobuf_array_contents[];
namespace foo {
@@ -114,7 +114,7 @@ class MyClass : public tensorflow::XlaCompiledCpuFunction {
// with dim indices specifying which value. No bounds checking is performed
// on dim indices.
- void set_arg0_data(void* data) {
+ void set_arg0_data(const void* data) {
set_arg_data(0, data);
}
float* arg0_data() {
@@ -132,7 +132,7 @@ class MyClass : public tensorflow::XlaCompiledCpuFunction {
arg_data(0)))[dim0][dim1];
}
- void set_arg_myfeed_data(void* data) {
+ void set_arg_myfeed_data(const void* data) {
set_arg_data(0, data);
}
float* arg_myfeed_data() {
@@ -150,7 +150,7 @@ class MyClass : public tensorflow::XlaCompiledCpuFunction {
arg_data(0)))[dim0][dim1];
}
- void set_arg1_data(void* data) {
+ void set_arg1_data(const void* data) {
set_arg_data(1, data);
}
tensorflow::int64* arg1_data() {
@@ -253,10 +253,10 @@ class MyClass : public tensorflow::XlaCompiledCpuFunction {
}
// Shape of the args and results.
- static const xla::ProgramShape* StaticProgramShape() {
- static const xla::ProgramShape* kShape = []() {
- xla::ProgramShape* proto = new xla::ProgramShape;
- proto->ParseFromArray(&__tfcompile_foo_bar_MyClass_ProgramShape_protobuf_array_contents[0], 52);
+ static const xla::ProgramShapeProto* StaticProgramShape() {
+ static const xla::ProgramShapeProto* kShape = []() {
+ xla::ProgramShapeProto* proto = new xla::ProgramShapeProto;
+ proto->ParseFromArray(&__tfcompile_foo_bar_MyClass_ProgramShapeProto_protobuf_array_contents[0], 52);
return proto;
}();
return kShape;
diff --git a/tensorflow/compiler/aot/codegen_test_o.golden b/tensorflow/compiler/aot/codegen_test_o.golden
index eb001c5d45bdfefc76629d7303d89f5480432235..ce8e5ec8c96a2c3696f14b8eea206d648182ecb5 100644
Binary files a/tensorflow/compiler/aot/codegen_test_o.golden and b/tensorflow/compiler/aot/codegen_test_o.golden differ
diff --git a/tensorflow/compiler/aot/compile.cc b/tensorflow/compiler/aot/compile.cc
index 2b5f97b34cd928d32eb220536342c715d91d45bb..9fc223bdc7c0e207ce2005cb86250aa77e709df8 100644
--- a/tensorflow/compiler/aot/compile.cc
+++ b/tensorflow/compiler/aot/compile.cc
@@ -56,17 +56,23 @@ Status CompileXla(xla::CompileOnlyClient* client,
return errors::Unknown("Couldn't get XLA program shape: ",
pshape_or.status().error_message());
}
- compile_result->program_shape = *pshape_or.ValueOrDie();
- xla::ProgramShape* pshape = &compile_result->program_shape;
- std::vector arg_layouts;
- arg_layouts.reserve(pshape->parameters_size());
+ compile_result->program_shape = pshape_or.ValueOrDie()->ToProto();
+ xla::ProgramShapeProto* pshape = &compile_result->program_shape;
+
+ // AotXlaComputationInstance::argument_layouts is a vector of Shape
+ // pointers. Accumulate the Shape objects themselves in a separate vector
+ // while building the vector of pointers.
+ std::vector arg_layout_ptrs(pshape->parameters_size());
+ std::vector arg_layouts(pshape->parameters_size());
for (int i = 0; i < pshape->parameters_size(); ++i) {
- arg_layouts.push_back(pshape->mutable_parameters(i));
+ arg_layouts[i] = xla::Shape(*pshape->mutable_parameters(i));
+ arg_layout_ptrs[i] = &arg_layouts[i];
}
xla::CompileOnlyClient::AotXlaComputationInstance instance;
instance.computation = &computation;
- instance.argument_layouts = std::move(arg_layouts);
- instance.result_layout = &pshape->result();
+ instance.argument_layouts = std::move(arg_layout_ptrs);
+ xla::Shape result_shape(pshape->result());
+ instance.result_layout = &result_shape;
xla::StatusOr>>
aot_or = client->CompileAheadOfTime({instance}, aot_opts);
if (!aot_or.ok()) {
diff --git a/tensorflow/compiler/aot/compile.h b/tensorflow/compiler/aot/compile.h
index e03c5b1aa77c1262ed903aae3072ef65f34d80a2..ee7bb26fabd2d897b85b62f38778ecbfe2238eb6 100644
--- a/tensorflow/compiler/aot/compile.h
+++ b/tensorflow/compiler/aot/compile.h
@@ -33,9 +33,9 @@ namespace tfcompile {
struct CompileResult {
// Contains object file and meta-info.
std::unique_ptr aot;
- xla::ProgramShape program_shape; // Static shape of args and results.
- string entry_point; // Name of generated function.
- int pointer_size = 0; // Size of a pointer in bytes.
+ xla::ProgramShapeProto program_shape; // Static shape of args and results.
+ string entry_point; // Name of generated function.
+ int pointer_size = 0; // Size of a pointer in bytes.
};
// CompileGraph compiles the graph_def into an object file containing a function
diff --git a/tensorflow/compiler/aot/tests/tfcompile_test.cc b/tensorflow/compiler/aot/tests/tfcompile_test.cc
index f10852c7850f61bfd8b99fa9f1648202d182085e..4dd79e5882d7da61be029735ef2b165908c599f9 100644
--- a/tensorflow/compiler/aot/tests/tfcompile_test.cc
+++ b/tensorflow/compiler/aot/tests/tfcompile_test.cc
@@ -526,13 +526,15 @@ TEST(TFCompileTest, ProgramShape) {
// muladd has the program shape defined.
MatMulAndAddComp muladd;
- const xla::ProgramShape* muladd_shape = muladd.ProgramShape();
+ const xla::ProgramShapeProto* muladd_shape = muladd.ProgramShape();
ASSERT_TRUE(muladd_shape != nullptr);
ASSERT_EQ(muladd_shape->parameters_size(), 2);
- EXPECT_TRUE(ShapeUtil::Compatible(muladd_shape->parameters(0), f32_2x2));
- EXPECT_TRUE(ShapeUtil::Compatible(muladd_shape->parameters(1), f32_2x2));
+ EXPECT_TRUE(
+ ShapeUtil::Compatible(xla::Shape(muladd_shape->parameters(0)), f32_2x2));
+ EXPECT_TRUE(
+ ShapeUtil::Compatible(xla::Shape(muladd_shape->parameters(1)), f32_2x2));
- const xla::Shape& muladd_result = muladd_shape->result();
+ const xla::Shape muladd_result(muladd_shape->result());
ASSERT_EQ(muladd_result.element_type(), xla::TUPLE);
ASSERT_EQ(ShapeUtil::TupleElementCount(muladd_result), 2);
const xla::Shape& muladd_result0 =
diff --git a/tensorflow/compiler/jit/BUILD b/tensorflow/compiler/jit/BUILD
index 682c0f0cb05c8c83acac28c8f3abf4f5e355e7c0..d8c88a9fca2db74265b4962e07a66ab214b1d994 100644
--- a/tensorflow/compiler/jit/BUILD
+++ b/tensorflow/compiler/jit/BUILD
@@ -23,7 +23,6 @@ package(
load("//tensorflow:tensorflow.bzl", "cc_header_only_library")
load("//tensorflow:tensorflow.bzl", "tf_cc_test")
load("@local_config_cuda//cuda:build_defs.bzl", "if_cuda")
-load("@local_config_cuda//cuda:build_defs.bzl", "if_cuda_is_configured")
load("//tensorflow:tensorflow.bzl", "tf_cuda_cc_test")
load("//tensorflow:tensorflow.bzl", "tf_custom_op_py_library")
@@ -38,7 +37,7 @@ cc_library(
":xla_cpu_device",
":xla_cpu_jit",
"//tensorflow/compiler/plugin",
- ] + if_cuda_is_configured([
+ ] + if_cuda([
":xla_gpu_device",
":xla_gpu_jit",
]),
@@ -51,6 +50,7 @@ cc_library(
deps = [
":jit_compilation_passes",
"//tensorflow/compiler/jit/kernels:xla_ops",
+ "//tensorflow/compiler/tf2xla/kernels:xla_cpu_only_ops",
"//tensorflow/compiler/tf2xla/kernels:xla_dummy_ops",
"//tensorflow/compiler/tf2xla/kernels:xla_ops",
"//tensorflow/compiler/xla/service:cpu_plugin",
@@ -76,6 +76,7 @@ cc_library(
srcs = ["xla_cpu_device.cc"],
visibility = [":friends"],
deps = [
+ ":create_xla_launch_op", # buildcleaner: keep
":flags",
":jit_compilation_passes",
":xla_device",
@@ -95,6 +96,7 @@ cc_library(
srcs = ["xla_gpu_device.cc"],
visibility = [":friends"],
deps = [
+ ":create_xla_launch_op", # buildcleaner: keep
":jit_compilation_passes",
":xla_device",
"//tensorflow/compiler/jit/kernels:xla_ops",
@@ -104,6 +106,7 @@ cc_library(
"//tensorflow/core:core_cpu_internal",
"//tensorflow/core:lib",
"@com_google_absl//absl/memory",
+ "@com_google_absl//absl/strings",
],
alwayslink = 1,
)
@@ -268,6 +271,7 @@ cc_library(
"//tensorflow/compiler/tf2xla:common",
"//tensorflow/compiler/tf2xla:dump_graph",
"//tensorflow/compiler/tf2xla:xla_compiler",
+ "//tensorflow/compiler/xla:debug_options_flags",
"//tensorflow/compiler/xla:statusor",
"//tensorflow/compiler/xla/client:client_library",
"//tensorflow/compiler/xla/client:local_client",
@@ -511,6 +515,7 @@ cc_library(
"//tensorflow/compiler/jit/ops:xla_ops",
"//tensorflow/compiler/tf2xla:dump_graph",
"//tensorflow/compiler/tf2xla:resource_operation_table",
+ "//tensorflow/compiler/tf2xla:side_effect_util",
"//tensorflow/compiler/tf2xla:tf2xla_util",
"//tensorflow/compiler/tf2xla:xla_compiler",
"//tensorflow/compiler/tf2xla/cc:xla_jit_ops",
@@ -609,6 +614,7 @@ tf_cc_test(
"//tensorflow/cc:cc_ops",
"//tensorflow/cc:cc_ops_internal",
"//tensorflow/cc:function_ops",
+ "//tensorflow/cc:functional_ops",
"//tensorflow/cc:ops",
"//tensorflow/cc:resource_variable_ops",
"//tensorflow/cc:scope",
@@ -621,6 +627,7 @@ tf_cc_test(
"//tensorflow/compiler/tf2xla/cc:xla_ops",
"//tensorflow/compiler/tf2xla/kernels:xla_dummy_ops",
"//tensorflow/compiler/tf2xla/kernels:xla_ops",
+ "//tensorflow/compiler/xla:test",
"//tensorflow/core:core_cpu",
"//tensorflow/core:framework",
"//tensorflow/core:framework_internal",
@@ -736,7 +743,10 @@ tf_custom_op_py_library(
visibility = [
":friends",
],
- deps = ["//tensorflow/compiler/jit/ops:xla_ops_wrapper_py"],
+ deps = [
+ "//tensorflow/compiler/jit/ops:xla_ops_grad",
+ "//tensorflow/compiler/jit/ops:xla_ops_wrapper_py",
+ ],
)
# This target can be used by XLA device plugins to prevent circular dependencies, and provides access to all of the required headers for building a device library.
diff --git a/tensorflow/compiler/jit/build_xla_ops_pass_test.cc b/tensorflow/compiler/jit/build_xla_ops_pass_test.cc
index 11df946cc186660242574c2644463a26ead44f1f..48a23a4c1711ac88a329723c46559112d5a39dbd 100644
--- a/tensorflow/compiler/jit/build_xla_ops_pass_test.cc
+++ b/tensorflow/compiler/jit/build_xla_ops_pass_test.cc
@@ -42,14 +42,8 @@ class BuildXlaOpsTest : public ::testing::Test {
.ok());
}
- void TearDown() override {
- for (Device* device : devices_) {
- delete device;
- }
- }
-
private:
- std::vector devices_;
+ std::vector> devices_;
};
using ::tensorflow::testing::FindNodeByName;
diff --git a/tensorflow/compiler/jit/create_xla_launch_op_test.cc b/tensorflow/compiler/jit/create_xla_launch_op_test.cc
index 73866607621cd745f6e640a14405daebf0dd9985..0f872a480f4d4843217f1df3452c4dc62531264e 100644
--- a/tensorflow/compiler/jit/create_xla_launch_op_test.cc
+++ b/tensorflow/compiler/jit/create_xla_launch_op_test.cc
@@ -59,8 +59,9 @@ class CreateXlaLaunchOpTest : public ::testing::Test {
SessionOptions options;
auto* device_count = options.config.mutable_device_count();
device_count->insert({"CPU", 1});
+ std::vector> devices;
TF_CHECK_OK(DeviceFactory::AddDevices(
- options, "/job:localhost/replica:0/task:0", &devices_));
+ options, "/job:localhost/replica:0/task:0", &devices));
FunctionDefLibrary proto;
for (const auto& fdef : flib) {
@@ -69,7 +70,7 @@ class CreateXlaLaunchOpTest : public ::testing::Test {
lib_def_ = absl::make_unique(
OpRegistry::Global(), proto);
OptimizerOptions opts;
- device_mgr_ = absl::make_unique(devices_);
+ device_mgr_ = absl::make_unique(std::move(devices));
pflr_ = absl::make_unique(
device_mgr_.get(), Env::Default(), TF_GRAPH_DEF_VERSION, lib_def_.get(),
opts, /*default_thread_pool=*/nullptr, /*cluster_flr=*/nullptr);
@@ -77,7 +78,6 @@ class CreateXlaLaunchOpTest : public ::testing::Test {
}
FunctionLibraryRuntime* flr_;
- std::vector devices_;
std::unique_ptr device_mgr_;
std::unique_ptr lib_def_;
std::unique_ptr pflr_;
diff --git a/tensorflow/compiler/jit/encapsulate_subgraphs_pass_test.cc b/tensorflow/compiler/jit/encapsulate_subgraphs_pass_test.cc
index de89be9a3555960dabe7bacd17226c15ae888ae6..7476d1dc51d3beebce087fd687d971b1465607a2 100644
--- a/tensorflow/compiler/jit/encapsulate_subgraphs_pass_test.cc
+++ b/tensorflow/compiler/jit/encapsulate_subgraphs_pass_test.cc
@@ -299,7 +299,7 @@ REGISTER_OP("XlaHostCompute")
.Attr("Toutputs: list(type) >= 0")
.Attr("ancestors: list(string) >= 0")
.Attr("key: string")
- .Attr("shape_inference_graph: string = ''")
+ .Attr("shape_inference_graph: func")
.Attr("shapes: list(shape) >= 0")
.SetShapeFn(::tensorflow::shape_inference::UnknownShape);
@@ -901,18 +901,22 @@ TEST(EncapsulateSubgraphsTest, OneFunctionOneOutside) {
{
GraphDefBuilder shape(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape.opts());
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT, DT_FLOAT}, shape.opts());
+ Node* recv = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT, DT_FLOAT},
+ shape.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Binary(ops::NodeOut(recv, 0), ops::NodeOut(recv, 1),
shape.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, shape.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ shape.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape, "F1_O1", &library_expected));
}
+ NameAttrList shape_inference_graph;
+ shape_inference_graph.set_name("_outside_compilation_shape_inference_F1_O1");
*library_expected.add_function() = test::function::XTimesTwo();
*library_expected.add_function() = FunctionDefHelper::Create(
"F1", {"a_0_arg:float", "b_0_arg:float"}, {"f_0_retval_retval:float"}, {},
@@ -931,8 +935,7 @@ TEST(EncapsulateSubgraphsTest, OneFunctionOneOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O1"},
+ {"shape_inference_graph", shape_inference_graph},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}},
{"c"}},
@@ -948,8 +951,9 @@ TEST(EncapsulateSubgraphsTest, OneFunctionOneOutside) {
Node* key_constant =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT, DT_FLOAT}, b2.opts());
+ Node* recv = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT, DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Binary(ops::NodeOut(recv, 0), ops::NodeOut(recv, 1),
b2.opts()
.WithName("E")
@@ -957,7 +961,8 @@ TEST(EncapsulateSubgraphsTest, OneFunctionOneOutside) {
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
Node* send = SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
- b2.opts().WithControlInput(e));
+ b2.opts().WithControlInput(e).WithAttr(
+ kXlaHasHostTransferAttrName, true));
Node* s = Sequencer(
b2.opts().WithName("F1_sequencer").WithControlInputs({recv, send}),
@@ -1022,14 +1027,16 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
{
GraphDefBuilder shape1(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape1.opts());
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT, DT_FLOAT}, shape1.opts());
+ Node* recv = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT, DT_FLOAT},
+ shape1.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Binary(ops::NodeOut(recv, 0), ops::NodeOut(recv, 1),
shape1.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, shape1.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ shape1.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape1, "F1_O1", &library_expected));
}
@@ -1037,25 +1044,31 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
{
GraphDefBuilder shape2(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape2.opts());
- Node* recv1 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT, DT_FLOAT}, shape2.opts());
+ Node* recv1 = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT, DT_FLOAT},
+ shape2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Binary(ops::NodeOut(recv1, 0), ops::NodeOut(recv1, 1),
shape2.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2",
- {DT_FLOAT, DT_FLOAT}, shape2.opts());
+ Node* recv2 = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O2", {DT_FLOAT, DT_FLOAT},
+ shape2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* h = Binary(ops::NodeOut(recv2, 1), e,
shape2.opts()
.WithName("H")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O2"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {h}, shape2.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {h},
+ shape2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape2, "F1_O2", &library_expected));
}
+ NameAttrList shape_inference_graph1, shape_inference_graph2;
+ shape_inference_graph1.set_name("_outside_compilation_shape_inference_F1_O1");
+ shape_inference_graph2.set_name("_outside_compilation_shape_inference_F1_O2");
*library_expected.add_function() = FunctionDefHelper::Create(
"F1", {"a_0_arg:float", "b_0_arg:float"}, {"i_0_retval_retval:float"}, {},
{
@@ -1076,8 +1089,7 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O2"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O2"},
+ {"shape_inference_graph", shape_inference_graph2},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O2"}},
{"F"}},
@@ -1088,8 +1100,7 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O1"},
+ {"shape_inference_graph", shape_inference_graph1},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}},
{"D"}},
@@ -1105,8 +1116,9 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
Node* key_constant =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv1 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT, DT_FLOAT}, b2.opts());
+ Node* recv1 = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT, DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Binary(ops::NodeOut(recv1, 0), ops::NodeOut(recv1, 1),
b2.opts()
.WithName("E")
@@ -1114,10 +1126,12 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
Node* send1 = SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
- b2.opts().WithControlInput(e));
+ b2.opts().WithControlInput(e).WithAttr(
+ kXlaHasHostTransferAttrName, true));
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2",
- {DT_FLOAT, DT_FLOAT}, b2.opts());
+ Node* recv2 = RecvAtHost(
+ ops::NodeOut(key_constant, 0), "F1", "O2", {DT_FLOAT, DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* g = Binary(e, ops::NodeOut(recv2, 0),
b2.opts()
.WithName("G")
@@ -1130,7 +1144,8 @@ TEST(EncapsulateSubgraphsTest, OneFunctionTwoOutside) {
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O2"));
Node* send2 =
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {h}, b2.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {h},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* s = Sequencer(b2.opts()
.WithName("F1_sequencer")
@@ -1212,7 +1227,7 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes",
absl::Span({shape_proto_expected})},
{"_outside_compilation_subgraph", "O1"}},
@@ -1235,7 +1250,7 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F2_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes",
absl::Span({shape_proto_expected})},
{"_outside_compilation_subgraph", "O1"}}},
@@ -1251,8 +1266,9 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutside) {
Node* key_constant1 =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv1 = RecvAtHost(ops::NodeOut(key_constant1, 0), "F1", "O1",
- {DT_FLOAT, DT_FLOAT}, b2.opts());
+ Node* recv1 = RecvAtHost(
+ ops::NodeOut(key_constant1, 0), "F1", "O1", {DT_FLOAT, DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Binary(ops::NodeOut(recv1, 0), ops::NodeOut(recv1, 1),
b2.opts()
.WithName("E")
@@ -1260,7 +1276,8 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutside) {
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
Node* send1 = SendFromHost(ops::NodeOut(key_constant1, 0), "F1", "O1", {e},
- b2.opts().WithControlInput(e));
+ b2.opts().WithControlInput(e).WithAttr(
+ kXlaHasHostTransferAttrName, true));
Node* s1 = Sequencer(
b2.opts().WithName("F1_sequencer").WithControlInputs({recv1, send1}),
"F1");
@@ -1272,15 +1289,17 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutside) {
Node* key_constant2 =
KeyPlaceholder("F2", b2.opts().WithName("F2_key_placeholder"));
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant2, 0), "F2", "O1",
- {DT_FLOAT}, b2.opts());
+ Node* recv2 =
+ RecvAtHost(ops::NodeOut(key_constant2, 0), "F2", "O1", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* h = Binary(ops::NodeOut(call1, 1), recv2,
b2.opts()
.WithName("H")
.WithAttr("_encapsulate", "F2")
.WithAttr("_outside", "O1"));
- Node* send2 = SendFromHost(ops::NodeOut(key_constant2, 0), "F2", "O1", {h},
- b2.opts());
+ Node* send2 =
+ SendFromHost(ops::NodeOut(key_constant2, 0), "F2", "O1", {h},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* s2 = Sequencer(
b2.opts().WithName("F2_sequencer").WithControlInputs({recv2, send2}),
@@ -1358,7 +1377,7 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutsideDependencyFromOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes",
absl::Span({shape_proto_expected})},
{"_outside_compilation_subgraph", "O1"}},
@@ -1380,7 +1399,7 @@ TEST(EncapsulateSubgraphsTest, TwoFunctionsTwoOutsideDependencyFromOutside) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F2_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes",
absl::Span({shape_proto_expected})},
{"_outside_compilation_subgraph", "O1"}}},
@@ -1489,7 +1508,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationNoInputs) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes",
absl::Span({shape_proto_expected})},
{"_outside_compilation_subgraph", "O1"}}},
@@ -1574,7 +1593,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationControlInput) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes",
absl::Span({shape_proto_expected})},
{"_outside_compilation_subgraph", "O1"}},
@@ -1657,7 +1676,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationNoOutputs) {
{"Toutputs", absl::Span({})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}}},
},
@@ -1739,7 +1758,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationControlOutput) {
{"Toutputs", absl::Span({})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}}},
},
@@ -1816,17 +1835,21 @@ TEST(EncapsulateSubgraphsTest,
{
GraphDefBuilder shape2(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape2.opts());
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2",
- {DT_FLOAT}, shape2.opts());
+ Node* recv2 =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2", {DT_FLOAT},
+ shape2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* g = Unary(ops::NodeOut(recv2, 0), shape2.opts()
.WithName("G")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O2"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {g}, shape2.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {g},
+ shape2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape2, "F1_O2", &library_expected));
}
+ NameAttrList shape_inference_graph;
+ shape_inference_graph.set_name("_outside_compilation_shape_inference_F1_O2");
*library_expected.add_function() = FunctionDefHelper::Create(
"F1", {"a_0_arg:float", "b_0_arg:float"}, {"h_0_retval_retval:float"}, {},
{
@@ -1843,8 +1866,7 @@ TEST(EncapsulateSubgraphsTest,
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O2"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O2"},
+ {"shape_inference_graph", shape_inference_graph},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O2"}}},
},
@@ -1863,15 +1885,17 @@ TEST(EncapsulateSubgraphsTest,
.WithAttr("_outside", "O1"));
Node* key_constant =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2",
- {DT_FLOAT}, b2.opts());
+ Node* recv =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* g = Unary(recv, b2.opts()
.WithName("G")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O2")
.WithControlInput(e));
Node* send =
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {g}, b2.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O2", {g},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* s1 = Sequencer(
b2.opts().WithName("F1_sequencer").WithControlInputs({recv, send}),
"F1");
@@ -1925,17 +1949,21 @@ TEST(EncapsulateSubgraphsTest,
{
GraphDefBuilder shape1(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape1.opts());
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT}, shape1.opts());
+ Node* recv2 =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT},
+ shape1.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Unary(ops::NodeOut(recv2, 0), shape1.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, shape1.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ shape1.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape1, "F1_O1", &library_expected));
}
+ NameAttrList shape_inference_graph;
+ shape_inference_graph.set_name("_outside_compilation_shape_inference_F1_O1");
*library_expected.add_function() = FunctionDefHelper::Create(
"F1", {"a_0_arg:float", "b_0_arg:float"}, {"h_0_retval_retval:float"}, {},
{
@@ -1952,8 +1980,7 @@ TEST(EncapsulateSubgraphsTest,
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O1"},
+ {"shape_inference_graph", shape_inference_graph},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}}},
},
@@ -1968,14 +1995,16 @@ TEST(EncapsulateSubgraphsTest,
Node* key_constant =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT}, b2.opts());
+ Node* recv =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Unary(recv, b2.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
Node* send =
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, b2.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
/*Node* g =*/Unary(a, b2.opts()
.WithName("G")
.WithAttr("_encapsulate", "F1")
@@ -2039,17 +2068,21 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationClusterDependency) {
{
GraphDefBuilder shape1(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape1.opts());
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT}, shape1.opts());
+ Node* recv2 =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT},
+ shape1.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Unary(ops::NodeOut(recv2, 0), shape1.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, shape1.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ shape1.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape1, "F1_O1", &library_expected));
}
+ NameAttrList shape_inference_graph;
+ shape_inference_graph.set_name("_outside_compilation_shape_inference_F1_O1");
*library_expected.add_function() = FunctionDefHelper::Create(
"F1", {"a_0_arg:float", "b_0_arg:float"}, {"h_0_retval_retval:float"}, {},
{{{"C"}, "UnaryTest", {"a_0_arg"}},
@@ -2063,8 +2096,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationClusterDependency) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O1"},
+ {"shape_inference_graph", shape_inference_graph},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}}},
{{"outside_compilation_O2_host_compute"},
@@ -2074,7 +2106,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationClusterDependency) {
{"Toutputs", absl::Span({})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O2"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O2"}},
{}},
@@ -2085,7 +2117,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationClusterDependency) {
{"Toutputs", absl::Span({})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O3"},
- {"shape_inference_graph", ""},
+ {"shape_inference_graph", NameAttrList()},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O3"}},
{}}},
@@ -2100,23 +2132,27 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationClusterDependency) {
Node* key_constant =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv1 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT}, b2.opts());
+ Node* recv1 =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = Unary(recv1, b2.opts()
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
Node* send =
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, b2.opts());
- Node* recv2 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2",
- {DT_FLOAT}, b2.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
+ Node* recv2 =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O2", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* g = Unary(recv2, b2.opts()
.WithName("G")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O2")
.WithControlInput(e));
- Node* recv3 = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O3",
- {DT_FLOAT}, b2.opts());
+ Node* recv3 =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O3", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
/*Node* i =*/Binary(recv3, e,
b2.opts()
.WithName("I")
@@ -2236,8 +2272,9 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationShapeInference) {
{
GraphDefBuilder shape(GraphDefBuilder::kFailImmediately);
Node* key_constant = KeyPlaceholder("F1", shape.opts());
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT}, shape.opts());
+ Node* recv =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT},
+ shape.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* a = InputShaped(shape.opts().WithName("A"));
Node* c = Unary(a, shape.opts().WithName("C"));
Node* e = BinaryUnknownShape(c, recv,
@@ -2245,11 +2282,14 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationShapeInference) {
.WithName("E")
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
- SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e}, shape.opts());
+ SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
+ shape.opts().WithAttr(kXlaHasHostTransferAttrName, true));
TF_EXPECT_OK(
AddGraphDefToFunctionLibrary(shape, "F1_O1", &library_expected));
}
+ NameAttrList shape_inference_graph;
+ shape_inference_graph.set_name("_outside_compilation_shape_inference_F1_O1");
*library_expected.add_function() = test::function::XTimesTwo();
*library_expected.add_function() = FunctionDefHelper::Create(
"F1", {"b_0_arg:float", "c_0_arg:float"}, {"f_0_retval_retval:float"}, {},
@@ -2267,8 +2307,7 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationShapeInference) {
{"Toutputs", absl::Span({DT_FLOAT})},
{"ancestors", absl::Span({})},
{"key", "host_compute_channel_F1_O1"},
- {"shape_inference_graph",
- "_outside_compilation_shape_inference_F1_O1"},
+ {"shape_inference_graph", shape_inference_graph},
{"shapes", absl::Span({})},
{"_outside_compilation_subgraph", "O1"}},
{"c"}},
@@ -2285,8 +2324,9 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationShapeInference) {
Node* key_constant =
KeyPlaceholder("F1", b2.opts().WithName("F1_key_placeholder"));
- Node* recv = RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1",
- {DT_FLOAT}, b2.opts());
+ Node* recv =
+ RecvAtHost(ops::NodeOut(key_constant, 0), "F1", "O1", {DT_FLOAT},
+ b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
Node* e = BinaryUnknownShape(c, ops::NodeOut(recv, 0),
b2.opts()
.WithName("E")
@@ -2294,7 +2334,8 @@ TEST(EncapsulateSubgraphsTest, OutsideCompilationShapeInference) {
.WithAttr("_encapsulate", "F1")
.WithAttr("_outside", "O1"));
Node* send = SendFromHost(ops::NodeOut(key_constant, 0), "F1", "O1", {e},
- b2.opts().WithControlInput(e));
+ b2.opts().WithControlInput(e).WithAttr(
+ kXlaHasHostTransferAttrName, true));
Node* s = Sequencer(
b2.opts().WithName("F1_sequencer").WithControlInputs({recv, send}),
diff --git a/tensorflow/compiler/jit/encapsulate_util.cc b/tensorflow/compiler/jit/encapsulate_util.cc
index 28ec37b1b9c8a1a306b5e778bac5b6ba01c2c997..1f4b9c90a4ff0b1166cdb7b5942771b350740ef3 100644
--- a/tensorflow/compiler/jit/encapsulate_util.cc
+++ b/tensorflow/compiler/jit/encapsulate_util.cc
@@ -86,7 +86,7 @@ Status ProcessControlEdges(Graph* g, const string& xla_computation_attr_name,
continue;
} else if (src_xla_computation && !dst_xla_computation) {
if (src_outside_compilation) {
- // Case 1d: outside compilation to host computation control edge.
+ // Case 1c: outside compilation to host computation control edge.
edges_to_remove.push_back(e);
TF_RETURN_IF_ERROR(AppendToListAttr(
@@ -94,7 +94,7 @@ Status ProcessControlEdges(Graph* g, const string& xla_computation_attr_name,
}
} else if (!src_xla_computation && dst_xla_computation) {
if (dst_outside_compilation) {
- // Case 1d: host computation control to outside compilation edge.
+ // Case 1c: host computation control to outside compilation edge.
edges_to_remove.push_back(e);
TF_RETURN_IF_ERROR(AppendToListAttr(
@@ -103,40 +103,24 @@ Status ProcessControlEdges(Graph* g, const string& xla_computation_attr_name,
} else { // src_xla_computation && dst_xla_computation
if (*src_xla_computation != *dst_xla_computation) {
if (src_outside_compilation && dst_outside_compilation) {
- // Case 1c: outside compilation to outside compilation control edge.
+ // Case 1b: outside compilation to outside compilation control edge.
edges_to_remove.push_back(e);
TF_RETURN_IF_ERROR(AppendToListAttr(
e->dst(), kXlaControlDependenciesAttrName, e->src()->name()));
} else if (src_outside_compilation && !dst_outside_compilation) {
- // Case 1b: outside compilation to another XLA computaition control
+ // Case 1a: outside compilation to another XLA computaition control
// edge.
TF_RETURN_IF_ERROR(AppendToListAttr(
e->src(), kXlaConnectedToOtherXlaComputationAttrName,
*dst_xla_computation));
} else if (!src_outside_compilation && dst_outside_compilation) {
- // Case 1b: another XLA computaition to outside compilation control
+ // Case 1a: another XLA computaition to outside compilation control
// edge.
TF_RETURN_IF_ERROR(AppendToListAttr(
e->dst(), kXlaConnectedFromOtherXlaComputationAttrName,
*src_xla_computation));
}
- } else { // *src_xla_computation == *dst_xla_computation
- if (src_outside_compilation && dst_outside_compilation) {
- if (*src_outside_compilation != *dst_outside_compilation) {
- // Case 1c: outside compilation to outside compilation control edge.
- edges_to_remove.push_back(e);
-
- TF_RETURN_IF_ERROR(AppendToListAttr(
- e->dst(), kXlaControlDependenciesAttrName, e->src()->name()));
- }
- } else if (src_outside_compilation && !dst_outside_compilation) {
- // Case 1a: outside compilation to its XLA computation control edge.
- ReplaceAttr(e->src(), kXlaConnectedToXlaComputationAttrName, true);
- } else if (!src_outside_compilation && dst_outside_compilation) {
- // Case 1a: XLA computation to outside compilation in it control edge.
- ReplaceAttr(e->dst(), kXlaConnectedFromXlaComputationAttrName, true);
- }
}
}
}
@@ -181,12 +165,6 @@ Status ProcessXlaToXlaDataEdges(Graph* g,
edges.push_back(EdgeInfo{e->dst_input(), e->dst()->id()});
VLOG(4) << "XLA -> XLA edge: " << e->DebugString();
}
- } else { // *src_xla_computation == *dst_xla_computation
- if (src_outside_compilation && dst_outside_compilation &&
- *src_outside_compilation != *dst_outside_compilation) {
- edges.push_back(EdgeInfo{e->dst_input(), e->dst()->id()});
- VLOG(4) << "XLA -> XLA edge: " << e->DebugString();
- }
}
}
@@ -263,7 +241,7 @@ Status ProcessDataEdgeBetweenOutsideCompilationAndHostComputation(
// Remove the edge from host to outside compilation. Add a placeholder as
// outside compilation node input.
- std::map placeholders;
+ std::map, Node*> placeholders;
for (int i = 0; i < edges.size(); i++) {
Node* dst = g->FindNodeId(edges[i].dst_node_id);
const Edge* e;
@@ -275,9 +253,10 @@ Status ProcessDataEdgeBetweenOutsideCompilationAndHostComputation(
// Find or create placeholder node.
string new_name =
edges[i].is_host_to_outside_compilation
- ? absl::StrCat(src->name(), "_host_to_oc_placeholder")
- : absl::StrCat(src->name(), "_oc_to_host_placeholder");
- auto iter = placeholders.find(new_name);
+ ? absl::StrCat(src->name(), "_host_to_oc_placeholder_", src_output)
+ : absl::StrCat(src->name(), "_oc_to_host_placeholder_", src_output);
+ auto placeholder_index = std::make_pair(src->name(), src_output);
+ auto iter = placeholders.find(placeholder_index);
Node* placeholder_node;
if (iter == placeholders.end()) {
NodeDefBuilder placeholder_builder(new_name, "Placeholder");
@@ -310,7 +289,7 @@ Status ProcessDataEdgeBetweenOutsideCompilationAndHostComputation(
Status s;
placeholder_node = g->AddNode(placeholder_def, &s);
TF_RETURN_IF_ERROR(s);
- placeholders[new_name] = placeholder_node;
+ placeholders[placeholder_index] = placeholder_node;
} else {
placeholder_node = iter->second;
}
@@ -594,14 +573,244 @@ Status AddControlDependencies(
return Status::OK();
}
+// Step 1 for `PreprocessEdgesBetweenOutsideCompilations`. See comments of
+// `PreprocessEdgesBetweenOutsideCompilations` for details.
+Status PreprocessControlEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name) {
+ // Gather edges to remove. We should not remove the edge while iterating.
+ std::vector edges_to_remove;
+ for (const Edge* e : g->edges()) {
+ if (!e->IsControlEdge()) {
+ continue;
+ }
+
+ auto src_outside_compilation =
+ GetStringAttr(*e->src(), outside_compilation_attr_name);
+ auto dst_outside_compilation =
+ GetStringAttr(*e->dst(), outside_compilation_attr_name);
+
+ if (src_outside_compilation && dst_outside_compilation) {
+ if (*src_outside_compilation != *dst_outside_compilation) {
+ // Case 1a: outside compilation to outside compilation control edge.
+ edges_to_remove.push_back(e);
+
+ TF_RETURN_IF_ERROR(AppendToListAttr(
+ e->dst(), kXlaControlDependenciesWithinXlaClusterAttrName,
+ e->src()->name()));
+ }
+ } else if (src_outside_compilation && !dst_outside_compilation) {
+ // Case 1b: outside compilation to its XLA computation control edge.
+ ReplaceAttr(e->src(), kXlaConnectedToXlaComputationAttrName, true);
+ } else if (!src_outside_compilation && dst_outside_compilation) {
+ // Case 1b: XLA computation to outside compilation in it control edge.
+ ReplaceAttr(e->dst(), kXlaConnectedFromXlaComputationAttrName, true);
+ }
+ }
+
+ for (auto e : edges_to_remove) {
+ g->RemoveEdge(e);
+ }
+ return Status::OK();
+}
+
+// Step 2 for `PreprocessEdgesBetweenOutsideCompilations`. See comments of
+// `PreprocessEdgesBetweenOutsideCompilations` for details.
+Status PreprocessDataEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name) {
+ // Gather edges between outside compilation and host computation. Notice that
+ // we do not store `Edge*` directly because we remove some nodes while adding
+ // Identity nodes, and those Edge pointers might be invalidated.
+ struct EdgeInfo {
+ int dst_input, dst_node_id;
+ };
+ std::vector edges;
+ for (const Edge* e : g->edges()) {
+ if (e->IsControlEdge()) {
+ continue;
+ }
+
+ auto src_outside_compilation =
+ GetStringAttr(*e->src(), outside_compilation_attr_name);
+ auto dst_outside_compilation =
+ GetStringAttr(*e->dst(), outside_compilation_attr_name);
+
+ if (src_outside_compilation && dst_outside_compilation &&
+ *src_outside_compilation != *dst_outside_compilation) {
+ edges.push_back(EdgeInfo{e->dst_input(), e->dst()->id()});
+ VLOG(4) << "Oc -> oc edge: " << e->DebugString();
+ }
+ }
+
+ // Remove the edge from host to outside compilation. Add a placeholder as
+ // outside compilation node input.
+ std::map, Node*> placeholders;
+ for (int i = 0; i < edges.size(); i++) {
+ Node* dst = g->FindNodeId(edges[i].dst_node_id);
+ const Edge* e;
+ TF_RETURN_IF_ERROR(dst->input_edge(edges[i].dst_input, &e));
+ Node* src = e->src();
+ int src_output = e->src_output(), dst_input = e->dst_input();
+ g->RemoveEdge(e);
+
+ // Find or create placeholder node.
+ string new_name =
+ absl::StrCat(src->name(), "_oc_to_oc_placeholder_", src_output);
+ auto placeholder_index = std::make_pair(src->name(), src_output);
+ auto iter = placeholders.find(placeholder_index);
+ Node* placeholder_node;
+ if (iter == placeholders.end()) {
+ NodeDefBuilder placeholder_builder(new_name, "Placeholder");
+ placeholder_builder.Attr("dtype", src->output_type(src_output));
+ string outside_compilation_attr;
+ TF_RETURN_IF_ERROR(GetNodeAttr(dst->attrs(),
+ outside_compilation_attr_name,
+ &outside_compilation_attr));
+ placeholder_builder.Attr(outside_compilation_attr_name,
+ outside_compilation_attr);
+ placeholder_builder.Attr(kOutsideCompilationOriginalNodeAttrName,
+ src->name());
+ placeholder_builder.Attr(kOutsideCompilationSrcOutputAttrName,
+ src_output);
+ NodeDef placeholder_def;
+ TF_RETURN_IF_ERROR(placeholder_builder.Finalize(&placeholder_def));
+ Status s;
+ placeholder_node = g->AddNode(placeholder_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ placeholders[placeholder_index] = placeholder_node;
+ } else {
+ placeholder_node = iter->second;
+ }
+ g->AddEdge(placeholder_node, 0, dst, dst_input);
+
+ // Replace `e->dst()` because its input node changed.
+ NodeDef new_def = dst->def();
+ *new_def.mutable_input(dst_input) = placeholder_node->name();
+ TF_ASSIGN_OR_RETURN(Node * dst_replace_node, ReplaceNode(g, dst, new_def));
+
+ // Other edge in `edges` might have `e->dst()` as src or dst
+ // node. Before removing `e->dst()`, replace those edges with
+ // corresponding edges for `dst_replace_node`.
+ for (int j = i + 1; j < edges.size(); j++) {
+ if (edges[j].dst_node_id == edges[i].dst_node_id) {
+ edges[j].dst_node_id = dst_replace_node->id();
+ }
+ }
+ }
+ return Status::OK();
+}
+
+// Step 1 for `PostprocessEdgesBetweenOutsideCompilations`. See comments of
+// `PostprocessEdgesBetweenOutsideCompilations` for details.
+Status PostprocessDataEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name) {
+ // Gather all outside compilation to outside compilation nodes.
+ std::vector placeholder_nodes;
+ for (Node* n : g->nodes()) {
+ if (n->type_string() == "Placeholder" &&
+ HasNodeAttr(n->def(), kOutsideCompilationOriginalNodeAttrName)) {
+ placeholder_nodes.push_back(n);
+ }
+ }
+
+ // Remove the placeholder nodes, and reconnect original edge.
+ auto node_name_index = g->BuildNodeNameIndex();
+ for (auto n : placeholder_nodes) {
+ string node_name;
+ int node_src_output;
+ TF_RETURN_IF_ERROR(GetNodeAttr(
+ n->attrs(), kOutsideCompilationOriginalNodeAttrName, &node_name));
+ TF_RETURN_IF_ERROR(GetNodeAttr(
+ n->attrs(), kOutsideCompilationSrcOutputAttrName, &node_src_output));
+ auto iter = node_name_index.find(node_name);
+ if (iter == node_name_index.end()) {
+ return errors::Internal(
+ "Cannot find original node for oc -> host placeholder node ",
+ node_name);
+ }
+
+ // Change all usage node to use the original node instead.
+ Node* original_node = iter->second;
+ std::vector control_edges;
+ std::vector data_edges;
+ for (auto e : n->out_edges()) {
+ if (e->IsControlEdge()) {
+ control_edges.push_back(e);
+ } else {
+ data_edges.push_back({e->dst(), e->src_output(), e->dst_input()});
+ }
+ }
+ for (const Edge* e : control_edges) {
+ g->AddControlEdge(original_node, e->dst());
+ g->RemoveEdge(e);
+ }
+ for (int i = 0; i < data_edges.size(); i++) {
+ Node* dst = data_edges[i].dst;
+ NodeDef new_def = dst->def();
+ int dst_input = data_edges[i].dst_input;
+ *new_def.mutable_input(dst_input) =
+ absl::StrCat(original_node->name(), ":", node_src_output);
+ TF_ASSIGN_OR_RETURN(Node * replace_node, ReplaceNode(g, dst, new_def));
+
+ const Edge* edge_to_replace = nullptr;
+ TF_RETURN_IF_ERROR(replace_node->input_edge(dst_input, &edge_to_replace));
+ g->RemoveEdge(edge_to_replace);
+ g->AddEdge(original_node, node_src_output, replace_node, dst_input);
+
+ // Other edges might have `dst` as dst node. Update those edges with
+ // `replace_node`.
+ for (int j = i + 1; j < data_edges.size(); j++) {
+ if (data_edges[j].dst == dst) {
+ data_edges[j].dst = replace_node;
+ }
+ }
+
+ // Other placeholder node might have `dst` as original node. Update
+ // `node_name_index` with `replace_node`.
+ node_name_index[replace_node->name()] = replace_node;
+ }
+
+ // Remove placeholder node.
+ g->RemoveNode(n);
+ }
+ return Status::OK();
+}
+
+// Step 2 for `PostprocessEdgesBetweenOutsideCompilations`. See comments of
+// `PostprocessEdgesBetweenOutsideCompilations` for details.
+Status PostprocessControlEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name) {
+ auto node_name_index = g->BuildNodeNameIndex();
+
+ // Reconnect outside compilation to outside compilation control edge.
+ for (Node* n : g->nodes()) {
+ std::vector control_deps;
+ Status s =
+ GetNodeAttr(n->attrs(), kXlaControlDependenciesWithinXlaClusterAttrName,
+ &control_deps);
+ if (!s.ok()) {
+ if (s.code() != error::NOT_FOUND) {
+ return s;
+ } else {
+ continue;
+ }
+ } else {
+ n->ClearAttr(kXlaControlDependenciesWithinXlaClusterAttrName);
+ for (const string& control_input : control_deps) {
+ auto iter = node_name_index.find(control_input);
+ if (iter == node_name_index.end()) {
+ return errors::Internal("Cannot find original node for ",
+ control_input);
+ }
+ g->AddControlEdge(iter->second, n);
+ }
+ }
+ }
+ return Status::OK();
+}
} // namespace
const char kXlaInferredShapesAttrName[] = "_xla_inferred_shapes";
-const char kXlaConnectedToXlaComputationAttrName[] =
- "_xla_connected_to_xla_computation";
-const char kXlaConnectedFromXlaComputationAttrName[] =
- "_xla_connected_from_xla_computation";
const char kXlaConnectedToOtherXlaComputationAttrName[] =
"_xla_connected_to_other_xla_computation";
const char kXlaConnectedFromOtherXlaComputationAttrName[] =
@@ -616,6 +825,15 @@ const char kHostToOutsideCompilationOriginalNodeAttrName[] =
"_xla_host_to_oc_node_name";
const char kHostToOutsideCompilationSrcOutputAttrName[] =
"_xla_host_to_oc_src_output";
+const char kXlaConnectedToXlaComputationAttrName[] =
+ "_xla_connected_to_xla_computation";
+const char kXlaConnectedFromXlaComputationAttrName[] =
+ "_xla_connected_from_xla_computation";
+const char kOutsideCompilationOriginalNodeAttrName[] =
+ "_xla_oc_to_oc_node_name";
+const char kOutsideCompilationSrcOutputAttrName[] = "_xla_oc_to_oc_src_output";
+const char kXlaControlDependenciesWithinXlaClusterAttrName[] =
+ "_xla_control_dependencies_within_xla_cluster";
Status PerformStaticShapeInferenceBeforeEncapsulation(
Graph* g, const string& xla_computation_attr_name,
@@ -699,4 +917,39 @@ Status PostprocessForEncapsulation(
return Status::OK();
}
+Status PreprocessEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name) {
+ // Remove edges from source node to outside compilation nodes, and edges
+ // from outside compilation nodes to sink node.
+ std::vector edges_to_remove;
+ for (const Edge* e : g->source_node()->out_edges()) {
+ if (HasNodeAttr(e->dst()->def(), outside_compilation_attr_name)) {
+ edges_to_remove.push_back(e);
+ }
+ }
+ for (const Edge* e : g->sink_node()->in_edges()) {
+ if (HasNodeAttr(e->src()->def(), outside_compilation_attr_name)) {
+ edges_to_remove.push_back(e);
+ }
+ }
+ for (auto e : edges_to_remove) {
+ g->RemoveEdge(e);
+ }
+
+ TF_RETURN_IF_ERROR(PreprocessControlEdgesBetweenOutsideCompilations(
+ g, outside_compilation_attr_name));
+ TF_RETURN_IF_ERROR(PreprocessDataEdgesBetweenOutsideCompilations(
+ g, outside_compilation_attr_name));
+ return Status::OK();
+}
+
+Status PostprocessEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name) {
+ TF_RETURN_IF_ERROR(PostprocessDataEdgesBetweenOutsideCompilations(
+ g, outside_compilation_attr_name));
+ TF_RETURN_IF_ERROR(PostprocessControlEdgesBetweenOutsideCompilations(
+ g, outside_compilation_attr_name));
+ return Status::OK();
+}
+
} // namespace tensorflow
diff --git a/tensorflow/compiler/jit/encapsulate_util.h b/tensorflow/compiler/jit/encapsulate_util.h
index 5e0c4bf6a0cc92d69209595e257989665404db6b..e363bc5754ac395bae262dc67a780a0173efaf5e 100644
--- a/tensorflow/compiler/jit/encapsulate_util.h
+++ b/tensorflow/compiler/jit/encapsulate_util.h
@@ -44,14 +44,6 @@ Status PerformStaticShapeInferenceBeforeEncapsulation(
Graph* g, const string& xla_computation_attr_name,
const string& outside_compilation_attr_name);
-// Attribute indicating that some ops in this node's XLA computation has control
-// dependency on this node. Attribute value will always be "true".
-extern const char kXlaConnectedToXlaComputationAttrName[];
-
-// Attribute indicating that this node has control dependency on some ops in
-// this node's XLA computation. Attribute value will always be "true".
-extern const char kXlaConnectedFromXlaComputationAttrName[];
-
// Attribute indicating that some ops in other XLA computation has control
// dependency on this node. Attribute value will be a list of string (XLA
// computation names).
@@ -81,6 +73,14 @@ extern const char kOutsideCompilationToHostOriginalNodeAttrName[];
// int (src_output for original edge).
extern const char kOutsideCompilationToHostSrcOutputAttrName[];
+// Attribute indicating that some ops in this node's XLA computation has control
+// dependency on this node. Attribute value will always be "true".
+extern const char kXlaConnectedToXlaComputationAttrName[];
+
+// Attribute indicating that this node has control dependency on some ops in
+// this node's XLA computation. Attribute value will always be "true".
+extern const char kXlaConnectedFromXlaComputationAttrName[];
+
// Attribute indicating that this is an Placeholder node added to act as a
// temporary input node for an host node. Attribute value will be string
// (original input node name).
@@ -91,19 +91,31 @@ extern const char kHostToOutsideCompilationOriginalNodeAttrName[];
// for original edge).
extern const char kHostToOutsideCompilationSrcOutputAttrName[];
-// Preprocesses the graph for encapsulation. It will perform the following
-// operations in order:
+// Attribute indicating that this is an Placeholder node added to act as a
+// temporary input node for an outside compilation node. Attribute value will be
+// string (original input node name).
+extern const char kOutsideCompilationOriginalNodeAttrName[];
+
+// Attribute indicating that this is an Placeholder node added to act as a
+// temporary input node for an outside compilation node. Attribute value will be
+// int (src_output for original edge).
+extern const char kOutsideCompilationSrcOutputAttrName[];
+
+// Attribute indicating that this node has control dependencies on some other
+// nodes within the same XLA cluster. Attribute value will be a list of string
+// (node names).
+extern const char kXlaControlDependenciesWithinXlaClusterAttrName[];
+
+// Preprocesses edges between different XLA clusters for encapsulation. It will
+// perform the following operations in order:
//
-// 1a. For control edges between outside compilation and its XLA computation,
-// add attr "kXlaConnected{From, To}XlaComputationAttrName = true" to the
-// outside compilation node.
-// 1b. For control edges between outside compilation and another XLA
+// 1a. For control edges between outside compilation and another XLA
// computation, add attr "kXlaConnected{From, To}OtherXlaComputationAttrName
// = XLA computation node name" to the outside compilation node.
-// 1c. For control edges between different outside compilations, remove the edge
-// and add attr "kXlaControlDependenciesAttrName = src node name" to dst
-// node.
-// 1d. For control edges between outside compilation and host computation,
+// 1b. For control edges between different outside compilations (in different
+// XLA computations), remove the edge and add attr
+// "kXlaControlDependenciesAttrName = src node name" to dst node.
+// 1c. For control edges between outside compilation and host computation,
// remove the edge and add attr "kXlaControlDependenciesAttrName = src node
// name" to dst node.
// 2. For data edges between different XLA computations, if either src or dst
@@ -146,26 +158,53 @@ struct XlaClusterInfo {
const std::map host_compute_core;
};
-// Postprocesses the graph for encapsulation. This function reverts what
-// `PreprocessForEncapsulation` did. It will perform the following operations in
-// order:
+// Postprocesses edges between different XLA clusters for encapsulation. This
+// function reverts what `PreprocessForEncapsulation` did. It will perform the
+// following operations in order:
//
// 1. Remove Placeholder nodes between outside compilation and host computation
// (created in `PreprocessForEncapsulation` step 3).
// 2. Remove Identity nodes created in `PreprocessForEncapsulation` step 2.
-// 3a. Reconnect control edges between different outside compilations (marked by
-// `PreprocessForEncapsulation` step 1c) and control edges between outside
-// compilation and host computation (marked by `PreprocessForEncapsulation`
-// step 1d).
-// 3b. Reconnect control edges between outside compilation and another XLA
-// computation (marked by `PreprocessForEncapsulation` step 1b).
-// Notice that control edges marked by `PreprocessForEncapsulation` step 1a are
-// not handled here. They are handled in `RewriteOutsideCompilationSubgraphFn`.
+// 3a. Reconnect control edges between outside compilation and another XLA
+// computation (marked by `PreprocessForEncapsulation` step 1a).
+// 3b. Reconnect control edges between different outside compilations (marked by
+// `PreprocessForEncapsulation` step 1b).
+// 3c. Reconnect control edges between outside compilation and host computation
+// (marked by `PreprocessForEncapsulation` step 1c).
Status PostprocessForEncapsulation(
Graph* g, const string& xla_computation_attr_name,
const string& outside_compilation_attr_name,
const std::unordered_map& clusters);
+// Preprocesses edges within the same XLA cluster. It will perform the following
+// operations in order:
+//
+// 0. Remove edges from source node to outside compilation nodes, and edges
+// from outside compilation nodes to sink node.
+// 1a. For edges between different outside compilation clusters, remove the edge
+// and add attr "kXlaControlDependenciesWithinXlaClusterAttrName = src node
+// name" to dst node.
+// 1b. For control edges between outside compilation and its XLA computation,
+// add attr "kXlaConnected{From, To}XlaComputationAttrName = true" to the
+// outside compilation node.
+// 2. For data edges between different outside compilations, remove the edge
+// and create a Placeholder node as dst node's input.
+Status PreprocessEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name);
+
+// Postprocesses edges within the same XLA cluster. This function reverts what
+// `PreprocessEdgesBetweenOutsideCompilations` did. It will perform the
+// following operations in order:
+//
+// 1. Remove Placeholder nodes between different outside compilations (created
+// in `PreprocessEdgesBetweenOutsideCompilations` step 2).
+// 2a. Reconnect control edges between different outside compilations (marked by
+// `PreprocessEdgesBetweenOutsideCompilations` step 1a).
+// Notice that control edges marked by
+// `PreprocessEdgesBetweenOutsideCompilations` step 1b are not handled here.
+// They are handled in `RewriteOutsideCompilationSubgraphFn`.
+Status PostprocessEdgesBetweenOutsideCompilations(
+ Graph* g, const string& outside_compilation_attr_name);
} // namespace tensorflow
#endif // TENSORFLOW_COMPILER_JIT_ENCAPSULATE_UTIL_H_
diff --git a/tensorflow/compiler/jit/encapsulate_util_test.cc b/tensorflow/compiler/jit/encapsulate_util_test.cc
index 7255df3112916b7abcc98ff8204efc8c02209b13..3b8b49cb92f3e453883a8e64e12ce3748a5173f6 100644
--- a/tensorflow/compiler/jit/encapsulate_util_test.cc
+++ b/tensorflow/compiler/jit/encapsulate_util_test.cc
@@ -107,28 +107,19 @@ TEST(PreprocessForEncapsulationTest, ControlEdges) {
identity4_node->AddAttr("_xla", "1");
identity4_node->AddAttr("_oc", "0");
identity5_node->AddAttr("_xla", "1");
- // Case 1a: control edges between outside compilation and its XLA computation.
- g.AddControlEdge(add_node, identity0_node);
- g.AddControlEdge(identity0_node, identity1_node);
- // Case 1b: control edges between outside compilation and another XLA
+ // Case 1a: control edges between outside compilation and another XLA
// computation.
g.AddControlEdge(identity0_node, identity3_node);
g.AddControlEdge(identity1_node, identity4_node);
- // Case 1c: control edges between different outside compilations.
+ // Case 1b: control edges between different outside compilations.
g.AddControlEdge(identity0_node, identity4_node);
- // Case 1d: control edges between outside compilation and host computation.
+ // Case 1c: control edges between outside compilation and host computation.
g.AddControlEdge(const0_node, identity0_node);
g.AddControlEdge(identity0_node, identity2_node);
TF_CHECK_OK(PreprocessForEncapsulation(&g, "_xla", "_oc"));
- // Case 1a: add attr "_xla_connected_{from/to}_xla_computation = true" to the
- // outside compilation node.
- EXPECT_TRUE(HasNodeAttr(identity0_node->def(),
- kXlaConnectedFromXlaComputationAttrName));
- EXPECT_TRUE(HasNodeAttr(identity0_node->def(),
- kXlaConnectedToXlaComputationAttrName));
- // Case 1b: add attr "_xla_control_deps_{from/to} = XLA computation node name"
+ // Case 1a: add attr "_xla_control_deps_{from/to} = XLA computation node name"
// to the outside compilation node.
std::vector attr;
TF_CHECK_OK(GetNodeAttr(identity0_node->def(),
@@ -140,13 +131,13 @@ TEST(PreprocessForEncapsulationTest, ControlEdges) {
kXlaConnectedFromOtherXlaComputationAttrName, &attr));
EXPECT_EQ(attr.size(), 1);
EXPECT_EQ(attr[0], "0");
- // Case 1c: add attr "_xla_control_deps = src node name" to dst node.
+ // Case 1b: add attr "_xla_control_deps = src node name" to dst node.
attr.clear();
TF_CHECK_OK(GetNodeAttr(identity4_node->def(),
kXlaControlDependenciesAttrName, &attr));
EXPECT_EQ(attr.size(), 1);
EXPECT_EQ(attr[0], "identity0");
- // Case 1d: add attr "_xla_control_deps = src node name" to dst node.
+ // Case 1c: add attr "_xla_control_deps = src node name" to dst node.
attr.clear();
TF_CHECK_OK(GetNodeAttr(identity0_node->def(),
kXlaControlDependenciesAttrName, &attr));
@@ -162,23 +153,33 @@ TEST(PreprocessForEncapsulationTest, ControlEdges) {
TEST(PreprocessForEncapsulationTest, DataEdges) {
// Build the graph:
// "const_0" and "const_1" in host computation
+ // "identityn0" = ("const_0", "const_1") in host computation 0
// "add0" = "const_0" + "const_1" in XLA computation 0
// "add1" = "add0" + "const_0" in XLA computation 0 & outside compilation 0
// "identity0" = "add1" in XLA computation 0
// "add2" = "add1" + "identity0" in host computation
// "add3" = "add1" + "add2" in XLA computation 1
- // "add4" = "identity0" + "add2" in XLA computation 1 & outside compilation 1
+ // "add4" = "identity0" + "add2" in XLA computation 1 & outside compilation 0
+ // "add5" = "identityn0"[0] + "identityn0"[1] in XLA computation 1 &
+ // outside compilation 0
+ // "identityn1" = ("identityn0"[0], "identityn0"[1]) in XLA computation 1 &
+ // outside compilation 0
// "identity1" = "add4" in XLA computation 1
// "identity2" = "identity1" in host computation
tensorflow::Scope s = tensorflow::Scope::NewRootScope();
Output const_0 = ops::Const(s.WithOpName("const_0"), 1, {});
Output const_1 = ops::Const(s.WithOpName("const_1"), 2, {});
+ auto identityn0 =
+ ops::IdentityN(s.WithOpName("identityn_0"), {const_0, const_1});
Output add0 = ops::Add(s.WithOpName("add0"), const_0, const_1);
Output add1 = ops::Add(s.WithOpName("add1"), add0, const_0);
Output identity0 = ops::Identity(s.WithOpName("identity0"), add1);
Output add2 = ops::Add(s.WithOpName("add2"), add1, identity0);
Output add3 = ops::Add(s.WithOpName("add3"), add1, add2);
Output add4 = ops::Add(s.WithOpName("add4"), identity0, add2);
+ Output add5 = ops::Add(s.WithOpName("add5"), identityn0[0], identityn0[1]);
+ auto identityn1 = ops::IdentityN(s.WithOpName("identityn_1"),
+ {identityn0[0], identityn0[1]});
Output identity1 = ops::Identity(s.WithOpName("identity1"), add4);
Output identity2 = ops::Identity(s.WithOpName("identity2"), add4);
Graph g(OpRegistry::Global());
@@ -189,6 +190,8 @@ TEST(PreprocessForEncapsulationTest, DataEdges) {
Node *add0_node = node_index["add0"], *add1_node = node_index["add1"],
*identity0_node = node_index["identity0"],
*add3_node = node_index["add3"], *add4_node = node_index["add4"],
+ *add5_node = node_index["add5"],
+ *identityn1_node = node_index["identityn_1"],
*identity1_node = node_index["identity1"];
add0_node->AddAttr("_xla", "0");
add1_node->AddAttr("_xla", "0");
@@ -197,6 +200,10 @@ TEST(PreprocessForEncapsulationTest, DataEdges) {
add3_node->AddAttr("_xla", "1");
add4_node->AddAttr("_xla", "1");
add4_node->AddAttr("_oc", "0");
+ add5_node->AddAttr("_xla", "1");
+ add5_node->AddAttr("_oc", "0");
+ identityn1_node->AddAttr("_xla", "1");
+ identityn1_node->AddAttr("_oc", "0");
identity1_node->AddAttr("_xla", "1");
TF_CHECK_OK(PreprocessForEncapsulation(&g, "_xla", "_oc"));
@@ -214,8 +221,9 @@ TEST(PreprocessForEncapsulationTest, DataEdges) {
EXPECT_NE(bridge_identity0_add4, nullptr);
// Step 3: add placeholder for edges between host computation and outside
// compilation.
- EXPECT_EQ(bridge_add1_add3->def().input(0), "add1_oc_to_host_placeholder");
- Node *add1_oc_to_host_placeholder = node_index["add1_oc_to_host_placeholder"];
+ EXPECT_EQ(bridge_add1_add3->def().input(0), "add1_oc_to_host_placeholder_0");
+ Node *add1_oc_to_host_placeholder =
+ node_index["add1_oc_to_host_placeholder_0"];
TF_CHECK_OK(GetNodeAttr(add1_oc_to_host_placeholder->attrs(),
kOutsideCompilationToHostOriginalNodeAttrName, &str));
EXPECT_EQ(str, "add1");
@@ -226,15 +234,34 @@ TEST(PreprocessForEncapsulationTest, DataEdges) {
add4_node = node_index["add4"];
ASSERT_NE(add4_node, nullptr);
EXPECT_EQ(add4_node->def().input(0),
- "bridge_identity0_add4_host_to_oc_placeholder");
+ "bridge_identity0_add4_host_to_oc_placeholder_0");
Node *identity0_host_to_oc_placeholder =
- node_index["bridge_identity0_add4_host_to_oc_placeholder"];
+ node_index["bridge_identity0_add4_host_to_oc_placeholder_0"];
TF_CHECK_OK(GetNodeAttr(identity0_host_to_oc_placeholder->attrs(),
kHostToOutsideCompilationOriginalNodeAttrName, &str));
EXPECT_EQ(str, "bridge_identity0_add4");
TF_CHECK_OK(GetNodeAttr(identity0_host_to_oc_placeholder->attrs(),
kHostToOutsideCompilationSrcOutputAttrName, &i));
EXPECT_EQ(i, 0);
+
+ // Check different placeholder nodes are created for different src_output.
+ Node *placeholder0 = node_index["identityn_0_host_to_oc_placeholder_0"],
+ *placeholder1 = node_index["identityn_0_host_to_oc_placeholder_1"];
+ EXPECT_NE(placeholder0, nullptr);
+ EXPECT_NE(placeholder1, nullptr);
+ // Check we only have 2 placeholder nodes created for "identityn_0".
+ int placeholder_count = 0;
+ for (Node *n : g.nodes()) {
+ if (HasNodeAttr(n->def(), kHostToOutsideCompilationOriginalNodeAttrName)) {
+ string attr;
+ TF_CHECK_OK(GetNodeAttr(
+ n->attrs(), kHostToOutsideCompilationOriginalNodeAttrName, &attr));
+ if (attr == "identityn_0") {
+ ++placeholder_count;
+ }
+ }
+ }
+ EXPECT_EQ(placeholder_count, 2);
}
TEST(PostprocessForEncapsulationTest, ControlEdges) {
diff --git a/tensorflow/compiler/jit/encapsulate_xla_computations_pass.cc b/tensorflow/compiler/jit/encapsulate_xla_computations_pass.cc
index 2ce6fa73fc448ca83fa392aa909cb385453eb8b6..d334100aa4a915a87fb05d371e0e3379a7ee05f2 100644
--- a/tensorflow/compiler/jit/encapsulate_xla_computations_pass.cc
+++ b/tensorflow/compiler/jit/encapsulate_xla_computations_pass.cc
@@ -195,8 +195,11 @@ Status RewriteSubgraph(const std::vector& arg_source_tensors,
e->dst()->attrs().Find(kXlaClusterAttr) == nullptr &&
e->dst()->type_string() != kXlaClusterOutput) {
return errors::InvalidArgument(
- "Undeclared output of XLA computation. A common cause of this error "
- "is variable initializers that depend on the XLA computation. Edge: ",
+ "Undeclared output of XLA computation. Some common causes of this "
+ "error are: 1) variable initializers that depend on the XLA "
+ "computation; 2) gradient computations that depend on the XLA "
+ "computation, which can be mitigated by moving gradient computations "
+ "inside XLA computation. Offending edge: ",
e->src()->name(), ":", e->src_output(), " -> ", e->dst()->name(), ":",
e->dst_input());
}
diff --git a/tensorflow/compiler/jit/extract_outside_compilation_pass.cc b/tensorflow/compiler/jit/extract_outside_compilation_pass.cc
index 8b3587c5087a0651c466f53f3709ba21e75dd273..feac98388469bf5fc6fba0ff305996ea93a6d261 100644
--- a/tensorflow/compiler/jit/extract_outside_compilation_pass.cc
+++ b/tensorflow/compiler/jit/extract_outside_compilation_pass.cc
@@ -20,8 +20,10 @@ limitations under the License.
#include "tensorflow/compiler/jit/encapsulate_subgraphs_pass.h"
#include "tensorflow/compiler/jit/encapsulate_util.h"
#include "tensorflow/compiler/tf2xla/dump_graph.h"
+#include "tensorflow/compiler/tf2xla/side_effect_util.h"
#include "tensorflow/compiler/tf2xla/tf2xla_util.h"
#include "tensorflow/core/common_runtime/function.h"
+#include "tensorflow/core/framework/function.h"
#include "tensorflow/core/framework/graph_to_functiondef.h"
#include "tensorflow/core/framework/node_def_builder.h"
#include "tensorflow/core/framework/node_def_util.h"
@@ -98,9 +100,12 @@ xla::StatusOr BuildRecvAtHostNode(
recv_at_host_builder.Attr("Toutputs", recv_at_host_dtypes);
// The correct device_ordinal will be inserted during replication in a
// subsequent rewrite.
- recv_at_host_builder.Attr("device_ordinal", 0);
+ AttrValue device_ordinal_value;
+ device_ordinal_value.set_placeholder("device_ordinal");
+ recv_at_host_builder.Attr("device_ordinal", device_ordinal_value);
recv_at_host_builder.Attr(
"key", absl::StrCat("host_compute_channel_", oc_cluster_name));
+ recv_at_host_builder.Attr(kXlaHasHostTransferAttrName, true);
recv_at_host_builder.Input(key_placeholder->name(), 0, DT_STRING);
TF_RETURN_IF_ERROR(recv_at_host_builder.Finalize(&recv_at_host_def));
Status s;
@@ -197,9 +202,12 @@ xla::StatusOr BuildSendFromHostNode(
send_from_host_builder.Attr("Tinputs", send_from_host_dtypes);
// The correct device_ordinal will be inserted during replication in a
// subsequent rewrite.
- send_from_host_builder.Attr("device_ordinal", 0);
+ AttrValue device_ordinal_value;
+ device_ordinal_value.set_placeholder("device_ordinal");
+ send_from_host_builder.Attr("device_ordinal", device_ordinal_value);
send_from_host_builder.Attr(
"key", absl::StrCat("host_compute_channel_", oc_cluster_name));
+ send_from_host_builder.Attr(kXlaHasHostTransferAttrName, true);
std::vector inputs(send_from_host_dtypes.size());
for (auto* n : ret_nodes) {
int index;
@@ -357,6 +365,47 @@ Status ReplaceOrRemoveOutsideCompilationCallNode(
return Status::OK();
}
+// Resets "device_ordinal" attr to placeholder value for related nodes
+// (XlaRecvAtHost nodes; XlaSendFromHost nodes; If nodes containing
+// XlaRecvAtHost/XlaSendFromHost).
+Status ResetDeviceOrdinalToPlaceholderValue(Graph* g) {
+ AttrValue device_ordinal_value;
+ device_ordinal_value.set_placeholder("device_ordinal");
+ for (Node* n : g->nodes()) {
+ if (!HasNodeAttr(n->def(), kXlaHasHostTransferAttrName)) {
+ continue;
+ }
+
+ if (n->type_string() == "_XlaRecvAtHost" ||
+ n->type_string() == "_XlaSendFromHost") {
+ n->ClearAttr("device_ordinal");
+ n->AddAttr("device_ordinal", device_ordinal_value);
+ } else if (n->type_string() == "If") {
+ for (const string& attr_name :
+ std::vector{"then_branch", "else_branch"}) {
+ NameAttrList branch_func;
+ TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), attr_name, &branch_func));
+ (*branch_func.mutable_attr())["device_ordinal"] = device_ordinal_value;
+ n->ClearAttr(attr_name);
+ n->AddAttr(attr_name, branch_func);
+ }
+ } else if (n->type_string() == "While") {
+ for (const string& attr_name : std::vector{"cond", "body"}) {
+ NameAttrList branch_func;
+ TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), attr_name, &branch_func));
+ (*branch_func.mutable_attr())["device_ordinal"] = device_ordinal_value;
+ n->ClearAttr(attr_name);
+ n->AddAttr(attr_name, branch_func);
+ }
+ } else {
+ return errors::Internal("Unknown node marked with ",
+ kXlaHasHostTransferAttrName, ": ",
+ n->DebugString());
+ }
+ }
+ return Status::OK();
+}
+
// For an XLA computation, builds host side graph given all outside compilation
// graphs inside it. The host side graph contains:
// 1) a "sequencer" node (we will add control edge between XlaRecvAtHost and
@@ -366,10 +415,10 @@ Status ReplaceOrRemoveOutsideCompilationCallNode(
// replace this node with compilation result node.
// 3) all outside compilation graphs.
Status ConstructHostGraph(
- const string& xla_cluster_name,
+ const string& xla_cluster_name, const string& outside_compilation_attr_name,
const std::vector& outside_compilation_host_graphs,
- FunctionLibraryDefinition* fld, std::unique_ptr* host_graph) {
- host_graph->reset(new Graph(fld));
+ FunctionLibraryDefinition* fld, const string& host_graph_func_name) {
+ Graph host_graph(fld);
// Create sequencer node in host graph.
NodeDefBuilder sequencer_builder(absl::StrCat(xla_cluster_name, "_sequencer"),
@@ -378,24 +427,34 @@ Status ConstructHostGraph(
NodeDef sequencer_def;
TF_RETURN_IF_ERROR(sequencer_builder.Finalize(&sequencer_def));
Status s;
- Node* sequencer = (*host_graph)->AddNode(sequencer_def, &s);
+ Node* sequencer = host_graph.AddNode(sequencer_def, &s);
TF_RETURN_IF_ERROR(s);
// Create key placeholder in host graph.
TF_ASSIGN_OR_RETURN(
Node * key_placeholder,
- AddHostComputeKeyPlaceholder(xla_cluster_name, host_graph->get()));
+ AddHostComputeKeyPlaceholder(xla_cluster_name, &host_graph));
// For each outside compilation graph, copy them to host graph with the
// following changes:
// a) Use key_placeholder in host graph instead of its own.
- // b) Add control edge from RecvAtHost/SendFromHost to sequencer.
+ // b) Add control edge from host transfer nodes (XlaRecvAtHost,
+ // XlaSendFromHost, If/While nodes containing
+ // XlaRecvAtHost/XlaSendFromHost) to sequencer node.
// c) Clear node_def.device(), so device placer won't get confused.
for (const string& host_func : outside_compilation_host_graphs) {
VLOG(4) << "Expanding host graph " << host_func;
+ // Temporarily use "0" as "device_ordinal". It will be reset to placeholder
+ // value after we expanded all host graphs. We cannot just use placeholder
+ // value here because FunctionDef instantiation does not allow placeholder
+ // value for attributes.
+ AttrValue device_ordinal_attr;
+ device_ordinal_attr.set_i(0);
+ protobuf::Map attrs;
+ attrs["device_ordinal"] = device_ordinal_attr;
FunctionBody* host_fbody = nullptr;
TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
- *fld->Find(host_func), AttrSlice(), fld,
+ *fld->Find(host_func), AttrSlice(&attrs), fld,
[&](const string& op, const OpDef** sig) {
return fld->LookUpOpDef(op, sig);
},
@@ -408,8 +467,8 @@ Status ConstructHostGraph(
FixupSourceAndSinkEdges(host_fbody->graph);
std::map node_map;
- node_map[host_fbody->graph->source_node()] = (*host_graph)->source_node();
- node_map[host_fbody->graph->sink_node()] = (*host_graph)->sink_node();
+ node_map[host_fbody->graph->source_node()] = host_graph.source_node();
+ node_map[host_fbody->graph->sink_node()] = host_graph.sink_node();
Status s;
ReverseDFS(
*host_fbody->graph, /*enter=*/nullptr,
@@ -431,7 +490,7 @@ Status ConstructHostGraph(
NodeDef copy_def = n->def();
// Change c).
copy_def.clear_device();
- copy = (*host_graph)->AddNode(copy_def, &s);
+ copy = host_graph.AddNode(copy_def, &s);
if (!s.ok()) {
return;
}
@@ -446,22 +505,23 @@ Status ConstructHostGraph(
e->src()->DebugString());
return;
}
- (*host_graph)
- ->AddEdge(node_map[e->src()], e->src_output(), copy,
- e->dst_input());
+ host_graph.AddEdge(node_map[e->src()], e->src_output(), copy,
+ e->dst_input());
}
// Change b).
- if (copy->type_string() == "_XlaRecvAtHost" ||
- copy->type_string() == "_XlaSendFromHost") {
- (*host_graph)->AddControlEdge(copy, sequencer);
+ if (HasNodeAttr(copy->def(), kXlaHasHostTransferAttrName)) {
+ host_graph.AddControlEdge(copy, sequencer);
}
},
NodeComparatorID());
+
if (!s.ok()) {
return s;
}
}
+ // Reset "device_ordinal" to placeholder value.
+ TF_RETURN_IF_ERROR(ResetDeviceOrdinalToPlaceholderValue(&host_graph));
// sequencer and key_placeholder might be dead nodes. Prune them if necessary.
// - sequencer should be pruned iff it has no input control edges from
@@ -470,17 +530,30 @@ Status ConstructHostGraph(
// - key_placeholder should be pruned iff there's no RecvAtHost/SendFromHost.
// We don't need to do anything special.
if (!sequencer->in_edges().empty()) {
- (*host_graph)->AddControlEdge(sequencer, (*host_graph)->sink_node());
+ host_graph.AddControlEdge(sequencer, host_graph.sink_node());
}
PruneForReverseReachability(
- host_graph->get(),
- std::unordered_set{(*host_graph)->sink_node()});
+ &host_graph, std::unordered_set{host_graph.sink_node()});
+
+ // Postprocess edges between different outside compilations.
+ TF_RETURN_IF_ERROR(PostprocessEdgesBetweenOutsideCompilations(
+ &host_graph, outside_compilation_attr_name));
if (VLOG_IS_ON(4)) {
dump_graph::DumpGraphToFile(
absl::StrCat("extract_outside_compilation_host_graph_for_",
xla_cluster_name),
- **host_graph, fld);
+ host_graph, fld);
+ }
+
+ FunctionDef host_graph_fdef;
+ TF_RETURN_IF_ERROR(
+ GraphToFunctionDef(host_graph, host_graph_func_name, &host_graph_fdef));
+ if (fld->Find(host_graph_func_name)) {
+ TF_RETURN_IF_ERROR(
+ fld->ReplaceFunction(host_graph_func_name, host_graph_fdef));
+ } else {
+ TF_RETURN_IF_ERROR(fld->AddFunctionDef(host_graph_fdef));
}
return Status::OK();
@@ -488,8 +561,28 @@ Status ConstructHostGraph(
// Expand XLA computation's outside compilation host side graph into main graph.
// Add a control edge between sequencer node and the XLA computation node.
-Status ExpandHostGraphIntoMainGraph(Graph* main_graph, Graph* host_graph,
+Status ExpandHostGraphIntoMainGraph(Graph* main_graph,
+ FunctionLibraryDefinition* fld,
+ const string& host_graph_func_name,
Node* xla_computation_node) {
+ // Temporarily use "0" as "device_ordinal". It will be rewritten with the
+ // correct value in a later pass. We cannot just use placeholder value here
+ // because FunctionDef instantiation does not allow placeholder value for
+ // attributes.
+ AttrValue device_ordinal_attr;
+ device_ordinal_attr.set_i(0);
+ protobuf::Map attrs;
+ attrs["device_ordinal"] = device_ordinal_attr;
+ FunctionBody* fbody = nullptr;
+ TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
+ *fld->Find(host_graph_func_name), AttrSlice(&attrs), fld,
+ [&](const string& op, const OpDef** sig) {
+ return fld->LookUpOpDef(op, sig);
+ },
+ &fbody));
+ std::unique_ptr fbody_deleter(fbody);
+ Graph* host_graph = fbody->graph;
+
// We use ReverseDFS() to copy nodes. Make sure all nodes are reverse
// reachable from sink node so all nodes will be copied.
// TODO(b/77601805): consolidate copy graph functions.
@@ -555,9 +648,14 @@ Status ExpandHostGraphIntoMainGraph(Graph* main_graph, Graph* host_graph,
Status RewriteShapeInferenceGraph(const string& shape_inference_graph_name,
Graph* host_graph,
FunctionLibraryDefinition* fld) {
+ // Use "0" as "device_ordinal". It does not matter for shape inference.
+ AttrValue device_ordinal_attr;
+ device_ordinal_attr.set_i(0);
+ protobuf::Map attrs;
+ attrs["device_ordinal"] = device_ordinal_attr;
FunctionBody* fbody = nullptr;
TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
- *fld->Find(shape_inference_graph_name), AttrSlice(), fld,
+ *fld->Find(shape_inference_graph_name), AttrSlice(&attrs), fld,
[&](const string& op, const OpDef** sig) {
return fld->LookUpOpDef(op, sig);
},
@@ -665,6 +763,567 @@ Status RewriteShapeInferenceGraph(const string& shape_inference_graph_name,
return Status::OK();
}
+// Builds XlaSendToHost node which sends cond predicate to host.
+xla::StatusOr BuildSendIfPredNode(const string& name,
+ const string& host_transfer_key,
+ Node* pred_node, Graph* g) {
+ NodeDefBuilder send_pred_builder(name, "XlaSendToHost");
+ send_pred_builder.Attr("Tinput", DT_BOOL);
+ send_pred_builder.Attr("key", absl::StrCat(host_transfer_key, "_dtoh_0"));
+ send_pred_builder.Attr(kXlaTokenInputNodesAttrName,
+ std::vector{kXlaTokenArgNodeName});
+ send_pred_builder.Input(pred_node->name(), 0, DT_BOOL);
+ NodeDef send_pred_def;
+ TF_RETURN_IF_ERROR(send_pred_builder.Finalize(&send_pred_def));
+ Status s;
+ Node* send_pred_node = g->AddNode(send_pred_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ g->AddEdge(pred_node, 0, send_pred_node, 0);
+ return send_pred_node;
+}
+
+// Replaces key placeholder node with an _Arg node.
+Status ReplaceKeyPlaceholderWithArgNode(const string& xla_cluster_name,
+ const string& func_name,
+ FunctionLibraryDefinition* fld) {
+ // Temporarily use "0" as "device_ordinal". It will be reset to placeholder
+ // value after rewriting.
+ AttrValue device_ordinal_attr;
+ device_ordinal_attr.set_i(0);
+ protobuf::Map attrs;
+ attrs["device_ordinal"] = device_ordinal_attr;
+ FunctionBody* fbody = nullptr;
+ TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
+ *fld->Find(func_name), AttrSlice(&attrs), fld,
+ [&](const string& op, const OpDef** sig) {
+ return fld->LookUpOpDef(op, sig);
+ },
+ &fbody));
+ std::unique_ptr fbody_deleter(fbody);
+ Graph* g = fbody->graph;
+
+ // Find or create the key placeholder node.
+ Node* key_placeholder = nullptr;
+ for (Node* n : g->nodes()) {
+ if (IsKeyPlaceholderNode(*n)) {
+ key_placeholder = n;
+ break;
+ }
+ }
+ if (!key_placeholder) {
+ TF_ASSIGN_OR_RETURN(key_placeholder,
+ AddHostComputeKeyPlaceholder(xla_cluster_name, g));
+ }
+
+ // Build the _Arg node, and replace key placeholder node with it.
+ NodeDefBuilder arg_builder("key_arg", FunctionLibraryDefinition::kArgOp);
+ arg_builder.Attr("T", DT_STRING);
+ arg_builder.Attr("index", 0);
+ NodeDef arg_def;
+ TF_RETURN_IF_ERROR(arg_builder.Finalize(&arg_def));
+ TF_RETURN_IF_ERROR(ReplaceNode(g, key_placeholder, arg_def).status());
+
+ // Reset "device_ordinal" to placeholder value.
+ TF_RETURN_IF_ERROR(ResetDeviceOrdinalToPlaceholderValue(g));
+
+ FunctionDef replace_fdef;
+ TF_RETURN_IF_ERROR(GraphToFunctionDef(*g, func_name, &replace_fdef));
+ TF_RETURN_IF_ERROR(fld->ReplaceFunction(func_name, replace_fdef));
+ return Status::OK();
+}
+
+// Builds host side graph for If node.
+Status BuildHostGraphForIfNode(const string& xla_cluster_attr_name,
+ const string& outside_compilation_attr_name,
+ const string& xla_cluster_name,
+ const string& if_node_name,
+ const string& host_transfer_key,
+ const string& host_graph_func_name,
+ FunctionLibraryDefinition* fld,
+ const string& then_branch_host_func_name,
+ const string& else_branch_host_func_name) {
+ Graph host_graph(fld);
+ string outside_compilation_name = absl::StrCat("oc_if_", if_node_name);
+ AttrValue device_ordinal_value;
+ device_ordinal_value.set_placeholder("device_ordinal");
+
+ // Step 1: add key placeholder node.
+ TF_ASSIGN_OR_RETURN(
+ Node * key_placeholder,
+ AddHostComputeKeyPlaceholder(xla_cluster_name, &host_graph));
+
+ // Step 2: build XlaRecvAtHost node to recv predicate.
+ NodeDefBuilder recv_pred_builder(
+ absl::StrCat("recv_oc_if_pred_", if_node_name), "_XlaRecvAtHost");
+ recv_pred_builder.Attr("Toutputs", std::vector{DT_BOOL});
+ recv_pred_builder.Attr("key", host_transfer_key);
+ recv_pred_builder.Attr("device_ordinal", device_ordinal_value);
+ recv_pred_builder.Attr(xla_cluster_attr_name, xla_cluster_name);
+ recv_pred_builder.Attr(outside_compilation_attr_name,
+ outside_compilation_name);
+ recv_pred_builder.Attr(kXlaHasHostTransferAttrName, true);
+ recv_pred_builder.Input(key_placeholder->name(), 0, DT_STRING);
+ NodeDef recv_pred_def;
+ TF_RETURN_IF_ERROR(recv_pred_builder.Finalize(&recv_pred_def));
+ Status s;
+ Node* recv_pred_node = host_graph.AddNode(recv_pred_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ host_graph.AddEdge(key_placeholder, 0, recv_pred_node, 0);
+
+ // Step 3: rewrite `{then, else}_branch_host_func_name`, replace key
+ // placeholder with an _Arg node.
+ TF_RETURN_IF_ERROR(ReplaceKeyPlaceholderWithArgNode(
+ xla_cluster_name, then_branch_host_func_name, fld));
+ TF_RETURN_IF_ERROR(ReplaceKeyPlaceholderWithArgNode(
+ xla_cluster_name, else_branch_host_func_name, fld));
+
+ // Step 4: build If node to choose between `{then, else}_branch_host_graph`.
+ NodeDefBuilder if_builder(absl::StrCat("oc_if_", if_node_name), "If");
+ if_builder.Attr("Tcond", DT_BOOL);
+ if_builder.Attr("Tin", std::vector{DT_STRING});
+ if_builder.Attr("Tout", std::vector{});
+ NameAttrList host_then_branch, host_else_branch;
+ host_then_branch.set_name(then_branch_host_func_name);
+ (*host_then_branch.mutable_attr())["device_ordinal"] = device_ordinal_value;
+ host_else_branch.set_name(else_branch_host_func_name);
+ (*host_else_branch.mutable_attr())["device_ordinal"] = device_ordinal_value;
+ if_builder.Attr("then_branch", host_then_branch);
+ if_builder.Attr("else_branch", host_else_branch);
+ if_builder.Attr(kXlaHasHostTransferAttrName, true);
+ if_builder.Attr(xla_cluster_attr_name, xla_cluster_name);
+ if_builder.Attr(outside_compilation_attr_name, outside_compilation_name);
+ if_builder.Input(recv_pred_node->name(), 0, DT_BOOL);
+ std::vector if_inputs{
+ {key_placeholder->name(), 0, DT_STRING}};
+ if_builder.Input(if_inputs);
+ NodeDef if_def;
+ TF_RETURN_IF_ERROR(if_builder.Finalize(&if_def));
+ Node* if_node = host_graph.AddNode(if_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ host_graph.AddEdge(recv_pred_node, 0, if_node, 0);
+ host_graph.AddEdge(key_placeholder, 0, if_node, 1);
+
+ // Convert `host_graph` to function, and add a "device_ordinal" attr.
+ FunctionDef oc_host_graph_fdef;
+ TF_RETURN_IF_ERROR(GraphToFunctionDef(host_graph, host_graph_func_name,
+ &oc_host_graph_fdef));
+ if (fld->Find(host_graph_func_name)) {
+ TF_RETURN_IF_ERROR(
+ fld->ReplaceFunction(host_graph_func_name, oc_host_graph_fdef));
+ } else {
+ TF_RETURN_IF_ERROR(fld->AddFunctionDef(oc_host_graph_fdef));
+ }
+
+ return Status::OK();
+}
+
+// Rewrites loop cond to add a node which sends loop cond to host.
+Status AddSendLoopPredToLoopCond(FunctionLibraryDefinition* fld,
+ const NameAttrList& loop_cond_func,
+ const string& while_node_name,
+ const string& host_transfer_key) {
+ // Instantiate the loop cond function.
+ FunctionBody* fbody = nullptr;
+ TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
+ *fld->Find(loop_cond_func.name()), AttrSlice(&loop_cond_func.attr()), fld,
+ [&](const string& op, const OpDef** sig) {
+ return fld->LookUpOpDef(op, sig);
+ },
+ &fbody));
+ std::unique_ptr fbody_deleter(fbody);
+ Graph* g = fbody->graph;
+
+ // Find the _Retval node and the loop cond node.
+ Node* ret_node = nullptr;
+ for (Node* n : g->nodes()) {
+ if (n->type_string() == "_Retval") {
+ if (ret_node) {
+ return errors::Internal("Multiple return node for loop cond function ",
+ loop_cond_func.name(), ": ",
+ ret_node->DebugString(), " and ",
+ n->DebugString());
+ } else {
+ ret_node = n;
+ }
+ }
+ }
+ if (!ret_node) {
+ return errors::Internal("No _Retval node for loop cond function ",
+ loop_cond_func.name());
+ }
+ Node* loop_cond;
+ TF_RETURN_IF_ERROR(ret_node->input_node(0, &loop_cond));
+
+ // Build the XlaSendToHost node.
+ NodeDefBuilder send_loop_cond_builder(
+ absl::StrCat("send_oc_while_cond_", while_node_name), "XlaSendToHost");
+ send_loop_cond_builder.Attr("Tinput", DT_BOOL);
+ send_loop_cond_builder.Attr("key",
+ absl::StrCat(host_transfer_key, "_dtoh_0"));
+ send_loop_cond_builder.Attr(kXlaTokenInputNodesAttrName,
+ std::vector{kXlaTokenArgNodeName});
+ send_loop_cond_builder.Input(loop_cond->name(), 0, DT_BOOL);
+ NodeDef send_loop_cond_def;
+ TF_RETURN_IF_ERROR(send_loop_cond_builder.Finalize(&send_loop_cond_def));
+ Status s;
+ Node* send_loop_cond_node = g->AddNode(send_loop_cond_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ g->AddEdge(loop_cond, 0, send_loop_cond_node, 0);
+
+ // Replace original function.
+ FunctionDef replace_fdef;
+ TF_RETURN_IF_ERROR(
+ GraphToFunctionDef(*g, loop_cond_func.name(), &replace_fdef));
+ TF_RETURN_IF_ERROR(fld->ReplaceFunction(loop_cond_func.name(), replace_fdef));
+
+ return Status::OK();
+}
+
+// Rewrites while loop cond function for host.
+Status RewriteHostWhileLoopCond(
+ const string& cond_host_func_name, const string& while_node_name,
+ const string& host_transfer_key, const string& xla_cluster_attr_name,
+ const string& xla_cluster_name, const string& outside_compilation_attr_name,
+ const string& outside_compilation_name, FunctionLibraryDefinition* fld) {
+ // Replace key placeholder node with _Arg node.
+ TF_RETURN_IF_ERROR(ReplaceKeyPlaceholderWithArgNode(
+ xla_cluster_name, cond_host_func_name, fld));
+
+ // Instantiate cond function.
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map attrs;
+ attrs["device_ordinal"] = device_ordinal_temp_value;
+ FunctionBody* cond_fbody = nullptr;
+ TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
+ *fld->Find(cond_host_func_name), AttrSlice(&attrs), fld,
+ [&](const string& op, const OpDef** sig) {
+ return fld->LookUpOpDef(op, sig);
+ },
+ &cond_fbody));
+ std::unique_ptr cond_fbody_deleter(cond_fbody);
+ Graph* cond_graph = cond_fbody->graph;
+ Node* key_arg = nullptr;
+ for (Node* n : cond_graph->nodes()) {
+ if (n->type_string() == "_Arg") {
+ key_arg = n;
+ }
+ }
+ if (!key_arg) {
+ return errors::Internal(
+ "No _Arg node found for host compute key in function ",
+ cond_host_func_name);
+ }
+
+ // Add an XlaRecvAtHost node to use as cond function return value.
+ // We don't need to set kXlaHasHostTransferAttrName for this node, because
+ // it's already added for the "While" node on the host.
+ NodeDefBuilder recv_pred_builder(
+ absl::StrCat("recv_oc_while_cond_", while_node_name), "_XlaRecvAtHost");
+ recv_pred_builder.Attr("Toutputs", std::vector{DT_BOOL});
+ recv_pred_builder.Attr("key", host_transfer_key);
+ AttrValue device_ordinal_value;
+ device_ordinal_value.set_placeholder("device_ordinal");
+ recv_pred_builder.Attr("device_ordinal", device_ordinal_value);
+ recv_pred_builder.Attr(xla_cluster_attr_name, xla_cluster_name);
+ recv_pred_builder.Attr(outside_compilation_attr_name,
+ outside_compilation_name);
+ recv_pred_builder.Input(key_arg->name(), 0, DT_STRING);
+ NodeDef recv_pred_def;
+ TF_RETURN_IF_ERROR(recv_pred_builder.Finalize(&recv_pred_def));
+ Status s;
+ Node* recv_pred_node = cond_graph->AddNode(recv_pred_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ cond_graph->AddEdge(key_arg, 0, recv_pred_node, 0);
+ NodeDefBuilder ret_builder(
+ absl::StrCat("recv_oc_while_cond_ret_", while_node_name), "_Retval");
+ ret_builder.Attr("T", DT_BOOL);
+ ret_builder.Attr("index", 0);
+ ret_builder.Input(recv_pred_node->name(), 0, DT_BOOL);
+ NodeDef ret_def;
+ TF_RETURN_IF_ERROR(ret_builder.Finalize(&ret_def));
+ Node* ret_node = cond_graph->AddNode(ret_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ cond_graph->AddEdge(recv_pred_node, 0, ret_node, 0);
+
+ // Reset device_ordinal to placeholder value.
+ TF_RETURN_IF_ERROR(ResetDeviceOrdinalToPlaceholderValue(cond_graph));
+
+ // Replace original function.
+ FunctionDef cond_replace_fdef;
+ TF_RETURN_IF_ERROR(
+ GraphToFunctionDef(*cond_graph, cond_host_func_name, &cond_replace_fdef));
+ TF_RETURN_IF_ERROR(
+ fld->ReplaceFunction(cond_host_func_name, cond_replace_fdef));
+
+ return Status::OK();
+}
+
+// Rewrites while loop body function for host.
+Status RewriteHostWhileLoopBody(
+ const string& body_host_func_name, const string& while_node_name,
+ const string& host_transfer_key, const string& xla_cluster_attr_name,
+ const string& xla_cluster_name, const string& outside_compilation_attr_name,
+ const string& outside_compilation_name, FunctionLibraryDefinition* fld) {
+ // Replace key placeholder node with _Arg node.
+ TF_RETURN_IF_ERROR(ReplaceKeyPlaceholderWithArgNode(
+ xla_cluster_name, body_host_func_name, fld));
+
+ // Instantiate body function.
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map attrs;
+ attrs["device_ordinal"] = device_ordinal_temp_value;
+ FunctionBody* body_fbody = nullptr;
+ TF_RETURN_IF_ERROR(FunctionDefToBodyHelper(
+ *fld->Find(body_host_func_name), AttrSlice(&attrs), fld,
+ [&](const string& op, const OpDef** sig) {
+ return fld->LookUpOpDef(op, sig);
+ },
+ &body_fbody));
+ std::unique_ptr body_fbody_deleter(body_fbody);
+ Graph* body_graph = body_fbody->graph;
+ Node* key_arg = nullptr;
+ for (Node* n : body_graph->nodes()) {
+ if (n->type_string() == "_Arg") {
+ key_arg = n;
+ }
+ }
+ if (!key_arg) {
+ return errors::Internal(
+ "No _Arg node found for host compute key in function ",
+ body_host_func_name);
+ }
+
+ // Add a _Retval node to loop body.
+ NodeDefBuilder ret_builder(
+ absl::StrCat("recv_oc_while_body_ret_", while_node_name), "_Retval");
+ ret_builder.Attr("T", DT_STRING);
+ ret_builder.Attr("index", 0);
+ ret_builder.Input(key_arg->name(), 0, DT_STRING);
+ NodeDef ret_def;
+ TF_RETURN_IF_ERROR(ret_builder.Finalize(&ret_def));
+ Status s;
+ Node* ret_node = body_graph->AddNode(ret_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ body_graph->AddEdge(key_arg, 0, ret_node, 0);
+
+ // Reset device_ordinal to placeholder value.
+ TF_RETURN_IF_ERROR(ResetDeviceOrdinalToPlaceholderValue(body_graph));
+
+ // Replace original function.
+ FunctionDef body_replace_fdef;
+ TF_RETURN_IF_ERROR(
+ GraphToFunctionDef(*body_graph, body_host_func_name, &body_replace_fdef));
+ TF_RETURN_IF_ERROR(
+ fld->ReplaceFunction(body_host_func_name, body_replace_fdef));
+
+ return Status::OK();
+}
+
+// Builds host side graph for while node.
+Status BuildHostGraphForWhileNode(
+ const string& xla_cluster_attr_name,
+ const string& outside_compilation_attr_name, const string& xla_cluster_name,
+ const string& while_node_name, const string& host_transfer_key,
+ const string& host_graph_func_name, FunctionLibraryDefinition* fld,
+ const string& cond_host_func_name, const string& body_host_func_name) {
+ Graph host_graph(fld);
+ string outside_compilation_name = absl::StrCat("oc_while_", while_node_name);
+
+ // Step 1: add key placeholder node.
+ TF_ASSIGN_OR_RETURN(
+ Node * key_placeholder,
+ AddHostComputeKeyPlaceholder(xla_cluster_name, &host_graph));
+
+ // Step 2: rewrite cond function.
+ TF_RETURN_IF_ERROR(RewriteHostWhileLoopCond(
+ cond_host_func_name, while_node_name, host_transfer_key,
+ xla_cluster_attr_name, xla_cluster_name, outside_compilation_attr_name,
+ outside_compilation_name, fld));
+
+ // Step 3: rewrite body function.
+ TF_RETURN_IF_ERROR(RewriteHostWhileLoopBody(
+ body_host_func_name, while_node_name, host_transfer_key,
+ xla_cluster_attr_name, xla_cluster_name, outside_compilation_attr_name,
+ outside_compilation_name, fld));
+
+ // Step 4: build While node.
+ NodeDefBuilder while_builder(absl::StrCat("oc_while_", while_node_name),
+ "While");
+ while_builder.Attr("T", std::vector{DT_STRING});
+ NameAttrList func;
+ AttrValue device_ordinal_value;
+ device_ordinal_value.set_placeholder("device_ordinal");
+ (*func.mutable_attr())["device_ordinal"] = device_ordinal_value;
+ func.set_name(cond_host_func_name);
+ while_builder.Attr("cond", func);
+ func.set_name(body_host_func_name);
+ while_builder.Attr("body", func);
+ while_builder.Attr(kXlaHasHostTransferAttrName, true);
+ while_builder.Attr(xla_cluster_attr_name, xla_cluster_name);
+ while_builder.Attr(outside_compilation_attr_name, outside_compilation_name);
+ std::vector while_inputs{
+ {key_placeholder->name(), 0, DT_STRING}};
+ while_builder.Input(while_inputs);
+ NodeDef while_def;
+ TF_RETURN_IF_ERROR(while_builder.Finalize(&while_def));
+ Status s;
+ Node* while_node = host_graph.AddNode(while_def, &s);
+ TF_RETURN_IF_ERROR(s);
+ host_graph.AddEdge(key_placeholder, 0, while_node, 0);
+
+ // Convert `host_graph` to function.
+ FunctionDef oc_host_graph_fdef;
+ TF_RETURN_IF_ERROR(GraphToFunctionDef(host_graph, host_graph_func_name,
+ &oc_host_graph_fdef));
+ if (fld->Find(host_graph_func_name)) {
+ TF_RETURN_IF_ERROR(
+ fld->ReplaceFunction(host_graph_func_name, oc_host_graph_fdef));
+ } else {
+ TF_RETURN_IF_ERROR(fld->AddFunctionDef(oc_host_graph_fdef));
+ }
+
+ return Status::OK();
+}
+
+Status ExtractOutsideCompilationForNodesWithAssociatedFunctions(
+ Graph* g, const string& xla_cluster_attr_name,
+ const string& outside_compilation_attr_name, const string& xla_cluster_name,
+ const std::map& host_compute_core,
+ FunctionLibraryDefinition* fld, std::vector* host_graphs,
+ std::vector* shape_inference_graphs,
+ bool* has_outside_compilation) {
+ std::vector if_nodes, while_nodes;
+ for (Node* n : g->nodes()) {
+ if (n->type_string() == "If") {
+ if_nodes.push_back(n);
+ } else if (n->type_string() == "While") {
+ while_nodes.push_back(n);
+ }
+ }
+
+ for (Node* n : if_nodes) {
+ // Instantiate "then_branch" and "else_branch".
+ NameAttrList then_branch, else_branch;
+ TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "then_branch", &then_branch));
+ TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "else_branch", &else_branch));
+
+ // Extract outside compilation for then_branch and else_branch.
+ bool then_branch_has_outside_compilation = false;
+ bool else_branch_has_outside_compilation = false;
+ string then_branch_host_func_name =
+ absl::StrCat("oc_then_branch_host_if_", n->name()),
+ else_branch_host_func_name =
+ absl::StrCat("oc_else_branch_host_if_", n->name());
+ string then_branch_xla_func_name = absl::StrCat(then_branch.name(), "_oc"),
+ else_branch_xla_func_name = absl::StrCat(else_branch.name(), "_oc");
+ TF_RETURN_IF_ERROR(ExtractOutsideCompilationForFunction(
+ xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
+ then_branch, then_branch_xla_func_name, then_branch_host_func_name,
+ host_compute_core, fld, shape_inference_graphs,
+ &then_branch_has_outside_compilation));
+ TF_RETURN_IF_ERROR(ExtractOutsideCompilationForFunction(
+ xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
+ else_branch, else_branch_xla_func_name, else_branch_host_func_name,
+ host_compute_core, fld, shape_inference_graphs,
+ &else_branch_has_outside_compilation));
+
+ // If then/else branch do not have outside compilation, nothing to do.
+ if (!then_branch_has_outside_compilation &&
+ !else_branch_has_outside_compilation) {
+ continue;
+ }
+
+ *has_outside_compilation = true;
+
+ // Change If node to call the new functions.
+ then_branch.set_name(then_branch_xla_func_name);
+ n->ClearAttr("then_branch");
+ n->AddAttr("then_branch", then_branch);
+ else_branch.set_name(else_branch_xla_func_name);
+ n->ClearAttr("else_branch");
+ n->AddAttr("else_branch", else_branch);
+
+ string host_transfer_key = absl::StrCat("oc_if_pred_", n->name());
+
+ // XLA computation: add a SendToHost node to send cond predicate.
+ Node* pred_node;
+ TF_RETURN_IF_ERROR(n->input_node(0, &pred_node));
+ TF_ASSIGN_OR_RETURN(
+ Node * send_pred_node,
+ BuildSendIfPredNode(absl::StrCat("send_oc_if_pred_", n->name()),
+ host_transfer_key, pred_node, g));
+ n->AddAttr(kXlaTokenInputNodesAttrName,
+ std::vector{send_pred_node->name()});
+
+ // Build host side graph for the "If" node.
+ string oc_host_graph_name = absl::StrCat("oc_if_host_graph_", n->name());
+ TF_RETURN_IF_ERROR(BuildHostGraphForIfNode(
+ xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
+ n->name(), host_transfer_key, oc_host_graph_name, fld,
+ then_branch_host_func_name, else_branch_host_func_name));
+ host_graphs->push_back(oc_host_graph_name);
+ }
+
+ for (Node* n : while_nodes) {
+ // Instantiate "cond" and "body".
+ NameAttrList cond, body;
+ TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "cond", &cond));
+ TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "body", &body));
+
+ // Extract outside compilation for cond and body.
+ bool cond_has_outside_compilation = false;
+ bool body_has_outside_compilation = false;
+ string cond_host_func_name = absl::StrCat("oc_cond_host_while_", n->name()),
+ body_host_func_name = absl::StrCat("oc_body_host_while_", n->name());
+ string cond_xla_func_name = absl::StrCat(cond.name(), "_oc"),
+ body_xla_func_name = absl::StrCat(body.name(), "_oc");
+ TF_RETURN_IF_ERROR(ExtractOutsideCompilationForFunction(
+ xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
+ cond, cond_xla_func_name, cond_host_func_name, host_compute_core, fld,
+ shape_inference_graphs, &cond_has_outside_compilation));
+ TF_RETURN_IF_ERROR(ExtractOutsideCompilationForFunction(
+ xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
+ body, body_xla_func_name, body_host_func_name, host_compute_core, fld,
+ shape_inference_graphs, &body_has_outside_compilation));
+
+ // If cond/body do not have outside compilation, nothing to do.
+ if (!cond_has_outside_compilation && !body_has_outside_compilation) {
+ continue;
+ }
+
+ *has_outside_compilation = true;
+
+ // Change While node to call the new functions.
+ cond.set_name(cond_xla_func_name);
+ n->ClearAttr("cond");
+ n->AddAttr("cond", cond);
+ body.set_name(body_xla_func_name);
+ n->ClearAttr("body");
+ n->AddAttr("body", body);
+
+ string host_transfer_key = absl::StrCat("oc_while_pred_", n->name());
+
+ // XLA computation: rewrite cond function to add a SendToHost node to send
+ // loop predicate.
+ TF_RETURN_IF_ERROR(
+ AddSendLoopPredToLoopCond(fld, cond, n->name(), host_transfer_key));
+ n->AddAttr(kXlaTokenInputNodesAttrName,
+ std::vector{kXlaTokenArgNodeName});
+
+ // Build host side graph for the "While" node.
+ string oc_host_graph_name = absl::StrCat("oc_while_host_graph_", n->name());
+ TF_RETURN_IF_ERROR(BuildHostGraphForWhileNode(
+ xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
+ n->name(), host_transfer_key, oc_host_graph_name, fld,
+ cond_host_func_name, body_host_func_name));
+ host_graphs->push_back(oc_host_graph_name);
+ }
+
+ return Status::OK();
+}
+
} // namespace
Status RewriteOutsideCompilationSubgraphFn::operator()(
@@ -751,12 +1410,15 @@ Status RewriteOutsideCompilationSubgraphFn::operator()(
// it with HostCompute node later.
AddNodeAttr("_outside_compilation_subgraph", old_name, node_def);
if (shapes) {
- AddNodeAttr("shape_inference_graph", "", node_def);
+ NameAttrList shape_inference_graph;
+ AddNodeAttr("shape_inference_graph", shape_inference_graph, node_def);
AddNodeAttr("shapes", *shapes, node_def);
} else {
string shape_inference_func_name =
absl::StrCat("_outside_compilation_shape_inference_", new_name);
- AddNodeAttr("shape_inference_graph", shape_inference_func_name, node_def);
+ NameAttrList shape_inference_graph;
+ shape_inference_graph.set_name(shape_inference_func_name);
+ AddNodeAttr("shape_inference_graph", shape_inference_graph, node_def);
AddNodeAttr("shapes", std::vector{}, node_def);
}
AddNodeAttr("ancestors", std::vector{}, node_def);
@@ -771,11 +1433,10 @@ Status ExtractOutsideCompilationForFunction(
const string& xla_cluster_attr_name,
const string& outside_compilation_attr_name, const string& xla_cluster_name,
const NameAttrList& func_name_attrs, const string& new_func_name,
+ const string& host_graph_func_name,
const std::map& host_compute_core,
- FunctionLibraryDefinition* fld, std::unique_ptr* host_graph,
- std::vector* shape_inference_graphs,
+ FunctionLibraryDefinition* fld, std::vector* shape_inference_graphs,
bool* has_outside_compilation) {
- // Early return if function does not have any outside compilation nodes.
const string& func_name = func_name_attrs.name();
const FunctionDef* fdef = fld->Find(func_name);
if (!fdef) {
@@ -788,9 +1449,8 @@ Status ExtractOutsideCompilationForFunction(
break;
}
}
- if (!has_outside_compilation) {
- return Status::OK();
- }
+ // We cannot early return here, because we might have outside compilation in
+ // If/While function body.
// Convert the function to graph.
FunctionBody* fbody = nullptr;
@@ -801,6 +1461,11 @@ Status ExtractOutsideCompilationForFunction(
},
&fbody));
std::unique_ptr fbody_deleter(fbody);
+
+ // Preprocess edges between different outside compilations. They will be
+ // restored in `ConstructHostGraph()`.
+ TF_RETURN_IF_ERROR(PreprocessEdgesBetweenOutsideCompilations(
+ fbody->graph, outside_compilation_attr_name));
if (VLOG_IS_ON(4)) {
dump_graph::DumpGraphToFile(
absl::StrCat("extract_outside_compilation_for_func_before_", func_name),
@@ -826,11 +1491,11 @@ Status ExtractOutsideCompilationForFunction(
// If we could not infer shapes for XlaSendFromHost inputs statically, we
// will set the "shape_inference_graph" attribute. In that case, copy
// outside compilation subgraph as shape inference graph in `fld`.
- string shape_inference_graph;
+ NameAttrList shape_inference_graph;
TF_RETURN_IF_ERROR(GetNodeAttr(n->attrs(), "shape_inference_graph",
&shape_inference_graph));
- if (!shape_inference_graph.empty()) {
- shape_inference_graphs->push_back(shape_inference_graph);
+ if (!shape_inference_graph.name().empty()) {
+ shape_inference_graphs->push_back(shape_inference_graph.name());
const FunctionDef* xla_fdef = fld->Find(n->name());
if (!xla_fdef) {
@@ -838,9 +1503,9 @@ Status ExtractOutsideCompilationForFunction(
}
FunctionDef shape_inference_fdef = *xla_fdef;
shape_inference_fdef.mutable_signature()->set_name(
- shape_inference_graph);
- if (fld->Find(shape_inference_graph)) {
- TF_RETURN_IF_ERROR(fld->ReplaceFunction(shape_inference_graph,
+ shape_inference_graph.name());
+ if (fld->Find(shape_inference_graph.name())) {
+ TF_RETURN_IF_ERROR(fld->ReplaceFunction(shape_inference_graph.name(),
shape_inference_fdef));
} else {
TF_RETURN_IF_ERROR(fld->AddFunctionDef(shape_inference_fdef));
@@ -858,11 +1523,17 @@ Status ExtractOutsideCompilationForFunction(
*graph_out, fld);
}
+ // Handle nodes with associated functions.
+ TF_RETURN_IF_ERROR(ExtractOutsideCompilationForNodesWithAssociatedFunctions(
+ graph_out.get(), xla_cluster_attr_name, outside_compilation_attr_name,
+ xla_cluster_name, host_compute_core, fld,
+ &outside_compilation_host_graphs, shape_inference_graphs,
+ has_outside_compilation));
+
// Construct host graph.
- if (!outside_compilation_host_graphs.empty()) {
- TF_RETURN_IF_ERROR(ConstructHostGraph(
- xla_cluster_name, outside_compilation_host_graphs, fld, host_graph));
- }
+ TF_RETURN_IF_ERROR(ConstructHostGraph(
+ xla_cluster_name, outside_compilation_attr_name,
+ outside_compilation_host_graphs, fld, host_graph_func_name));
// Remove the outside compilation graphs from function library.
for (const string& func : outside_compilation_host_graphs) {
@@ -899,14 +1570,15 @@ Status ExtractOutsideCompilation(
auto const& host_compute_core = iter.second.host_compute_core;
bool has_outside_compilation;
- std::unique_ptr host_graph;
+ string host_graph_func_name = absl::StrCat("oc_host_graph_", n->name());
TF_RETURN_IF_ERROR(ExtractOutsideCompilationForFunction(
xla_cluster_attr_name, outside_compilation_attr_name, xla_cluster_name,
- func_name_attrs, func_name_attrs.name(), host_compute_core, fld,
- &host_graph, &shape_inference_graphs, &has_outside_compilation));
- if (host_graph) {
- TF_RETURN_IF_ERROR(ExpandHostGraphIntoMainGraph(g, host_graph.get(), n));
- }
+ func_name_attrs, func_name_attrs.name(), host_graph_func_name,
+ host_compute_core, fld, &shape_inference_graphs,
+ &has_outside_compilation));
+ TF_RETURN_IF_ERROR(
+ ExpandHostGraphIntoMainGraph(g, fld, host_graph_func_name, n));
+ TF_RETURN_IF_ERROR(fld->RemoveFunction(host_graph_func_name));
}
if (VLOG_IS_ON(4)) {
diff --git a/tensorflow/compiler/jit/extract_outside_compilation_pass.h b/tensorflow/compiler/jit/extract_outside_compilation_pass.h
index 2a4f07cca213d999202024294f5d8f94527059c3..e07e7c5dd0cd42ddd4d643d8b36583c82056bbb2 100644
--- a/tensorflow/compiler/jit/extract_outside_compilation_pass.h
+++ b/tensorflow/compiler/jit/extract_outside_compilation_pass.h
@@ -88,9 +88,10 @@ Status ExtractOutsideCompilationForFunction(
const string& xla_cluster_attr_name,
const string& outside_compilation_attr_name, const string& xla_cluster_name,
const NameAttrList& func_name_attrs, const string& new_func_name,
+ const string& host_graph_func_name,
const std::map& host_compute_core,
- FunctionLibraryDefinition* fld, std::unique_ptr* host_graph,
- std::vector* shape_inference_graphs, bool* has_outside_compilation);
+ FunctionLibraryDefinition* fld, std::vector* shape_inference_graphs,
+ bool* has_outside_compilation);
// Rewrites XLA computation in `clusters` to replace outside compilation nodes
// with XlaHostCompute, and moves those outside compilations into `g`. If shapes
diff --git a/tensorflow/compiler/jit/extract_outside_compilation_pass_test.cc b/tensorflow/compiler/jit/extract_outside_compilation_pass_test.cc
index c5bd64f004ef98853955372680277e04c16bdc9e..0887fbcde9c3c901748ce810e857b1e1e082a8b0 100644
--- a/tensorflow/compiler/jit/extract_outside_compilation_pass_test.cc
+++ b/tensorflow/compiler/jit/extract_outside_compilation_pass_test.cc
@@ -19,8 +19,10 @@ limitations under the License.
#include "tensorflow/cc/framework/scope.h"
#include "tensorflow/cc/ops/array_ops.h"
#include "tensorflow/cc/ops/function_ops.h"
+#include "tensorflow/cc/ops/functional_ops.h"
#include "tensorflow/cc/ops/standard_ops.h"
#include "tensorflow/compiler/jit/encapsulate_util.h"
+#include "tensorflow/compiler/xla/test.h"
#include "tensorflow/core/common_runtime/function.h"
#include "tensorflow/core/framework/common_shape_fns.h"
#include "tensorflow/core/framework/function.h"
@@ -109,10 +111,10 @@ TEST(RewriteOutsideCompilationSubgraphFnTest, Basic) {
}
EXPECT_TRUE(has_control_edge_to_send_from_host);
// Verify step 7: necessary attrs added to call_node_def.
- string shape_inference_graph;
+ NameAttrList shape_inference_graph;
TF_CHECK_OK(GetNodeAttr(AttrSlice(&call_node_def.attr()),
"shape_inference_graph", &shape_inference_graph));
- EXPECT_EQ(shape_inference_graph,
+ EXPECT_EQ(shape_inference_graph.name(),
"_outside_compilation_shape_inference_cluster_0");
}
@@ -249,27 +251,26 @@ TEST(ExtractOutsideCompilationForFunctionTest, Basic) {
protobuf::Map attrs;
std::map host_compute_core = {{"0", 1}, {"1", 0}};
- std::unique_ptr host_graph;
std::vector shape_inference_graphs;
bool has_outside_compilation;
NameAttrList name_attrs;
name_attrs.set_name("cluster");
*name_attrs.mutable_attr() = attrs;
TF_CHECK_OK(ExtractOutsideCompilationForFunction(
- "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten",
- host_compute_core, &fld, &host_graph, &shape_inference_graphs,
+ "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten", "host_graph",
+ host_compute_core, &fld, &shape_inference_graphs,
&has_outside_compilation));
// Get rewritten XLA computation function.
- FunctionBody *fbody = nullptr;
- TF_CHECK_OK(FunctionDefToBodyHelper(*fld.Find("cluster_rewritten"),
- AttrSlice(), &fld,
- [&](const string &op, const OpDef **sig) {
- return fld.LookUpOpDef(op, sig);
- },
- &fbody));
- std::unique_ptr fbody_deleter(fbody);
- auto node_name_index = fbody->graph->BuildNodeNameIndex();
+ FunctionBody *xla_fbody = nullptr;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("cluster_rewritten"), AttrSlice(), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &xla_fbody));
+ std::unique_ptr xla_fbody_deleter(xla_fbody);
+ auto node_name_index = xla_fbody->graph->BuildNodeNameIndex();
// Check XlaHostCompute nodes.
Node *host_compute_0 = node_name_index["outside_compilation_0_host_compute"];
@@ -290,23 +291,33 @@ TEST(ExtractOutsideCompilationForFunctionTest, Basic) {
TF_CHECK_OK(GetNodeAttr(host_compute_1->attrs(), "shapes", &shapes));
EXPECT_EQ(shapes.size(), 1);
EXPECT_EQ(shapes[0].dim_size(), 1);
- // Check XlaHostCompute nodes' "shape_inference_graph" attr. "0" should have a
- // non-empty value, and "1" should have an empty value.
- string shape_inference_graph;
+ // Check XlaHostCompute nodes' "shape_inference_graph" attr. Both should have
+ // empty values.
+ NameAttrList shape_inference_graph;
TF_CHECK_OK(GetNodeAttr(host_compute_0->attrs(), "shape_inference_graph",
&shape_inference_graph));
- EXPECT_EQ(shape_inference_graph,
- "_outside_compilation_shape_inference_cluster_0");
+ EXPECT_EQ(shape_inference_graph.name(), "");
TF_CHECK_OK(GetNodeAttr(host_compute_1->attrs(), "shape_inference_graph",
&shape_inference_graph));
- EXPECT_EQ(shape_inference_graph, "");
+ EXPECT_EQ(shape_inference_graph.name(), "");
// Check `shape_inference_graphs`.
- EXPECT_EQ(shape_inference_graphs.size(), 1);
- EXPECT_EQ(shape_inference_graphs[0],
- "_outside_compilation_shape_inference_cluster_0");
-
- // Check `host_graph`: verify we have key placeholder and sequencer.
+ EXPECT_EQ(shape_inference_graphs.size(), 0);
+
+ // Check host graph: verify we have key placeholder and sequencer.
+ FunctionBody *host_fbody = nullptr;
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map host_func_attrs;
+ host_func_attrs["device_ordinal"] = device_ordinal_temp_value;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("host_graph"), AttrSlice(&host_func_attrs), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &host_fbody));
+ std::unique_ptr host_fbody_deleter(host_fbody);
+ Graph *host_graph = host_fbody->graph;
Node *key_placeholder = nullptr, *sequencer = nullptr;
for (Node *n : host_graph->nodes()) {
if (n->type_string() == "Placeholder" &&
@@ -333,8 +344,8 @@ TEST(ExtractOutsideCompilationForFunctionTest, Basic) {
send_recv_nodes.push_back(n);
}
}
- EXPECT_EQ(num_send_from_host, 2);
- EXPECT_EQ(num_recv_at_host, 2);
+ EXPECT_EQ(num_send_from_host, 1);
+ EXPECT_EQ(num_recv_at_host, 1);
for (Node *n : send_recv_nodes) {
Node *input_node;
TF_CHECK_OK(n->input_node(n->num_inputs() - 1, &input_node));
@@ -368,25 +379,37 @@ TEST(ExtractOutsideCompilationForFunctionTest, NoHostGraph) {
protobuf::Map attrs;
std::map host_compute_core = {{"0", 1}, {"1", 0}};
- std::unique_ptr host_graph;
std::vector shape_inference_graphs;
bool has_outside_compilation;
NameAttrList name_attrs;
name_attrs.set_name("cluster");
*name_attrs.mutable_attr() = attrs;
TF_CHECK_OK(ExtractOutsideCompilationForFunction(
- "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten",
- host_compute_core, &fld, &host_graph, &shape_inference_graphs,
+ "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten", "host_graph",
+ host_compute_core, &fld, &shape_inference_graphs,
&has_outside_compilation));
- // Check `host_graph` is empty.
- EXPECT_FALSE(host_graph);
+ // Check host graph is empty.
+ FunctionBody *host_fbody = nullptr;
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map host_func_attrs;
+ host_func_attrs["device_ordinal"] = device_ordinal_temp_value;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("host_graph"), AttrSlice(&host_func_attrs), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &host_fbody));
+ std::unique_ptr host_fbody_deleter(host_fbody);
+ Graph *host_graph = host_fbody->graph;
+ EXPECT_EQ(host_graph->num_nodes(), 2);
}
TEST(ExtractOutsideCompilationForFunctionTest, XlaHostComputeRemoved) {
// Build the XLA computation func.
// "const0"
- // "const1" (outside compilation clsuter "0")
+ // "const1" (outside compilation cluster "0")
FunctionDefLibrary fdl;
{
tensorflow::Scope s = tensorflow::Scope::NewRootScope();
@@ -404,31 +427,43 @@ TEST(ExtractOutsideCompilationForFunctionTest, XlaHostComputeRemoved) {
protobuf::Map attrs;
std::map host_compute_core = {{"0", 1}, {"1", 0}};
- std::unique_ptr host_graph;
std::vector shape_inference_graphs;
bool has_outside_compilation;
NameAttrList name_attrs;
name_attrs.set_name("cluster");
*name_attrs.mutable_attr() = attrs;
TF_CHECK_OK(ExtractOutsideCompilationForFunction(
- "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten",
- host_compute_core, &fld, &host_graph, &shape_inference_graphs,
+ "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten", "host_graph",
+ host_compute_core, &fld, &shape_inference_graphs,
&has_outside_compilation));
// Check rewritten XLA graph: verify that we have no XlaHostCompute.
- FunctionBody *fbody = nullptr;
- TF_CHECK_OK(FunctionDefToBodyHelper(*fld.Find("cluster_rewritten"),
- AttrSlice(), &fld,
- [&](const string &op, const OpDef **sig) {
- return fld.LookUpOpDef(op, sig);
- },
- &fbody));
- std::unique_ptr fbody_deleter(fbody);
- for (Node *n : fbody->graph->nodes()) {
+ FunctionBody *xla_fbody = nullptr;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("cluster_rewritten"), AttrSlice(), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &xla_fbody));
+ std::unique_ptr xla_fbody_deleter(xla_fbody);
+ for (Node *n : xla_fbody->graph->nodes()) {
EXPECT_NE(n->type_string(), "XlaHostCompute");
}
- // Check `host_graph`: verify we have no placeholder, but we have "const1".
+ // Check host graph: verify we have no placeholder, but we have "const1".
+ FunctionBody *host_fbody = nullptr;
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map host_func_attrs;
+ host_func_attrs["device_ordinal"] = device_ordinal_temp_value;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("host_graph"), AttrSlice(&host_func_attrs), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &host_fbody));
+ std::unique_ptr host_fbody_deleter(host_fbody);
+ Graph *host_graph = host_fbody->graph;
int num_key_placeholders = 0;
for (Node *n : host_graph->nodes()) {
if (n->type_string() == "Placeholder" &&
@@ -441,4 +476,301 @@ TEST(ExtractOutsideCompilationForFunctionTest, XlaHostComputeRemoved) {
EXPECT_NE(node_name_index.find("const1"), node_name_index.end());
}
+REGISTER_OP("XlaSendToHost")
+ .Input("input: Tinput")
+ .Attr("Tinput: type")
+ .Attr("key: string")
+ .SetIsStateful();
+
+REGISTER_OP("XlaRecvFromHost")
+ .Output("output: Toutput")
+ .Attr("Toutput: type")
+ .Attr("shape: shape")
+ .Attr("key: string")
+ .SetIsStateful();
+
+TEST(ExtractOutsideCompilationForFunctionTest, OutsideCompilationInIf) {
+ // Build the XLA computation func.
+ // "const0" (bool)
+ // "const1" (int32)
+ // "if0" (pred = "const0", input = "const1", then_branch = "true_fn",
+ // else_branch = "false_fn")
+ FunctionDefLibrary fdl;
+ {
+ tensorflow::Scope s = tensorflow::Scope::NewRootScope();
+ Output arg = ops::_Arg(s.WithOpName("arg"), DT_INT32, 0);
+ Output identity = ops::Identity(s.WithOpName("identity_true_fn"), arg);
+ ops::_Retval retval(s.WithOpName("retval"), identity, 0);
+ std::unique_ptr g(new Graph(OpRegistry::Global()));
+ TF_CHECK_OK(s.ToGraph(g.get()));
+ auto node_name_image = g->BuildNodeNameIndex();
+ node_name_image["identity_true_fn"]->AddAttr("_oc", "0");
+ PartialTensorShape shape({2});
+ node_name_image["identity_true_fn"]->AddAttr(
+ kXlaInferredShapesAttrName, std::vector{shape});
+
+ FunctionDef *true_fn_fdef = fdl.add_function();
+ TF_CHECK_OK(GraphToFunctionDef(*g, "true_fn", true_fn_fdef));
+ }
+ {
+ tensorflow::Scope s = tensorflow::Scope::NewRootScope();
+ Output arg = ops::_Arg(s.WithOpName("arg"), DT_INT32, 0);
+ Output identity = ops::Identity(s.WithOpName("identity_false_fn"), arg);
+ ops::_Retval retval(s.WithOpName("retval"), identity, 0);
+ std::unique_ptr g(new Graph(OpRegistry::Global()));
+ TF_CHECK_OK(s.ToGraph(g.get()));
+ auto node_name_image = g->BuildNodeNameIndex();
+ node_name_image["identity_false_fn"]->AddAttr("_oc", "0");
+ PartialTensorShape shape({2});
+ node_name_image["identity_false_fn"]->AddAttr(
+ kXlaInferredShapesAttrName, std::vector{shape});
+
+ FunctionDef *false_fn_fdef = fdl.add_function();
+ TF_CHECK_OK(GraphToFunctionDef(*g, "false_fn", false_fn_fdef));
+ }
+ {
+ tensorflow::Scope s = tensorflow::Scope::NewRootScope();
+ Output cond = ops::Const(s.WithOpName("const0"), true, {2});
+ Output input = ops::Const(s.WithOpName("const1"), 1, {2});
+ NameAttrList true_fn;
+ true_fn.set_name("true_fn");
+ NameAttrList false_fn;
+ false_fn.set_name("false_fn");
+ auto if_op = ops::If(s.WithOpName("if"), cond,
+ std::initializer_list{cond, input}, {DT_INT32},
+ true_fn, false_fn);
+ ops::_Retval retval(s.WithOpName("retval"), if_op.output[0], 0);
+ std::unique_ptr g(new Graph(OpRegistry::Global()));
+ TF_CHECK_OK(s.ToGraph(g.get()));
+
+ FunctionDef *xla_fdef = fdl.add_function();
+ TF_CHECK_OK(GraphToFunctionDef(*g, "cluster", xla_fdef));
+ }
+ FunctionLibraryDefinition fld(OpRegistry::Global(), fdl);
+
+ protobuf::Map attrs;
+ std::map host_compute_core;
+ std::vector shape_inference_graphs;
+ bool has_outside_compilation;
+ NameAttrList name_attrs;
+ name_attrs.set_name("cluster");
+ *name_attrs.mutable_attr() = attrs;
+ TF_CHECK_OK(ExtractOutsideCompilationForFunction(
+ "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten", "host_graph",
+ host_compute_core, &fld, &shape_inference_graphs,
+ &has_outside_compilation));
+
+ // Check host graph.
+ {
+ FunctionBody *host_fbody = nullptr;
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map host_func_attrs;
+ host_func_attrs["device_ordinal"] = device_ordinal_temp_value;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("host_graph"), AttrSlice(&host_func_attrs), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &host_fbody));
+ std::unique_ptr host_fbody_deleter(host_fbody);
+ Graph *host_graph = host_fbody->graph;
+ auto node_name_index = host_graph->BuildNodeNameIndex();
+
+ // Verify we have XlaRecvAtHost to receive "If" predicate.
+ Node *recv_if_pred_node = node_name_index["recv_oc_if_pred_if"];
+ EXPECT_NE(recv_if_pred_node, nullptr);
+
+ // Verify we have an "If" to choose outside compilation between then_branch
+ // and else_branch, and it has `recv_if_pred_node` as cond input.
+ Node *if_oc_node = node_name_index["oc_if_if"];
+ EXPECT_NE(if_oc_node, nullptr);
+ Node *if_oc_node_cond_input;
+ TF_CHECK_OK(if_oc_node->input_node(0, &if_oc_node_cond_input));
+ EXPECT_EQ(if_oc_node_cond_input, recv_if_pred_node);
+
+ // Check that then_branch outside compilation has node "identity_true_fn".
+ const FunctionDef *true_def = fld.Find("oc_then_branch_host_if_if");
+ EXPECT_NE(true_def, nullptr);
+ bool has_identity_true_fn_node = false;
+ for (const auto &node_def : true_def->node_def()) {
+ if (node_def.name() == "identity_true_fn") {
+ has_identity_true_fn_node = true;
+ break;
+ }
+ }
+ EXPECT_TRUE(has_identity_true_fn_node);
+
+ // Check that else_branch outside compilation has node "identity_false_fn".
+ const FunctionDef *false_def = fld.Find("oc_else_branch_host_if_if");
+ EXPECT_NE(false_def, nullptr);
+ bool has_identity_false_fn_node = false;
+ for (const auto &node_def : false_def->node_def()) {
+ if (node_def.name() == "identity_false_fn") {
+ has_identity_false_fn_node = true;
+ break;
+ }
+ }
+ EXPECT_TRUE(has_identity_false_fn_node);
+ }
+
+ // Check XLA graph.
+ {
+ FunctionBody *xla_fbody = nullptr;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("cluster_rewritten"), AttrSlice(), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &xla_fbody));
+ std::unique_ptr xla_fbody_deleter(xla_fbody);
+ Graph *xla_graph = xla_fbody->graph;
+ auto node_name_index = xla_graph->BuildNodeNameIndex();
+
+ // Check that we have XlaSendToHost to send cond predicate to host.
+ Node *send_if_pred_node = node_name_index["send_oc_if_pred_if"];
+ EXPECT_NE(send_if_pred_node, nullptr);
+
+ // Check that the "If" node now has `send_if_pred_node` as attribute
+ // _xla_token_input_nodes.
+ Node *if_node = node_name_index["if"];
+ EXPECT_NE(if_node, nullptr);
+ std::vector token_inputs;
+ TF_CHECK_OK(
+ GetNodeAttr(if_node->def(), "_xla_token_input_nodes", &token_inputs));
+ EXPECT_THAT(token_inputs, ::testing::ElementsAre("send_oc_if_pred_if"));
+ }
+}
+
+TEST(ExtractOutsideCompilationForFunctionTest, OutsideCompilationInWhile) {
+ // Build the XLA computation func.
+ // "const0" (bool)
+ // "while0" (input = "const0", cond = "cond_fn", body = "body_fn")
+ FunctionDefLibrary fdl;
+ {
+ tensorflow::Scope s = tensorflow::Scope::NewRootScope();
+ Output arg = ops::_Arg(s.WithOpName("arg"), DT_BOOL, 0);
+ Output identity = ops::Identity(s.WithOpName("identity_cond_fn"), arg);
+ ops::_Retval retval(s.WithOpName("retval"), identity, 0);
+ std::unique_ptr g(new Graph(OpRegistry::Global()));
+ TF_CHECK_OK(s.ToGraph(g.get()));
+ auto node_name_image = g->BuildNodeNameIndex();
+ node_name_image["identity_cond_fn"]->AddAttr("_oc", "0");
+ PartialTensorShape shape({2});
+ node_name_image["identity_cond_fn"]->AddAttr(
+ kXlaInferredShapesAttrName, std::vector{shape});
+
+ FunctionDef *cond_fn_fdef = fdl.add_function();
+ TF_CHECK_OK(GraphToFunctionDef(*g, "cond_fn", cond_fn_fdef));
+ }
+ {
+ tensorflow::Scope s = tensorflow::Scope::NewRootScope();
+ Output arg = ops::_Arg(s.WithOpName("arg"), DT_BOOL, 0);
+ Output identity = ops::Identity(s.WithOpName("identity_body_fn"), arg);
+ ops::_Retval retval(s.WithOpName("retval"), identity, 0);
+ std::unique_ptr g(new Graph(OpRegistry::Global()));
+ TF_CHECK_OK(s.ToGraph(g.get()));
+ auto node_name_image = g->BuildNodeNameIndex();
+ node_name_image["identity_body_fn"]->AddAttr("_oc", "0");
+ PartialTensorShape shape({2});
+ node_name_image["identity_body_fn"]->AddAttr(
+ kXlaInferredShapesAttrName, std::vector{shape});
+
+ FunctionDef *body_fn_fdef = fdl.add_function();
+ TF_CHECK_OK(GraphToFunctionDef(*g, "body_fn", body_fn_fdef));
+ }
+ {
+ tensorflow::Scope s = tensorflow::Scope::NewRootScope();
+ Output input = ops::Const(s.WithOpName("const0"), true, {2});
+ NameAttrList cond_fn;
+ cond_fn.set_name("cond_fn");
+ NameAttrList body_fn;
+ body_fn.set_name("body_fn");
+ auto while_op =
+ ops::While(s.WithOpName("while"), std::initializer_list{input},
+ cond_fn, body_fn);
+ ops::_Retval retval(s.WithOpName("retval"), while_op.output[0], 0);
+ std::unique_ptr g(new Graph(OpRegistry::Global()));
+ TF_CHECK_OK(s.ToGraph(g.get()));
+
+ FunctionDef *xla_fdef = fdl.add_function();
+ TF_CHECK_OK(GraphToFunctionDef(*g, "cluster", xla_fdef));
+ }
+ FunctionLibraryDefinition fld(OpRegistry::Global(), fdl);
+
+ protobuf::Map attrs;
+ std::map host_compute_core;
+ std::vector shape_inference_graphs;
+ bool has_outside_compilation;
+ NameAttrList name_attrs;
+ name_attrs.set_name("cluster");
+ *name_attrs.mutable_attr() = attrs;
+ TF_CHECK_OK(ExtractOutsideCompilationForFunction(
+ "_xla", "_oc", "cluster", name_attrs, "cluster_rewritten", "host_graph",
+ host_compute_core, &fld, &shape_inference_graphs,
+ &has_outside_compilation));
+
+ // Check host graph.
+ {
+ FunctionBody *host_fbody = nullptr;
+ AttrValue device_ordinal_temp_value;
+ device_ordinal_temp_value.set_i(0);
+ protobuf::Map host_func_attrs;
+ host_func_attrs["device_ordinal"] = device_ordinal_temp_value;
+ TF_CHECK_OK(FunctionDefToBodyHelper(
+ *fld.Find("host_graph"), AttrSlice(&host_func_attrs), &fld,
+ [&](const string &op, const OpDef **sig) {
+ return fld.LookUpOpDef(op, sig);
+ },
+ &host_fbody));
+ std::unique_ptr host_fbody_deleter(host_fbody);
+ Graph *host_graph = host_fbody->graph;
+ auto node_name_index = host_graph->BuildNodeNameIndex();
+
+ // Verify we have an "While" to execute outside compilation.
+ Node *while_oc_node = node_name_index["oc_while_while"];
+ EXPECT_NE(while_oc_node, nullptr);
+
+ // Check that cond outside compilation has node "identity_cond_fn".
+ const FunctionDef *cond_def = fld.Find("oc_cond_host_while_while");
+ EXPECT_NE(cond_def, nullptr);
+ bool has_identity_cond_fn_node = false;
+ for (const auto &node_def : cond_def->node_def()) {
+ if (node_def.name() == "identity_cond_fn") {
+ has_identity_cond_fn_node = true;
+ break;
+ }
+ }
+ EXPECT_TRUE(has_identity_cond_fn_node);
+
+ // Check that body outside compilation has node "identity_body_fn".
+ const FunctionDef *body_def = fld.Find("oc_body_host_while_while");
+ EXPECT_NE(body_def, nullptr);
+ bool has_identity_body_fn_node = false;
+ for (const auto &node_def : body_def->node_def()) {
+ if (node_def.name() == "identity_body_fn") {
+ has_identity_body_fn_node = true;
+ break;
+ }
+ }
+ EXPECT_TRUE(has_identity_body_fn_node);
+ }
+
+ // Check XLA graph.
+ {
+ // Verify that rewritten cond fn has XlaSendToHost to send loop predicate to
+ // host.
+ const FunctionDef *cond_def = fld.Find("cond_fn_oc");
+ EXPECT_NE(cond_def, nullptr);
+ bool has_send_oc_while_cond_node = false;
+ for (const auto &node_def : cond_def->node_def()) {
+ if (node_def.name() == "send_oc_while_cond_while") {
+ has_send_oc_while_cond_node = true;
+ break;
+ }
+ }
+ EXPECT_TRUE(has_send_oc_while_cond_node);
+ }
+}
+
} // namespace tensorflow
diff --git a/tensorflow/compiler/jit/mark_for_compilation_pass.cc b/tensorflow/compiler/jit/mark_for_compilation_pass.cc
index 60b962d2e8889c287c103078be8a96c2aa32278d..6618e3a58ab7b6374ed775cd6e4e18a6a4975588 100644
--- a/tensorflow/compiler/jit/mark_for_compilation_pass.cc
+++ b/tensorflow/compiler/jit/mark_for_compilation_pass.cc
@@ -72,6 +72,11 @@ struct OperationFilter {
// to resort to a dummy implementation. Currently Assert and CheckNumerics ops
// have dummy XLA implementations.
bool allow_dummy_ops;
+
+ // Whether ops that produce or consume DT_VARIANT values are allowed. We
+ // don't auto-cluster these ops because we don't yet support live-in or
+ // live-out DT_VARIANT values.
+ bool allow_ops_producing_or_consuming_variant;
};
bool IsDummyImplOp(absl::string_view op_name) {
@@ -81,7 +86,13 @@ bool IsDummyImplOp(absl::string_view op_name) {
bool IsStatefulRandomOp(absl::string_view op_name) {
return op_name == "RandomUniform" || op_name == "RandomShuffle" ||
op_name == "RandomUniformInt" || op_name == "RandomStandardNormal" ||
- op_name == "TruncatedNormal";
+ op_name == "TruncatedNormal" || op_name == "Multinomial";
+}
+
+bool OpProducesOrConsumesVariant(const Node& node) {
+ auto is_variant = [](DataType dtype) { return dtype == DT_VARIANT; };
+ return absl::c_any_of(node.input_types(), is_variant) ||
+ absl::c_any_of(node.output_types(), is_variant);
}
bool HasXLAKernel(const Node& node, const DeviceType& jit_device_type) {
@@ -246,6 +257,10 @@ bool IsCompilableCall(const NodeDef& call_def,
if (!op_filter.allow_dummy_ops && IsDummyImplOp(node->type_string())) {
return false;
}
+ if (!op_filter.allow_ops_producing_or_consuming_variant &&
+ OpProducesOrConsumesVariant(*node)) {
+ return false;
+ }
if (!HasXLAKernel(*node, jit_device_type) &&
!IsCompilableCall(node->def(), jit_device_type, op_filter, depth + 1,
lib_runtime)) {
@@ -470,16 +485,15 @@ Status FindCompilationCandidates(
XlaOpRegistry::GetCompilationDevice(device_type.type(), ®istration));
DeviceType jit_device_type(registration->compilation_device_name);
+ bool always_auto_cluster = registration->autoclustering_policy ==
+ XlaOpRegistry::AutoclusteringPolicy::kAlways;
+
OperationFilter op_filter;
op_filter.allow_resource_ops = registration->compile_resource_ops;
- op_filter.allow_stateful_rng_ops =
- (registration->autoclustering_policy ==
- XlaOpRegistry::AutoclusteringPolicy::kAlways);
- op_filter.allow_control_trigger =
- (registration->autoclustering_policy ==
- XlaOpRegistry::AutoclusteringPolicy::kAlways);
- op_filter.allow_dummy_ops = (registration->autoclustering_policy ==
- XlaOpRegistry::AutoclusteringPolicy::kAlways);
+ op_filter.allow_stateful_rng_ops = always_auto_cluster;
+ op_filter.allow_control_trigger = always_auto_cluster;
+ op_filter.allow_dummy_ops = always_auto_cluster;
+ op_filter.allow_ops_producing_or_consuming_variant = always_auto_cluster;
if (!HasXLAKernel(*node, jit_device_type) &&
!IsCompilableCall(node->def(), jit_device_type, op_filter, 0,
@@ -503,6 +517,12 @@ Status FindCompilationCandidates(
<< node->type_string() << ")";
continue;
}
+ if (!op_filter.allow_ops_producing_or_consuming_variant &&
+ OpProducesOrConsumesVariant(*node)) {
+ VLOG(2) << "Rejecting " << node->name()
+ << ": produces or consumes DT_VARIANT";
+ continue;
+ }
if (!op_filter.allow_resource_ops &&
(HasResourceOutput(*node) || IsNonResourceVarResourceOp(*node))) {
@@ -639,6 +659,7 @@ bool IsCompilable(FunctionLibraryRuntime* flr, const NodeDef& ndef) {
op_filter.allow_stateful_rng_ops = true;
op_filter.allow_control_trigger = true;
op_filter.allow_dummy_ops = true;
+ op_filter.allow_ops_producing_or_consuming_variant = true;
return IsCompilableCall(ndef, jit_device_type, op_filter, 0, flr);
}
diff --git a/tensorflow/compiler/jit/mark_for_compilation_pass_test.cc b/tensorflow/compiler/jit/mark_for_compilation_pass_test.cc
index 24d78c077268f83cebbdafddc1a658ae8dc6b8d8..bf2c5508ea9e987e80093f4c2e15d3ff5191126f 100644
--- a/tensorflow/compiler/jit/mark_for_compilation_pass_test.cc
+++ b/tensorflow/compiler/jit/mark_for_compilation_pass_test.cc
@@ -22,6 +22,7 @@ limitations under the License.
#include "tensorflow/cc/ops/array_ops.h"
#include "tensorflow/cc/ops/control_flow_ops_internal.h"
#include "tensorflow/cc/ops/function_ops.h"
+#include "tensorflow/cc/ops/list_ops.h"
#include "tensorflow/cc/ops/resource_variable_ops.h"
#include "tensorflow/cc/ops/sendrecv_ops.h"
#include "tensorflow/cc/ops/standard_ops.h"
@@ -1147,5 +1148,80 @@ TEST(XlaCompilationTest, DontAutoClusterDummyOps) {
EXPECT_EQ(clusters["test/check"], "");
}
+TEST(XlaCompilationTest, DontAutoClusterOpsProducingVariant) {
+ Scope root = Scope::NewRootScope().ExitOnError();
+ Output a = ops::Placeholder(root.WithOpName("test/a"), DT_INT64);
+ Output b = ops::Placeholder(root.WithOpName("test/b"), DT_INT64);
+
+ Output cast_a = ops::Cast(root.WithOpName("test/cast_a"), a, DT_INT32);
+ Output cast_b = ops::Cast(root.WithOpName("test/cast_b"), b, DT_INT32);
+
+ Output tensor_list_reserve = ops::TensorListReserve(
+ root.WithOpName("test/tensor_list_reserve"), cast_a, cast_b, DT_FLOAT);
+
+ std::unique_ptr graph(new Graph(OpRegistry::Global()));
+ TF_ASSERT_OK(root.ToGraph(graph.get()));
+
+ TF_ASSERT_OK(MarkForCompilationPassTestHelper::MarkForCompilation(&graph));
+
+ std::unordered_map clusters = GetClusters(*graph);
+ EXPECT_EQ(clusters["test/tensor_list_reserve"], "");
+}
+
+TEST(XlaCompilationTest, DontAutoClusterOpsConsumingVariant) {
+ Scope root = Scope::NewRootScope().ExitOnError();
+ Output dummy_input =
+ ops::Placeholder(root.WithOpName("test/dummy_input"), DT_INT64);
+ Output variant_input =
+ ops::Placeholder(root.WithOpName("test/variant_input"), DT_VARIANT);
+
+ // Create one more node so that we don't avoid creating a cluster solely
+ // because it would be trivial.
+ Output dummy_cast =
+ ops::Cast(root.WithOpName("test/dummy_cast"), dummy_input, DT_INT32);
+
+ Output tensor_list_element_shape = ops::TensorListElementShape(
+ root.WithOpName("test/tensor_list_element_shape"), variant_input,
+ DT_INT32);
+
+ root.graph()->AddControlEdge(dummy_cast.node(),
+ tensor_list_element_shape.node());
+
+ std::unique_ptr graph(new Graph(OpRegistry::Global()));
+ TF_ASSERT_OK(root.ToGraph(graph.get()));
+
+ TF_ASSERT_OK(MarkForCompilationPassTestHelper::MarkForCompilation(&graph));
+
+ std::unordered_map clusters = GetClusters(*graph);
+ EXPECT_EQ(clusters["test/tensor_list_element_shape"], "");
+}
+
+TEST(XlaCompilationTest, ClusterOpsProducingVariantIfOnXlaDevice) {
+ Scope root = Scope::NewRootScope().ExitOnError();
+ Output a = ops::Placeholder(root.WithOpName("test/a"), DT_INT64);
+ Output b = ops::Placeholder(root.WithOpName("test/b"), DT_INT64);
+
+ Output cast_a = ops::Cast(root.WithOpName("test/cast_a"), a, DT_INT32);
+ Output cast_b = ops::Cast(root.WithOpName("test/cast_b"), b, DT_INT32);
+
+ Output tensor_list_reserve = ops::TensorListReserve(
+ root.WithOpName("test/tensor_list_reserve"), cast_a, cast_b, DT_FLOAT);
+
+ std::unique_ptr graph(new Graph(OpRegistry::Global()));
+ TF_ASSERT_OK(root.ToGraph(graph.get()));
+
+ string xla_cpu_device = "/job:worker/replica:0/task:0/device:XLA_CPU:0";
+ for (Node* n : graph->nodes()) {
+ if (absl::StartsWith(n->name(), /*prefix=*/"test/")) {
+ n->set_assigned_device_name(xla_cpu_device);
+ }
+ }
+
+ TF_ASSERT_OK(MarkForCompilationPassTestHelper::MarkForCompilation(&graph));
+
+ std::unordered_map clusters = GetClusters(*graph);
+ EXPECT_NE(clusters["test/tensor_list_reserve"], "");
+}
+
} // namespace
} // namespace tensorflow
diff --git a/tensorflow/compiler/jit/mark_for_compilation_pass_test_helper.cc b/tensorflow/compiler/jit/mark_for_compilation_pass_test_helper.cc
index d56d0f8ccfcdab40003be38059228cb255921b64..64a3301745790132fe3149bf8fb52d6c45ecc3c1 100644
--- a/tensorflow/compiler/jit/mark_for_compilation_pass_test_helper.cc
+++ b/tensorflow/compiler/jit/mark_for_compilation_pass_test_helper.cc
@@ -34,15 +34,9 @@ namespace tensorflow {
//
// It may be worth refactoring out XlaOpRegistry::RegisterCompilationDevice to
// make this more direct, but probably not worth it solely for this test.
- std::vector devices;
+ std::vector> devices;
TF_RETURN_IF_ERROR(DeviceFactory::AddDevices(*session_options, "", &devices));
- auto delete_devices = gtl::MakeCleanup([&] {
- for (Device* d : devices) {
- delete d;
- }
- });
-
GraphOptimizationPassOptions opt_options;
opt_options.graph = graph;
opt_options.session_options = session_options;
diff --git a/tensorflow/compiler/jit/ops/BUILD b/tensorflow/compiler/jit/ops/BUILD
index f72224545b25bc7100e0b6788e6fbf0a7ca63dad..64409d9334751e0edfce9091a4e5697dd2c712c5 100644
--- a/tensorflow/compiler/jit/ops/BUILD
+++ b/tensorflow/compiler/jit/ops/BUILD
@@ -18,3 +18,9 @@ tf_gen_op_wrapper_py(
out = "xla_ops.py",
deps = ["//tensorflow/compiler/jit/ops:xla_ops"],
)
+
+py_library(
+ name = "xla_ops_grad",
+ srcs = ["xla_ops_grad.py"],
+ deps = ["//tensorflow/python:framework_ops"],
+)
diff --git a/tensorflow/contrib/estimator/python/estimator/dnn.py b/tensorflow/compiler/jit/ops/xla_ops_grad.py
similarity index 62%
rename from tensorflow/contrib/estimator/python/estimator/dnn.py
rename to tensorflow/compiler/jit/ops/xla_ops_grad.py
index 10f657df8de64cc96f0cf04f434a77df66629dca..2d31d8dc714307a48932d061fb1af643940a0872 100644
--- a/tensorflow/contrib/estimator/python/estimator/dnn.py
+++ b/tensorflow/compiler/jit/ops/xla_ops_grad.py
@@ -1,3 +1,4 @@
+"""Gradients for XLA ops."""
# Copyright 2018 The TensorFlow Authors. All Rights Reserved.
#
# Licensed under the Apache License, Version 2.0 (the "License");
@@ -12,21 +13,17 @@
# See the License for the specific language governing permissions and
# limitations under the License.
# ==============================================================================
-"""dnn python module.
-
-Importing from tensorflow.python.estimator is unsupported
-and will soon break!
-"""
-# pylint: disable=unused-import,g-bad-import-order,g-import-not-at-top,wildcard-import
from __future__ import absolute_import
from __future__ import division
from __future__ import print_function
-from tensorflow_estimator.contrib.estimator.python.estimator import dnn
+from tensorflow.python.framework import ops
-# Include attrs that start with single underscore.
-_HAS_DYNAMIC_ATTRIBUTES = True
-dnn.__all__ = [s for s in dir(dnn) if not s.startswith('__')]
-from tensorflow_estimator.contrib.estimator.python.estimator.dnn import *
+@ops.RegisterGradient("XlaClusterOutput")
+def _XlaClusterOutputGrad(_, grad):
+ del grad # unused
+ raise RuntimeError("Gradient computation of graph in xla.compile() is "
+ "prohibited because it can cause performance degradation."
+ "Please move gradient computation inside xla.compile().")
diff --git a/tensorflow/compiler/jit/partially_decluster_pass_test.cc b/tensorflow/compiler/jit/partially_decluster_pass_test.cc
index 1fc5da5071f7aa6f6dd6636aacd60e33c12431a6..38a54cc5efae35ad77b6dc8039c653e920cfc071 100644
--- a/tensorflow/compiler/jit/partially_decluster_pass_test.cc
+++ b/tensorflow/compiler/jit/partially_decluster_pass_test.cc
@@ -386,7 +386,7 @@ TEST(PartiallyDeclusterPassTest, DontDeclusterXlaDeviceOps) {
TF_ASSERT_OK(s.ToGraph(graph.get()));
// This is needed to register the XLA_GPU device.
- std::vector devices;
+ std::vector> devices;
TF_ASSERT_OK(DeviceFactory::AddDevices(
SessionOptions(), "/job:localhost/replica:0/task:0", &devices));
@@ -400,10 +400,6 @@ TEST(PartiallyDeclusterPassTest, DontDeclusterXlaDeviceOps) {
TF_ASSERT_OK(PartiallyDecluster(&graph));
EXPECT_EQ(GetXlaClusterForNode(*n), "cluster_0");
-
- for (Device* d : devices) {
- delete d;
- }
}
TEST(PartiallyDeclusterPassTest, DontDeclusterNonTensorFlowOps) {
diff --git a/tensorflow/compiler/jit/xla_compile_on_demand_op.cc b/tensorflow/compiler/jit/xla_compile_on_demand_op.cc
index 1fe612d43d10030675cf307b109e4dcc89cb2d79..c7e8d61d280a33a83c3386d8ef801018634d31ec 100644
--- a/tensorflow/compiler/jit/xla_compile_on_demand_op.cc
+++ b/tensorflow/compiler/jit/xla_compile_on_demand_op.cc
@@ -142,11 +142,22 @@ Status XlaCompileOnDemandOp::Compile(
TF_RETURN_IF_ERROR(ctx->allocate_temp(
device_tensor.dtype(), device_tensor.shape(), &host_tensor, attrs));
Notification n;
+ Status status;
ctx->op_device_context()->CopyDeviceTensorToCPU(
&device_tensor, "ConstantArgument",
reinterpret_cast(ctx->device()), &host_tensor,
- [&](Status status) { n.Notify(); });
+ [&](Status s) {
+ status = s;
+ n.Notify();
+ });
n.WaitForNotification();
+ if (!status.ok()) {
+ LOG(ERROR) << "Copying tensor of shape "
+ << device_tensor.shape().DebugString() << " from "
+ << ctx->device()->name() << "to CPU failed with "
+ << status.ToString();
+ return status;
+ }
constant_arguments[i] = host_tensor;
}
}
@@ -189,6 +200,7 @@ Status XlaCompileOnDemandOp::Compile(
std::map variable_args = GetVariables(ctx);
std::vector args;
+
TF_RETURN_IF_ERROR(XlaComputationLaunchContext::BuildXlaCompilerArguments(
constant_arguments, variable_args, ctx, &args));
diff --git a/tensorflow/compiler/jit/xla_cpu_device.cc b/tensorflow/compiler/jit/xla_cpu_device.cc
index 9006dd514b166ad8291d2d437305e53de2a093a4..e9770647e7ba96cc1db026d12d5f11f52ce98d35 100644
--- a/tensorflow/compiler/jit/xla_cpu_device.cc
+++ b/tensorflow/compiler/jit/xla_cpu_device.cc
@@ -31,12 +31,12 @@ namespace tensorflow {
class XlaCpuDeviceFactory : public DeviceFactory {
public:
Status CreateDevices(const SessionOptions& options, const string& name_prefix,
- std::vector* devices) override;
+ std::vector>* devices) override;
};
-Status XlaCpuDeviceFactory::CreateDevices(const SessionOptions& session_options,
- const string& name_prefix,
- std::vector* devices) {
+Status XlaCpuDeviceFactory::CreateDevices(
+ const SessionOptions& session_options, const string& name_prefix,
+ std::vector>* devices) {
XlaDeviceFlags* flags = GetXlaDeviceFlags();
bool compile_on_demand = flags->tf_xla_compile_on_demand;
@@ -64,7 +64,18 @@ Status XlaCpuDeviceFactory::CreateDevices(const SessionOptions& session_options,
options.compilation_device_name = DEVICE_CPU_XLA_JIT;
options.use_multiple_streams = false;
auto device = absl::make_unique(session_options, options);
- devices->push_back(device.release());
+
+ // Setting GpuDeviceInfo because eager runtime relies on the device
+ // context in tensorflow_gpu_device_info(). Also,
+ // tensorflow_gpu_device_info() == nullptr is used as an IsCPU test.
+ // We need XlaCpuDevice to be treated not as CPU because it allocates
+ // XlaTensors, not regular Tensors.
+ Status status = device->UseGpuDeviceInfo();
+ if (!status.ok()) {
+ errors::AppendToMessage(&status, "while setting up ", DEVICE_GPU_XLA_JIT);
+ return status;
+ }
+ devices->push_back(std::move(device));
return Status::OK();
}
diff --git a/tensorflow/compiler/jit/xla_device.cc b/tensorflow/compiler/jit/xla_device.cc
index 738bac54cca450857b506681a6d8fe54fbffb86c..4201ff91a89b1bee370e6a43337c51abe3bf974a 100644
--- a/tensorflow/compiler/jit/xla_device.cc
+++ b/tensorflow/compiler/jit/xla_device.cc
@@ -410,6 +410,31 @@ Status XlaDevice::Sync() {
return Status::OK();
}
+void XlaDevice::Sync(const DoneCallback& done) {
+ VLOG(1) << "XlaDevice::Sync (asynchronous)";
+ std::shared_ptr stream;
+ {
+ mutex_lock lock(mu_);
+ stream = stream_;
+ }
+ if (!stream) {
+ done(Status::OK());
+ return;
+ }
+
+ stream->ThenEnqueueOnBackgroundThread(
+ [this, stream, done](se::StreamExecutor*) {
+ tracing::ScopedActivity activity("XlaDevice::Sync::Callback",
+ /*is_expensive=*/true);
+ mutex_lock lock(mu_);
+ while (outstanding_asynchronous_operations_ > 0) {
+ outstanding_asynchronous_operations_cv_.wait(lock);
+ }
+ done(stream->ok() ? Status::OK()
+ : errors::Internal("XlaDevice::Sync() failed."));
+ });
+}
+
Status XlaDevice::MakeTensorFromProto(const TensorProto& tensor_proto,
const AllocatorAttributes alloc_attrs,
Tensor* tensor) {
diff --git a/tensorflow/compiler/jit/xla_device.h b/tensorflow/compiler/jit/xla_device.h
index dc8f49a9c975a25cbae0e980749db5a1daf039e4..c8bb276cdb9673fdcba4cc15a9f33ecd3ae96dbb 100644
--- a/tensorflow/compiler/jit/xla_device.h
+++ b/tensorflow/compiler/jit/xla_device.h
@@ -135,6 +135,7 @@ class XlaDevice : public LocalDevice {
void ComputeAsync(AsyncOpKernel* op_kernel, OpKernelContext* context,
AsyncOpKernel::DoneCallback done) override;
Status Sync() override;
+ void Sync(const DoneCallback& done) override;
Status FillContextMap(const Graph* graph,
DeviceContextMap* device_context_map) override
diff --git a/tensorflow/compiler/jit/xla_device_context.cc b/tensorflow/compiler/jit/xla_device_context.cc
index 6e6532731e64bd42ee56aa719748988f321e0f17..1f3afe8822d441a5ce37617fe18d7767e9bc72e4 100644
--- a/tensorflow/compiler/jit/xla_device_context.cc
+++ b/tensorflow/compiler/jit/xla_device_context.cc
@@ -79,6 +79,13 @@ XlaDeviceContext::XlaDeviceContext(
}
}
+void XlaDeviceContext::CopyTensorInSameDevice(const Tensor* input_tensor,
+ Device* device,
+ Tensor* output_tensor,
+ StatusCallback done) const {
+ done(errors::Unimplemented("XLA->XLA same-device copies not implemented."));
+}
+
void XlaDeviceContext::CopyCPUTensorToDevice(const Tensor* cpu_tensor,
Device* device,
Tensor* device_tensor,
diff --git a/tensorflow/compiler/jit/xla_device_context.h b/tensorflow/compiler/jit/xla_device_context.h
index 1e18df197a2dd65590c5181b4dae4481dca36641..e45db989fac720df6c3458c93a6b8dbb0919f930 100644
--- a/tensorflow/compiler/jit/xla_device_context.h
+++ b/tensorflow/compiler/jit/xla_device_context.h
@@ -62,6 +62,9 @@ class XlaDeviceContext : public DeviceContext {
void CopyDeviceTensorToCPU(const Tensor* device_tensor,
absl::string_view tensor_name, Device* device,
Tensor* cpu_tensor, StatusCallback done) override;
+ void CopyTensorInSameDevice(const Tensor* input_tensor, Device* device,
+ Tensor* output_tensor,
+ StatusCallback done) const override;
xla::LocalClient* client() const { return client_; }
se::Stream* stream() const { return stream_.get(); }
diff --git a/tensorflow/compiler/jit/xla_device_ops.h b/tensorflow/compiler/jit/xla_device_ops.h
index adf0f994b84d9fbf918a5b2478aa7d106853e038..927f983ba9ef23c8509523f42366c0c89c29db9f 100644
--- a/tensorflow/compiler/jit/xla_device_ops.h
+++ b/tensorflow/compiler/jit/xla_device_ops.h
@@ -203,6 +203,8 @@ class XlaAssignVariableOp : public OpKernel {
.HostMemory("output") \
.TypeConstraint("T"), \
ArgOp); \
+ REGISTER_KERNEL_BUILDER( \
+ Name(kArgOp).Device(DEVICE).TypeConstraint("T"), ArgOp); \
\
REGISTER_KERNEL_BUILDER(Name(kRetOp) \
.Device(DEVICE) \
diff --git a/tensorflow/compiler/jit/xla_gpu_device.cc b/tensorflow/compiler/jit/xla_gpu_device.cc
index 441970169581d53e0d8683b98d26712445b170ea..0191315a66f4d331e54fadc9dc6a073a05fd67ef 100644
--- a/tensorflow/compiler/jit/xla_gpu_device.cc
+++ b/tensorflow/compiler/jit/xla_gpu_device.cc
@@ -16,7 +16,10 @@ limitations under the License.
// Registers the XLA_GPU device, which is an XlaDevice instantiation that runs
// operators using XLA via the XLA "CUDA" (GPU) backend.
+#include
#include "absl/memory/memory.h"
+#include "absl/strings/numbers.h"
+#include "absl/strings/str_split.h"
#include "tensorflow/compiler/jit/kernels/xla_ops.h"
#include "tensorflow/compiler/jit/xla_device.h"
#include "tensorflow/compiler/jit/xla_device_ops.h"
@@ -29,12 +32,12 @@ namespace tensorflow {
class XlaGpuDeviceFactory : public DeviceFactory {
public:
Status CreateDevices(const SessionOptions& options, const string& name_prefix,
- std::vector