[TF2XLA] Remove unused and broken call to enable XLA, remove unneeded build macro
Additionally, "common" bits of the XLA compiler are now linked in by default, so we don't even need the build option to check for linkage. PiperOrigin-RevId: 328776929 Change-Id: Ie59796a8c0f0bb400b38397b49853ed23912bdaf
This commit is contained in:
parent
27a2c42be8
commit
fbb7dab9fd
tensorflow
@ -1704,66 +1704,5 @@ TEST_F(CApiFunctionTest, GetFunctionsFromGraph) {
|
||||
TF_DeleteFunction(func1);
|
||||
}
|
||||
|
||||
// This test only works when the TF build includes XLA compiler. One way to set
|
||||
// this up is via bazel build option "--define with_xla_support=true".
|
||||
//
|
||||
// FIXME: generalize the macro name TENSORFLOW_EAGER_USE_XLA to
|
||||
// something like TENSORFLOW_CAPI_USE_XLA.
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
TEST_F(CApiFunctionTest, StatelessIf_XLA) {
|
||||
TF_Function* func;
|
||||
const std::string funcName = "BranchFunc";
|
||||
DefineFunction(funcName.c_str(), &func);
|
||||
TF_GraphCopyFunction(host_graph_, func, nullptr, s_);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
|
||||
TF_Operation* feed = Placeholder(host_graph_, s_);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
|
||||
TF_Operation* true_cond = ScalarConst(true, host_graph_, s_);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
|
||||
TF_OperationDescription* desc =
|
||||
TF_NewOperation(host_graph_, "StatelessIf", "IfNode");
|
||||
TF_AddInput(desc, {true_cond, 0});
|
||||
TF_Output inputs[] = {{feed, 0}};
|
||||
TF_AddInputList(desc, inputs, TF_ARRAYSIZE(inputs));
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
TF_SetAttrType(desc, "Tcond", TF_BOOL);
|
||||
TF_DataType inputType = TF_INT32;
|
||||
TF_SetAttrTypeList(desc, "Tin", &inputType, 1);
|
||||
TF_SetAttrTypeList(desc, "Tout", &inputType, 1);
|
||||
TF_SetAttrFuncName(desc, "then_branch", funcName.data(), funcName.size());
|
||||
TF_SetAttrFuncName(desc, "else_branch", funcName.data(), funcName.size());
|
||||
TF_SetDevice(desc, "/device:XLA_CPU:0");
|
||||
auto op = TF_FinishOperation(desc, s_);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
ASSERT_NE(op, nullptr);
|
||||
|
||||
// Create a session for this graph.
|
||||
CSession csession(host_graph_, s_, /*use_XLA*/ true);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
|
||||
// Run the graph.
|
||||
csession.SetInputs({{feed, Int32Tensor(17)}});
|
||||
csession.SetOutputs({op});
|
||||
csession.Run(s_);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
TF_Tensor* out = csession.output_tensor(0);
|
||||
ASSERT_TRUE(out != nullptr);
|
||||
EXPECT_EQ(TF_INT32, TF_TensorType(out));
|
||||
EXPECT_EQ(0, TF_NumDims(out)); // scalar
|
||||
ASSERT_EQ(sizeof(int32), TF_TensorByteSize(out));
|
||||
int32* output_contents = static_cast<int32*>(TF_TensorData(out));
|
||||
EXPECT_EQ(-17, *output_contents);
|
||||
|
||||
// Clean up
|
||||
csession.CloseAndDelete(s_);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
|
||||
|
||||
TF_DeleteFunction(func);
|
||||
}
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
|
||||
} // namespace
|
||||
} // namespace tensorflow
|
||||
|
@ -6,7 +6,6 @@ load(
|
||||
"tf_copts",
|
||||
"tf_cuda_cc_test",
|
||||
"tf_cuda_library",
|
||||
"tfe_xla_copts",
|
||||
)
|
||||
load(
|
||||
"//tensorflow/core/platform:build_config.bzl",
|
||||
@ -31,7 +30,7 @@ tf_cuda_library(
|
||||
"c_api_unified_experimental.h",
|
||||
],
|
||||
hdrs = ["c_api.h"],
|
||||
copts = tf_copts() + tfe_xla_copts(),
|
||||
copts = tf_copts(),
|
||||
visibility = ["//visibility:public"],
|
||||
deps = select({
|
||||
"//tensorflow:android": [
|
||||
@ -72,13 +71,6 @@ tf_cuda_library(
|
||||
"//tensorflow/core:protos_all_cc",
|
||||
"//tensorflow/core/profiler/lib:traceme",
|
||||
],
|
||||
}) + select({
|
||||
"//tensorflow:with_xla_support": [
|
||||
"//tensorflow/compiler/tf2xla:xla_compiler",
|
||||
"//tensorflow/compiler/jit",
|
||||
"//tensorflow/compiler/jit:xla_device",
|
||||
],
|
||||
"//conditions:default": [],
|
||||
}) + [
|
||||
"@com_google_absl//absl/memory",
|
||||
"//tensorflow/core/common_runtime/eager:eager_operation",
|
||||
@ -228,7 +220,6 @@ tf_cuda_cc_test(
|
||||
"gradients_test.cc",
|
||||
],
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
linkstatic = tf_kernel_tests_linkstatic(),
|
||||
tags = tf_cuda_tests_tags() + ["nomac"],
|
||||
deps = [
|
||||
@ -290,7 +281,6 @@ tf_cuda_cc_test(
|
||||
"mnist_gradients_test.cc",
|
||||
],
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
linkstatic = tf_kernel_tests_linkstatic(),
|
||||
tags = tf_cuda_tests_tags() + [
|
||||
"nomac",
|
||||
@ -553,7 +543,6 @@ tf_cuda_cc_test(
|
||||
"c_api_debug_test.cc",
|
||||
"c_api_test.cc",
|
||||
],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
tags = [
|
||||
"noguitar", # TODO(b/155445984): flaky
|
||||
#"guitar",
|
||||
@ -608,7 +597,6 @@ tf_cuda_cc_test(
|
||||
],
|
||||
# TODO(b/136478427): Figure out how to correctly shut the server down
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
tags = [
|
||||
"no_windows",
|
||||
],
|
||||
@ -641,7 +629,6 @@ tf_cuda_cc_test(
|
||||
],
|
||||
# TODO(b/136478427): Figure out how to correctly shut the server down
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
tags = [
|
||||
"no_windows",
|
||||
],
|
||||
@ -660,7 +647,6 @@ tf_cuda_cc_test(
|
||||
],
|
||||
# TODO(b/136478427): Figure out how to correctly shut the server down
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
tags = [
|
||||
"no_windows",
|
||||
"noasan", # leaks gRPC server instances
|
||||
@ -694,7 +680,6 @@ tf_cuda_cc_test(
|
||||
],
|
||||
# TODO(b/136478427): Figure out how to correctly shut the server down
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
tags = [
|
||||
"no_windows",
|
||||
],
|
||||
@ -729,7 +714,7 @@ tf_cuda_library(
|
||||
"c_api_experimental.h",
|
||||
"c_api_unified_experimental.h",
|
||||
],
|
||||
copts = tf_copts() + tfe_xla_copts(),
|
||||
copts = tf_copts(),
|
||||
visibility = ["//visibility:public"],
|
||||
deps = select({
|
||||
"//tensorflow:android": [
|
||||
@ -801,7 +786,6 @@ tf_cuda_cc_test(
|
||||
"c_api_experimental_test.cc",
|
||||
],
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
linkstatic = tf_kernel_tests_linkstatic(),
|
||||
tags = tf_cuda_tests_tags() + ["nomac"],
|
||||
deps = [
|
||||
@ -825,7 +809,6 @@ tf_cuda_cc_test(
|
||||
"c_api_unified_experimental_test.cc",
|
||||
],
|
||||
args = ["--heap_check=local"],
|
||||
extra_copts = tfe_xla_copts(),
|
||||
linkstatic = tf_kernel_tests_linkstatic(),
|
||||
tags = tf_cuda_tests_tags() + ["nomac"],
|
||||
deps = [
|
||||
|
@ -51,9 +51,6 @@ limitations under the License.
|
||||
#include "tensorflow/core/protobuf/device_filters.pb.h"
|
||||
#include "tensorflow/core/protobuf/error_codes.pb.h"
|
||||
#include "tensorflow/core/util/device_name_utils.h"
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
#include "tensorflow/compiler/tf2xla/xla_op_registry.h"
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
#include "tensorflow/core/common_runtime/copy_tensor.h"
|
||||
#include "tensorflow/core/common_runtime/device_factory.h"
|
||||
#include "tensorflow/core/common_runtime/device_mgr.h"
|
||||
@ -1165,18 +1162,6 @@ const char* TFE_OpGetDevice(const TFE_Op* op, TF_Status* status) {
|
||||
return tensorflow::unwrap(op)->DeviceName().c_str();
|
||||
}
|
||||
|
||||
void TFE_OpSetXLACompilation(TFE_Op* op, unsigned char enable) {
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
tensorflow::Status s = tensorflow::unwrap(op)->SetUseXla(enable);
|
||||
if (!s.ok()) {
|
||||
LOG(ERROR) << "Could not enable XLA compilation for op: " << s;
|
||||
}
|
||||
#else
|
||||
LOG(WARNING) << "This call is a no-op, as the TensorFlow library is not "
|
||||
"built with XLA support.";
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
}
|
||||
|
||||
void TFE_OpAddInput(TFE_Op* op, TFE_TensorHandle* input, TF_Status* status) {
|
||||
status->status = tensorflow::unwrap(op)->AddInput(tensorflow::unwrap(input));
|
||||
}
|
||||
|
@ -264,13 +264,6 @@ TF_CAPI_EXPORT extern void TFE_OpSetDevice(TFE_Op* op, const char* device_name,
|
||||
TF_CAPI_EXPORT extern const char* TFE_OpGetDevice(const TFE_Op* op,
|
||||
TF_Status* status);
|
||||
|
||||
// When 'enable' is set to 1, and if TensorFlow library is built with XLA
|
||||
// support, a subsequent TFE_Execute() call on `op` will run the op via XLA.
|
||||
//
|
||||
// If the library is not built with XLA support, this call would be a no-op.
|
||||
TF_CAPI_EXPORT extern void TFE_OpSetXLACompilation(TFE_Op* op,
|
||||
unsigned char enable);
|
||||
|
||||
TF_CAPI_EXPORT extern void TFE_OpAddInput(TFE_Op* op, TFE_TensorHandle* input,
|
||||
TF_Status* status);
|
||||
|
||||
|
@ -22,9 +22,6 @@ limitations under the License.
|
||||
#include "tensorflow/c/tf_status_internal.h"
|
||||
#include "tensorflow/core/common_runtime/eager/tensor_handle.h"
|
||||
#include "tensorflow/core/platform/status.h"
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
#include "tensorflow/compiler/jit/xla_device.h"
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
|
||||
using tensorflow::string;
|
||||
|
||||
@ -64,87 +61,6 @@ TF_CAPI_EXPORT extern TFE_TensorDebugInfo* TFE_TensorHandleTensorDebugInfo(
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
auto* device = absl::get<tensorflow::Device*>(handle->device());
|
||||
|
||||
// If tensor resides on an XLA device, use XLA device's PaddedShapeFn.
|
||||
auto* xla_device = dynamic_cast<tensorflow::XlaDevice*>(device);
|
||||
if (xla_device != nullptr) {
|
||||
tensorflow::XlaDevice::PaddedShapeFn shape_fn =
|
||||
xla_device->metadata().padded_shape_fn();
|
||||
xla::Shape padded_shape;
|
||||
status->status = shape_fn(*tensor, &padded_shape);
|
||||
if (!status->status.ok()) {
|
||||
return nullptr;
|
||||
}
|
||||
if (VLOG_IS_ON(3)) {
|
||||
std::vector<tensorflow::int64> shape_to_log =
|
||||
TensorShapeAsVector(*handle, &status->status);
|
||||
if (!status->status.ok()) {
|
||||
// Ignore the status here as we are simply logging.
|
||||
status->status = tensorflow::Status::OK();
|
||||
} else {
|
||||
VLOG(3) << "Fully padded shape of ["
|
||||
<< absl::StrJoin(shape_to_log, ", ") << "] is "
|
||||
<< padded_shape.DebugString();
|
||||
}
|
||||
}
|
||||
|
||||
if (padded_shape.IsTuple()) {
|
||||
if (xla::ShapeUtil::TupleElementCount(padded_shape) != 2) {
|
||||
// Currently, the only case of XlaTensor containing a tuple shape is to
|
||||
// represent 64 bit ints, doubles, and complex numbers (we don't support
|
||||
// 64bit complex numbers).
|
||||
status->status = tensorflow::errors::InvalidArgument(
|
||||
"XlaTensors should only contain tuples of size 2. Shape: ",
|
||||
padded_shape.DebugString());
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
// shape0 is not a const& because we will assign it to padded_shape below.
|
||||
// It is illegal to assign a part of a message to itself.
|
||||
xla::Shape shape0 = xla::ShapeUtil::GetTupleElementShape(padded_shape, 0);
|
||||
const xla::Shape& shape1 =
|
||||
xla::ShapeUtil::GetTupleElementShape(padded_shape, 1);
|
||||
if (shape0.IsTuple() || shape1.IsTuple()) {
|
||||
status->status = tensorflow::errors::InvalidArgument(
|
||||
"XlaTensors should not contain nested tuples. Shape: ",
|
||||
padded_shape.DebugString());
|
||||
return nullptr;
|
||||
}
|
||||
if (!xla::ShapeUtil::Equal(shape0, shape1)) {
|
||||
status->status = tensorflow::errors::InvalidArgument(
|
||||
"Subshapes of XlaTensors should be the same. Shape: ",
|
||||
padded_shape.DebugString());
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
// Since the only case we handle here are two equal subshapes, we
|
||||
// simply return one of them. The caller will interpret it as this
|
||||
// shape directly storing the 64bit types. This approximation is good
|
||||
// enough for this API's debugging use case.
|
||||
padded_shape = shape0;
|
||||
}
|
||||
|
||||
int rank = padded_shape.dimensions_size();
|
||||
std::vector<tensorflow::int64> dev_dims;
|
||||
dev_dims.reserve(rank);
|
||||
if (rank == 1) {
|
||||
// Rank 1 tensors might not have padded_shape.layout.minor_to_major set,
|
||||
dev_dims.push_back(padded_shape.dimensions(0));
|
||||
} else {
|
||||
for (int i = rank - 1; i >= 0; --i) {
|
||||
tensorflow::int64 dim_index = padded_shape.layout().minor_to_major(i);
|
||||
dev_dims.push_back(padded_shape.dimensions(dim_index));
|
||||
}
|
||||
}
|
||||
status->status = tensorflow::Status::OK();
|
||||
return new TFE_TensorDebugInfo(dev_dims);
|
||||
}
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
|
||||
// If the tensor is not an XLA tensor, the device shape is
|
||||
// the same as regular tensor shape.
|
||||
std::vector<tensorflow::int64> dev_dims =
|
||||
TensorShapeAsVector(*handle, &status->status);
|
||||
if (!status->status.ok()) {
|
||||
|
@ -316,86 +316,6 @@ TEST(CAPI, Function_ident_CPU) {
|
||||
TF_DeleteStatus(status);
|
||||
}
|
||||
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
TEST(CAPI, Function_ident_XLA_CPU) {
|
||||
// First create a simple identity function.
|
||||
TF_Graph* function_graph = TF_NewGraph();
|
||||
TF_OperationDescription* arg_descr =
|
||||
TF_NewOperation(function_graph, "Placeholder", "arg");
|
||||
TF_SetAttrType(arg_descr, "dtype", TF_INT32);
|
||||
TF_Status* status = TF_NewStatus();
|
||||
TF_Operation* arg = TF_FinishOperation(arg_descr, status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TF_OperationDescription* id_descr =
|
||||
TF_NewOperation(function_graph, "Identity", "id");
|
||||
TF_SetAttrType(id_descr, "T", TF_INT32);
|
||||
TF_AddInput(id_descr, {arg, 0});
|
||||
TF_Operation* id = TF_FinishOperation(id_descr, status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TF_Output input{arg, 0};
|
||||
TF_Output output{id, 0};
|
||||
TF_Function* fn =
|
||||
TF_GraphToFunction(function_graph, "ident", 0, 1, &id, 1, &input, 1,
|
||||
&output, nullptr, nullptr, "test", status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TF_DeleteGraph(function_graph);
|
||||
TFE_ContextOptions* opts = TFE_NewContextOptions();
|
||||
TFE_Context* ctx = TFE_NewContext(opts, status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TFE_DeleteContextOptions(opts);
|
||||
TFE_ContextAddFunction(ctx, fn, status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TF_DeleteFunction(fn);
|
||||
|
||||
for (bool async : {false, true, false}) {
|
||||
TFE_Executor* old_executor = TFE_ContextGetExecutorForThread(ctx);
|
||||
TFE_Executor* executor = TFE_NewExecutor(async);
|
||||
TFE_ContextSetExecutorForThread(ctx, executor);
|
||||
CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK);
|
||||
TF_Tensor* t =
|
||||
TF_AllocateTensor(TF_INT32, nullptr, 0, 1 * sizeof(tensorflow::int32));
|
||||
*reinterpret_cast<tensorflow::int32*>(TF_TensorData(t)) = 42;
|
||||
TFE_TensorHandle* h = TFE_NewTensorHandle(t, status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TF_DeleteTensor(t);
|
||||
|
||||
TFE_Op* op = TFE_NewOp(ctx, "ident", status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TFE_OpAddInput(op, h, status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
|
||||
// Now run it via XLA.
|
||||
TFE_OpSetXLACompilation(op, true);
|
||||
|
||||
std::vector<TFE_TensorHandle*> result;
|
||||
result.push_back(nullptr);
|
||||
int num_retvals = 1;
|
||||
TFE_Execute(op, result.data(), &num_retvals, status);
|
||||
TFE_DeleteOp(op);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
ASSERT_EQ(num_retvals, 1);
|
||||
|
||||
TF_Tensor* r = TFE_TensorHandleResolve(result[0], status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
EXPECT_EQ(*reinterpret_cast<tensorflow::int32*>(TF_TensorData(r)), 42);
|
||||
TFE_ContextSetExecutorForThread(ctx, old_executor);
|
||||
TFE_ExecutorWaitForAllPendingNodes(executor, status);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
TFE_DeleteExecutor(executor);
|
||||
TFE_DeleteExecutor(old_executor);
|
||||
TFE_DeleteTensorHandle(h);
|
||||
TF_DeleteTensor(r);
|
||||
TFE_DeleteTensorHandle(result[0]);
|
||||
}
|
||||
TFE_ContextRemoveFunction(ctx, "ident", status);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TFE_DeleteContext(ctx);
|
||||
ASSERT_TRUE(TF_GetCode(status) == TF_OK) << TF_Message(status);
|
||||
TF_DeleteStatus(status);
|
||||
}
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
|
||||
void Executor_MatMul_CPU(bool async) {
|
||||
TF_Status* status = TF_NewStatus();
|
||||
TFE_ContextOptions* opts = TFE_NewContextOptions();
|
||||
|
@ -876,89 +876,6 @@ TEST(CAPI, Execute_Min_CPU) {
|
||||
TF_DeleteStatus(status);
|
||||
}
|
||||
|
||||
#ifdef TENSORFLOW_EAGER_USE_XLA
|
||||
void Execute_MatMul_XLA_CPU(bool async) {
|
||||
TF_Status* status = TF_NewStatus();
|
||||
TFE_ContextOptions* opts = TFE_NewContextOptions();
|
||||
TFE_ContextOptionsSetAsync(opts, static_cast<unsigned char>(async));
|
||||
TFE_Context* ctx = TFE_NewContext(opts, status);
|
||||
CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
TFE_DeleteContextOptions(opts);
|
||||
|
||||
TFE_TensorHandle* m = TestMatrixTensorHandle(ctx);
|
||||
TFE_Op* matmul = MatMulOp(ctx, m, m);
|
||||
|
||||
TFE_OpSetXLACompilation(matmul, true);
|
||||
|
||||
TFE_TensorHandle* retvals[1] = {nullptr};
|
||||
int num_retvals = 1;
|
||||
TFE_Execute(matmul, &retvals[0], &num_retvals, status);
|
||||
// Running a primitive TF operator via XLA is not yet supported.
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
|
||||
TFE_DeleteOp(matmul);
|
||||
TFE_DeleteTensorHandle(m);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
|
||||
EXPECT_EQ(1, num_retvals);
|
||||
|
||||
TF_Tensor* t = TFE_TensorHandleResolve(retvals[0], status);
|
||||
TFE_DeleteTensorHandle(retvals[0]);
|
||||
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_DeleteContext(ctx);
|
||||
TF_DeleteStatus(status);
|
||||
}
|
||||
TEST(CAPI, Execute_MatMul_XLA_CPU) { Execute_MatMul_XLA_CPU(false); }
|
||||
TEST(CAPI, Execute_MatMul_XLA_CPUAsync) { Execute_MatMul_XLA_CPU(true); }
|
||||
|
||||
void Execute_Min_XLA_CPU(bool async) {
|
||||
TF_Status* status = TF_NewStatus();
|
||||
TFE_ContextOptions* opts = TFE_NewContextOptions();
|
||||
TFE_ContextOptionsSetAsync(opts, static_cast<unsigned char>(async));
|
||||
TFE_Context* ctx = TFE_NewContext(opts, status);
|
||||
CHECK_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
TFE_DeleteContextOptions(opts);
|
||||
|
||||
TFE_TensorHandle* input = TestMatrixTensorHandle(ctx);
|
||||
TFE_TensorHandle* axis = TestAxisTensorHandle(ctx);
|
||||
TFE_Op* minOp = MinOp(ctx, input, axis);
|
||||
|
||||
TFE_OpSetXLACompilation(minOp, true);
|
||||
|
||||
TFE_TensorHandle* retvals[1] = {nullptr};
|
||||
int num_retvals = 1;
|
||||
TFE_Execute(minOp, &retvals[0], &num_retvals, status);
|
||||
EXPECT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
TFE_DeleteOp(minOp);
|
||||
TFE_DeleteTensorHandle(input);
|
||||
TFE_DeleteTensorHandle(axis);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
ASSERT_EQ(1, num_retvals);
|
||||
|
||||
TF_Tensor* t = TFE_TensorHandleResolve(retvals[0], status);
|
||||
TFE_DeleteTensorHandle(retvals[0]);
|
||||
ASSERT_EQ(TF_OK, TF_GetCode(status)) << TF_Message(status);
|
||||
float output[2] = {0};
|
||||
EXPECT_EQ(sizeof(output), TF_TensorByteSize(t));
|
||||
memcpy(&output[0], TF_TensorData(t), TF_TensorByteSize(t));
|
||||
TF_DeleteTensor(t);
|
||||
EXPECT_EQ(1, output[0]);
|
||||
EXPECT_EQ(3, output[1]);
|
||||
TFE_DeleteContext(ctx);
|
||||
TF_DeleteStatus(status);
|
||||
}
|
||||
TEST(CAPI, Execute_Min_XLA_CPU) { Execute_Min_XLA_CPU(false); }
|
||||
TEST(CAPI, Execute_Min_XLA_CPUAsync) { Execute_Min_XLA_CPU(true); }
|
||||
#endif // TENSORFLOW_EAGER_USE_XLA
|
||||
|
||||
void ExecuteWithTracing(bool async) {
|
||||
TF_Status* status = TF_NewStatus();
|
||||
TFE_ContextOptions* opts = TFE_NewContextOptions();
|
||||
|
@ -47,9 +47,6 @@ class ImmediateExecutionOperation : public AbstractOperation {
|
||||
virtual Status InputLength(const char* input_name, int* length) = 0;
|
||||
virtual Status OutputLength(const char* output_name, int* length) = 0;
|
||||
|
||||
// Experimental
|
||||
virtual Status SetUseXla(bool enable) = 0;
|
||||
|
||||
// Set stack trace to be used for potential async error reporting.
|
||||
virtual void SetStackTrace(AbstractStackTrace stack_trace) = 0;
|
||||
|
||||
|
@ -2,7 +2,6 @@ load(
|
||||
"//tensorflow:tensorflow.bzl",
|
||||
"tf_copts",
|
||||
"tf_cuda_library",
|
||||
"tfe_xla_copts",
|
||||
)
|
||||
|
||||
package(
|
||||
@ -20,7 +19,7 @@ tf_cuda_library(
|
||||
srcs = [
|
||||
"c_api_unified_experimental_mlir.cc",
|
||||
],
|
||||
copts = tf_copts() + tfe_xla_copts(),
|
||||
copts = tf_copts(),
|
||||
deps = [
|
||||
"//tensorflow/c:c_api",
|
||||
"//tensorflow/c:tensor_interface",
|
||||
|
@ -277,11 +277,6 @@ Status EagerOperation::AddInputList(
|
||||
return InferInputListAttrs(inputs.size());
|
||||
}
|
||||
|
||||
Status EagerOperation::SetUseXla(bool enable) {
|
||||
use_xla_ = enable;
|
||||
return Status::OK();
|
||||
}
|
||||
|
||||
Status EagerOperation::Reset(
|
||||
const char* op, const char* device_name, bool remote,
|
||||
EagerExecutor* executor,
|
||||
@ -313,7 +308,6 @@ Status EagerOperation::Reset(
|
||||
"registered in the binary running in this process.");
|
||||
}
|
||||
attrs_.Reset(op);
|
||||
use_xla_ = false;
|
||||
stack_trace_.reset();
|
||||
is_function_ = is_function;
|
||||
cancellation_manager_ = nullptr;
|
||||
|
@ -120,8 +120,6 @@ class EagerOperation : public ImmediateExecutionOperation {
|
||||
Status InputLength(const char* input_name, int* length) override;
|
||||
Status OutputLength(const char* output_name, int* length) override;
|
||||
|
||||
Status SetUseXla(bool enable) override;
|
||||
|
||||
void SetStackTrace(AbstractStackTrace stack_trace) override {
|
||||
stack_trace_ = stack_trace;
|
||||
}
|
||||
@ -227,7 +225,6 @@ class EagerOperation : public ImmediateExecutionOperation {
|
||||
// updated accordingly.
|
||||
VariantDevice device_;
|
||||
|
||||
bool use_xla_ = false;
|
||||
absl::optional<AbstractStackTrace> stack_trace_;
|
||||
bool is_function_; // Conceptually const, but can't be because of Reset
|
||||
bool colocation_exempt_;
|
||||
|
@ -352,12 +352,6 @@ def tf_copts(
|
||||
def tf_openmp_copts():
|
||||
return (if_mkl_lnx_x64(["-fopenmp"]) + if_mkldnn_threadpool(["-fno-openmp"]))
|
||||
|
||||
def tfe_xla_copts():
|
||||
return select({
|
||||
"//tensorflow:with_xla_support": ["-DTENSORFLOW_EAGER_USE_XLA"],
|
||||
"//conditions:default": [],
|
||||
})
|
||||
|
||||
def tf_opts_nortti():
|
||||
return [
|
||||
"-fno-rtti",
|
||||
|
Loading…
Reference in New Issue
Block a user