diff -ruN tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/translate/export_graphdef.cc tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/translate/export_graphdef.cc --- tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/translate/export_graphdef.cc 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/translate/export_graphdef.cc 2020-01-29 22:26:46.168800847 +0900 @@ -221,7 +221,7 @@ AttrValue index_attr; index_attr.set_i(index); (*node_def->mutable_attr())["index"] = index_attr; - return node_def; + return std::move(node_def); } StatusOr> Exporter::GetReturnNode( @@ -240,7 +240,7 @@ AttrValue index_attr; index_attr.set_i(index); (*node_def->mutable_attr())["index"] = index_attr; - return node_def; + return std::move(node_def); } Status Exporter::AddEdgeBetweenNodes(mlir::Value* src, Node* dst_node, @@ -511,7 +511,7 @@ // Fixes the edges between the inserted nodes and special "_SOURCE" and // "_SINK". FixupSourceAndSinkEdges(graph.get()); - return graph; + return std::move(graph); } Status Exporter::ConvertLibFunction(const ExporterConfigs& configs, @@ -638,7 +638,7 @@ node_def.clear_experimental_debug_info(); } } - return graphdef; + return std::move(graphdef); } } // namespace tensorflow diff -ruN tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/translate/export_tf_dialect_op.cc tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/translate/export_tf_dialect_op.cc --- tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/translate/export_tf_dialect_op.cc 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/translate/export_tf_dialect_op.cc 2020-01-29 22:26:46.169800843 +0900 @@ -95,7 +95,7 @@ return errors::Internal("Falied to populate derived attrs: " + status.ToString()); } - return node_def; + return std::move(node_def); } } // namespace tensorflow diff -ruN tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/translate/import_model.cc tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/translate/import_model.cc --- tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/translate/import_model.cc 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/translate/import_model.cc 2020-01-29 22:26:46.169800843 +0900 @@ -1548,7 +1548,7 @@ TF_RETURN_IF_ERROR(importer.ImporterBase::Convert( "main", func_type, arg_nodes, ret_nodes, control_ret_nodes, attrs)); - return module; + return std::move(module); } StatusOr GraphDefImporter::InferMainFunctionType( diff -ruN tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/utils/export_utils.cc tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/utils/export_utils.cc --- tensorflow-1.15.2/tensorflow/compiler/mlir/tensorflow/utils/export_utils.cc 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/compiler/mlir/tensorflow/utils/export_utils.cc 2020-01-29 22:26:46.170800838 +0900 @@ -221,7 +221,7 @@ if (node_def->op() == "If") UpdateCompositeIfOp(node_def.get()); if (node_def->op() == "While") UpdateCompositeWhileOp(node_def.get()); - return node_def; + return std::move(node_def); } Status ConvertAttributes(const llvm::ArrayRef attrs, diff -ruN tensorflow-1.15.2/tensorflow/core/kernels/tridiagonal_solve_op_gpu.cu.cc tensorflow-1.15.2-cuda/tensorflow/core/kernels/tridiagonal_solve_op_gpu.cu.cc --- tensorflow-1.15.2/tensorflow/core/kernels/tridiagonal_solve_op_gpu.cu.cc 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/core/kernels/tridiagonal_solve_op_gpu.cu.cc 2020-01-29 22:26:46.170800838 +0900 @@ -43,8 +43,11 @@ __global__ void SolveForSizeOneOrTwoKernel(const int m, const Scalar* diags, const Scalar* rhs, const int num_rhs, Scalar* x, bool* not_invertible) { + Scalar z(0); if (m == 1) { - if (diags[1] == Scalar(0)) { + bool b = true; + for(int i = 0; i < sizeof(z); i++) b &= ((char*)&diags[1])[i] == ((char*)&z)[i]; + if (b) { *not_invertible = true; return; } @@ -53,7 +56,9 @@ } } else { Scalar det = diags[2] * diags[3] - diags[0] * diags[5]; - if (det == Scalar(0)) { + bool b = true; + for(int i = 0; i < sizeof(z); i++) b &= ((char*)&det)[i] == ((char*)&z)[i]; + if (b) { *not_invertible = true; return; } diff -ruN tensorflow-1.15.2/tensorflow/stream_executor/cuda/cuda_gpu_executor.cc tensorflow-1.15.2-cuda/tensorflow/stream_executor/cuda/cuda_gpu_executor.cc --- tensorflow-1.15.2/tensorflow/stream_executor/cuda/cuda_gpu_executor.cc 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/stream_executor/cuda/cuda_gpu_executor.cc 2020-01-29 22:26:46.170800838 +0900 @@ -195,7 +195,7 @@ _NSGetExecutablePath(nullptr, &buffer_size); char unresolved_path[buffer_size]; _NSGetExecutablePath(unresolved_path, &buffer_size); - CHECK_ERR(realpath(unresolved_path, exe_path) ? 1 : -1); + CHECK(realpath(unresolved_path, exe_path)); #else #if defined(PLATFORM_WINDOWS) HMODULE hModule = GetModuleHandle(NULL); diff -ruN tensorflow-1.15.2/tensorflow/tensorflow.bzl tensorflow-1.15.2-cuda/tensorflow/tensorflow.bzl --- tensorflow-1.15.2/tensorflow/tensorflow.bzl 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/tensorflow.bzl 2020-01-29 22:26:46.171800833 +0900 @@ -365,7 +365,7 @@ ) def _make_search_paths(prefix, levels_to_root): - return ",".join( + return "-rpath,/usr/local/cuda/lib/,-rpath,/usr/local/cuda/lib64/," + ",".join( [ "-rpath,%s/%s" % (prefix, "/".join([".."] * search_level)) for search_level in range(levels_to_root + 1) diff -ruN tensorflow-1.15.2/tensorflow/tools/pip_package/BUILD tensorflow-1.15.2-cuda/tensorflow/tools/pip_package/BUILD --- tensorflow-1.15.2/tensorflow/tools/pip_package/BUILD 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/tools/pip_package/BUILD 2020-01-29 22:26:46.171800833 +0900 @@ -200,7 +200,6 @@ "//conditions:default": [], }) + if_cuda([ "@cub_archive//:LICENSE.TXT", - "@local_config_nccl//:LICENSE", ]) + if_mkl([ "//third_party/mkl:LICENSE", "//third_party/mkl_dnn:LICENSE", diff -ruN tensorflow-1.15.2/tensorflow/workspace.bzl tensorflow-1.15.2-cuda/tensorflow/workspace.bzl --- tensorflow-1.15.2/tensorflow/workspace.bzl 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/tensorflow/workspace.bzl 2020-01-29 22:26:46.172800829 +0900 @@ -157,6 +157,7 @@ tf_http_archive( name = "com_google_absl", build_file = clean_dep("//third_party:com_google_absl.BUILD"), + patch_file = clean_dep("//third_party:absl.patch"), sha256 = "acd93f6baaedc4414ebd08b33bebca7c7a46888916101d8c0b8083573526d070", strip_prefix = "abseil-cpp-43ef2148c0936ebf7cb4be6b19927a9d9d145b8f", urls = [ diff -ruN tensorflow-1.15.2/third_party/absl.patch tensorflow-1.15.2-cuda/third_party/absl.patch --- tensorflow-1.15.2/third_party/absl.patch 1970-01-01 09:00:00.000000000 +0900 +++ tensorflow-1.15.2-cuda/third_party/absl.patch 2019-09-11 11:15:18.073174802 +0900 @@ -0,0 +1,22 @@ +diff -ruN abseil-cpp-daf381e8535a1f1f1b8a75966a74e7cca63dee89/absl/container/internal/compressed_tuple.h abseil-cpp-daf381e8535a1f1f1b8a75966a74e7cca63dee89-patch/absl/container/internal/compressed_tuple.h +--- abseil-cpp-daf381e8535a1f1f1b8a75966a74e7cca63dee89/absl/container/internal/compressed_tuple.h 2019-05-17 04:48:51.000000000 +0900 ++++ abseil-cpp-daf381e8535a1f1f1b8a75966a74e7cca63dee89-patch/absl/container/internal/compressed_tuple.h 2019-06-01 18:08:31.079239816 +0900 +@@ -167,14 +167,16 @@ + + template + ElemT&& get() && { ++ using s = internal_compressed_tuple::Storage; + return std::move(*this) +- .internal_compressed_tuple::template Storage::get(); ++ .s::get(); + } + + template + constexpr const ElemT&& get() const&& { ++ using s = internal_compressed_tuple::Storage; + return absl::move(*this) +- .internal_compressed_tuple::template Storage::get(); ++ .s::get(); + } + }; + diff -ruN tensorflow-1.15.2/third_party/eigen3/gpu_packet_math.patch tensorflow-1.15.2-cuda/third_party/eigen3/gpu_packet_math.patch --- tensorflow-1.15.2/third_party/eigen3/gpu_packet_math.patch 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/third_party/eigen3/gpu_packet_math.patch 2020-01-29 22:26:59.093742824 +0900 @@ -150,3 +150,14 @@ +} // end namespace Eigen + +#endif // HIP_VECTOR_COMPATIBILITY_H +--- a/unsupported/Eigen/CXX11/src/Tensor/TensorRandom.h ++++ b/unsupported/Eigen/CXX11/src/Tensor/TensorRandom.h +@@ -15,7 +15,7 @@ + + namespace { + +-EIGEN_DEVICE_FUNC uint64_t get_random_seed() { ++static EIGEN_DEVICE_FUNC uint64_t get_random_seed() { + #if defined(EIGEN_GPU_COMPILE_PHASE) + // We don't support 3d kernels since we currently only use 1 and + // 2d kernels. diff -ruN tensorflow-1.15.2/third_party/protobuf/protobuf.patch tensorflow-1.15.2-cuda/third_party/protobuf/protobuf.patch --- tensorflow-1.15.2/third_party/protobuf/protobuf.patch 2020-01-26 12:57:51.000000000 +0900 +++ tensorflow-1.15.2-cuda/third_party/protobuf/protobuf.patch 2020-01-29 22:26:46.172800829 +0900 @@ -20,3 +20,38 @@ includes = ["src/"], visibility = ["//visibility:public"], ) +diff --git a/src/google/protobuf/map.h b/src/google/protobuf/map.h +index 40f35e92..1af91f9c 100644 +--- a/src/google/protobuf/map.h ++++ b/src/google/protobuf/map.h +@@ -424,10 +424,10 @@ + node_ = NULL; + for (bucket_index_ = start_bucket; bucket_index_ < m_->num_buckets_; + bucket_index_++) { +- if (m_->TableEntryIsNonEmptyList(bucket_index_)) { ++ if ((*m_).TableEntryIsNonEmptyList(bucket_index_)) { + node_ = static_cast(m_->table_[bucket_index_]); + break; +- } else if (m_->TableEntryIsTree(bucket_index_)) { ++ } else if ((*m_).TableEntryIsTree(bucket_index_)) { + Tree* tree = static_cast(m_->table_[bucket_index_]); + GOOGLE_DCHECK(!tree->empty()); + node_ = NodePtrFromKeyPtr(*tree->begin()); +@@ -485,7 +485,7 @@ + if (m_->table_[bucket_index_] == static_cast(node_)) return true; + // Less common: the bucket is a linked list with node_ somewhere in it, + // but not at the head. +- if (m_->TableEntryIsNonEmptyList(bucket_index_)) { ++ if ((*m_).TableEntryIsNonEmptyList(bucket_index_)) { + Node* l = static_cast(m_->table_[bucket_index_]); + while ((l = l->next) != NULL) { + if (l == node_) { +@@ -499,7 +499,7 @@ + // find-like method that compares Node* instead of const Key&. + iterator_base i(m_->find(*KeyPtrFromNodePtr(node_), it)); + bucket_index_ = i.bucket_index_; +- return m_->TableEntryIsList(bucket_index_); ++ return (*m_).TableEntryIsList(bucket_index_); + } + + Node* node_;