From 9f5be553f6e475b323587ce0a5cda61491eb62f7 Mon Sep 17 00:00:00 2001 From: jackyko Date: Tue, 23 Jan 2018 18:31:39 +0800 Subject: [PATCH 001/596] tpu contrib fix --- tensorflow/contrib/cmake/tf_core_framework.cmake | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tensorflow/contrib/cmake/tf_core_framework.cmake b/tensorflow/contrib/cmake/tf_core_framework.cmake index 24d7fb82a2..129c208ecd 100644 --- a/tensorflow/contrib/cmake/tf_core_framework.cmake +++ b/tensorflow/contrib/cmake/tf_core_framework.cmake @@ -126,7 +126,9 @@ endfunction() file(GLOB_RECURSE tf_protos_cc_srcs RELATIVE ${tensorflow_source_dir} "${tensorflow_source_dir}/tensorflow/core/*.proto" "${tensorflow_source_dir}/tensorflow/contrib/boosted_trees/proto/*.proto" + "${tensorflow_source_dir}/tensorflow/contrib/tpu/proto/*.proto" ) + RELATIVE_PROTOBUF_GENERATE_CPP(PROTO_SRCS PROTO_HDRS ${tensorflow_source_dir} ${tf_protos_cc_srcs} ) -- GitLab From 1c6b52927680e8bb1b1f0ebe98c5a7a7033f9af4 Mon Sep 17 00:00:00 2001 From: jackyko Date: Thu, 25 Jan 2018 14:16:53 +0800 Subject: [PATCH 002/596] improve cmake gpu build --- tensorflow/contrib/cmake/CMakeLists.txt | 12 ++++++++---- tensorflow/contrib/cmake/external/grpc.cmake | 11 ++++++----- tensorflow/contrib/cmake/tf_cc_ops.cmake | 2 +- tensorflow/contrib/cmake/tf_core_framework.cmake | 8 ++++++-- tensorflow/core/platform/default/gpu/cupti_wrapper.h | 8 ++++---- 5 files changed, 25 insertions(+), 16 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 817e96f5da..5f54ba5e2d 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -291,6 +291,8 @@ if (tensorflow_ENABLE_GPU) set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};--include-path ${PROJECT_BINARY_DIR}/$\{build_configuration\};--expt-relaxed-constexpr) set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-ftz=true) # Flush denormals to zero set(CUDA_INCLUDE ${CUDA_TOOLKIT_TARGET_DIR} ${CUDA_TOOLKIT_TARGET_DIR}/extras/CUPTI/include) + + message(STATUS ${CUDA_INCLUDE}) include_directories(${CUDA_INCLUDE}) if (WIN32) add_definitions(-DGOOGLE_CUDA=1 -DTF_EXTRA_CUDA_CAPABILITIES=3.0,3.5,5.2) @@ -399,10 +401,12 @@ else(tensorflow_ENABLE_GPU) msvcp_dll_name=msvcp140.dll) endif(tensorflow_ENABLE_GPU) -# Find python executable -include(FindPythonInterp) -if(NOT ${PYTHONINTERP_FOUND}) - message(FATAL_ERROR "CMake was unable to find a python interpreter.") +if(tensorflow_BUILD_PYTHON_BINDINGS) + # Find python executable + include(FindPythonInterp) + if(NOT ${PYTHONINTERP_FOUND}) + message(FATAL_ERROR "CMake was unable to find a python interpreter.") + endif() endif() # Let's get to work! diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 28adb4fe84..28b85e0a19 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -21,14 +21,16 @@ set(GRPC_TAG 730b778632e79cc3c96ad237f282d687ee325ce7) if(WIN32) set(grpc_STATIC_LIBRARIES - ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/Release/grpc++_unsecure.lib - ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/Release/grpc_unsecure.lib - ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/Release/gpr.lib) + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/$(Configuration)/grpc++_unsecure.lib + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/$(Configuration)/grpc_unsecure.lib + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/$(Configuration)/gpr.lib) else() set(grpc_STATIC_LIBRARIES ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgrpc++_unsecure.a ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgrpc_unsecure.a - ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgpr.a) + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgpr.a + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/third_party/cares/cares/lib/libcares.a + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/third_party/zlib/libz.a) endif() add_definitions(-DGRPC_ARES=0) @@ -49,7 +51,6 @@ ExternalProject_Add(grpc -DPROTOBUF_INCLUDE_DIRS:STRING=${PROTOBUF_INCLUDE_DIRS} -DPROTOBUF_LIBRARIES:STRING=${protobuf_STATIC_LIBRARIES} -DZLIB_ROOT:STRING=${ZLIB_INSTALL} - -DgRPC_SSL_PROVIDER:STRING=NONE ) # grpc/src/core/ext/census/tracing.c depends on the existence of openssl/rand.h. diff --git a/tensorflow/contrib/cmake/tf_cc_ops.cmake b/tensorflow/contrib/cmake/tf_cc_ops.cmake index f3cf3e7044..1791dad48c 100644 --- a/tensorflow/contrib/cmake/tf_cc_ops.cmake +++ b/tensorflow/contrib/cmake/tf_cc_ops.cmake @@ -149,7 +149,7 @@ add_library(tf_cc OBJECT ${tf_cc_srcs}) add_dependencies(tf_cc tf_cc_framework tf_cc_ops) if (WIN32) - set (pywrap_tensorflow_lib "${CMAKE_CURRENT_BINARY_DIR}/${CMAKE_BUILD_TYPE}/pywrap_tensorflow_internal.lib") + set (pywrap_tensorflow_lib "${CMAKE_CURRENT_BINARY_DIR}/$(Configuration)/pywrap_tensorflow_internal.lib") else (WIN32) set (pywrap_tensorflow_lib "${CMAKE_CURRENT_BINARY_DIR}/libpywrap_tensorflow_internal.so") endif (WIN32) diff --git a/tensorflow/contrib/cmake/tf_core_framework.cmake b/tensorflow/contrib/cmake/tf_core_framework.cmake index 129c208ecd..91f8f51e5b 100644 --- a/tensorflow/contrib/cmake/tf_core_framework.cmake +++ b/tensorflow/contrib/cmake/tf_core_framework.cmake @@ -203,14 +203,18 @@ file(GLOB tf_core_platform_srcs "${tensorflow_source_dir}/tensorflow/core/framework/resource_handle.h" "${tensorflow_source_dir}/tensorflow/core/framework/resource_handle.cc") if (NOT tensorflow_ENABLE_GPU) - file(GLOB tf_core_platform_gpu_srcs + file(GLOB tf_core_platform_gpu_srcs_exclude "${tensorflow_source_dir}/tensorflow/core/platform/cuda_libdevice_path.*" "${tensorflow_source_dir}/tensorflow/core/platform/default/cuda_libdevice_path.*") - list(REMOVE_ITEM tf_core_platform_srcs ${tf_core_platform_gpu_srcs}) + list(REMOVE_ITEM tf_core_platform_srcs ${tf_core_platform_gpu_srcs_exclude}) else() file(GLOB tf_core_platform_srcs_exclude "${tensorflow_source_dir}/tensorflow/core/platform/default/device_tracer.cc") list(REMOVE_ITEM tf_core_platform_srcs ${tf_core_platform_srcs_exclude}) + # file(GLOB tf_core_platform_gpu_srcs + # "${tensorflow_source_dir}/tensorflow/core/platform/default/gpu/*.h" + # "${tensorflow_source_dir}/tensorflow/core/platform/default/gpu/*.cc") + # list(APPEND tf_core_platform_srcs ${tf_core_platform_gpu_srcs}) endif() file(GLOB tf_core_platform_exclude_srcs diff --git a/tensorflow/core/platform/default/gpu/cupti_wrapper.h b/tensorflow/core/platform/default/gpu/cupti_wrapper.h index acd889e474..f1e3af12cf 100644 --- a/tensorflow/core/platform/default/gpu/cupti_wrapper.h +++ b/tensorflow/core/platform/default/gpu/cupti_wrapper.h @@ -20,11 +20,11 @@ limitations under the License. #include #include -#if defined(WIN32) +// #if defined(WIN32) #include "extras/CUPTI/include/cupti.h" -#else -#include "cuda/extras/CUPTI/include/cupti.h" -#endif +// #else +// #include "cuda/extras/CUPTI/include/cupti.h" +// #endif namespace perftools { namespace gputools { namespace profiler { -- GitLab From f8ebe607de0f71a4eace98cd110a4e931332aa6b Mon Sep 17 00:00:00 2001 From: jackyko Date: Thu, 25 Jan 2018 14:19:22 +0800 Subject: [PATCH 003/596] improve cmake option with python build --- tensorflow/contrib/cmake/CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 5f54ba5e2d..8e736f4def 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -292,7 +292,6 @@ if (tensorflow_ENABLE_GPU) set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-ftz=true) # Flush denormals to zero set(CUDA_INCLUDE ${CUDA_TOOLKIT_TARGET_DIR} ${CUDA_TOOLKIT_TARGET_DIR}/extras/CUPTI/include) - message(STATUS ${CUDA_INCLUDE}) include_directories(${CUDA_INCLUDE}) if (WIN32) add_definitions(-DGOOGLE_CUDA=1 -DTF_EXTRA_CUDA_CAPABILITIES=3.0,3.5,5.2) -- GitLab From 27ddbee4d59874a796020678f8eb45cd4b5c5ff1 Mon Sep 17 00:00:00 2001 From: jackyko Date: Thu, 25 Jan 2018 14:21:38 +0800 Subject: [PATCH 004/596] better variable name --- tensorflow/contrib/cmake/tf_core_framework.cmake | 4 ---- 1 file changed, 4 deletions(-) diff --git a/tensorflow/contrib/cmake/tf_core_framework.cmake b/tensorflow/contrib/cmake/tf_core_framework.cmake index 91f8f51e5b..e39b7f9289 100644 --- a/tensorflow/contrib/cmake/tf_core_framework.cmake +++ b/tensorflow/contrib/cmake/tf_core_framework.cmake @@ -211,10 +211,6 @@ else() file(GLOB tf_core_platform_srcs_exclude "${tensorflow_source_dir}/tensorflow/core/platform/default/device_tracer.cc") list(REMOVE_ITEM tf_core_platform_srcs ${tf_core_platform_srcs_exclude}) - # file(GLOB tf_core_platform_gpu_srcs - # "${tensorflow_source_dir}/tensorflow/core/platform/default/gpu/*.h" - # "${tensorflow_source_dir}/tensorflow/core/platform/default/gpu/*.cc") - # list(APPEND tf_core_platform_srcs ${tf_core_platform_gpu_srcs}) endif() file(GLOB tf_core_platform_exclude_srcs -- GitLab From 418657e25288945cbccf11eb4bbd4e1bcfc4b2ad Mon Sep 17 00:00:00 2001 From: jackyko Date: Fri, 26 Jan 2018 10:22:45 +0800 Subject: [PATCH 005/596] gcc bug fix --- tensorflow/contrib/cmake/tf_python.cmake | 17 ++++++++++++++++- .../core/platform/default/gpu/cupti_wrapper.h | 8 ++++---- 2 files changed, 20 insertions(+), 5 deletions(-) diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index 8862390d2b..c8b6ced79c 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -585,13 +585,28 @@ target_include_directories(pywrap_tensorflow_internal PUBLIC ${NUMPY_INCLUDE_DIR} ) -target_link_libraries(pywrap_tensorflow_internal PRIVATE +if(CMAKE_COMPILER_IS_GNUCC AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 5.0) + # There is a bug in GCC 5 resulting in undefined reference to a __cpu_model function when + # linking to the tensorflow library. Adding the following libraries fixes it. + # See issue on github: https://github.com/tensorflow/tensorflow/issues/9593 + target_link_libraries(pywrap_tensorflow_internal PRIVATE ${tf_core_gpu_kernels_lib} ${tensorflow_EXTERNAL_LIBRARIES} tf_protos_cc tf_python_protos_cc ${PYTHON_LIBRARIES} + gcc_s + gcc ) +else() + target_link_libraries(pywrap_tensorflow_internal PRIVATE + ${tf_core_gpu_kernels_lib} + ${tensorflow_EXTERNAL_LIBRARIES} + tf_protos_cc + tf_python_protos_cc + ${PYTHON_LIBRARIES} +) +endif() if(WIN32) diff --git a/tensorflow/core/platform/default/gpu/cupti_wrapper.h b/tensorflow/core/platform/default/gpu/cupti_wrapper.h index f1e3af12cf..acd889e474 100644 --- a/tensorflow/core/platform/default/gpu/cupti_wrapper.h +++ b/tensorflow/core/platform/default/gpu/cupti_wrapper.h @@ -20,11 +20,11 @@ limitations under the License. #include #include -// #if defined(WIN32) +#if defined(WIN32) #include "extras/CUPTI/include/cupti.h" -// #else -// #include "cuda/extras/CUPTI/include/cupti.h" -// #endif +#else +#include "cuda/extras/CUPTI/include/cupti.h" +#endif namespace perftools { namespace gputools { namespace profiler { -- GitLab From b5fb75bb77d671bfd77fb3eb5f7ac3be5604e96b Mon Sep 17 00:00:00 2001 From: jackyko Date: Fri, 26 Jan 2018 11:39:47 +0800 Subject: [PATCH 006/596] fix build issue with cmake --- tensorflow/core/platform/default/gpu/cupti_wrapper.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tensorflow/core/platform/default/gpu/cupti_wrapper.h b/tensorflow/core/platform/default/gpu/cupti_wrapper.h index acd889e474..fd5968c827 100644 --- a/tensorflow/core/platform/default/gpu/cupti_wrapper.h +++ b/tensorflow/core/platform/default/gpu/cupti_wrapper.h @@ -20,10 +20,10 @@ limitations under the License. #include #include -#if defined(WIN32) -#include "extras/CUPTI/include/cupti.h" -#else +#if defined(PLATFORM_GOOGLE) #include "cuda/extras/CUPTI/include/cupti.h" +#else +#include "extras/CUPTI/include/cupti.h" #endif namespace perftools { namespace gputools { -- GitLab From cf55a3eb11df0ad5de84202ba9805977c5301ca5 Mon Sep 17 00:00:00 2001 From: jackyko Date: Fri, 26 Jan 2018 16:15:52 +0800 Subject: [PATCH 007/596] add missing ops --- tensorflow/contrib/cmake/tf_core_ops.cmake | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/tensorflow/contrib/cmake/tf_core_ops.cmake b/tensorflow/contrib/cmake/tf_core_ops.cmake index 138993db35..15b0227335 100644 --- a/tensorflow/contrib/cmake/tf_core_ops.cmake +++ b/tensorflow/contrib/cmake/tf_core_ops.cmake @@ -15,7 +15,7 @@ set(tf_op_lib_names "audio_ops" "array_ops" - "batch_ops" + "batch_ops" "bitwise_ops" "candidate_sampling_ops" "checkpoint_ops" @@ -23,11 +23,12 @@ set(tf_op_lib_names "ctc_ops" "data_flow_ops" "dataset_ops" + "function_ops" "functional_ops" "image_ops" "io_ops" "linalg_ops" - "list_ops" + "list_ops" "lookup_ops" "logging_ops" "math_ops" @@ -46,8 +47,9 @@ set(tf_op_lib_names "state_ops" "stateless_random_ops" "string_ops" - "summary_ops" + "summary_ops" "training_ops" + "word2vec_ops" ) foreach(tf_op_lib_name ${tf_op_lib_names}) -- GitLab From ab59f29dc8f930443881a1e27f7080d8ba87b119 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Fri, 26 Jan 2018 21:52:02 +0800 Subject: [PATCH 008/596] fix for cmake gui build --- tensorflow/contrib/cmake/tf_cc_ops.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/tf_cc_ops.cmake b/tensorflow/contrib/cmake/tf_cc_ops.cmake index f3cf3e7044..1791dad48c 100644 --- a/tensorflow/contrib/cmake/tf_cc_ops.cmake +++ b/tensorflow/contrib/cmake/tf_cc_ops.cmake @@ -149,7 +149,7 @@ add_library(tf_cc OBJECT ${tf_cc_srcs}) add_dependencies(tf_cc tf_cc_framework tf_cc_ops) if (WIN32) - set (pywrap_tensorflow_lib "${CMAKE_CURRENT_BINARY_DIR}/${CMAKE_BUILD_TYPE}/pywrap_tensorflow_internal.lib") + set (pywrap_tensorflow_lib "${CMAKE_CURRENT_BINARY_DIR}/$(Configuration)/pywrap_tensorflow_internal.lib") else (WIN32) set (pywrap_tensorflow_lib "${CMAKE_CURRENT_BINARY_DIR}/libpywrap_tensorflow_internal.so") endif (WIN32) -- GitLab From 3b316fc9d9468d957d4866e0397d4455bd3516bb Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sat, 27 Jan 2018 12:26:21 +0800 Subject: [PATCH 009/596] add version --- tensorflow/contrib/cmake/CMakeLists.txt | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 25d5dedb0e..c26c7cd3b7 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.5) # Project -project(tensorflow C CXX) +project(tensorflow VERSION 1.5.0 LANGUAGES C CXX) # Set C++14 as standard for the whole project set(CMAKE_CXX_STANDARD 14) @@ -134,9 +134,6 @@ if(WIN32) set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} /D_ITERATOR_DEBUG_LEVEL=0") set(CMAKE_CXX_FLAGS_MINSIZEREL "${CMAKE_CXX_FLAGS_MINSIZEREL} /D_ITERATOR_DEBUG_LEVEL=0") set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "${CMAKE_CXX_FLAGS_RELWITHDEBINFO} /D_ITERATOR_DEBUG_LEVEL=0") - - # Try to avoid flaky failures due to failed generation of generate.stamp files. - set(CMAKE_SUPPRESS_REGENERATION ON) endif() if ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU") -- GitLab From d6b7077d2dfd2d93410fc77ac3c14064f73ae933 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sat, 27 Jan 2018 13:53:04 +0800 Subject: [PATCH 010/596] tensorflow can be used by cmake find_package function --- .../contrib/cmake/TensorflowConfig.cmake.in | 16 ++++ .../cmake/TensorflowConfigVersion.cmake.in | 11 +++ tensorflow/contrib/cmake/tf_shared_lib.cmake | 80 +++++++++++++++++-- 3 files changed, 99 insertions(+), 8 deletions(-) create mode 100644 tensorflow/contrib/cmake/TensorflowConfig.cmake.in create mode 100644 tensorflow/contrib/cmake/TensorflowConfigVersion.cmake.in diff --git a/tensorflow/contrib/cmake/TensorflowConfig.cmake.in b/tensorflow/contrib/cmake/TensorflowConfig.cmake.in new file mode 100644 index 0000000000..cc04db6e95 --- /dev/null +++ b/tensorflow/contrib/cmake/TensorflowConfig.cmake.in @@ -0,0 +1,16 @@ +# - Config file for the Tensorflow package +# It defines the following variables +# TENSORFLOW_INCLUDE_DIRS - include directories for FooBar +# TENSORFLOW_LIBRARIES - libraries to link against + +# Compute paths +get_filename_component(TENSORFLOW_CMAKE_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH) +set(TENSORFLOW_INCLUDE_DIRS "@CONF_INCLUDE_DIRS@") + +# Our library dependencies (contains definitions for IMPORTED targets) +if(NOT TENSORFLOW_BINARY_DIR) + include("${TENSORFLOW_CMAKE_DIR}/TensorflowTargets.cmake") +endif() + +# These are IMPORTED targets created by TensorflowTargets.cmake +set(TENSORFLOW_LIBRARIES tensorflow) \ No newline at end of file diff --git a/tensorflow/contrib/cmake/TensorflowConfigVersion.cmake.in b/tensorflow/contrib/cmake/TensorflowConfigVersion.cmake.in new file mode 100644 index 0000000000..2a9609ddb9 --- /dev/null +++ b/tensorflow/contrib/cmake/TensorflowConfigVersion.cmake.in @@ -0,0 +1,11 @@ +set(PACKAGE_VERSION "@TENSORFLOW_VERSION@") + +# Check whether the requested PACKAGE_FIND_VERSION is compatible +if("${PACKAGE_VERSION}" VERSION_LESS "${PACKAGE_FIND_VERSION}") + set(PACKAGE_VERSION_COMPATIBLE FALSE) +else() + set(PACKAGE_VERSION_COMPATIBLE TRUE) + if ("${PACKAGE_VERSION}" VERSION_EQUAL "${PACKAGE_FIND_VERSION}") + set(PACKAGE_VERSION_EXACT TRUE) + endif() +endif() \ No newline at end of file diff --git a/tensorflow/contrib/cmake/tf_shared_lib.cmake b/tensorflow/contrib/cmake/tf_shared_lib.cmake index 571d2b0dec..a4a25dd557 100644 --- a/tensorflow/contrib/cmake/tf_shared_lib.cmake +++ b/tensorflow/contrib/cmake/tf_shared_lib.cmake @@ -91,6 +91,27 @@ if(CMAKE_COMPILER_IS_GNUCC AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER 5.0) target_link_libraries(tensorflow PRIVATE gcc_s gcc) endif() +# Offer the user the choice of overriding the installation directories +set(INSTALL_LIB_DIR lib CACHE PATH "Installation directory for libraries") +set(INSTALL_BIN_DIR bin CACHE PATH "Installation directory for executables") +set(INSTALL_INCLUDE_DIR include CACHE PATH + "Installation directory for header files") +if(WIN32 AND NOT CYGWIN) + set(DEF_INSTALL_CMAKE_DIR cmake) +else() + set(DEF_INSTALL_CMAKE_DIR lib/cmake) +endif() +set(INSTALL_CMAKE_DIR ${DEF_INSTALL_CMAKE_DIR} CACHE PATH + "Installation directory for CMake files") + +# Make relative paths absolute (needed later on) +foreach(p LIB BIN INCLUDE CMAKE) + set(var INSTALL_${p}_DIR) + if(NOT IS_ABSOLUTE "${${var}}") + set(${var} "${CMAKE_INSTALL_PREFIX}/${${var}}") + endif() +endforeach() + if(WIN32) add_dependencies(tensorflow tensorflow_static) endif(WIN32) @@ -99,14 +120,57 @@ target_include_directories(tensorflow PUBLIC $ $) -install(TARGETS tensorflow EXPORT tensorflow_export - RUNTIME DESTINATION bin - LIBRARY DESTINATION lib - ARCHIVE DESTINATION lib) +# Add all targets to build-tree export set +export(TARGETS tensorflow + FILE ${PROJECT_BINARY_DIR}/TensorflowTargets.cmake) + +# Export the package for use from the build-tree +export(PACKAGE Tensorflow) + +# Create the TensorflowConfig.cmake and TensorflowConfigVersion files +file(RELATIVE_PATH REL_INCLUDE_DIR "${INSTALL_CMAKE_DIR}" + "${INSTALL_INCLUDE_DIR}") +# for the build tree +set(CONF_INCLUDE_DIRS "${tensorflow_source_dir}" + "${PROJECT_BINARY_DIR}" + "${CMAKE_CURRENT_BINARY_DIR}/protobuf/src/protobuf/src" + "${CMAKE_CURRENT_BINARY_DIR}/nsync/install/include" # Please if there is a better directory + "${CMAKE_CURRENT_BINARY_DIR}/eigen/src/eigen/Eigen/" + "${CMAKE_CURRENT_BINARY_DIR}/external/eigen_archive/" + "${tensorflow_source_dir}/third_party/eigen3/" + "${CMAKE_CURRENT_BINARY_DIR}/eigen/src/eigen/unsupported/Eigen/") +configure_file(TensorflowConfig.cmake.in + "${PROJECT_BINARY_DIR}/TensorflowConfig.cmake" @ONLY) +# for the install tree, yet to be complete +set(CONF_INCLUDE_DIRS "\${TENSORFLOW_CMAKE_DIR}/${REL_INCLUDE_DIR}") +configure_file(TensorflowConfig.cmake.in + "${PROJECT_BINARY_DIR}/${CMAKE_FILES_DIRECTORY}/TensorflowConfig.cmake" @ONLY) +# for both +configure_file(TensorflowConfigVersion.cmake.in + "${PROJECT_BINARY_DIR}/TensorflowConfigVersion.cmake" @ONLY) + +# install(TARGETS tensorflow EXPORT tensorflow_export +# RUNTIME DESTINATION ${INSTALL_BIN_DIR} +# LIBRARY DESTINATION ${INSTALL_LIB_DIR} +# ARCHIVE DESTINATION ${INSTALL_LIB_DIR}) + +# install(EXPORT tensorflow_export +# FILE TensorflowConfig.cmake +# DESTINATION ${INSTALL_CMAKE_DIR}) -install(EXPORT tensorflow_export - FILE TensorflowConfig.cmake - DESTINATION lib/cmake) +install(FILES + "${PROJECT_BINARY_DIR}/${CMAKE_FILES_DIRECTORY}/TensorflowConfig.cmake" + "${PROJECT_BINARY_DIR}/TensorflowConfigVersion.cmake" + DESTINATION "${INSTALL_CMAKE_DIR}" COMPONENT dev) + +# install the export set for use with the install-tree +install(EXPORT TensorflowTargets + DESTINATION ${INSTALL_CMAKE_DIR}) + +install(TARGETS tensorflow EXPORT TensorflowTargets + RUNTIME DESTINATION ${INSTALL_BIN_DIR} + LIBRARY DESTINATION ${INSTALL_LIB_DIR} + ARCHIVE DESTINATION ${INSTALL_LIB_DIR}) # install necessary headers # tensorflow headers @@ -144,4 +208,4 @@ install(DIRECTORY ${tensorflow_source_dir}/third_party/eigen3/ DESTINATION include/third_party/eigen3) # unsupported Eigen directory install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/eigen/src/eigen/unsupported/Eigen/ - DESTINATION include/unsupported/Eigen) + DESTINATION include/unsupported/Eigen) \ No newline at end of file -- GitLab From 6dd6cc95087066274eccc0205043f1c657c8a558 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sun, 28 Jan 2018 17:26:22 +0800 Subject: [PATCH 011/596] readme add cmake gui build and c++ interface --- tensorflow/contrib/cmake/README.md | 160 ++++++++++++++++++++++++++--- 1 file changed, 147 insertions(+), 13 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index 8f85a75ee4..f6497c2138 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -2,9 +2,9 @@ TensorFlow CMake build ====================== This directory contains CMake files for building TensorFlow on Microsoft -Windows. [CMake](https://cmake.org) is a cross-platform tool that can +Windows and Linux. [CMake](https://cmake.org) is a cross-platform tool that can generate build scripts for multiple build systems, including Microsoft -Visual Studio. +Visual Studio and GCC. The method suppose could work on MacOS, but not tested yet. **N.B.** We provide Linux build instructions primarily for the purpose of testing the build. We recommend using the standard Bazel-based build on @@ -13,12 +13,13 @@ Linux. Current Status -------------- -CMake can be used to build TensorFlow on Windows. See the [getting started documentation](https://www.tensorflow.org/install/install_windows) -for instructions on how to install a pre-built TensorFlow package on Windows. +CMake can be used to build TensorFlow on all platforms. See the [getting started documentation](https://www.tensorflow.org/install/install_windows) +for instructions on how to install a pre-built TensorFlow package on Windows and Linux. Procedure in MacOS is similar to the Linux build. ### Current known limitations * It is not possible to load a custom Op library. * GCS file system is not supported. +* Debug build is not available since python no longer release debug library. ## Building with CMake @@ -35,7 +36,7 @@ bindings. * [SWIG](http://www.swig.org/download.html) * Additional pre-requisites for Microsoft Windows: - - Visual Studio 2015 + - Visual Studio 2015 (latest version of MSVC 2017 is not suppored by CUDA yet, try it on your own) - Python 3.5 - NumPy 1.11.0 or later @@ -47,12 +48,12 @@ bindings. ### Known-good configurations * Microsoft Windows 10 - - Microsoft Visual Studio Enterprise 2015 with Visual C++ 2015 + - Microsoft Visual Studio Enterprise/ Community 2015 with Visual C++ 2015 - [Anaconda 4.1.1 (Python 3.5 64-bit)](https://www.anaconda.com/download/) - [Git for Windows version 2.9.2.windows.1](https://git-scm.com/download/win) - [swigwin-3.0.10](http://www.swig.org/download.html) - - [NVidia CUDA Toolkit 8.0](https://developer.nvidia.com/cuda-downloads) - - [NVidia CUDNN 5.1](https://developer.nvidia.com/cudnn) + - [NVidia CUDA Toolkit 9.0](https://developer.nvidia.com/cuda-downloads) + - [NVidia CUDNN 7](https://developer.nvidia.com/cudnn) - [CMake 3.6](https://cmake.org/files/v3.6/cmake-3.6.3-win64-x64.msi) * Ubuntu 14.04 @@ -60,8 +61,8 @@ bindings. - Docker 1.9.1 (for automated testing) ### Current known limitations - - The Python package supports **Python 3.5 only**, because that is the only - version for which standard Python binaries exist and those binaries are + - The Python package supports **Python 3.5/3.6 only**, because these are the only + versions for which standard Python binaries exist and those binaries are compatible with the TensorFlow runtime. (On Windows, the standard Python binaries for versions earlier than 3.5 were compiled with older compilers that do not have all of the features (e.g. C++11 support) needed to compile @@ -98,8 +99,141 @@ We are actively working on improving CMake and Windows support, and addressing these limitations. We would appreciate pull requests that implement missing ops or APIs. +CMake GUI build (all platforms) +================================== +Install from CMake GUI would be a convenient way to generate C++ build projects. The software supports Windows, MacOS and Linux, while the posix platform provides an extra ccmake binary to run command line GUI. Both working principal of cmake, ccmake and cmake-gui are the same, the only difference is by providing suitable interface for project configuration and dependency setting. + +0. Pre-buid checklist: + * Compiler (GCC for Linux, MSVC for Windows) + * Make sure compiler directory has been set to system path + * CUDA 9.0 (GPU build) + * CUDNN (GPU build) + * NCCL (GPU build on Linux) + * SWIG (python binding) +1. Start CMake GUI +2. Click on `Browse Source` and direct to the the folder `/tensorflow/contrib/cmake` +3. Click on `Browse Build` and spectify a location that you want tensorflow to be build +4. Click on `Configure`, a new window will be prompted out, specify the generator mode for the project generation. For Windows, choose `Visual Studio Win64`, for Linux, choose `Unix Makefiles`, then press `Finish`. Wait for a moment, the default project dependecy would automatically generate. +5. There are a few options that you can customize your own build. **The setting here is crucial for a sucessful build, please check all items carefully.** + * `tensorflow_BUILD_ALL_KERNELS` should alway be `on` + * `tensorflow_BUILD_CC_EXAMPLE` is default to be `on`. This can help you to test build (optional) + * `tensorflow_BUILD_CONTRIB_KERNELS` is default to be `on`, but it won't affect tensorflow function, turn it to `off` if you want a slim build. (optional) + * `tensorflow_BUILD_PYTHON_BINDING` is default to be `on`. Set to `off` if you don't need python interaface. If SWIG is not in system path, you need set it manually. (optional) + * `tensorflow_BUILD_SHARED_LIB` is default to be `off`. Set to `on` if you want the c++ interface. (optional) + * `tensorflow_ENABLE_GPU` is default to be `off`. Set to `on` if you want GPU support. It will search CUDA and CUDNN dependecies if you have set them to system path, otherwise CMake would prompt error and request you to set it manually. (optional) + * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows you may turn it to `off`. + * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` + * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` + * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` + * `CMAKE_INSTALL_PREFIX` is the location where the final package will be installed. You may change it to you own preferred path (optional) + + ##### Current known bug + * `WIN_CPU_SIMD_OPTIONS` cannot be turned on now. Going to be fix +6. After changing the configuration in step 5, press `Configure` again +7. If not error is found, press `Generate` + +#### Windows + +1. Open `tensorflow.sln` in the build folder (Windows). Change build type from `Debug` to `Release`. Choose `Build`->`Build Solution`. This may take more than hours of compilation. If everything is alright, the output window would show no error. + + ##### Python + In solution explorer, right click on `tf_python_build_pip_package` -> `build`. It will generate the wheel file in `/tf_python/dist`. Install with following command: + + ```pip install --upgrade tensorflow-.whl``` + + ***The wheel name varies depends on you config. Change to your own wheel filename.*** + + Remind that some pip installation requires administrator right command prompt. + + ##### C++ + You can directly use the build folder tree for C++ interface with cmake. If you want to do installation for api releasing, right click on `Install` -> `build`. The headers and library will be installed in the directory specify by `CMAKE_INSTALL_PREFIX` during configuration. + +2. For smaller RAM computer, it is noticed that out of heap space error appears. Change to command prompt build is an alternative to do step 1. + + Open `VS2015 x64 Native Tools Command Prompt`. You can open it by press `Start`, then type the binary name. Use `VS2017 x64 Native Tools Command Prompt` if you are using MSVC 2017. + + ##### Python + Directly build python wheel package by following command: + ```MSBuild /p:Configuration=Release ``` + + Remember to change `` to the actual path of the file, it can be found at the root of build directory + + Install the wheel file generated as instructed by step 1. + + ##### C++ interface + Build from VS native toolchain with following command: + ```MSBuild /p:Configuration=Release ``` + + Headers are discretely located in the build folders. Tensorflow library can be found at `/Release`, namely `tensorflow.dll` and `tensorflow.lib`. + + * Build to install for api release (optional): + ```MSBuild /p:Configuration=Release ``` + + Remember to change `` and `` to the actual path of the file, it can be found at the root of build directory. + +#### Linux/MacOS (command line GNU build) + +1. Open the terminal, change working directory to the one specified in step 3. + +2. Type the following command: + + ```make -sj all``` + + ##### Python + + ``` + make -sj tf_python_build_pip_package + cd tf_python + pip install --upgrade tensorflow-.whl + ``` + + ##### C++ interface + + ```make -sj install``` + + Where `` is the threads used for the compilation, change to any integer less or equal to your computer's maxiumum thread number. + + Headers are discretely located in the build folders. Tensorflow library can be found at ``, namely `tensorflow.so` (Linux) or `tensorflow.dylib` (MacOS). + +#### Start a Tensorflow C++ project with CMake +Here we assume that you have basic knowledge on gathering dependency with `CMakeLists.txt`. Here we introduce how the C++ api works with [official hello world tutorial](https://www.tensorflow.org/api_guides/cc/guide). + +1. Create a new working directory and create a new text file named `CMakeLists.txt` and the c++ file `main.cxx` +2. Fill in the `main.cxx` with the code provided in [official c++ api basic](https://www.tensorflow.org/api_guides/cc/guide). +3. Fill in the `CMakeLists.txt` with following code: + ``` cmake + cmake_minimum_required (VERSION 2.6) + project (tf_hello) + + # Tensorflow + find_package(Tensorflow REQUIRED) + include_directories(${TENSORFLOW_INCLUDE_DIRS}) + + # compiler setting required by tensorflow, to be tested on all compilers + # currently only tested on MSVC and GCC + if (${CMAKE_CXX_COMPILER_ID} STREQUAL MSVC) + add_definitions(-DCOMPILER_MSVC) + elseif (${CMAKE_CXX_COMPILER_ID} STREQUAL GNU) + if (${CMAKE_CXX_COMPILER_VERSION} VERSION_LESS "3") + add_definitions(-DCOMPILER_GCC3) + else() + add_definitions(-D__GNUC__) + endif() + else() + message(ERROR " compiler ${CMAKE_CXX_COMPILER_ID} not supported by this CMakeList.txt, under development") + endif() + + add_executable(tf_hello main.cxx) + target_link_libraries(tf_hello ${TENSORFLOW_LIBRARIES}) + ``` +4. Configure the folder with cmake-gui, an error should be prompted out, requesting you to locate the folder containing `TensorflowConfig.cmake`. This file can be found at `` or `` (for those have build install in previous steps). + +5. Configure again, generate the project. +6. Compile the project with `Release` config (Windows). For Linux users, just compile the project. +7. Copy the `tensorflow.dll`(Windows)/`tensorflow.so`(Linux) from build directory to the build folder containing `tf_hello` binary. +8. Run `tf_hello` binary -Step-by-step Windows build +Step-by-step Windows build (command prompt) ========================== 1. Install the pre-requisites detailed above, and set up your environment. @@ -122,7 +256,7 @@ Step-by-step Windows build For example: ``` - D:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\bin + D:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\bin D:\local\cuda\bin ``` @@ -261,4 +395,4 @@ $ cd tensorflow $ tensorflow/tools/ci_build/ci_build.sh CMAKE tensorflow/tools/ci_build/builds/cmake.sh ``` -That's it. Dependencies included. +That's it. Dependencies included. \ No newline at end of file -- GitLab From b18673c9921f3f692c4c8f4fb881769d392f656d Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sun, 28 Jan 2018 17:41:28 +0800 Subject: [PATCH 012/596] readme grpc update --- tensorflow/contrib/cmake/README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index f6497c2138..633d61a80e 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -121,7 +121,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * `tensorflow_BUILD_PYTHON_BINDING` is default to be `on`. Set to `off` if you don't need python interaface. If SWIG is not in system path, you need set it manually. (optional) * `tensorflow_BUILD_SHARED_LIB` is default to be `off`. Set to `on` if you want the c++ interface. (optional) * `tensorflow_ENABLE_GPU` is default to be `off`. Set to `on` if you want GPU support. It will search CUDA and CUDNN dependecies if you have set them to system path, otherwise CMake would prompt error and request you to set it manually. (optional) - * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows you may turn it to `off`. + * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows you need to turn it to `off`. * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` @@ -333,7 +333,7 @@ Step-by-step Windows build (command prompt) * `-Dtensorflow_ENABLE_GRPC_SUPPORT=(ON|OFF)`. Defaults to `ON`. Include gRPC support and the distributed client and server code in the TensorFlow - runtime. + runtime. * `-Dtensorflow_ENABLE_SSL_SUPPORT=(ON|OFF)`. Defaults to `OFF`. Include SSL support (for making secure HTTP requests) in the TensorFlow runtime. -- GitLab From 51b8978fe2697881d004d736460fde5ef0eb8a2c Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sun, 28 Jan 2018 20:46:58 +0800 Subject: [PATCH 013/596] fix for cmake build --- tensorflow/core/util/cuda_device_functions.h | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/tensorflow/core/util/cuda_device_functions.h b/tensorflow/core/util/cuda_device_functions.h index f787687f66..a994398c42 100644 --- a/tensorflow/core/util/cuda_device_functions.h +++ b/tensorflow/core/util/cuda_device_functions.h @@ -28,12 +28,21 @@ limitations under the License. #include #include +#ifdef GOOGLE_PLATFORM #include "cuda/include/cuda.h" #include "cuda/include/device_functions.h" +#else +#include "cuda.h" +#include "device_functions.h" +#endif #include "tensorflow/core/platform/types.h" #if CUDA_VERSION >= 7050 +#ifdef GOOGLE_PLATFORM #include "cuda/include/cuda_fp16.h" +#else +#include "cuda_fp16.h" +#endif #endif // CUDA_VERSION >= 7050 namespace tensorflow { -- GitLab From f6911068262751adba24a26344fed8394d042a07 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sun, 28 Jan 2018 20:55:20 +0800 Subject: [PATCH 014/596] cmake cuda fix --- tensorflow/core/util/cuda_device_functions.h | 4 ++-- tensorflow/core/util/cuda_launch_config.h | 4 ++++ 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/util/cuda_device_functions.h b/tensorflow/core/util/cuda_device_functions.h index a994398c42..66129d2425 100644 --- a/tensorflow/core/util/cuda_device_functions.h +++ b/tensorflow/core/util/cuda_device_functions.h @@ -28,7 +28,7 @@ limitations under the License. #include #include -#ifdef GOOGLE_PLATFORM +#ifdef PLATFORM_GOOGLE #include "cuda/include/cuda.h" #include "cuda/include/device_functions.h" #else @@ -38,7 +38,7 @@ limitations under the License. #include "tensorflow/core/platform/types.h" #if CUDA_VERSION >= 7050 -#ifdef GOOGLE_PLATFORM +#ifdef PLATFORM_GOOGLE #include "cuda/include/cuda_fp16.h" #else #include "cuda_fp16.h" diff --git a/tensorflow/core/util/cuda_launch_config.h b/tensorflow/core/util/cuda_launch_config.h index 3ea33ee6cf..7e552c7dce 100644 --- a/tensorflow/core/util/cuda_launch_config.h +++ b/tensorflow/core/util/cuda_launch_config.h @@ -21,7 +21,11 @@ limitations under the License. #include #include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" +#ifdef PLATFORM_GOOGLE #include "cuda/include/cuda.h" +#else +#include "cuda.h" +#endif #include "tensorflow/core/framework/op_kernel.h" #include "tensorflow/core/platform/logging.h" #include "tensorflow/core/platform/stream_executor.h" -- GitLab From 0e5faafaa3977584efd7ab6aab114f8c0689c247 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Mon, 29 Jan 2018 00:24:09 +0800 Subject: [PATCH 015/596] update protobuf to 3.5.1.1 --- tensorflow/contrib/cmake/external/protobuf.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/external/protobuf.cmake b/tensorflow/contrib/cmake/external/protobuf.cmake index aedb793d2a..31a13b980e 100644 --- a/tensorflow/contrib/cmake/external/protobuf.cmake +++ b/tensorflow/contrib/cmake/external/protobuf.cmake @@ -16,7 +16,7 @@ include (ExternalProject) set(PROTOBUF_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/protobuf/src/protobuf/src) set(PROTOBUF_URL https://github.com/google/protobuf.git) -set(PROTOBUF_TAG b04e5cba356212e4e8c66c61bbe0c3a20537c5b9) +set(PROTOBUF_TAG 860bd12fec5c69e6529565165532b3d5108a7d97) if(WIN32) set(protobuf_STATIC_LIBRARIES -- GitLab From cf101a9cd64dc9f101a27c45ae2a52e4630b6fb7 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Mon, 29 Jan 2018 00:50:21 +0800 Subject: [PATCH 016/596] windows SMID option imporve --- tensorflow/contrib/cmake/CMakeLists.txt | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index c26c7cd3b7..a7e6e8262d 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -31,10 +31,15 @@ option(tensorflow_BUILD_PYTHON_TESTS "Build python unit tests " OFF) option(tensorflow_BUILD_MORE_PYTHON_TESTS "Build more python unit tests for contrib packages" OFF) option(tensorflow_BUILD_SHARED_LIB "Build TensorFlow as a shared library" OFF) option(tensorflow_OPTIMIZE_FOR_NATIVE_ARCH "Enable compiler optimizations for the native processor architecture (if available)" ON) -option(tensorflow_WIN_CPU_SIMD_OPTIONS "Enables CPU SIMD instructions") option(tensorflow_ENABLE_SNAPPY_SUPPORT "Enable SNAPPY compression support" ON) option(tensorflow_DISABLE_EIGEN_FORCEINLINE "Disable forceinline, to speed up build on windows." OFF) +# enable SIMD instructions with MSVC, as AVX and SSE +if (WIN32) +SET(tensorflow_WIN_CPU_SIMD_OPTIONS "/arch:AVX" CACHE STRING "Enables CPU SIMD instructions") +SET_PROPERTY(CACHE tensorflow_WIN_CPU_SIMD_OPTIONS PROPERTY STRINGS /arch:AVX) +endif() + # GPU, CUDA and cuDNN options option(tensorflow_ENABLE_GPU "Enable GPU support" OFF) set(tensorflow_CUDA_VERSION "9.0" CACHE STRING "CUDA version to build against") -- GitLab From 9b8ecce73d4ad7d5753ecabe9c511566b31f120a Mon Sep 17 00:00:00 2001 From: JackyKo Date: Mon, 29 Jan 2018 00:52:45 +0800 Subject: [PATCH 017/596] GRPC updated --- tensorflow/contrib/cmake/README.md | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index 633d61a80e..dc75250265 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -121,14 +121,12 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * `tensorflow_BUILD_PYTHON_BINDING` is default to be `on`. Set to `off` if you don't need python interaface. If SWIG is not in system path, you need set it manually. (optional) * `tensorflow_BUILD_SHARED_LIB` is default to be `off`. Set to `on` if you want the c++ interface. (optional) * `tensorflow_ENABLE_GPU` is default to be `off`. Set to `on` if you want GPU support. It will search CUDA and CUDNN dependecies if you have set them to system path, otherwise CMake would prompt error and request you to set it manually. (optional) - * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows you need to turn it to `off`. + * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows this need to be `on` for gpu build. * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` * `CMAKE_INSTALL_PREFIX` is the location where the final package will be installed. You may change it to you own preferred path (optional) - ##### Current known bug - * `WIN_CPU_SIMD_OPTIONS` cannot be turned on now. Going to be fix 6. After changing the configuration in step 5, press `Configure` again 7. If not error is found, press `Generate` -- GitLab From 9284b3cc7782dd7d7f9459f7beec7cddec1c3f9a Mon Sep 17 00:00:00 2001 From: JackyKo Date: Mon, 29 Jan 2018 01:21:32 +0800 Subject: [PATCH 018/596] python build in linux failed --- tensorflow/contrib/cmake/README.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index dc75250265..a18a0c8e37 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -179,6 +179,10 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. ##### Python + **Important Note** CMake generated python wheel for Linux/MacOs is currently under development. Please use bazel build. + + Follow code is an expected Linux/MacOS python package build after development work is completed. + ``` make -sj tf_python_build_pip_package cd tf_python -- GitLab From 4219092017be9523cb2be7106111b1a4a1016d0c Mon Sep 17 00:00:00 2001 From: jackyko Date: Mon, 29 Jan 2018 11:33:04 +0800 Subject: [PATCH 019/596] grpc require compiler update --- tensorflow/contrib/cmake/README.md | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index a18a0c8e37..ba5f4e58c5 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -35,6 +35,12 @@ bindings. * [SWIG](http://www.swig.org/download.html) +* [Perl](https://www.perl.org/get.html) + +* [Go](https://golang.org/) + +* [NASM](http://www.nasm.us/) + * Additional pre-requisites for Microsoft Windows: - Visual Studio 2015 (latest version of MSVC 2017 is not suppored by CUDA yet, try it on your own) - Python 3.5 @@ -104,12 +110,16 @@ CMake GUI build (all platforms) Install from CMake GUI would be a convenient way to generate C++ build projects. The software supports Windows, MacOS and Linux, while the posix platform provides an extra ccmake binary to run command line GUI. Both working principal of cmake, ccmake and cmake-gui are the same, the only difference is by providing suitable interface for project configuration and dependency setting. 0. Pre-buid checklist: + The following binary/libraries should be setted in system path, otherwise you need to set manualy via cmake. * Compiler (GCC for Linux, MSVC for Windows) * Make sure compiler directory has been set to system path * CUDA 9.0 (GPU build) * CUDNN (GPU build) * NCCL (GPU build on Linux) * SWIG (python binding) + * Perl (GPU build on Windows, required by grpc) + * Go (GPU build on Windows, required by grpc) + * NASM (GPU build on Windows, required by grpc) 1. Start CMake GUI 2. Click on `Browse Source` and direct to the the folder `/tensorflow/contrib/cmake` 3. Click on `Browse Build` and spectify a location that you want tensorflow to be build @@ -121,7 +131,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * `tensorflow_BUILD_PYTHON_BINDING` is default to be `on`. Set to `off` if you don't need python interaface. If SWIG is not in system path, you need set it manually. (optional) * `tensorflow_BUILD_SHARED_LIB` is default to be `off`. Set to `on` if you want the c++ interface. (optional) * `tensorflow_ENABLE_GPU` is default to be `off`. Set to `on` if you want GPU support. It will search CUDA and CUDNN dependecies if you have set them to system path, otherwise CMake would prompt error and request you to set it manually. (optional) - * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows this need to be `on` for gpu build. + * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows this need to be `on` for gpu build. Reminded that Perl, Go and NASM is required for this option in windows. * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` @@ -135,6 +145,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. 1. Open `tensorflow.sln` in the build folder (Windows). Change build type from `Debug` to `Release`. Choose `Build`->`Build Solution`. This may take more than hours of compilation. If everything is alright, the output window would show no error. ##### Python + In solution explorer, right click on `tf_python_build_pip_package` -> `build`. It will generate the wheel file in `/tf_python/dist`. Install with following command: ```pip install --upgrade tensorflow-.whl``` @@ -144,6 +155,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. Remind that some pip installation requires administrator right command prompt. ##### C++ + You can directly use the build folder tree for C++ interface with cmake. If you want to do installation for api releasing, right click on `Install` -> `build`. The headers and library will be installed in the directory specify by `CMAKE_INSTALL_PREFIX` during configuration. 2. For smaller RAM computer, it is noticed that out of heap space error appears. Change to command prompt build is an alternative to do step 1. @@ -151,7 +163,9 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. Open `VS2015 x64 Native Tools Command Prompt`. You can open it by press `Start`, then type the binary name. Use `VS2017 x64 Native Tools Command Prompt` if you are using MSVC 2017. ##### Python + Directly build python wheel package by following command: + ```MSBuild /p:Configuration=Release ``` Remember to change `` to the actual path of the file, it can be found at the root of build directory -- GitLab From 4dd3d50cf8f1dcec7f4a4268b50bcea7c395a9a2 Mon Sep 17 00:00:00 2001 From: jackyko Date: Mon, 29 Jan 2018 13:34:04 +0800 Subject: [PATCH 020/596] add perl, NASM and Go request --- tensorflow/contrib/cmake/external/grpc.cmake | 25 ++++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 28b85e0a19..03205f58ca 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -14,6 +14,29 @@ # ============================================================================== include (ExternalProject) +if (WIN32) + # perl + find_package(Perl REQUIRED) + + # nasm + if(NOT NASM_COMPILER) + find_program(NASM_COMPILER nasm + "$ENV{ProgramFiles}/NASM" DOC "path to NASM.exe") + endif() + if(NOT NASM_COMPILER) + message(FATAL_ERROR "NASM not found!" DOC "path to go.exe") + endif() + + # go + if(NOT GO) + find_program(GO go + "$ENV{ProgramFiles}/go") + endif() + if(NOT GO) + message(FATAL_ERROR "Go not found!") + endif() +endif() + set(GRPC_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/include) set(GRPC_URL https://github.com/grpc/grpc.git) set(GRPC_BUILD ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc) @@ -51,6 +74,8 @@ ExternalProject_Add(grpc -DPROTOBUF_INCLUDE_DIRS:STRING=${PROTOBUF_INCLUDE_DIRS} -DPROTOBUF_LIBRARIES:STRING=${protobuf_STATIC_LIBRARIES} -DZLIB_ROOT:STRING=${ZLIB_INSTALL} + -DCMAKE_ASM_NASM_COMPILER:STRING=${NASM_COMPILER} + -DGO_EXECUTABLE:STRING=${GO} ) # grpc/src/core/ext/census/tracing.c depends on the existence of openssl/rand.h. -- GitLab From 6b6d24015f16bd150571920c013b5fd6f3d07f4b Mon Sep 17 00:00:00 2001 From: jackyko Date: Mon, 29 Jan 2018 13:41:33 +0800 Subject: [PATCH 021/596] search path for go update --- tensorflow/contrib/cmake/external/grpc.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 03205f58ca..944c49e48a 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -30,7 +30,7 @@ if (WIN32) # go if(NOT GO) find_program(GO go - "$ENV{ProgramFiles}/go") + "$ENV{ProgramFiles}/Go/bin") endif() if(NOT GO) message(FATAL_ERROR "Go not found!") -- GitLab From 4ea80c7e96ab059360faec42193ec67ef24fdbec Mon Sep 17 00:00:00 2001 From: jackyko Date: Mon, 29 Jan 2018 16:43:10 +0800 Subject: [PATCH 022/596] version update to 1.5.0 --- tensorflow/tools/pip_package/setup.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/tools/pip_package/setup.py b/tensorflow/tools/pip_package/setup.py index 62df6453fb..306bfab161 100644 --- a/tensorflow/tools/pip_package/setup.py +++ b/tensorflow/tools/pip_package/setup.py @@ -29,7 +29,7 @@ from setuptools.dist import Distribution # This version string is semver compatible, but incompatible with pip. # For pip, we will remove all '-' characters from this string, and use the # result for pip. -_VERSION = '1.5.0-rc1' +_VERSION = '1.5.0' REQUIRED_PACKAGES = [ 'absl-py >= 0.1.6', -- GitLab From 8e337006a72f6316a4afa7f1173f819b6baca1ba Mon Sep 17 00:00:00 2001 From: jackyko Date: Mon, 29 Jan 2018 16:45:50 +0800 Subject: [PATCH 023/596] install package add abseil for grpc support --- tensorflow/tools/pip_package/setup.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tensorflow/tools/pip_package/setup.py b/tensorflow/tools/pip_package/setup.py index 306bfab161..083d002c83 100644 --- a/tensorflow/tools/pip_package/setup.py +++ b/tensorflow/tools/pip_package/setup.py @@ -37,9 +37,10 @@ REQUIRED_PACKAGES = [ 'gast >= 0.2.0', 'numpy >= 1.12.1', 'six >= 1.10.0', - 'protobuf >= 3.4.0', + 'protobuf >= 3.5.0', 'tensorflow-tensorboard >= 0.4.0', 'termcolor >= 1.1.0', + 'absl-py >= 0.1.9' ] project_name = 'tensorflow' -- GitLab From 4a35d10279b8b28d56a83a5ea05879d2cac16bb5 Mon Sep 17 00:00:00 2001 From: jackyko Date: Mon, 29 Jan 2018 17:38:17 +0800 Subject: [PATCH 024/596] linux gpu build fix --- tensorflow/contrib/cmake/CMakeLists.txt | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index a7e6e8262d..3ec630b3d6 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -401,8 +401,12 @@ if (tensorflow_ENABLE_GPU) cudnn_version_number=${tensorflow_CUDNN_VERSION}) endif(WIN32) else(tensorflow_ENABLE_GPU) - set(tensorflow_BUILD_INFO_FLAGS --build_config cpu --key_value - msvcp_dll_name=msvcp140.dll) + if(WIN32) + set(tensorflow_BUILD_INFO_FLAGS --build_config cpu --key_value + msvcp_dll_name=msvcp140.dll) + else() + set(tensorflow_BUILD_INFO_FLAGS --build_config cpu) + endif() endif(tensorflow_ENABLE_GPU) if(tensorflow_BUILD_PYTHON_BINDINGS) -- GitLab From 547731fe8ffcc994870dda4d70eeaf8a0686427c Mon Sep 17 00:00:00 2001 From: jackyko Date: Tue, 30 Jan 2018 11:16:05 +0800 Subject: [PATCH 025/596] add linux support --- tensorflow/contrib/cmake/external/grpc.cmake | 55 +++++++++++++------- 1 file changed, 36 insertions(+), 19 deletions(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 944c49e48a..f3119aa749 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -14,28 +14,45 @@ # ============================================================================== include (ExternalProject) -if (WIN32) - # perl - find_package(Perl REQUIRED) +# perl +find_package(Perl REQUIRED) - # nasm - if(NOT NASM_COMPILER) - find_program(NASM_COMPILER nasm - "$ENV{ProgramFiles}/NASM" DOC "path to NASM.exe") - endif() - if(NOT NASM_COMPILER) - message(FATAL_ERROR "NASM not found!" DOC "path to go.exe") +# nasm +if(NOT CMAKE_ASM_NASM_COMPILER) + if (WIN32) + find_program(CMAKE_ASM_NASM_COMPILER + NAMES nasm + PATHS "$ENV{ProgramFiles}/NASM" + DOC "path to NASM compiler") + else() + find_program(CMAKE_ASM_NASM_COMPILER + NAMES nasm + PATHS "/usr/bin" + DOC "path to NASM compiler") endif() +endif() + +if(NOT NASM_COMPILER) + message(FATAL_ERROR "NASM not found!") +endif() - # go - if(NOT GO) - find_program(GO go - "$ENV{ProgramFiles}/Go/bin") - endif() - if(NOT GO) - message(FATAL_ERROR "Go not found!") +# go +if(NOT GO_EXECUTABLE) + if (WIN32) + find_program(GO_EXECUTABLE + NAMES go + PATHS "$ENV{ProgramFiles}/Go/bin" + DOC "path to go compiler") + else() + find_program(GO_EXECUTABLE + NAMES go + PATHS "/usr/bin" "/usr/local/bin" "/usr/local/go" + DOC "path to go compiler") endif() endif() +if(NOT GO_EXECUTABLE) + message(FATAL_ERROR "Go not found!") +endif() set(GRPC_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/include) set(GRPC_URL https://github.com/grpc/grpc.git) @@ -74,8 +91,8 @@ ExternalProject_Add(grpc -DPROTOBUF_INCLUDE_DIRS:STRING=${PROTOBUF_INCLUDE_DIRS} -DPROTOBUF_LIBRARIES:STRING=${protobuf_STATIC_LIBRARIES} -DZLIB_ROOT:STRING=${ZLIB_INSTALL} - -DCMAKE_ASM_NASM_COMPILER:STRING=${NASM_COMPILER} - -DGO_EXECUTABLE:STRING=${GO} + -DCMAKE_ASM_NASM_COMPILER:STRING=${CMAKE_ASM_NASM_COMPILER} + -DGO_EXECUTABLE:STRING=${GO_EXECUTABLE} ) # grpc/src/core/ext/census/tracing.c depends on the existence of openssl/rand.h. -- GitLab From 70127d3891ac0dcf291727eee91dfca6266bd6ee Mon Sep 17 00:00:00 2001 From: jackyko Date: Tue, 30 Jan 2018 11:17:51 +0800 Subject: [PATCH 026/596] nasm cmake fix --- tensorflow/contrib/cmake/external/grpc.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index f3119aa749..4a2989a945 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -32,7 +32,7 @@ if(NOT CMAKE_ASM_NASM_COMPILER) endif() endif() -if(NOT NASM_COMPILER) +if(NOT CMAKE_ASM_NASM_COMPILER) message(FATAL_ERROR "NASM not found!") endif() -- GitLab From eb77281b907075ef7cc3699f9222b839527a3b6b Mon Sep 17 00:00:00 2001 From: jackyko Date: Tue, 30 Jan 2018 14:03:09 +0800 Subject: [PATCH 027/596] change to version 3.5.1 --- tensorflow/contrib/cmake/external/protobuf.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/external/protobuf.cmake b/tensorflow/contrib/cmake/external/protobuf.cmake index 31a13b980e..269f1c5c89 100644 --- a/tensorflow/contrib/cmake/external/protobuf.cmake +++ b/tensorflow/contrib/cmake/external/protobuf.cmake @@ -16,7 +16,7 @@ include (ExternalProject) set(PROTOBUF_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/protobuf/src/protobuf/src) set(PROTOBUF_URL https://github.com/google/protobuf.git) -set(PROTOBUF_TAG 860bd12fec5c69e6529565165532b3d5108a7d97) +set(PROTOBUF_TAG 106ffc04be1abf3ff3399f54ccf149815b287dd9) if(WIN32) set(protobuf_STATIC_LIBRARIES -- GitLab From ae797531d6ea09b6fe8a30cada74be5dffbbac12 Mon Sep 17 00:00:00 2001 From: jackyko Date: Tue, 30 Jan 2018 14:05:37 +0800 Subject: [PATCH 028/596] change to 3.5.0 --- tensorflow/contrib/cmake/external/protobuf.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/external/protobuf.cmake b/tensorflow/contrib/cmake/external/protobuf.cmake index 269f1c5c89..d6fdb37406 100644 --- a/tensorflow/contrib/cmake/external/protobuf.cmake +++ b/tensorflow/contrib/cmake/external/protobuf.cmake @@ -16,7 +16,7 @@ include (ExternalProject) set(PROTOBUF_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/protobuf/src/protobuf/src) set(PROTOBUF_URL https://github.com/google/protobuf.git) -set(PROTOBUF_TAG 106ffc04be1abf3ff3399f54ccf149815b287dd9) +set(PROTOBUF_TAG 2761122b810fe8861004ae785cc3ab39f384d342) if(WIN32) set(protobuf_STATIC_LIBRARIES -- GitLab From 81f3f47a55aa70b12dccdaffafe7bf1f53f42818 Mon Sep 17 00:00:00 2001 From: jackyko Date: Tue, 30 Jan 2018 18:48:58 +0800 Subject: [PATCH 029/596] grpc build according to protobuf version --- tensorflow/contrib/cmake/external/grpc.cmake | 70 ++++++++++--------- .../contrib/cmake/external/protobuf.cmake | 11 ++- .../ci_build/windows/cpu/cmake/run_build.bat | 3 +- .../ci_build/windows/gpu/cmake/run_build.bat | 3 +- 4 files changed, 50 insertions(+), 37 deletions(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 4a2989a945..0de0433619 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -17,41 +17,43 @@ include (ExternalProject) # perl find_package(Perl REQUIRED) -# nasm -if(NOT CMAKE_ASM_NASM_COMPILER) - if (WIN32) - find_program(CMAKE_ASM_NASM_COMPILER - NAMES nasm - PATHS "$ENV{ProgramFiles}/NASM" - DOC "path to NASM compiler") - else() - find_program(CMAKE_ASM_NASM_COMPILER - NAMES nasm - PATHS "/usr/bin" - DOC "path to NASM compiler") - endif() -endif() - -if(NOT CMAKE_ASM_NASM_COMPILER) - message(FATAL_ERROR "NASM not found!") -endif() +if(${PROTOBUF_VERSION} STREQUAL "3.5.0") + # nasm + if(NOT CMAKE_ASM_NASM_COMPILER) + if (WIN32) + find_program(CMAKE_ASM_NASM_COMPILER + NAMES nasm + PATHS "$ENV{ProgramFiles}/NASM" + DOC "path to NASM compiler") + else() + find_program(CMAKE_ASM_NASM_COMPILER + NAMES nasm + PATHS "/usr/bin" + DOC "path to NASM compiler") + endif() + endif() -# go -if(NOT GO_EXECUTABLE) - if (WIN32) - find_program(GO_EXECUTABLE - NAMES go - PATHS "$ENV{ProgramFiles}/Go/bin" - DOC "path to go compiler") - else() - find_program(GO_EXECUTABLE - NAMES go - PATHS "/usr/bin" "/usr/local/bin" "/usr/local/go" - DOC "path to go compiler") - endif() -endif() -if(NOT GO_EXECUTABLE) - message(FATAL_ERROR "Go not found!") + if(NOT CMAKE_ASM_NASM_COMPILER) + message(FATAL_ERROR "NASM not found!") + endif() + + # go + if(NOT GO_EXECUTABLE) + if (WIN32) + find_program(GO_EXECUTABLE + NAMES go + PATHS "$ENV{ProgramFiles}/Go/bin" + DOC "path to go compiler") + else() + find_program(GO_EXECUTABLE + NAMES go + PATHS "/usr/bin" "/usr/local/bin" "/usr/local/go" + DOC "path to go compiler") + endif() + endif() + if(NOT GO_EXECUTABLE) + message(FATAL_ERROR "Go not found!") + endif() endif() set(GRPC_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/include) diff --git a/tensorflow/contrib/cmake/external/protobuf.cmake b/tensorflow/contrib/cmake/external/protobuf.cmake index d6fdb37406..785ed4c739 100644 --- a/tensorflow/contrib/cmake/external/protobuf.cmake +++ b/tensorflow/contrib/cmake/external/protobuf.cmake @@ -16,7 +16,16 @@ include (ExternalProject) set(PROTOBUF_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/protobuf/src/protobuf/src) set(PROTOBUF_URL https://github.com/google/protobuf.git) -set(PROTOBUF_TAG 2761122b810fe8861004ae785cc3ab39f384d342) + +# enable choose protobuf versions +SET(PROTOBUF_VERSION "3.5.0" CACHE STRING "Protobuf version") +SET_PROPERTY(CACHE PROTOBUF_VERSION PROPERTY STRINGS "3.4.0" "3.5.0") + +if(${PROTOBUF_VERSION} STREQUAL "3.5.0") + set(PROTOBUF_TAG 2761122b810fe8861004ae785cc3ab39f384d342) +else() + set(PROTOBUF_TAG b04e5cba356212e4e8c66c61bbe0c3a20537c5b9) +endif() if(WIN32) set(protobuf_STATIC_LIBRARIES diff --git a/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat b/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat index 957729bb37..11eb9f9771 100644 --- a/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat +++ b/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat @@ -34,9 +34,10 @@ IF DEFINED DISABLE_FORCEINLINE (ECHO DISABLE_FORCEINLINE is set to %DISABLE_FORC SET CMAKE_DIR=%REPO_ROOT%\tensorflow\contrib\cmake SET MSBUILD_EXE="C:\Program Files (x86)\MSBuild\14.0\Bin\msbuild.exe" +SET PROTOBUF_VERSION="3.4.0" :: Run cmake to create Visual Studio Project files. -%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% +%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% -DPROTOBUF_VERSION=%PROTOBUF_VERSION% :: Run msbuild in the resulting VS project files to build a pip package. %MSBUILD_EXE% /p:Configuration=Release /maxcpucount:32 tf_python_build_pip_package.vcxproj diff --git a/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat b/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat index 5a362de399..df70cae55d 100644 --- a/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat +++ b/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat @@ -35,9 +35,10 @@ IF DEFINED DISABLE_FORCEINLINE (ECHO DISABLE_FORCEINLINE is set to %DISABLE_FORC SET CMAKE_DIR=%REPO_ROOT%\tensorflow\contrib\cmake SET MSBUILD_EXE="C:\Program Files (x86)\MSBuild\14.0\Bin\msbuild.exe" +SET PROTOBUF_VERSION="3.4.0" :: Run cmake to create Visual Studio Project files. -%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_ENABLE_GPU=ON -DCUDNN_HOME=%CUDNN_HOME% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% +%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_ENABLE_GPU=ON -DCUDNN_HOME=%CUDNN_HOME% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% -DPROTOBUF_VERSION=%PROTOBUF_VERSION% :: Run msbuild in the resulting VS project files to build a pip package. %MSBUILD_EXE% /p:Configuration=Release /maxcpucount:32 tf_python_build_pip_package.vcxproj -- GitLab From 946e0772432acc29663a5a4bf5738d0f429b6622 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Tue, 30 Jan 2018 23:00:36 +0800 Subject: [PATCH 030/596] ssl change back to none --- tensorflow/contrib/cmake/external/grpc.cmake | 79 ++++++++++---------- 1 file changed, 39 insertions(+), 40 deletions(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 0de0433619..cd866f74c9 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -14,47 +14,45 @@ # ============================================================================== include (ExternalProject) -# perl -find_package(Perl REQUIRED) +# # perl +# find_package(Perl REQUIRED) -if(${PROTOBUF_VERSION} STREQUAL "3.5.0") - # nasm - if(NOT CMAKE_ASM_NASM_COMPILER) - if (WIN32) - find_program(CMAKE_ASM_NASM_COMPILER - NAMES nasm - PATHS "$ENV{ProgramFiles}/NASM" - DOC "path to NASM compiler") - else() - find_program(CMAKE_ASM_NASM_COMPILER - NAMES nasm - PATHS "/usr/bin" - DOC "path to NASM compiler") - endif() - endif() +# # nasm +# if(NOT CMAKE_ASM_NASM_COMPILER) +# if (WIN32) +# find_program(CMAKE_ASM_NASM_COMPILER +# NAMES nasm +# PATHS "$ENV{ProgramFiles}/NASM" +# DOC "path to NASM compiler") +# else() +# find_program(CMAKE_ASM_NASM_COMPILER +# NAMES nasm +# PATHS "/usr/bin" +# DOC "path to NASM compiler") +# endif() +# endif() - if(NOT CMAKE_ASM_NASM_COMPILER) - message(FATAL_ERROR "NASM not found!") - endif() +# if(NOT CMAKE_ASM_NASM_COMPILER) +# message(FATAL_ERROR "NASM not found!") +# endif() - # go - if(NOT GO_EXECUTABLE) - if (WIN32) - find_program(GO_EXECUTABLE - NAMES go - PATHS "$ENV{ProgramFiles}/Go/bin" - DOC "path to go compiler") - else() - find_program(GO_EXECUTABLE - NAMES go - PATHS "/usr/bin" "/usr/local/bin" "/usr/local/go" - DOC "path to go compiler") - endif() - endif() - if(NOT GO_EXECUTABLE) - message(FATAL_ERROR "Go not found!") - endif() -endif() +# # go +# if(NOT GO_EXECUTABLE) +# if (WIN32) +# find_program(GO_EXECUTABLE +# NAMES go +# PATHS "$ENV{ProgramFiles}/Go/bin" +# DOC "path to go compiler") +# else() +# find_program(GO_EXECUTABLE +# NAMES go +# PATHS "/usr/bin" "/usr/local/bin" "/usr/local/go" +# DOC "path to go compiler") +# endif() +# endif() +# if(NOT GO_EXECUTABLE) +# message(FATAL_ERROR "Go not found!") +# endif() set(GRPC_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/include) set(GRPC_URL https://github.com/grpc/grpc.git) @@ -93,8 +91,9 @@ ExternalProject_Add(grpc -DPROTOBUF_INCLUDE_DIRS:STRING=${PROTOBUF_INCLUDE_DIRS} -DPROTOBUF_LIBRARIES:STRING=${protobuf_STATIC_LIBRARIES} -DZLIB_ROOT:STRING=${ZLIB_INSTALL} - -DCMAKE_ASM_NASM_COMPILER:STRING=${CMAKE_ASM_NASM_COMPILER} - -DGO_EXECUTABLE:STRING=${GO_EXECUTABLE} + -DgRPC_SSL_PROVIDER:STRING=NONE + # -DCMAKE_ASM_NASM_COMPILER:STRING=${CMAKE_ASM_NASM_COMPILER} + # -DGO_EXECUTABLE:STRING=${GO_EXECUTABLE} ) # grpc/src/core/ext/census/tracing.c depends on the existence of openssl/rand.h. -- GitLab From 2cd3a0edc2b9119f3434105ace28411731d403bc Mon Sep 17 00:00:00 2001 From: JackyKo Date: Tue, 30 Jan 2018 23:10:18 +0800 Subject: [PATCH 031/596] readme update --- tensorflow/contrib/cmake/README.md | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index ba5f4e58c5..da72ea320f 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -35,11 +35,11 @@ bindings. * [SWIG](http://www.swig.org/download.html) -* [Perl](https://www.perl.org/get.html) +* [Perl](https://www.perl.org/get.html) (optional, for full grpc build) -* [Go](https://golang.org/) +* [Go](https://golang.org/) (optional, for full grpc build) -* [NASM](http://www.nasm.us/) +* [NASM](http://www.nasm.us/)/[yasm](http://yasm.tortall.net/) (optional, for full grpc build) * Additional pre-requisites for Microsoft Windows: - Visual Studio 2015 (latest version of MSVC 2017 is not suppored by CUDA yet, try it on your own) @@ -117,9 +117,9 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * CUDNN (GPU build) * NCCL (GPU build on Linux) * SWIG (python binding) - * Perl (GPU build on Windows, required by grpc) - * Go (GPU build on Windows, required by grpc) - * NASM (GPU build on Windows, required by grpc) + * Perl (GPU build on Windows, required by full function grpc, optional) + * Go (GPU build on Windows, required by full function grpc, optional) + * NASM (GPU build on Windows, required by full function grpc, optional) 1. Start CMake GUI 2. Click on `Browse Source` and direct to the the folder `/tensorflow/contrib/cmake` 3. Click on `Browse Build` and spectify a location that you want tensorflow to be build @@ -131,7 +131,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * `tensorflow_BUILD_PYTHON_BINDING` is default to be `on`. Set to `off` if you don't need python interaface. If SWIG is not in system path, you need set it manually. (optional) * `tensorflow_BUILD_SHARED_LIB` is default to be `off`. Set to `on` if you want the c++ interface. (optional) * `tensorflow_ENABLE_GPU` is default to be `off`. Set to `on` if you want GPU support. It will search CUDA and CUDNN dependecies if you have set them to system path, otherwise CMake would prompt error and request you to set it manually. (optional) - * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. In Windows this need to be `on` for gpu build. Reminded that Perl, Go and NASM is required for this option in windows. + * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. This need to be `on` for a gpu build. Reminded that Perl, Go and NASM is required for this option if you want to build grpc with offical SSL support. * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` -- GitLab From f3335e8b1e69d641162f4873c0f532778946f147 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Tue, 30 Jan 2018 23:15:40 +0800 Subject: [PATCH 032/596] Update run_build.bat --- tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat b/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat index 11eb9f9771..957729bb37 100644 --- a/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat +++ b/tensorflow/tools/ci_build/windows/cpu/cmake/run_build.bat @@ -34,10 +34,9 @@ IF DEFINED DISABLE_FORCEINLINE (ECHO DISABLE_FORCEINLINE is set to %DISABLE_FORC SET CMAKE_DIR=%REPO_ROOT%\tensorflow\contrib\cmake SET MSBUILD_EXE="C:\Program Files (x86)\MSBuild\14.0\Bin\msbuild.exe" -SET PROTOBUF_VERSION="3.4.0" :: Run cmake to create Visual Studio Project files. -%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% -DPROTOBUF_VERSION=%PROTOBUF_VERSION% +%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% :: Run msbuild in the resulting VS project files to build a pip package. %MSBUILD_EXE% /p:Configuration=Release /maxcpucount:32 tf_python_build_pip_package.vcxproj -- GitLab From 3d5ffc90cd3c1bb750046b3b0bc3c85e4f317e33 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Tue, 30 Jan 2018 23:16:38 +0800 Subject: [PATCH 033/596] Update run_build.bat --- tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat b/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat index df70cae55d..5a362de399 100644 --- a/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat +++ b/tensorflow/tools/ci_build/windows/gpu/cmake/run_build.bat @@ -35,10 +35,9 @@ IF DEFINED DISABLE_FORCEINLINE (ECHO DISABLE_FORCEINLINE is set to %DISABLE_FORC SET CMAKE_DIR=%REPO_ROOT%\tensorflow\contrib\cmake SET MSBUILD_EXE="C:\Program Files (x86)\MSBuild\14.0\Bin\msbuild.exe" -SET PROTOBUF_VERSION="3.4.0" :: Run cmake to create Visual Studio Project files. -%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_ENABLE_GPU=ON -DCUDNN_HOME=%CUDNN_HOME% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% -DPROTOBUF_VERSION=%PROTOBUF_VERSION% +%CMAKE_EXE% %CMAKE_DIR% -A x64 -DSWIG_EXECUTABLE=%SWIG_EXE% -DPYTHON_EXECUTABLE=%PY_EXE% -DCMAKE_BUILD_TYPE=Release -DPYTHON_LIBRARIES=%PY_LIB% -Dtensorflow_BUILD_PYTHON_TESTS=%BUILD_PYTHON_TESTS% -Dtensorflow_BUILD_CC_TESTS=%BUILD_CC_TESTS% -Dtensorflow_ENABLE_GPU=ON -DCUDNN_HOME=%CUDNN_HOME% -Dtensorflow_TF_NIGHTLY=%TF_NIGHTLY% -Dtensorflow_DISABLE_EIGEN_FORCEINLINE=%DISABLE_FORCEINLINE% :: Run msbuild in the resulting VS project files to build a pip package. %MSBUILD_EXE% /p:Configuration=Release /maxcpucount:32 tf_python_build_pip_package.vcxproj -- GitLab From d6b17ce09bddfec1670935bffb06eecddb7c3ba4 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Tue, 30 Jan 2018 23:25:36 +0800 Subject: [PATCH 034/596] README ipdate --- tensorflow/contrib/cmake/README.md | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index da72ea320f..ac7b7bcb04 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -39,7 +39,7 @@ bindings. * [Go](https://golang.org/) (optional, for full grpc build) -* [NASM](http://www.nasm.us/)/[yasm](http://yasm.tortall.net/) (optional, for full grpc build) +* [NASM](http://www.nasm.us/)/[YASM](http://yasm.tortall.net/) (optional, for full grpc build) * Additional pre-requisites for Microsoft Windows: - Visual Studio 2015 (latest version of MSVC 2017 is not suppored by CUDA yet, try it on your own) @@ -117,9 +117,9 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * CUDNN (GPU build) * NCCL (GPU build on Linux) * SWIG (python binding) - * Perl (GPU build on Windows, required by full function grpc, optional) - * Go (GPU build on Windows, required by full function grpc, optional) - * NASM (GPU build on Windows, required by full function grpc, optional) + * Perl (required by grpc for ssl support, optional) + * Go (required by grpc for ssl support, optional) + * NASM/YASM (required by grpc for ssl support, optional) 1. Start CMake GUI 2. Click on `Browse Source` and direct to the the folder `/tensorflow/contrib/cmake` 3. Click on `Browse Build` and spectify a location that you want tensorflow to be build @@ -131,7 +131,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * `tensorflow_BUILD_PYTHON_BINDING` is default to be `on`. Set to `off` if you don't need python interaface. If SWIG is not in system path, you need set it manually. (optional) * `tensorflow_BUILD_SHARED_LIB` is default to be `off`. Set to `on` if you want the c++ interface. (optional) * `tensorflow_ENABLE_GPU` is default to be `off`. Set to `on` if you want GPU support. It will search CUDA and CUDNN dependecies if you have set them to system path, otherwise CMake would prompt error and request you to set it manually. (optional) - * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. This need to be `on` for a gpu build. Reminded that Perl, Go and NASM is required for this option if you want to build grpc with offical SSL support. + * `tensorflow_ENABLE_GRPC_SUPPORT` is default to be `on`. For Linux build, this option must always be `on`. This need to be `on` for a gpu build. Reminded that Perl, Go and NASM/YASM are required for this option if you want to build grpc with offical SSL support. * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` @@ -152,7 +152,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. ***The wheel name varies depends on you config. Change to your own wheel filename.*** - Remind that some pip installation requires administrator right command prompt. + Reminded that some pip installation requires administrator right command prompt. ##### C++ -- GitLab From 6d84ada6eed5ad6e362ef4bcbf768ebf095753b2 Mon Sep 17 00:00:00 2001 From: JackyKo Date: Tue, 30 Jan 2018 23:26:45 +0800 Subject: [PATCH 035/596] README update --- tensorflow/contrib/cmake/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index ac7b7bcb04..c2faafa9f2 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -135,7 +135,7 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * `tensorflow_ENABLE_POSITION_INDEPENDENT_CODE` should always be `on` * `tensorflow_ENABLE_SNAPPY_SUPPORT` should always be `on` * `tensorflow_OPTIMIZE_FOR_NATIVE_ARCH` should always be `on` - * `CMAKE_INSTALL_PREFIX` is the location where the final package will be installed. You may change it to you own preferred path (optional) + * `CMAKE_INSTALL_PREFIX` is the location where the final package will be installed. You may change it to your own preferred path (optional) 6. After changing the configuration in step 5, press `Configure` again 7. If not error is found, press `Generate` -- GitLab From 8b377c00cfb40f5c364b8c7f96262ae37e9bb19c Mon Sep 17 00:00:00 2001 From: JackyKo Date: Tue, 30 Jan 2018 23:36:28 +0800 Subject: [PATCH 036/596] typo fix --- tensorflow/contrib/cmake/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index c2faafa9f2..e935abb5b2 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -42,7 +42,7 @@ bindings. * [NASM](http://www.nasm.us/)/[YASM](http://yasm.tortall.net/) (optional, for full grpc build) * Additional pre-requisites for Microsoft Windows: - - Visual Studio 2015 (latest version of MSVC 2017 is not suppored by CUDA yet, try it on your own) + - Visual Studio 2015 (latest version of MSVC 2017 is not supported by CUDA yet, try it on your own risk) - Python 3.5 - NumPy 1.11.0 or later -- GitLab From a74e6b6b13e79937b546f8cc0a5d172d5c787f98 Mon Sep 17 00:00:00 2001 From: jackyko Date: Wed, 31 Jan 2018 14:39:49 +0800 Subject: [PATCH 037/596] remove perl, go, asm requirement --- tensorflow/contrib/cmake/external/grpc.cmake | 40 -------------------- 1 file changed, 40 deletions(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index cd866f74c9..6ac087892a 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -14,46 +14,6 @@ # ============================================================================== include (ExternalProject) -# # perl -# find_package(Perl REQUIRED) - -# # nasm -# if(NOT CMAKE_ASM_NASM_COMPILER) -# if (WIN32) -# find_program(CMAKE_ASM_NASM_COMPILER -# NAMES nasm -# PATHS "$ENV{ProgramFiles}/NASM" -# DOC "path to NASM compiler") -# else() -# find_program(CMAKE_ASM_NASM_COMPILER -# NAMES nasm -# PATHS "/usr/bin" -# DOC "path to NASM compiler") -# endif() -# endif() - -# if(NOT CMAKE_ASM_NASM_COMPILER) -# message(FATAL_ERROR "NASM not found!") -# endif() - -# # go -# if(NOT GO_EXECUTABLE) -# if (WIN32) -# find_program(GO_EXECUTABLE -# NAMES go -# PATHS "$ENV{ProgramFiles}/Go/bin" -# DOC "path to go compiler") -# else() -# find_program(GO_EXECUTABLE -# NAMES go -# PATHS "/usr/bin" "/usr/local/bin" "/usr/local/go" -# DOC "path to go compiler") -# endif() -# endif() -# if(NOT GO_EXECUTABLE) -# message(FATAL_ERROR "Go not found!") -# endif() - set(GRPC_INCLUDE_DIRS ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/include) set(GRPC_URL https://github.com/grpc/grpc.git) set(GRPC_BUILD ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc) -- GitLab From b83f1b17b72d0a2313b025e4c0e94561dd54d950 Mon Sep 17 00:00:00 2001 From: jackyko Date: Wed, 31 Jan 2018 14:57:41 +0800 Subject: [PATCH 038/596] nccl header update --- tensorflow/contrib/cmake/CMakeLists.txt | 5 +++++ tensorflow/contrib/nccl/kernels/nccl_manager.h | 5 +++++ tensorflow/contrib/nccl/kernels/nccl_ops.cc | 4 ++++ 3 files changed, 14 insertions(+) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 3ec630b3d6..c03c53e305 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -63,6 +63,11 @@ if (NOT WIN32) # option's default value is OFF. Fill it with real default values set(tensorflow_CUDNN_INCLUDE /usr/include) endif (NOT tensorflow_CUDNN_INCLUDE) + option(tensorflow_NCCL_INCLUDE "nccl.h headher install path" /usr/include/) + if (NOT tensorflow_NCCL_INCLUDE) + # option's default value is OFF. Fill it with real default values + set(tensorflow_NCCL_INCLUDE /usr/include) + endif (NOT tensorflow_NCCL_INCLUDE) option(tensorflow_PATH_CUDNN_STATIC_LIB "Override PATH_STATIC_LIB for libcudnn_static.a" ${tensorflow_PATH_STATIC_LIB}) if (NOT tensorflow_PATH_CUDNN_STATIC_LIB) # option's default value is OFF. Fill it with real default values diff --git a/tensorflow/contrib/nccl/kernels/nccl_manager.h b/tensorflow/contrib/nccl/kernels/nccl_manager.h index bb219e0edc..7726f5813a 100644 --- a/tensorflow/contrib/nccl/kernels/nccl_manager.h +++ b/tensorflow/contrib/nccl/kernels/nccl_manager.h @@ -20,7 +20,12 @@ limitations under the License. #include #include +#ifdef (PLATFORM_GOOGLE) #include "src/nccl.h" +#else +#include "nccl.h" +#endif + #include "tensorflow/core/common_runtime/gpu/gpu_event_mgr.h" #include "tensorflow/core/framework/tensor.h" #include "tensorflow/core/platform/mutex.h" diff --git a/tensorflow/contrib/nccl/kernels/nccl_ops.cc b/tensorflow/contrib/nccl/kernels/nccl_ops.cc index 266d4f6f0d..7e59e70bbc 100644 --- a/tensorflow/contrib/nccl/kernels/nccl_ops.cc +++ b/tensorflow/contrib/nccl/kernels/nccl_ops.cc @@ -17,7 +17,11 @@ limitations under the License. #include +#ifdef (PLATFORM_GOOGLE) #include "src/nccl.h" +#else +#include "nccl.h" +#endif #include "tensorflow/contrib/nccl/kernels/nccl_manager.h" #include "tensorflow/core/framework/op_kernel.h" -- GitLab From ffa7c0a459153d9d7c8f1a390fb708dca8519ed4 Mon Sep 17 00:00:00 2001 From: jackyko Date: Wed, 31 Jan 2018 15:45:38 +0800 Subject: [PATCH 039/596] macro define fix --- tensorflow/contrib/nccl/kernels/nccl_manager.h | 2 +- tensorflow/contrib/nccl/kernels/nccl_ops.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/nccl/kernels/nccl_manager.h b/tensorflow/contrib/nccl/kernels/nccl_manager.h index 7726f5813a..0fe0476a2d 100644 --- a/tensorflow/contrib/nccl/kernels/nccl_manager.h +++ b/tensorflow/contrib/nccl/kernels/nccl_manager.h @@ -20,7 +20,7 @@ limitations under the License. #include #include -#ifdef (PLATFORM_GOOGLE) +#ifdef PLATFORM_GOOGLE #include "src/nccl.h" #else #include "nccl.h" diff --git a/tensorflow/contrib/nccl/kernels/nccl_ops.cc b/tensorflow/contrib/nccl/kernels/nccl_ops.cc index 7e59e70bbc..03d4dc50a4 100644 --- a/tensorflow/contrib/nccl/kernels/nccl_ops.cc +++ b/tensorflow/contrib/nccl/kernels/nccl_ops.cc @@ -17,7 +17,7 @@ limitations under the License. #include -#ifdef (PLATFORM_GOOGLE) +#ifdef PLATFORM_GOOGLE #include "src/nccl.h" #else #include "nccl.h" -- GitLab From 9800b2e903d7ae486008509987b736e777cb1ead Mon Sep 17 00:00:00 2001 From: jackyko Date: Thu, 1 Feb 2018 16:25:16 +0800 Subject: [PATCH 040/596] fix for windows api --- tensorflow/core/platform/default/logging.h | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/tensorflow/core/platform/default/logging.h b/tensorflow/core/platform/default/logging.h index 40c260f236..d456f631f2 100644 --- a/tensorflow/core/platform/default/logging.h +++ b/tensorflow/core/platform/default/logging.h @@ -187,6 +187,11 @@ string* MakeCheckOpString(const T1& v1, const T2& v2, const char* exprtext) { return comb.NewString(); } +#ifdef COMPILER_MSVC +#undef max; +#undef min; +#endif + // Helper functions for CHECK_OP macro. // The (int, int) specialization works around the issue that the compiler // will not instantiate the template version of the function on values of -- GitLab From 5cb510ebc01efc9c680df34fd6175ea7d5bd307b Mon Sep 17 00:00:00 2001 From: JackyKo Date: Sun, 4 Feb 2018 20:19:24 +0800 Subject: [PATCH 041/596] readme update --- tensorflow/contrib/cmake/README.md | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index e935abb5b2..eb915bd9cb 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -35,11 +35,11 @@ bindings. * [SWIG](http://www.swig.org/download.html) -* [Perl](https://www.perl.org/get.html) (optional, for full grpc build) +* [Perl](https://www.perl.org/get.html) (optional, for SSL support build) -* [Go](https://golang.org/) (optional, for full grpc build) +* [Go](https://golang.org/) (optional, for SSL support build) -* [NASM](http://www.nasm.us/)/[YASM](http://yasm.tortall.net/) (optional, for full grpc build) +* [NASM](http://www.nasm.us/)/[YASM](http://yasm.tortall.net/) (optional, for SSL support build) * Additional pre-requisites for Microsoft Windows: - Visual Studio 2015 (latest version of MSVC 2017 is not supported by CUDA yet, try it on your own risk) @@ -117,8 +117,8 @@ Install from CMake GUI would be a convenient way to generate C++ build projects. * CUDNN (GPU build) * NCCL (GPU build on Linux) * SWIG (python binding) - * Perl (required by grpc for ssl support, optional) - * Go (required by grpc for ssl support, optional) + * Perl (required if you need ssl support, optional) + * Go (required if you need ssl support, optional) * NASM/YASM (required by grpc for ssl support, optional) 1. Start CMake GUI 2. Click on `Browse Source` and direct to the the folder `/tensorflow/contrib/cmake` -- GitLab From 21a9efc4cddbce661073544db31a63639686310a Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Tue, 28 Nov 2017 05:28:49 -0800 Subject: [PATCH 042/596] Add complex dtypes support for `tf.squared_difference` This fix tries to address the issue raised in 14932 where complex dtypes are not supported for `tf.squared_difference`, which is different from the doc string in `math_ops.cc` (see `BINARY_FEWER`). This fix adds the complex64 and complex128 support in kernel, and adds additional test cases. This fix fixes 14932. Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_op_squared_difference.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/kernels/cwise_op_squared_difference.cc b/tensorflow/core/kernels/cwise_op_squared_difference.cc index 78fefc69c7..d0ff271df6 100644 --- a/tensorflow/core/kernels/cwise_op_squared_difference.cc +++ b/tensorflow/core/kernels/cwise_op_squared_difference.cc @@ -16,8 +16,8 @@ limitations under the License. #include "tensorflow/core/kernels/cwise_ops_common.h" namespace tensorflow { -REGISTER5(BinaryOp, CPU, "SquaredDifference", functor::squared_difference, - float, Eigen::half, double, int32, int64); +REGISTER7(BinaryOp, CPU, "SquaredDifference", functor::squared_difference, + float, Eigen::half, double, int32, int64, complex64, complex128); #if GOOGLE_CUDA REGISTER4(BinaryOp, GPU, "SquaredDifference", functor::squared_difference, float, Eigen::half, double, int64); -- GitLab From 4f5e66aca388ee13e925d173a82644eed9d5a760 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Tue, 28 Nov 2017 05:32:08 -0800 Subject: [PATCH 043/596] Add test cases for complex dtypes support with `tf.squared_difference` Signed-off-by: Yong Tang --- tensorflow/python/ops/math_ops_test.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tensorflow/python/ops/math_ops_test.py b/tensorflow/python/ops/math_ops_test.py index d314124ccd..7078ac99c8 100644 --- a/tensorflow/python/ops/math_ops_test.py +++ b/tensorflow/python/ops/math_ops_test.py @@ -203,7 +203,9 @@ class SquaredDifferenceTest(test_util.TensorFlowTestCase): @test_util.run_in_graph_and_eager_modes() def testSquaredDifference(self): - for dtype in [np.int32, np.float16]: + for dtype in [np.float16, np.float32, np.float64, + np.int32, np.int64, + np.complex64, np.complex128]: x = np.array([[1, 2, 3], [4, 5, 6]], dtype=dtype) y = np.array([-3, -2, -1], dtype=dtype) z = (x - y) * (x - y) -- GitLab From f7bb3741549e791f687fa8289fb281717eae7426 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 24 Feb 2018 19:05:52 +0000 Subject: [PATCH 044/596] Add additional test to cover squared difference for complex where imag parts are not 0. Signed-off-by: Yong Tang --- tensorflow/python/ops/math_ops_test.py | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/tensorflow/python/ops/math_ops_test.py b/tensorflow/python/ops/math_ops_test.py index 7078ac99c8..3224e40db2 100644 --- a/tensorflow/python/ops/math_ops_test.py +++ b/tensorflow/python/ops/math_ops_test.py @@ -213,6 +213,16 @@ class SquaredDifferenceTest(test_util.TensorFlowTestCase): z_tf = self.evaluate(math_ops.squared_difference(x, y)) self.assertAllClose(z, z_tf) + @test_util.run_in_graph_and_eager_modes() + def testComplexSquaredDifference(self): + for dtype in [np.complex64, np.complex128]: + x = np.array([[1+3j, 2+2j, 3+1j], [4-1j, 5-2j, 6-3j]], dtype=dtype) + y = np.array([-3+1j, -2+2j, -1+3j], dtype=dtype) + z = (x - y) * (x - y) + with test_util.device(use_gpu=True): + z_tf = self.evaluate(math_ops.squared_difference(x, y)) + self.assertAllClose(z, z_tf) + @test_util.with_c_api class ApproximateEqualTest(test_util.TensorFlowTestCase): -- GitLab From ab5cdb187d96e3a865724c3d41671dd253288456 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 24 Feb 2018 21:22:26 +0000 Subject: [PATCH 045/596] Enable squared_difference complex on CPU only Signed-off-by: Yong Tang --- tensorflow/python/ops/math_ops_test.py | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/tensorflow/python/ops/math_ops_test.py b/tensorflow/python/ops/math_ops_test.py index 3224e40db2..533a00e737 100644 --- a/tensorflow/python/ops/math_ops_test.py +++ b/tensorflow/python/ops/math_ops_test.py @@ -204,8 +204,7 @@ class SquaredDifferenceTest(test_util.TensorFlowTestCase): @test_util.run_in_graph_and_eager_modes() def testSquaredDifference(self): for dtype in [np.float16, np.float32, np.float64, - np.int32, np.int64, - np.complex64, np.complex128]: + np.int32, np.int64]: x = np.array([[1, 2, 3], [4, 5, 6]], dtype=dtype) y = np.array([-3, -2, -1], dtype=dtype) z = (x - y) * (x - y) @@ -219,7 +218,7 @@ class SquaredDifferenceTest(test_util.TensorFlowTestCase): x = np.array([[1+3j, 2+2j, 3+1j], [4-1j, 5-2j, 6-3j]], dtype=dtype) y = np.array([-3+1j, -2+2j, -1+3j], dtype=dtype) z = (x - y) * (x - y) - with test_util.device(use_gpu=True): + with test_util.device(use_gpu=False): z_tf = self.evaluate(math_ops.squared_difference(x, y)) self.assertAllClose(z, z_tf) -- GitLab From d7f53eee873cf675eccd7a3f3d5966b8fa398887 Mon Sep 17 00:00:00 2001 From: Ka Long Date: Fri, 13 Apr 2018 10:29:47 +0800 Subject: [PATCH 046/596] Update CMakeLists.txt --- tensorflow/contrib/cmake/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 251f3d2e0c..75a7664f1d 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.5) # Project -project(tensorflow VERSION 1.5.0 LANGUAGES C CXX) +project(tensorflow VERSION 1.7.0 LANGUAGES C CXX) # Set C++14 as standard for the whole project set(CMAKE_CXX_STANDARD 14) -- GitLab From 346ec501531e6af0bad0da52d2db4a81733fac27 Mon Sep 17 00:00:00 2001 From: Jacky Date: Mon, 23 Apr 2018 15:02:21 +0800 Subject: [PATCH 047/596] turn off png_test --- tensorflow/contrib/cmake/external/png.cmake | 1 + 1 file changed, 1 insertion(+) diff --git a/tensorflow/contrib/cmake/external/png.cmake b/tensorflow/contrib/cmake/external/png.cmake index ad2af01bc0..d2aa99efdf 100644 --- a/tensorflow/contrib/cmake/external/png.cmake +++ b/tensorflow/contrib/cmake/external/png.cmake @@ -58,6 +58,7 @@ ExternalProject_Add(png -DCMAKE_VERBOSE_MAKEFILE:BOOL=OFF -DCMAKE_INSTALL_PREFIX:STRING=${png_INSTALL} -DZLIB_ROOT:STRING=${ZLIB_INSTALL} + -DPNG_TESTS:BOOL=OFF ) ## put png includes in the directory where they are expected -- GitLab From 522eebf29e2c7fdb46b9bb5af5900cc864b19fba Mon Sep 17 00:00:00 2001 From: Jacky Date: Tue, 24 Apr 2018 14:41:23 +0800 Subject: [PATCH 048/596] new gpu framework modules added --- tensorflow/contrib/cmake/tf_core_framework.cmake | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tensorflow/contrib/cmake/tf_core_framework.cmake b/tensorflow/contrib/cmake/tf_core_framework.cmake index 8cb5bc4f1b..ece815d205 100644 --- a/tensorflow/contrib/cmake/tf_core_framework.cmake +++ b/tensorflow/contrib/cmake/tf_core_framework.cmake @@ -139,16 +139,19 @@ set(tf_proto_text_srcs "tensorflow/core/example/example.proto" "tensorflow/core/example/feature.proto" "tensorflow/core/framework/allocation_description.proto" + "tensorflow/core/framework/api_def.proto" "tensorflow/core/framework/attr_value.proto" "tensorflow/core/framework/cost_graph.proto" "tensorflow/core/framework/device_attributes.proto" "tensorflow/core/framework/function.proto" "tensorflow/core/framework/graph.proto" "tensorflow/core/framework/graph_transfer_info.proto" + "tensorflow/core/framework/iterator.proto" "tensorflow/core/framework/kernel_def.proto" "tensorflow/core/framework/log_memory.proto" "tensorflow/core/framework/node_def.proto" "tensorflow/core/framework/op_def.proto" + "tensorflow/core/framework/reader_base.proto" "tensorflow/core/framework/remote_fused_graph_execute_info.proto" "tensorflow/core/framework/resource_handle.proto" "tensorflow/core/framework/step_stats.proto" @@ -158,6 +161,7 @@ set(tf_proto_text_srcs "tensorflow/core/framework/tensor_shape.proto" "tensorflow/core/framework/tensor_slice.proto" "tensorflow/core/framework/types.proto" + "tensorflow/core/framework/variable.proto" "tensorflow/core/framework/versions.proto" "tensorflow/core/lib/core/error_codes.proto" "tensorflow/core/protobuf/cluster.proto" -- GitLab From 973b80039920d25d93dc38aff4ab45b98b6b4872 Mon Sep 17 00:00:00 2001 From: Jacky Date: Tue, 24 Apr 2018 15:47:26 +0800 Subject: [PATCH 049/596] core_ops updated --- tensorflow/contrib/cmake/tf_core_ops.cmake | 8 +++++--- tensorflow/contrib/cmake/tf_python.cmake | 21 ++++++++++++++------- 2 files changed, 19 insertions(+), 10 deletions(-) diff --git a/tensorflow/contrib/cmake/tf_core_ops.cmake b/tensorflow/contrib/cmake/tf_core_ops.cmake index c356fb1494..42c621700b 100644 --- a/tensorflow/contrib/cmake/tf_core_ops.cmake +++ b/tensorflow/contrib/cmake/tf_core_ops.cmake @@ -13,13 +13,14 @@ # limitations under the License. # ============================================================================== set(tf_op_lib_names - "audio_ops" "array_ops" + "audio_ops" "batch_ops" "bitwise_ops" "boosted_trees_ops" "candidate_sampling_ops" "checkpoint_ops" + "collective_ops" "control_flow_ops" "ctc_ops" "cudnn_rnn_ops" @@ -33,8 +34,8 @@ set(tf_op_lib_names "io_ops" "linalg_ops" "list_ops" - "lookup_ops" "logging_ops" + "lookup_ops" "manip_ops" "math_ops" "nn_ops" @@ -44,10 +45,11 @@ set(tf_op_lib_names "remote_fused_graph_ops" "resource_variable_ops" "rpc_ops" + "scoped_allocator_ops" "script_ops" "sdca_ops" - "set_ops" "sendrecv_ops" + "set_ops" "sparse_ops" "spectral_ops" "state_ops" diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index 46d36d227f..1a630a2bd4 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -315,15 +315,14 @@ function(GENERATE_PYTHON_OP_LIB tf_python_op_lib_name) ${GENERATE_PYTHON_OP_LIB_DESTINATION} PARENT_SCOPE) endfunction() -GENERATE_PYTHON_OP_LIB("audio_ops") GENERATE_PYTHON_OP_LIB("array_ops") +GENERATE_PYTHON_OP_LIB("audio_ops") GENERATE_PYTHON_OP_LIB("batch_ops") GENERATE_PYTHON_OP_LIB("bitwise_ops") GENERATE_PYTHON_OP_LIB("boosted_trees_ops") -GENERATE_PYTHON_OP_LIB("math_ops") -GENERATE_PYTHON_OP_LIB("functional_ops") GENERATE_PYTHON_OP_LIB("candidate_sampling_ops") GENERATE_PYTHON_OP_LIB("checkpoint_ops") +GENERATE_PYTHON_OP_LIB("collective_ops") GENERATE_PYTHON_OP_LIB("control_flow_ops" ADDITIONAL_LIBRARIES $) GENERATE_PYTHON_OP_LIB("ctc_ops") @@ -334,14 +333,18 @@ GENERATE_PYTHON_OP_LIB("decode_proto_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/proto/python/ops/gen_decode_proto_op.py) GENERATE_PYTHON_OP_LIB("encode_proto_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/proto/python/ops/gen_encode_proto_op.py) +GENERATE_PYTHON_OP_LIB("function_ops") +GENERATE_PYTHON_OP_LIB("functional_ops") GENERATE_PYTHON_OP_LIB("image_ops") GENERATE_PYTHON_OP_LIB("io_ops") GENERATE_PYTHON_OP_LIB("linalg_ops") GENERATE_PYTHON_OP_LIB("list_ops") GENERATE_PYTHON_OP_LIB("logging_ops") GENERATE_PYTHON_OP_LIB("lookup_ops") -GENERATE_PYTHON_OP_LIB("nn_ops") GENERATE_PYTHON_OP_LIB("manip_ops") +GENERATE_PYTHON_OP_LIB("math_ops") +GENERATE_PYTHON_OP_LIB("nn_ops") +GENERATE_PYTHON_OP_LIB("no_op") GENERATE_PYTHON_OP_LIB("parsing_ops") GENERATE_PYTHON_OP_LIB("random_ops") GENERATE_PYTHON_OP_LIB("remote_fused_graph_ops" @@ -349,17 +352,21 @@ GENERATE_PYTHON_OP_LIB("remote_fused_graph_ops" GENERATE_PYTHON_OP_LIB("resource_variable_ops") GENERATE_PYTHON_OP_LIB("rpc_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/rpc/python/ops/gen_rpc_op.py) +GENERATE_PYTHON_OP_LIB("scoped_allocator_ops") GENERATE_PYTHON_OP_LIB("script_ops") GENERATE_PYTHON_OP_LIB("sdca_ops") +GENERATE_PYTHON_OP_LIB("sendrecv_ops") GENERATE_PYTHON_OP_LIB("set_ops") -GENERATE_PYTHON_OP_LIB("state_ops") GENERATE_PYTHON_OP_LIB("sparse_ops") GENERATE_PYTHON_OP_LIB("spectral_ops") +GENERATE_PYTHON_OP_LIB("state_ops") +GENERATE_PYTHON_OP_LIB("stateless_random_ops") GENERATE_PYTHON_OP_LIB("string_ops") GENERATE_PYTHON_OP_LIB("summary_ops") GENERATE_PYTHON_OP_LIB("user_ops") GENERATE_PYTHON_OP_LIB("training_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/python/training/gen_training_ops.py) +GENERATE_PYTHON_OP_LIB("word2vec_ops") GENERATE_PYTHON_OP_LIB("contrib_boosted_trees_model_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/boosted_trees/python/ops/gen_model_ops.py) @@ -422,8 +429,8 @@ GENERATE_PYTHON_OP_LIB("contrib_text_skip_gram_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/text/python/ops/gen_skip_gram_ops.py) GENERATE_PYTHON_OP_LIB("contrib_bigquery_reader_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/cloud/python/ops/gen_bigquery_reader_ops.py) -GENERATE_PYTHON_OP_LIB("stateless_random_ops" - DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/stateless/gen_stateless_random_ops.py) +# GENERATE_PYTHON_OP_LIB("stateless_random_ops" +# DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/stateless/gen_stateless_random_ops.py) GENERATE_PYTHON_OP_LIB("debug_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/python/debug/ops/gen_debug_ops.py) -- GitLab From e28cbfe68e11f97c89e0472b7fecd187485d0916 Mon Sep 17 00:00:00 2001 From: Jacky Date: Tue, 24 Apr 2018 19:23:20 +0800 Subject: [PATCH 050/596] version number update --- tensorflow/contrib/cmake/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 8c6d8f2e72..5396e434b1 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.5) # Project -project(tensorflow VERSION 1.7.0 LANGUAGES C CXX) +project(tensorflow VERSION 1.8.0 LANGUAGES C CXX) # Set C++14 as standard for the whole project set(CMAKE_CXX_STANDARD 14) @@ -69,7 +69,7 @@ if (NOT WIN32) # option's default value is OFF. Fill it with real default values set(tensorflow_CUDNN_INCLUDE /usr/include) endif (NOT tensorflow_CUDNN_INCLUDE) - option(tensorflow_NCCL_INCLUDE "nccl.h headher install path" /usr/include/) + option(tensorflow_NCCL_INCLUDE "nccl.h header install path" /usr/include/) if (NOT tensorflow_NCCL_INCLUDE) # option's default value is OFF. Fill it with real default values set(tensorflow_NCCL_INCLUDE /usr/include) -- GitLab From 664f3414aee08d0ec71e13c7329a6712c15de2da Mon Sep 17 00:00:00 2001 From: Jacky Date: Tue, 24 Apr 2018 19:24:18 +0800 Subject: [PATCH 051/596] typo correct --- tensorflow/contrib/cmake/README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tensorflow/contrib/cmake/README.md b/tensorflow/contrib/cmake/README.md index be3dfdf043..f5d17dbafc 100644 --- a/tensorflow/contrib/cmake/README.md +++ b/tensorflow/contrib/cmake/README.md @@ -4,7 +4,7 @@ TensorFlow CMake build This directory contains CMake files for building TensorFlow on Microsoft Windows and Linux. [CMake](https://cmake.org) is a cross-platform tool that can generate build scripts for multiple build systems, including Microsoft -Visual Studio and GCC. The method suppose could work on MacOS, but not tested yet. +Visual Studio and GCC. "The method has not been tested on Mac OS X. **N.B.** We provide Linux build instructions primarily for the purpose of testing the build. We recommend using the standard Bazel-based build on @@ -14,12 +14,12 @@ Current Status -------------- CMake can be used to build TensorFlow on all platforms. See the [getting started documentation](https://www.tensorflow.org/install/install_windows) -for instructions on how to install a pre-built TensorFlow package on Windows and Linux. Procedure in MacOS is similar to the Linux build. +for instructions on how to install a pre-built TensorFlow package on Windows and Linux. The procedure in MacOS is similar to the Linux build. ### Current known limitations * It is not possible to load a custom Op library. * GCS file system is not supported. -* Debug build is not available since python no longer release debug library. +* Debug build is not available since Python for Windows is no longer distributed with a debug library. ## Building with CMake -- GitLab From 149714f6845692f588329c54d20f93f10b1641c6 Mon Sep 17 00:00:00 2001 From: Jacky Date: Wed, 25 Apr 2018 17:34:33 +0800 Subject: [PATCH 052/596] refine under code review --- tensorflow/contrib/cmake/external/grpc.cmake | 3 +-- tensorflow/contrib/cmake/tf_python.cmake | 2 -- tensorflow/core/util/cuda_device_functions.h | 16 +++++++++------- 3 files changed, 10 insertions(+), 11 deletions(-) diff --git a/tensorflow/contrib/cmake/external/grpc.cmake b/tensorflow/contrib/cmake/external/grpc.cmake index 56774bf8aa..2258a81170 100644 --- a/tensorflow/contrib/cmake/external/grpc.cmake +++ b/tensorflow/contrib/cmake/external/grpc.cmake @@ -28,6 +28,7 @@ else() set(grpc_STATIC_LIBRARIES ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgrpc++_unsecure.a ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgrpc_unsecure.a + ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libaddress_sorting.a ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/libgpr.a ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/third_party/cares/cares/lib/libcares.a ${CMAKE_CURRENT_BINARY_DIR}/grpc/src/grpc/third_party/zlib/libz.a) @@ -53,8 +54,6 @@ ExternalProject_Add(grpc -DPROTOBUF_LIBRARIES:STRING=${protobuf_STATIC_LIBRARIES} -DZLIB_ROOT:STRING=${ZLIB_INSTALL} -DgRPC_SSL_PROVIDER:STRING=NONE - # -DCMAKE_ASM_NASM_COMPILER:STRING=${CMAKE_ASM_NASM_COMPILER} - # -DGO_EXECUTABLE:STRING=${GO_EXECUTABLE} ) # grpc/src/core/ext/census/tracing.c depends on the existence of openssl/rand.h. diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index 1a630a2bd4..bf5c54d087 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -429,8 +429,6 @@ GENERATE_PYTHON_OP_LIB("contrib_text_skip_gram_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/text/python/ops/gen_skip_gram_ops.py) GENERATE_PYTHON_OP_LIB("contrib_bigquery_reader_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/cloud/python/ops/gen_bigquery_reader_ops.py) -# GENERATE_PYTHON_OP_LIB("stateless_random_ops" -# DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/stateless/gen_stateless_random_ops.py) GENERATE_PYTHON_OP_LIB("debug_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/python/debug/ops/gen_debug_ops.py) diff --git a/tensorflow/core/util/cuda_device_functions.h b/tensorflow/core/util/cuda_device_functions.h index a8d67aa501..846a323a6d 100644 --- a/tensorflow/core/util/cuda_device_functions.h +++ b/tensorflow/core/util/cuda_device_functions.h @@ -28,20 +28,22 @@ limitations under the License. #include #include -#ifdef PLATFORM_GOOGLE -#include "cuda/include/cuda.h" -#include "cuda/include/device_functions.h" -#else +#include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" + +#ifdef _WIN32 #include "cuda.h" #include "device_functions.h" +#else +#include "cuda/include/cuda.h" +#include "cuda/include/device_functions.h" #endif #include "tensorflow/core/platform/types.h" #if CUDA_VERSION >= 7050 -#ifdef PLATFORM_GOOGLE -#include "cuda/include/cuda_fp16.h" -#else +#ifdef _WIN32 #include "cuda_fp16.h" +#else +#include "cuda/include/cuda_fp16.h" #endif #endif // CUDA_VERSION >= 7050 -- GitLab From 7f21edc8239c1f90ce438e1a14dc5560e01d4f65 Mon Sep 17 00:00:00 2001 From: Jacky Date: Wed, 25 Apr 2018 17:38:49 +0800 Subject: [PATCH 053/596] add suppress regeneration back --- tensorflow/contrib/cmake/CMakeLists.txt | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 5396e434b1..d7aab58b65 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -166,8 +166,13 @@ if(WIN32) set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} /D_ITERATOR_DEBUG_LEVEL=0") set(CMAKE_CXX_FLAGS_MINSIZEREL "${CMAKE_CXX_FLAGS_MINSIZEREL} /D_ITERATOR_DEBUG_LEVEL=0") set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "${CMAKE_CXX_FLAGS_RELWITHDEBINFO} /D_ITERATOR_DEBUG_LEVEL=0") + + # Try to avoid flaky failures due to failed generation of generate.stamp files. + set(CMAKE_SUPPRESS_REGENERATION ON) endif() + + if ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-exceptions -std=c++11") endif() -- GitLab From a210139765d916baa516dbaf571a435d588cb2e4 Mon Sep 17 00:00:00 2001 From: Jacky Date: Wed, 25 Apr 2018 18:01:47 +0800 Subject: [PATCH 054/596] update accroding to review --- tensorflow/contrib/cmake/CMakeLists.txt | 1 - tensorflow/core/platform/default/logging.h | 7 ++++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index d7aab58b65..c9e2002977 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -172,7 +172,6 @@ if(WIN32) endif() - if ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-exceptions -std=c++11") endif() diff --git a/tensorflow/core/platform/default/logging.h b/tensorflow/core/platform/default/logging.h index 916d3caa7e..6305cb2823 100644 --- a/tensorflow/core/platform/default/logging.h +++ b/tensorflow/core/platform/default/logging.h @@ -187,9 +187,10 @@ string* MakeCheckOpString(const T1& v1, const T2& v2, const char* exprtext) { return comb.NewString(); } -#ifdef COMPILER_MSVC -#undef max; -#undef min; +// for MSVC build, the max and min function maybe defined in other macros +#ifdef _WIN32 +#undef max +#undef min #endif // Helper functions for CHECK_OP macro. -- GitLab From 446707346fbbee5cf228aab0c25a2cf08de9e7dc Mon Sep 17 00:00:00 2001 From: Jacky Date: Thu, 26 Apr 2018 16:43:21 +0800 Subject: [PATCH 055/596] header fix --- tensorflow/contrib/nccl/kernels/nccl_manager.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/nccl/kernels/nccl_manager.h b/tensorflow/contrib/nccl/kernels/nccl_manager.h index 90cdea47c2..e81a529227 100644 --- a/tensorflow/contrib/nccl/kernels/nccl_manager.h +++ b/tensorflow/contrib/nccl/kernels/nccl_manager.h @@ -23,8 +23,8 @@ limitations under the License. #ifdef WIN32 #include "nccl.h" #else -"third_party/nccl/nccl.h" -#else +#include "third_party/nccl/nccl.h" +#endif #include "tensorflow/core/common_runtime/gpu/gpu_event_mgr.h" #include "tensorflow/core/framework/tensor.h" -- GitLab From 961b9086d6c435d078913fbc7293a91ee819bd38 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Sat, 28 Apr 2018 11:07:45 +0800 Subject: [PATCH 056/596] modified for latest eager runtime --- tensorflow/c/python_api.cc | 2 +- tensorflow/contrib/cmake/CMakeLists.txt | 3 +- tensorflow/contrib/cmake/tf_c.cmake | 27 ++++----- tensorflow/contrib/cmake/tf_core_cpu.cmake | 4 +- .../contrib/cmake/tf_core_eager_runtime.cmake | 57 +++++++++++++++++++ tensorflow/contrib/cmake/tf_python.cmake | 2 + tensorflow/contrib/cmake/tf_shared_lib.cmake | 4 ++ 7 files changed, 80 insertions(+), 19 deletions(-) create mode 100644 tensorflow/contrib/cmake/tf_core_eager_runtime.cmake diff --git a/tensorflow/c/python_api.cc b/tensorflow/c/python_api.cc index e18fdf6c57..cb9038978d 100644 --- a/tensorflow/c/python_api.cc +++ b/tensorflow/c/python_api.cc @@ -16,7 +16,7 @@ limitations under the License. #include "tensorflow/c/python_api.h" #include "tensorflow/c/c_api_internal.h" -#include "tensorflow/python/framework/cpp_shape_inference.pb.h" +#include "tensorflow/python/framework/cpp_shape_inference.h" namespace tensorflow { diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 981f50bc95..3016a7e2bd 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -539,6 +539,7 @@ include(tf_cc_ops.cmake) include(tf_c.cmake) include(tf_grappler.cmake) include(tf_core_profiler.cmake) +include(tf_core_eager_runtime.cmake) if(tensorflow_BUILD_CC_EXAMPLE) include(tf_tutorials.cmake) include(tf_label_image_example.cmake) @@ -552,4 +553,4 @@ if(tensorflow_BUILD_SHARED_LIB) endif() if(tensorflow_BUILD_CC_TESTS OR tensorflow_BUILD_PYTHON_TESTS) include(tf_tests.cmake) -endif() +endif() \ No newline at end of file diff --git a/tensorflow/contrib/cmake/tf_c.cmake b/tensorflow/contrib/cmake/tf_c.cmake index c6a15f2ca0..0479a90713 100644 --- a/tensorflow/contrib/cmake/tf_c.cmake +++ b/tensorflow/contrib/cmake/tf_c.cmake @@ -19,11 +19,6 @@ set(tf_c_srcs "${tensorflow_source_dir}/tensorflow/c/c_api.cc" "${tensorflow_source_dir}/tensorflow/c/c_api.h" "${tensorflow_source_dir}/tensorflow/c/c_api_function.cc" - "${tensorflow_source_dir}/tensorflow/c/eager/c_api.cc" - "${tensorflow_source_dir}/tensorflow/c/eager/c_api.h" - "${tensorflow_source_dir}/tensorflow/c/eager/tape.h" - "${tensorflow_source_dir}/tensorflow/c/eager/runtime.cc" - "${tensorflow_source_dir}/tensorflow/c/eager/runtime.h" "${tensorflow_source_dir}/tensorflow/c/checkpoint_reader.cc" "${tensorflow_source_dir}/tensorflow/c/checkpoint_reader.h" "${tensorflow_source_dir}/tensorflow/c/tf_status_helper.cc" @@ -38,13 +33,15 @@ add_dependencies( tf_core_lib tf_protos_cc) -add_library(tf_c_python_api OBJECT - "${tensorflow_source_dir}/tensorflow/c/python_api.cc" - "${tensorflow_source_dir}/tensorflow/c/python_api.h" -) -add_dependencies( - tf_c_python_api - tf_c - tf_core_lib - tf_core_framework - tf_protos_cc) +if(tensorflow_BUILD_PYTHON_BINDINGS) + add_library(tf_c_python_api OBJECT + "${tensorflow_source_dir}/tensorflow/c/python_api.cc" + "${tensorflow_source_dir}/tensorflow/c/python_api.h" + ) + add_dependencies( + tf_c_python_api + tf_c + tf_core_lib + tf_core_framework + tf_protos_cc) +endif() diff --git a/tensorflow/contrib/cmake/tf_core_cpu.cmake b/tensorflow/contrib/cmake/tf_core_cpu.cmake index 1562b6e0a3..d8884d464f 100644 --- a/tensorflow/contrib/cmake/tf_core_cpu.cmake +++ b/tensorflow/contrib/cmake/tf_core_cpu.cmake @@ -20,8 +20,6 @@ file(GLOB_RECURSE tf_core_cpu_srcs "${tensorflow_source_dir}/tensorflow/cc/saved_model/*.cc" "${tensorflow_source_dir}/tensorflow/core/common_runtime/*.h" "${tensorflow_source_dir}/tensorflow/core/common_runtime/*.cc" - "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*.h" - "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*.cc" "${tensorflow_source_dir}/tensorflow/core/debug/*.h" "${tensorflow_source_dir}/tensorflow/core/debug/*.cc" "${tensorflow_source_dir}/tensorflow/core/distributed_runtime/server_lib.h" @@ -41,6 +39,8 @@ file(GLOB_RECURSE tf_core_cpu_exclude_srcs "${tensorflow_source_dir}/tensorflow/core/*test*.h" "${tensorflow_source_dir}/tensorflow/core/*test*.cc" "${tensorflow_source_dir}/tensorflow/core/*main.cc" + "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*.cc" + "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*.h" "${tensorflow_source_dir}/tensorflow/core/common_runtime/gpu/*.cc" "${tensorflow_source_dir}/tensorflow/core/common_runtime/gpu_device_factory.cc" "${tensorflow_source_dir}/tensorflow/core/common_runtime/direct_session.cc" diff --git a/tensorflow/contrib/cmake/tf_core_eager_runtime.cmake b/tensorflow/contrib/cmake/tf_core_eager_runtime.cmake new file mode 100644 index 0000000000..78e4c0d303 --- /dev/null +++ b/tensorflow/contrib/cmake/tf_core_eager_runtime.cmake @@ -0,0 +1,57 @@ +# 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. +# ============================================================================== +######################################################## +# tf_core_eager_runtime library +######################################################## +file(GLOB_RECURSE tf_core_eager_runtime_srcs + "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*.cc" + "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*.h" +) + +file(GLOB_RECURSE tf_core_eager_runtime_exclude_srcs + "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*test*.h" + "${tensorflow_source_dir}/tensorflow/core/common_runtime/eager/*test*.cc" +) + +list(REMOVE_ITEM tf_core_eager_runtime_srcs ${tf_core_eager_runtime_exclude_srcs}) + +add_library(tf_core_eager_runtime OBJECT ${tf_core_eager_runtime_srcs}) +add_dependencies( + tf_core_eager_runtime + tf_c + tf_core_lib) + + +file(GLOB_RECURSE tf_c_eager_srcs + "${tensorflow_source_dir}/tensorflow/c/eager/*.cc" + "${tensorflow_source_dir}/tensorflow/c/eager/*.h" +) + +file(GLOB_RECURSE tf_c_eager_exlclude_srcs + "${tensorflow_source_dir}/tensorflow/c/eager/*test*.h" + "${tensorflow_source_dir}/tensorflow/c/eager/*test*.cc" +) + +list(REMOVE_ITEM tf_c_eager_srcs ${tf_c_eager_exlclude_srcs}) + +add_library(tf_c_eager OBJECT ${tf_c_eager_srcs}) +add_dependencies( + tf_c_eager + tf_core_eager_runtime + tf_c + tf_cc_framework + tf_cc_while_loop + tf_core_lib + tf_protos_cc) \ No newline at end of file diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index bf5c54d087..fd0b3bd38c 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -524,6 +524,7 @@ if(WIN32) $ $ $ + $ $ $ $ @@ -581,6 +582,7 @@ add_library(pywrap_tensorflow_internal SHARED $ $ $ + $ $ $ $ diff --git a/tensorflow/contrib/cmake/tf_shared_lib.cmake b/tensorflow/contrib/cmake/tf_shared_lib.cmake index 41f8391b6f..a6f4ac2f18 100644 --- a/tensorflow/contrib/cmake/tf_shared_lib.cmake +++ b/tensorflow/contrib/cmake/tf_shared_lib.cmake @@ -23,6 +23,8 @@ if(WIN32) # we need. # add_library(tensorflow_static STATIC + $ + $ $ $ $ @@ -65,6 +67,8 @@ endif(WIN32) # tensorflow is a shared library containing all of the # TensorFlow runtime and the standard ops and kernels. add_library(tensorflow SHARED + $ + $ $ $ $ -- GitLab From a1a036742ef614e28064b0839c8b17d1cb9025a0 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Sat, 28 Apr 2018 11:15:35 +0800 Subject: [PATCH 057/596] update pywrap dependecy with eager --- tensorflow/contrib/cmake/tf_python.cmake | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index fd0b3bd38c..02f56e9bad 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -519,6 +519,7 @@ if(WIN32) add_library(pywrap_tensorflow_internal_static STATIC ${pywrap_tensorflow_internal_src} $ + $ $ $ $ @@ -577,6 +578,7 @@ endif(WIN32) add_library(pywrap_tensorflow_internal SHARED ${pywrap_tensorflow_internal_src} $ + $ $ $ $ -- GitLab From 3165f6efe06d45caadbfb568ec6c0e4a6e9b789b Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Sat, 28 Apr 2018 12:14:37 +0800 Subject: [PATCH 058/596] add numpy dependency --- tensorflow/c/python_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/c/python_api.cc b/tensorflow/c/python_api.cc index cb9038978d..e18fdf6c57 100644 --- a/tensorflow/c/python_api.cc +++ b/tensorflow/c/python_api.cc @@ -16,7 +16,7 @@ limitations under the License. #include "tensorflow/c/python_api.h" #include "tensorflow/c/c_api_internal.h" -#include "tensorflow/python/framework/cpp_shape_inference.h" +#include "tensorflow/python/framework/cpp_shape_inference.pb.h" namespace tensorflow { -- GitLab From 7dc284ec17ba76258e14ac4e7f322863aaaa6546 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Sat, 28 Apr 2018 12:16:08 +0800 Subject: [PATCH 059/596] add numpy dependency --- tensorflow/contrib/cmake/tf_c.cmake | 34 +++++++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/tensorflow/contrib/cmake/tf_c.cmake b/tensorflow/contrib/cmake/tf_c.cmake index 0479a90713..9415a6ceb2 100644 --- a/tensorflow/contrib/cmake/tf_c.cmake +++ b/tensorflow/contrib/cmake/tf_c.cmake @@ -12,6 +12,34 @@ # See the License for the specific language governing permissions and # limitations under the License. # ============================================================================== + +# 1. Resolve the installed version of Python (for Python.h and python). +# TODO(mrry): Parameterize the build script to enable Python 3 building. +if(NOT PYTHON_INCLUDE_DIR) + set(PYTHON_NOT_FOUND false) + exec_program("${PYTHON_EXECUTABLE}" + ARGS "-c \"import distutils.sysconfig; print(distutils.sysconfig.get_python_inc())\"" + OUTPUT_VARIABLE PYTHON_INCLUDE_DIR + RETURN_VALUE PYTHON_NOT_FOUND) + if(${PYTHON_NOT_FOUND}) + message(FATAL_ERROR + "Cannot get Python include directory. Is distutils installed?") + endif(${PYTHON_NOT_FOUND}) +endif(NOT PYTHON_INCLUDE_DIR) + +# 2. Resolve the installed version of NumPy (for numpy/arrayobject.h). +if(NOT NUMPY_INCLUDE_DIR) + set(NUMPY_NOT_FOUND false) + exec_program("${PYTHON_EXECUTABLE}" + ARGS "-c \"import numpy; print(numpy.get_include())\"" + OUTPUT_VARIABLE NUMPY_INCLUDE_DIR + RETURN_VALUE NUMPY_NOT_FOUND) + if(${NUMPY_NOT_FOUND}) + message(FATAL_ERROR + "Cannot get NumPy include directory: Is NumPy installed?") + endif(${NUMPY_NOT_FOUND}) +endif(NOT NUMPY_INCLUDE_DIR) + ######################################################## # tf_c_framework library ######################################################## @@ -38,6 +66,12 @@ if(tensorflow_BUILD_PYTHON_BINDINGS) "${tensorflow_source_dir}/tensorflow/c/python_api.cc" "${tensorflow_source_dir}/tensorflow/c/python_api.h" ) + + target_include_directories(tf_c_python_api PUBLIC + ${PYTHON_INCLUDE_DIR} + ${NUMPY_INCLUDE_DIR} + ) + add_dependencies( tf_c_python_api tf_c -- GitLab From 7950d197767ef24a1525b809a310b82020f665ba Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Wed, 1 Aug 2018 13:43:35 -0700 Subject: [PATCH 060/596] MKL DNN: Adding support of fusing Pad and Conv2D in MKL DNN optimized code --- tensorflow/core/graph/mkl_layout_pass.cc | 298 +++++++++++++++++- tensorflow/core/graph/mkl_layout_pass_test.cc | 70 ++++ tensorflow/core/kernels/BUILD | 28 ++ tensorflow/core/kernels/mkl_conv_ops.cc | 96 +++++- tensorflow/core/kernels/mkl_conv_ops.h | 28 +- tensorflow/core/ops/nn_ops.cc | 49 +++ 6 files changed, 554 insertions(+), 15 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index c22e0a3872..d0abe5da35 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -2432,6 +2432,9 @@ class MklLayoutRewritePass : public GraphOptimizationPass { csinfo_.mkl_conv2d_with_bias = "_MklConv2DWithBias"; csinfo_.mkl_conv2d_grad_filter_with_bias = "_MklConv2DBackpropFilterWithBias"; + csinfo_.mkl_pad_with_conv2d = "_MklPadWithConv2D"; + csinfo_.pad = "Pad"; + csinfo_.pad_with_conv2d = "__MklDummyPadWithConv2D"; csinfo_.relu = "Relu"; csinfo_.relu_grad = "ReluGrad"; csinfo_.tanh = "Tanh"; @@ -2508,6 +2511,8 @@ class MklLayoutRewritePass : public GraphOptimizationPass { rinfo_.push_back({csinfo_.mul, mkl_op_registry::GetMklOpName(csinfo_.mul), CopyAttrsDataType, AlwaysRewrite}); + rinfo_.push_back({csinfo_.pad_with_conv2d, csinfo_.mkl_pad_with_conv2d, + CopyAttrsPadWithConv2D, AlwaysRewrite}); rinfo_.push_back({csinfo_.relu, mkl_op_registry::GetMklOpName(csinfo_.relu), CopyAttrsDataType, AlwaysRewrite}); rinfo_.push_back({csinfo_.relu_grad, @@ -2546,6 +2551,10 @@ class MklLayoutRewritePass : public GraphOptimizationPass { minfo_.push_back({csinfo_.conv2d_grad_filter, csinfo_.bias_add_grad, csinfo_.conv2d_grad_filter_with_bias, GetConv2DBackpropFilterOrBiasAddGrad}); + minfo_.push_back({csinfo_.pad, csinfo_.conv2d, + csinfo_.pad_with_conv2d, GetPadOrConv2D}); + //TODO : Need to check if pad is with zero or not + // if is zero then replace, if not then do not replace } // Standard interface to run pass @@ -2628,7 +2637,10 @@ class MklLayoutRewritePass : public GraphOptimizationPass { string mkl_conv2d_grad_filter; string mkl_conv2d_grad_filter_with_bias; string mkl_conv2d_with_bias; + string mkl_pad_with_conv2d; string mul; + string pad; + string pad_with_conv2d; string relu; string relu_grad; string tanh; @@ -2734,6 +2746,7 @@ class MklLayoutRewritePass : public GraphOptimizationPass { // Helper function to merge different nodes Status MergeConv2DWithBiasAdd(std::unique_ptr* g, Node* m, Node* n); + Status MergePadWithConv2D(std::unique_ptr* g, Node* m, Node* n); Status MergeConv2DBackpropFilterWithBiasAddGrad(std::unique_ptr* g, Node* m, Node* n); @@ -2771,6 +2784,59 @@ class MklLayoutRewritePass : public GraphOptimizationPass { return n; } + // Find Pad or Conv2D node that can be merged with input node 'm'. + // If input 'm' is Pad, then check if there exists Conv2D node that can be + // merged with 'm'. If input 'm' is Conv2D, then check if there exists BiasAdd + // node that can be merged with 'm'. + static Node* GetPadOrConv2D(const Node* m) { + CHECK_NOTNULL(m); + Node* n = nullptr; + + if (m->type_string() == csinfo_.pad) { + // If m is Pad, then Conv2D is the output of Pad. + for (const Edge* e : m->out_edges()) { + if (!e->IsControlEdge() && + e->dst()->type_string() == csinfo_.conv2d) { + n = e->dst(); + break; + } + } + } else { + CHECK_EQ(m->type_string(), csinfo_.conv2d); + // If m is conv2D, Go over all input edges + // and search for Pad Node. + for (const Edge* e : m->in_edges()) { + if (!e->IsControlEdge() && + e->src()->type_string() == csinfo_.pad) { + n = e->src(); + break; + } + } + } + // Check if only VALID type of padding is used + // or not. + if (n != nullptr) { + const Node* conv_node; + if (m->type_string() == csinfo_.conv2d) + conv_node = m; + else + conv_node = n; + string padding; + TF_CHECK_OK(GetNodeAttr(conv_node->def(), "padding", &padding)); + if (padding != "VALID") + // Then do not merge. + // Only VALID type of padding in conv op can be + // merged with Pad op. + n = nullptr; + } + if (n == nullptr) { + VLOG(1) << "MklLayoutRewritePass: Could not find matching " + << "Pad and Conv2D node for merging. Input node: " + << m->DebugString(); + } + + return n; + } // Find Conv2DBackpropFilter or BiasAddGrad node that can be merged with input // node 'm'. If input 'm' is Conv2DBackpropFilter, then check if there exists // BiasAddGrad node that can be merged with 'm'. If input 'm' is BiasAddGrad, @@ -3090,6 +3156,9 @@ class MklLayoutRewritePass : public GraphOptimizationPass { static void CopyAttrsDataType(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsFusedBatchNorm(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsLRN(const Node* orig_node, NodeBuilder* nb); + static void CopyAttrsPadWithConv2D(const Node* orig_node, NodeBuilder* nb); + static void CopyAttrsFromPadAndConv2D(const Node* orig_node1, const Node* orig_node2, + NodeBuilder* nb); static void CopyAttrsPooling(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsReshape(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsSplit(const Node* orig_node, NodeBuilder* nb); @@ -3289,6 +3358,8 @@ int MklLayoutRewritePass::SetUpContiguousInputs( // 2nd input (slot 1) of _MklConv2D and _MklConv2DWithBias. for (const Edge* e : filter_node->out_edges()) { if ((e->dst()->type_string() == csinfo_.mkl_conv2d || + // add check for mkl_pad_with_conv2d + e->dst()->type_string() == csinfo_.mkl_pad_with_conv2d || e->dst()->type_string() == csinfo_.mkl_conv2d_with_bias) && e->dst_input() == kConv2DFilterInputSlotIdx /* filter is 2nd input of Conv2D and _MklConv2D. */) { @@ -3598,6 +3669,65 @@ void MklLayoutRewritePass::CopyAttrsConv2D(const Node* orig_node, nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu); } +//used in rinfo when replacing __MklDummyPadWithConv2D by _MklPadWithConv2D +void MklLayoutRewritePass::CopyAttrsPadWithConv2D(const Node* orig_node, + NodeBuilder* nb) { + DataType Tpaddings; + DataType T; + string data_format; + string padding; + std::vector strides; + bool use_cudnn_on_gpu; + + // Get all attributes from old node 1. + TF_CHECK_OK(GetNodeAttr(orig_node->def(), "T", &T)); + TF_CHECK_OK(GetNodeAttr(orig_node->def(), "strides", &strides)); + TF_CHECK_OK(GetNodeAttr(orig_node->def(), "padding", &padding)); + TF_CHECK_OK(GetNodeAttr(orig_node->def(), "data_format", &data_format)); + TF_CHECK_OK( + GetNodeAttr(orig_node->def(), "use_cudnn_on_gpu", &use_cudnn_on_gpu)); + TF_CHECK_OK(GetNodeAttr(orig_node->def(), "Tpaddings", &Tpaddings)); + + // Add attributes to new node. + nb->Attr("T", T); + nb->Attr("strides", strides); + nb->Attr("padding", padding); + nb->Attr("data_format", data_format); + nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu); + nb->Attr("Tpaddings", Tpaddings); +} + +//used with MergePadWithConv2D +void MklLayoutRewritePass::CopyAttrsFromPadAndConv2D(const Node* orig_node1, + const Node* orig_node2, NodeBuilder* nb) { + DataType Tpaddings; + DataType T; + string data_format; + string padding; + std::vector strides; + bool use_cudnn_on_gpu; + + // Get all attributes from old node 1. + TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "T", &T)); + TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "strides", &strides)); + TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "padding", &padding)); + TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "data_format", &data_format)); + TF_CHECK_OK( + GetNodeAttr(orig_node1->def(), "use_cudnn_on_gpu", &use_cudnn_on_gpu)); + // Get all attributes from old node 2. + TF_CHECK_OK(GetNodeAttr(orig_node2->def(), "Tpaddings", &Tpaddings)); + + // Add attributes to new node. + nb->Attr("T", T); + nb->Attr("strides", strides); + nb->Attr("padding", padding); + nb->Attr("data_format", data_format); + nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu); + + + // Add attributes to new node. + nb->Attr("Tpaddings", Tpaddings); +} void MklLayoutRewritePass::CopyAttrsAddN(const Node* orig_node, NodeBuilder* nb) { DataType T; @@ -3824,7 +3954,7 @@ Status MklLayoutRewritePass::MergeConv2DWithBiasAdd(std::unique_ptr* g, // If 'm' is BiasAdd, then 'n' is Conv2D. Since Conv2D feeds BiasAdd, // BiasAdd is successor node, and Conv2D predecessor node. Node* pred = m->type_string() == csinfo_.bias_add ? n : m; - Node* succ = m->type_string() == csinfo_.bias_add ? m : n; + Node* succ = m->type_string() == csinfo_.bias_add ? m : n; // 1. Get all attributes from input nodes. DataType T_pred, T_succ; @@ -3963,6 +4093,161 @@ Status MklLayoutRewritePass::MergeConv2DWithBiasAdd(std::unique_ptr* g, return Status::OK(); } +Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, + Node* m, Node* n) { + CHECK_EQ(((m->type_string() == csinfo_.pad && + n->type_string() == csinfo_.conv2d)) || + ((n->type_string() == csinfo_.pad && + m->type_string() == csinfo_.conv2d)), + true); + + // Conv2D is successor node, and Pad predecessor node. + Node* pred = m->type_string() == csinfo_.pad ? m : n; + Node* succ = m->type_string() == csinfo_.pad ? n : m; + + // 1. Get all attributes from input nodes. + DataType T_pred, T_succ; + string padding; + std::vector strides; + std::vector dilations; + string data_format_pred, data_format_succ; + bool use_cudnn_on_gnu; + TF_CHECK_OK(GetNodeAttr(pred->def(), "T", &T_pred)); + TF_CHECK_OK(GetNodeAttr(succ->def(), "T", &T_succ)); + TF_CHECK_OK(GetNodeAttr(succ->def(), "padding", &padding)); + TF_CHECK_OK(GetNodeAttr(succ->def(), "strides", &strides)); + TF_CHECK_OK(GetNodeAttr(succ->def(), "dilations", &dilations)); + // data format for pad is not available and not necessary, thus + // we dont need to match data format + // TF_CHECK_OK(GetNodeAttr(pred->def(), "data_format", &data_format_pred)); + TF_CHECK_OK(GetNodeAttr(succ->def(), "data_format", &data_format_succ)); + TF_CHECK_OK(GetNodeAttr(succ->def(), "use_cudnn_on_gpu", &use_cudnn_on_gnu)); + // We check to ensure that data formats of both succ and pred are same. + // We expect them to be same, so we can enforce this as assert. + // But assert can be too strict, so we enforce this as a check. + // If the check fails, then we do not merge two nodes. + // We also do same check for devices. + // if (data_format_pred != data_format_succ || T_pred != T_succ || + if (T_pred != T_succ || + pred->assigned_device_name() != succ->assigned_device_name() || + pred->def().device() != succ->def().device()) { + return Status(error::Code::INVALID_ARGUMENT, + "data_format or T attribute or devices of Conv2D and " + "Pad do not match. Will skip node merge optimization"); + } + + const int succ_num = succ->num_inputs(); + gtl::InlinedVector succ_control_edges; + gtl::InlinedVector, 4> succ_in(succ_num); + FillInputs(succ, &succ_control_edges, &succ_in); + + const int pred_num = pred->num_inputs(); + gtl::InlinedVector pred_control_edges; + gtl::InlinedVector, 4> pred_in(pred_num); + FillInputs(pred, &pred_control_edges, &pred_in); + + // We need to ensure that Pad only feeds to Conv2D (some other operator is + // not expecting output of Pad). If this is not the case, then we cannot + // merge Conv2D with Pad. + const int kFirstOutputSlot = 0; + for (const Edge* e : pred->out_edges()) { + if (e->src_output() == kFirstOutputSlot && e->dst() != succ) { + return Status(error::Code::INVALID_ARGUMENT, + "Pad does not feed to Conv2D, or " + "it feeds Conv2D but has multiple outputs. " + "Will skip node merge optimization"); + } + } + + // 2. Get inputs from both the nodes. ( ? ? Explanation of the following) + // Find the 2 inputs from the Pad and the Filter input from the Conv2D. + // Get operand 0, 1 of conv2D. + CHECK_EQ(pred->in_edges().size(), 2); // Pad must have 2 inputs. + // Get operand 1 of add_bias??? + // Conv2D must have 2 inputs: pad output and Filter + CHECK_EQ(succ->in_edges().size(), 2); + + // We will use the node name of Conv2D as the name of new node + // Build new node. We use same name as original node, but change the op + // name. + NodeBuilder nb(succ->name(), csinfo_.pad_with_conv2d); + nb.Input(pred_in[0].first, pred_in[0].second); // In1 (input data) of Pad + // pred_in[1] will be 2nd Tensorflow tensor for Conv2D. + nb.Input(succ_in[1].first, succ_in[1].second); // In2 (filter) of conv2d + // In1 of Conv2D is same as output of Pad. + // Thus, only need to add In2 of Conv2D + nb.Input(pred_in[1].first, pred_in[1].second); // In2 (paddings) of Pad + + // Copy attributes from Pad and conv2D to PadWithConv2D. + CopyAttrsFromPadAndConv2D(const_cast(succ), const_cast(pred), + &nb); + + // Copy the device assigned to old node to new node. + nb.Device(succ->def().device()); + + // Create node. + Node* new_node; + TF_CHECK_OK(nb.Finalize(&**g, &new_node)); + CHECK_NOTNULL(new_node); + + // Incoming data edges from 'pred' node and 'succ' node to new 'new_node' + // node are already copied in BuildNode. + // We handle control edges now. + for (const Edge* e : pred->in_edges()) { + if (e->IsControlEdge()) { + // Allow duplicate while adding control edge as it would fail (return + // NULL) if we try to add duplicate edge. + CHECK_NOTNULL((*g)->AddControlEdge(e->src(), new_node, true)); + } + } + for (const Edge* e : succ->in_edges()) { + if (e->IsControlEdge()) { + // Allow duplicate while adding control edge as it would fail (return + // NULL) if we try to add duplicate edge. + CHECK_NOTNULL((*g)->AddControlEdge(e->src(), new_node, true)); + } + } + + // Incoming edges are fixed, we will fix the outgoing edges now. + // First, we will fix outgoing control edges from 'pred' node. + for (const Edge* e : pred->out_edges()) { + if (e->IsControlEdge()) { + // Allow duplicate while adding control edge as it would fail (return + // NULL) if we try to add duplicate edge. + CHECK_NOTNULL((*g)->AddControlEdge(new_node, e->dst(), true)); + } + } + + // Second, we will fix outgoing control and data edges from 'succ' node. + for (const Edge* e : succ->out_edges()) { + if (e->IsControlEdge()) { + // Allow duplicate while adding control edge as it would fail (return + // NULL) if we try to add duplicate edge. + CHECK_NOTNULL((*g)->AddControlEdge(new_node, e->dst(), true)); + } else { + // Conv2D has only 1 output (at slot 0) and merged node also has only 1 + // output (at slot 0). + const int kPadWithConv2DOutputSlot = 0; + CHECK_NOTNULL((*g)->AddEdge(new_node, kPadWithConv2DOutputSlot, e->dst(), + e->dst_input())); + } + } + + // Copy device assigned to old node to new node. + // It's ok to use pred or succ as we have enforced a check that + // both have same device assigned. + new_node->set_assigned_device_name(pred->assigned_device_name()); + + VLOG(1) << "MklLayoutRewritePass: Merged old node:" << pred->DebugString() + << ", and node: " << succ->DebugString() + << ", into node:" << new_node->DebugString(); + + (*g)->RemoveNode(succ); + (*g)->RemoveNode(pred); + + return Status::OK(); +} + Status MklLayoutRewritePass::MergeConv2DBackpropFilterWithBiasAddGrad( std::unique_ptr* g, Node* m, Node* n) { CHECK_EQ(((m->type_string() == csinfo_.bias_add_grad && @@ -4096,6 +4381,12 @@ Status MklLayoutRewritePass::MergeNode(std::unique_ptr* g, Node* m, m->type_string() == csinfo_.conv2d))) { return this->MergeConv2DWithBiasAdd(g, m, n); } + if (((m->type_string() == csinfo_.pad && + n->type_string() == csinfo_.conv2d)) || + ((n->type_string() == csinfo_.pad && + m->type_string() == csinfo_.conv2d))) { + return this->MergePadWithConv2D(g, m, n); + } if (((m->type_string() == csinfo_.bias_add_grad && n->type_string() == csinfo_.conv2d_grad_filter)) || @@ -4207,9 +4498,10 @@ MklLayoutRewritePass::CheckForNodeRewrite(const Node* n) const { } // We make an exception for __MklDummyConv2DWithBias and - // __MklConv2DBackpropFilterWithBias since their names do not match Mkl node - // names. + // __MklConv2DBackpropFilterWithBias, __MklDummyPadWithConv2D since their names + // do not match Mkl node names. if (n->type_string() != csinfo_.conv2d_with_bias && + n->type_string() != csinfo_.pad_with_conv2d && n->type_string() != csinfo_.conv2d_grad_filter_with_bias && !mkl_op_registry::IsMklOp(mkl_op_registry::GetMklOpName(n->type_string()), T)) { diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index a41f5861af..020e3c9168 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -2012,6 +2012,76 @@ TEST_F(MklLayoutPassTest, Basic) { "A->C;A->D;B->C:1;B->D:1"); } +// Test set 0: Pad + Conv2D; padding is VALID +// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// D=input(filter), E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) +// After layout pass +// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'D' op: 'Input'}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'VALID' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'D'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" + "A:control->DMT/_0:control;A:control->DMT/_1:control;" + "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;Y->Z:1"); +} + +// Test set 0: Pad + Conv2D; padding is SAME +// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// D=input(filter), E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) +// After layout pass - No merging +TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Negative) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'D' op: 'Input'}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'SAME' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'D'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);C(Pad);D(Input);DMT/_0(Const);DMT/_1(Const);" + "E(_MklConv2D);Y(Input);Z(Zeta)|A->C;B->C:1;C->E;" + "C:control->DMT/_0:control;C:control->DMT/_1:control;" + "D->E:1;DMT/_0->E:2;DMT/_1->E:3;E->Z;Y->Z:1"); +} + // Test set 1: Conv2D + AddBias // C=Conv2D(A,B); E=BiasAdd(C,D); Z=Zeta(E,Y) diff --git a/tensorflow/core/kernels/BUILD b/tensorflow/core/kernels/BUILD index 6126e8b7ba..f14542068d 100644 --- a/tensorflow/core/kernels/BUILD +++ b/tensorflow/core/kernels/BUILD @@ -39,6 +39,7 @@ load( "cc_header_only_library", "if_not_windows", "if_override_eigen_strong_inline", + "tf_cc_test_mkl", ) load("@local_config_sycl//sycl:build_defs.bzl", "if_sycl") load("//tensorflow:tensorflow.bzl", "tf_cuda_cc_test") @@ -1129,6 +1130,7 @@ tf_cuda_cc_test( ], ) + tf_cc_test( name = "decode_wav_op_test", size = "small", @@ -6124,6 +6126,7 @@ tf_mkl_kernel_library( ] + if_mkl(["@mkl_dnn"]), ) + tf_mkl_kernel_library( name = "mkl_tfconv_op", prefix = "mkl_tfconv", @@ -6269,6 +6272,31 @@ tf_mkl_kernel_library( ], ) +tf_cc_test_mkl( + name = "mkl_fused_ops_test", + size = "small", + srcs = ["mkl_fused_ops_test.cc"], + linkstatic = 1, + deps = [ + ":mkl_conv_op", + ":mkl_tfconv_op", + ":conv_ops", + ":image", + ":ops_testutil", + ":ops_util", + "//tensorflow/cc:cc_ops", + "//tensorflow/core:core_cpu", + "//tensorflow/core:framework", + "//tensorflow/core:framework_internal", + "//tensorflow/core:lib", + "//tensorflow/core:protos_all_cc", + "//tensorflow/core:tensorflow", + "//tensorflow/core:test", + "//tensorflow/core:test_main", + "//tensorflow/core:testlib", + ] +) + # NOTE(lespeholt): This rule is deprecated, please use: # tensorflow/core/util/batch_util.h cc_library( diff --git a/tensorflow/core/kernels/mkl_conv_ops.cc b/tensorflow/core/kernels/mkl_conv_ops.cc index 62396eeb8b..d4ec831cf2 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.cc +++ b/tensorflow/core/kernels/mkl_conv_ops.cc @@ -753,9 +753,31 @@ class MklConv2DOp : public OpKernel { TensorFormat data_format_; }; + +#define REGISTER_MKL_CPU(T) \ + REGISTER_KERNEL_BUILDER(Name("_MklConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConv2DOp); \ + REGISTER_KERNEL_BUILDER(Name("_MklConv2DWithBias") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConv2DOp); \ + REGISTER_KERNEL_BUILDER(Name("__MklDummyConv2DWithBias") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklDummyOp); + +TF_CALL_float(REGISTER_MKL_CPU); +#undef REGISTER_MKL_CPU + #else -template +template class MklConv2DOp : public OpKernel { public: ~MklConv2DOp() {} @@ -814,6 +836,11 @@ class MklConv2DOp : public OpKernel { dilations, strides; memory::dims dst_dims_tf_order, dst_dims_mkl_order; + // If pad with conv2d fusion is enabled + if (padEnabled) { + PadWithConvFusion(context, padding_left, padding_right); + } + // Get shapes of input tensors in MKL-DNN order MklDnnConvUtil conv_utl(context, strides_, padding_, data_format_, dilations_); @@ -822,7 +849,7 @@ class MklConv2DOp : public OpKernel { conv_utl.GetConvFwdSizesInMklOrder( src_tf_shape, filter_tf_shape, &src_dims, &filter_dims, &strides, &dilations, &dst_dims_tf_order, &dst_dims_mkl_order, - &padding_left, &padding_right); + &padding_left, &padding_right, padEnabled); if (!context->status().ok()) return; // Check for corner case - if there is nothing to compute, return. @@ -869,7 +896,6 @@ class MklConv2DOp : public OpKernel { // MKLDNN dilation starts from 0. dilations[kDilationH] -= 1; dilations[kDilationW] -= 1; - // get a conv2d fwd from primitive pool MklConv2DFwdPrimitive* conv2d_fwd = nullptr; if (biasEnabled) { @@ -937,13 +963,53 @@ class MklConv2DOp : public OpKernel { errors::Aborted("Operation received an exception:", error_msg)); } } + + void PadWithConvFusion(OpKernelContext* context, memory::dims &padding_left, + memory::dims &padding_right){ + const Tensor& paddings_tf = MklGetInput(context, 2); + OP_REQUIRES(context, paddings_tf.dims() == 2, + errors::InvalidArgument("paddings must be 2-dimensional: ", + paddings_tf.shape().DebugString())); + Tpadding* paddings = nullptr; + // To get individual pad, need to flatten the tensor + paddings = static_cast(const_cast + (paddings_tf.flat().data())); + // For NHWC format: + // paddings[0], paddings[1], paddings[6], paddings[7] should be zero + // if the paddings_tf is [ [0, 0] [1,2] [3,4] [0,0] ] + // paddings = {0, 0, 1, 2, 3, 4, 0, 0} ; flat method is row major + // then, values are: top = 1, bottom =2, left=3, right=4 + // For NCHW format, + // paddings[0], paddings[1], paddings[2], paddings[3] should be zero + // similar explanation as NHWC format will apply. + string data_format = ToString(data_format_); + if(data_format == "NHWC"){ + pad_top = paddings[2]; + pad_bottom = paddings[3]; + pad_left = paddings[4]; + pad_right = paddings[5]; + } + else if (data_format == "NCHW"){ + pad_top = paddings[4]; + pad_bottom = paddings[5]; + pad_left = paddings[6]; + pad_right = paddings[7]; + } + // Create padding arrays for MKL DNN convolutions. + // MKL-DNN uses asymetric padding. + padding_left = {static_cast(pad_top), static_cast(pad_left)}; + padding_right = {static_cast(pad_bottom), static_cast(pad_right)}; + } private: std::vector strides_; std::vector dilations_; + int64 pad_top, pad_left; + int64 pad_bottom, pad_right; Padding padding_; TensorFormat data_format_; const int kInputIndex_Src = 0, kInputIndex_Filter = 1, kInputIndex_Bias = 2; + const int kInputIndex_Pad = 2; const int kOutputIndex_Dst = 0, kOutputIndex_Filter = 1; const int kDilationH = 0, kDilationW = 1; engine cpu_engine = engine(engine::cpu, 0); @@ -1036,26 +1102,44 @@ class MklConv2DOp : public OpKernel { } }; -#endif #define REGISTER_MKL_CPU(T) \ REGISTER_KERNEL_BUILDER(Name("_MklConv2D") \ .Device(DEVICE_CPU) \ .TypeConstraint("T") \ .Label(mkl_op_registry::kMklOpLabel), \ - MklConv2DOp); \ + MklConv2DOp); \ REGISTER_KERNEL_BUILDER(Name("_MklConv2DWithBias") \ .Device(DEVICE_CPU) \ .TypeConstraint("T") \ .Label(mkl_op_registry::kMklOpLabel), \ - MklConv2DOp); \ + MklConv2DOp); \ REGISTER_KERNEL_BUILDER(Name("__MklDummyConv2DWithBias") \ .Device(DEVICE_CPU) \ .TypeConstraint("T") \ .Label(mkl_op_registry::kMklOpLabel), \ + MklDummyOp); \ + REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .TypeConstraint("Tpaddings") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConv2DOp); \ + REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .TypeConstraint("Tpaddings") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConv2DOp); \ + REGISTER_KERNEL_BUILDER(Name("__MklDummyPadWithConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .TypeConstraint("Tpaddings") \ + .Label(mkl_op_registry::kMklOpLabel), \ MklDummyOp); TF_CALL_float(REGISTER_MKL_CPU); +#endif } // namespace tensorflow #endif // INTEL_MKL diff --git a/tensorflow/core/kernels/mkl_conv_ops.h b/tensorflow/core/kernels/mkl_conv_ops.h index 3f154ff33b..c6487a4512 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.h +++ b/tensorflow/core/kernels/mkl_conv_ops.h @@ -232,7 +232,7 @@ class MklDnnConvUtil { const memory::dims& strides, const memory::dims& dilations, memory::dims* output_dims_tf_order, memory::dims* output_dims_mkl_order, memory::dims* pad_l, - memory::dims* pad_r) { + memory::dims* pad_r, bool padEnabled=false) { CHECK_NOTNULL(output_dims_tf_order); CHECK_NOTNULL(output_dims_mkl_order); CHECK_NOTNULL(pad_l); @@ -268,7 +268,19 @@ class MklDnnConvUtil { GetWindowedOutputSizeVerboseV2(input_cols, filter_cols, dilation_cols, stride_cols, padding_, &out_cols, &pad_left, &pad_right)); - + // If padEnabled, i.e., pad and conv op are fused, then + // all pads are already passed from pad op through + // *pad_l and *pad_r + if(padEnabled) { + pad_top = static_cast((*pad_l)[0]); + pad_left = static_cast((*pad_l)[1]); + pad_bottom = static_cast((*pad_r)[0]); + pad_right = static_cast((*pad_r)[1]); + // update the out_rows and out_cols based on all + // sides of the pads coming from pad op. + out_rows = out_rows + (pad_top + pad_bottom ) / stride_rows; + out_cols = out_cols + (pad_left + pad_right ) / stride_cols; + } // Tensorflow output is in data_format order. (NHWC or NCHW) TensorShape out_shape = ShapeFromFormat(data_format_, out_batch, out_rows, out_cols, out_depth); @@ -283,8 +295,12 @@ class MklDnnConvUtil { *output_dims_mkl_order = mkldnn_sizes; // Now handle padding. MKL-DNN uses asymetric padding. - *pad_l = {static_cast(pad_top), static_cast(pad_left)}; - *pad_r = {static_cast(pad_bottom), static_cast(pad_right)}; + // But, if padEnabled, i.e., pad and conv op are fused, + // then, *pad_l and *pad_r are already set from pad op + if(!padEnabled) { + *pad_l = {static_cast(pad_top), static_cast(pad_left)}; + *pad_r = {static_cast(pad_bottom), static_cast(pad_right)}; + } } // Calculate output and pad size of forward Convolution operator. @@ -325,7 +341,7 @@ class MklDnnConvUtil { memory::dims* strides, memory::dims *dilations, memory::dims* output_dims_tf_order, memory::dims* output_dims_mkl_order, memory::dims* pad_l, - memory::dims* pad_r) { + memory::dims* pad_r, bool padEnabled=false) { CHECK_NOTNULL(input_dims); CHECK_NOTNULL(filter_dims); CHECK_NOTNULL(strides); @@ -344,7 +360,7 @@ class MklDnnConvUtil { GetOutputAndPadSizeInMklOrder(input_shape, filter_shape, *strides, *dilations, output_dims_tf_order, output_dims_mkl_order, - pad_l, pad_r); + pad_l, pad_r, padEnabled); if (!context_->status().ok()) return; } }; diff --git a/tensorflow/core/ops/nn_ops.cc b/tensorflow/core/ops/nn_ops.cc index f947d4c30d..8bb22a8372 100644 --- a/tensorflow/core/ops/nn_ops.cc +++ b/tensorflow/core/ops/nn_ops.cc @@ -1573,6 +1573,55 @@ NOTE Do not invoke this operator directly in Python. Graph rewrite pass is expected to invoke these operators. )doc"); +REGISTER_OP("__MklDummyPadWithConv2D") + .Input("input: T") + .Input("filter: T") + .Input("paddings: Tpaddings") + .Output("output: T") + .Attr("T: {half, float, double}") + .Attr("strides: list(int)") + .Attr("use_cudnn_on_gpu: bool = true") + .Attr(GetPaddingAttrString()) + .Attr(GetConvnetDataFormatAttrString()) + .Attr("dilations: list(int) = [1, 1, 1, 1]") + .Attr("Tpaddings: {int32, int64} = DT_INT32") + .SetShapeFn(shape_inference::Conv2DShape) + .Doc(R"doc( +Dummy node that enables fusing Pad and Conv2D operator for MKL. This node +does not perform anything. It is just created as an intermediate output of +merging Pad and Conv2D. + +NOTE Do not invoke this operator directly in Python. Graph rewrite pass is +expected to invoke these operators. +)doc"); + +REGISTER_OP("_MklPadWithConv2D") + .Input("input: T") + .Input("filter: T") + .Input("paddings: Tpaddings") + .Input("mkl_input: uint8") + .Input("mkl_filter: uint8") + .Input("mkl_paddings: uint8") + .Output("output: T") + .Output("filter_output: T") + .Output("mkl_output: uint8") + .Output("mkl_filter_output: uint8") + .Attr("T: {half, float, double}") + .Attr("strides: list(int)") + .Attr("use_cudnn_on_gpu: bool = true") + .Attr(GetPaddingAttrString()) + .Attr(GetConvnetDataFormatAttrString()) + .Attr("dilations: list(int) = [1, 1, 1, 1]") + .Attr("Tpaddings: {int32, int64} = DT_INT32") + .SetShapeFn(shape_inference::Conv2DShape) + .Doc(R"doc( +MKL version of Pad and Conv2D operator. Uses MKL DNN APIs to perform +Pad and 2D convolution to the output of convolution. + +NOTE Do not invoke this operator directly in Python. Graph rewrite pass is +expected to invoke these operators. +)doc"); + REGISTER_OP("_MklConv2DBackpropFilter") .Input("input: T") .Input("filter_sizes: int32") -- GitLab From dd63093a599081accfe2a2d2ca8c029d413a15d7 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Thu, 2 Aug 2018 08:43:06 -0700 Subject: [PATCH 061/596] adding unit test for pad+conv2d fusion op --- tensorflow/core/kernels/mkl_fused_ops_test.cc | 164 ++++++++++++++++++ 1 file changed, 164 insertions(+) create mode 100644 tensorflow/core/kernels/mkl_fused_ops_test.cc diff --git a/tensorflow/core/kernels/mkl_fused_ops_test.cc b/tensorflow/core/kernels/mkl_fused_ops_test.cc new file mode 100644 index 0000000000..216e8d0206 --- /dev/null +++ b/tensorflow/core/kernels/mkl_fused_ops_test.cc @@ -0,0 +1,164 @@ +/* Copyright 2015 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. +==============================================================================*/ +#ifdef INTEL_MKL +#ifndef INTEL_MKL_ML // We don't support fusion in MKL ML +#include "tensorflow/cc/ops/const_op.h" +#include "tensorflow/cc/ops/image_ops.h" +#include "tensorflow/cc/ops/nn_ops.h" +#include "tensorflow/cc/ops/standard_ops.h" +#include "tensorflow/core/common_runtime/kernel_benchmark_testlib.h" +#include "tensorflow/core/framework/fake_input.h" +#include "tensorflow/core/framework/node_def_builder.h" +#include "tensorflow/core/framework/tensor.h" +#include "tensorflow/core/framework/types.pb.h" +#include "tensorflow/core/kernels/conv_ops_gpu.h" +#include "tensorflow/core/kernels/ops_testutil.h" +#include "tensorflow/core/kernels/ops_util.h" +#include "tensorflow/core/platform/test.h" +#include "tensorflow/core/platform/test_benchmark.h" +#include "tensorflow/core/platform/types.h" +#include "tensorflow/core/public/session.h" + +namespace tensorflow { + +// Helper class for converting MKL tesnors to TF tensor and comparing to +// expected values + +const uint8 dummy_tensor[] = {0, 0, 0, 0, 0, 0, 0, 0}; +const TensorShape dummy_shape({8}); + +class ConvMklToTF : public OpsTestBase { + public: + template + void ConvertAndCompare(DataType dtype, const Tensor& first, + const Tensor& second, const Tensor& expected) { + // Create an MKL to TF conversion node and execute it + TF_EXPECT_OK(NodeDefBuilder("mkl_to_tf_op", "_MklToTf") + .Input(FakeInput(dtype)) // Input + .Input(FakeInput(DT_UINT8)) // Mkl second tensor + .Attr("T", dtype) + .Attr("_kernel", "MklOp") + .Finalize(node_def())); + TF_EXPECT_OK(InitOp()); + AddInputFromArray(first.shape(), first.flat()); + AddInputFromArray(second.shape(), second.flat()); + TF_ASSERT_OK(RunOpKernel()); + + const Tensor& output = *GetOutput(0); + test::ExpectTensorNear(expected, output, 1e-5); + } + void TestBody(){}; +}; + +// Testing fusion of pad and convolution + +class FusedPadConvOpTest : public OpsTestBase { + public: + template + void Run(DataType dtype, Tensor& image, Tensor& filter, Tensor& padding, + Tensor& expected, const string data_format) { + const int stride = 1; + + // Create a fused pad+conv2d node + TF_EXPECT_OK(NodeDefBuilder("fused_pad_conv_op", "_MklPadWithConv2D") + .Input(FakeInput(dtype)) // Input + .Input(FakeInput(dtype)) // Filter + .Input(FakeInput(DT_INT32)) // Padding + .Input(FakeInput(DT_UINT8)) // MKl second tensor + .Input(FakeInput(DT_UINT8)) // MKl second tensor + .Input(FakeInput(DT_UINT8)) // MKl second tensor + .Attr("padding", "VALID") + .Attr("data_format", data_format) + .Attr("T", dtype) + .Attr("strides", {1, stride, stride, 1}) + .Attr("_kernel", "MklOp") + .Finalize(node_def())); + TF_EXPECT_OK(InitOp()); + + // Setting up inputs and execute + AddInputFromArray(image.shape(), image.flat()); + AddInputFromArray(filter.shape(), filter.flat()); + AddInputFromArray(padding.shape(), padding.flat()); + AddInputFromArray(dummy_shape, dummy_tensor); + AddInputFromArray(dummy_shape, dummy_tensor); + AddInputFromArray(dummy_shape, dummy_tensor); + TF_ASSERT_OK(RunOpKernel()); + + // Compare output to expected results + const Tensor& first = *GetOutput(0); + const Tensor& second = *GetOutput(2); + ConvMklToTF conv_comp; + conv_comp.ConvertAndCompare(dtype, first, second, expected); + } +}; + +TEST_F(FusedPadConvOpTest, PaddingConvTest) { + const int depth = 1; + const int image_width = 4; + const int image_height = 3; + const int image_batch_count = 1; + Tensor image(DT_FLOAT, {image_batch_count, image_height, image_width, depth}); + test::FillValues(&image, {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12}); + + const int filter_size = 3; + const int filter_count = 1; + Tensor filter(DT_FLOAT, {filter_size, filter_size, depth, filter_count}); + test::FillValues(&filter, {1, 4, 7, 2, 5, 8, 3, 6, 9}); + + const int padding_height = 4; + const int padding_width = 2; + Tensor padding(DT_INT32, {padding_height, padding_width}); + test::FillValues(&padding, {0, 0, 3, 4, 1, 2, 0, 0}); + + Tensor expected(DT_FLOAT, TensorShape({1, 8, 5, 1})); + test::FillValues( + &expected, + {0, 0, 0, 0, 0, 24, 42, 60, 33, 12, 105, 150, 183, 95, + 32, 235, 312, 357, 178, 56, 187, 234, 261, 121, 32, 106, 126, 138, + 59, 12, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}); + + Run(DT_FLOAT, image, filter, padding, expected, "NHWC"); +} + +TEST_F(FusedPadConvOpTest, PaddingConvTestNchw) { + const int depth = 1; + const int image_width = 4; + const int image_height = 3; + const int image_batch_count = 1; + Tensor image(DT_FLOAT, {image_batch_count, depth, image_height, image_width}); + test::FillValues(&image, {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12}); + + const int filter_size = 3; + const int filter_count = 1; + Tensor filter(DT_FLOAT, {filter_size, filter_size, depth, filter_count}); + test::FillValues(&filter, {1, 4, 7, 2, 5, 8, 3, 6, 9}); + + const int padding_height = 4; + const int padding_width = 2; + Tensor padding(DT_INT32, {padding_height, padding_width}); + test::FillValues(&padding, {0, 0, 0, 0, 3, 4, 1, 2}); + + Tensor expected(DT_FLOAT, TensorShape({1, 1, 8, 5})); + test::FillValues( + &expected, + {0, 0, 0, 0, 0, 24, 42, 60, 33, 12, 105, 150, 183, 95, + 32, 235, 312, 357, 178, 56, 187, 234, 261, 121, 32, 106, 126, 138, + 59, 12, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}); + + Run(DT_FLOAT, image, filter, padding, expected, "NCHW"); +} +} // namespace tensorflow +#endif // INTEL_MKL_ML +#endif // INTEL_MKL -- GitLab From 7f94025fe72369117bf32d69156f0bd947402c96 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Fri, 10 Aug 2018 20:00:27 -0700 Subject: [PATCH 062/596] Addressing the reviews for fused PAD and Conv2d PR --- tensorflow/core/graph/mkl_layout_pass.cc | 75 +++++++++---------- tensorflow/core/kernels/BUILD | 22 +++++- tensorflow/core/kernels/mkl_conv_ops.cc | 8 +- tensorflow/core/kernels/mkl_conv_ops.h | 12 ++- tensorflow/core/kernels/mkl_fused_ops_test.cc | 8 +- 5 files changed, 69 insertions(+), 56 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index d0abe5da35..1e85b50d99 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -2786,50 +2786,48 @@ class MklLayoutRewritePass : public GraphOptimizationPass { // Find Pad or Conv2D node that can be merged with input node 'm'. // If input 'm' is Pad, then check if there exists Conv2D node that can be - // merged with 'm'. If input 'm' is Conv2D, then check if there exists BiasAdd + // merged with 'm'. If input 'm' is Conv2D, then check if there exists Pad // node that can be merged with 'm'. static Node* GetPadOrConv2D(const Node* m) { CHECK_NOTNULL(m); Node* n = nullptr; + const Node* conv_node; if (m->type_string() == csinfo_.pad) { // If m is Pad, then Conv2D is the output of Pad. for (const Edge* e : m->out_edges()) { if (!e->IsControlEdge() && e->dst()->type_string() == csinfo_.conv2d) { n = e->dst(); + conv_node = n; break; } } } else { CHECK_EQ(m->type_string(), csinfo_.conv2d); - // If m is conv2D, Go over all input edges + // If m is conv2D, Go over all input edges // and search for Pad Node. for (const Edge* e : m->in_edges()) { if (!e->IsControlEdge() && e->src()->type_string() == csinfo_.pad) { n = e->src(); + conv_node = m; break; } } } - // Check if only VALID type of padding is used - // or not. + // Check if only VALID type of padding is used + // or not. if (n != nullptr) { - const Node* conv_node; - if (m->type_string() == csinfo_.conv2d) - conv_node = m; - else - conv_node = n; string padding; TF_CHECK_OK(GetNodeAttr(conv_node->def(), "padding", &padding)); - if (padding != "VALID") - // Then do not merge. - // Only VALID type of padding in conv op can be + if (padding != "VALID") + // Then do not merge. + // Only VALID type of padding in conv op can be // merged with Pad op. n = nullptr; } - if (n == nullptr) { + else { VLOG(1) << "MklLayoutRewritePass: Could not find matching " << "Pad and Conv2D node for merging. Input node: " << m->DebugString(); @@ -3669,7 +3667,7 @@ void MklLayoutRewritePass::CopyAttrsConv2D(const Node* orig_node, nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu); } -//used in rinfo when replacing __MklDummyPadWithConv2D by _MklPadWithConv2D +// Used in rinfo when replacing __MklDummyPadWithConv2D by _MklPadWithConv2D void MklLayoutRewritePass::CopyAttrsPadWithConv2D(const Node* orig_node, NodeBuilder* nb) { DataType Tpaddings; @@ -3677,11 +3675,13 @@ void MklLayoutRewritePass::CopyAttrsPadWithConv2D(const Node* orig_node, string data_format; string padding; std::vector strides; + std::vector dilations; bool use_cudnn_on_gpu; - // Get all attributes from old node 1. + // Get all attributes from old node. TF_CHECK_OK(GetNodeAttr(orig_node->def(), "T", &T)); TF_CHECK_OK(GetNodeAttr(orig_node->def(), "strides", &strides)); + TF_CHECK_OK(GetNodeAttr(orig_node->def(), "dilations", &dilations)); TF_CHECK_OK(GetNodeAttr(orig_node->def(), "padding", &padding)); TF_CHECK_OK(GetNodeAttr(orig_node->def(), "data_format", &data_format)); TF_CHECK_OK( @@ -3691,13 +3691,14 @@ void MklLayoutRewritePass::CopyAttrsPadWithConv2D(const Node* orig_node, // Add attributes to new node. nb->Attr("T", T); nb->Attr("strides", strides); + nb->Attr("dilations", dilations); nb->Attr("padding", padding); nb->Attr("data_format", data_format); nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu); nb->Attr("Tpaddings", Tpaddings); } -//used with MergePadWithConv2D +// Used with MergePadWithConv2D void MklLayoutRewritePass::CopyAttrsFromPadAndConv2D(const Node* orig_node1, const Node* orig_node2, NodeBuilder* nb) { DataType Tpaddings; @@ -3705,11 +3706,13 @@ void MklLayoutRewritePass::CopyAttrsFromPadAndConv2D(const Node* orig_node1, string data_format; string padding; std::vector strides; + std::vector dilations; bool use_cudnn_on_gpu; // Get all attributes from old node 1. TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "T", &T)); TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "strides", &strides)); + TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "dilations", &dilations)); TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "padding", &padding)); TF_CHECK_OK(GetNodeAttr(orig_node1->def(), "data_format", &data_format)); TF_CHECK_OK( @@ -3720,12 +3723,10 @@ void MklLayoutRewritePass::CopyAttrsFromPadAndConv2D(const Node* orig_node1, // Add attributes to new node. nb->Attr("T", T); nb->Attr("strides", strides); + nb->Attr("dilations", dilations); nb->Attr("padding", padding); nb->Attr("data_format", data_format); nb->Attr("use_cudnn_on_gpu", use_cudnn_on_gpu); - - - // Add attributes to new node. nb->Attr("Tpaddings", Tpaddings); } void MklLayoutRewritePass::CopyAttrsAddN(const Node* orig_node, @@ -3954,7 +3955,7 @@ Status MklLayoutRewritePass::MergeConv2DWithBiasAdd(std::unique_ptr* g, // If 'm' is BiasAdd, then 'n' is Conv2D. Since Conv2D feeds BiasAdd, // BiasAdd is successor node, and Conv2D predecessor node. Node* pred = m->type_string() == csinfo_.bias_add ? n : m; - Node* succ = m->type_string() == csinfo_.bias_add ? m : n; + Node* succ = m->type_string() == csinfo_.bias_add ? m : n; // 1. Get all attributes from input nodes. DataType T_pred, T_succ; @@ -4095,11 +4096,10 @@ Status MklLayoutRewritePass::MergeConv2DWithBiasAdd(std::unique_ptr* g, Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, Node* m, Node* n) { - CHECK_EQ(((m->type_string() == csinfo_.pad && + CHECK(((m->type_string() == csinfo_.pad && n->type_string() == csinfo_.conv2d)) || ((n->type_string() == csinfo_.pad && - m->type_string() == csinfo_.conv2d)), - true); + m->type_string() == csinfo_.conv2d))); // Conv2D is successor node, and Pad predecessor node. Node* pred = m->type_string() == csinfo_.pad ? m : n; @@ -4117,22 +4117,18 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, TF_CHECK_OK(GetNodeAttr(succ->def(), "padding", &padding)); TF_CHECK_OK(GetNodeAttr(succ->def(), "strides", &strides)); TF_CHECK_OK(GetNodeAttr(succ->def(), "dilations", &dilations)); - // data format for pad is not available and not necessary, thus - // we dont need to match data format - // TF_CHECK_OK(GetNodeAttr(pred->def(), "data_format", &data_format_pred)); + // Data format for pad is not available and not necessary, thus + // dont need to match data format for Pad TF_CHECK_OK(GetNodeAttr(succ->def(), "data_format", &data_format_succ)); TF_CHECK_OK(GetNodeAttr(succ->def(), "use_cudnn_on_gpu", &use_cudnn_on_gnu)); - // We check to ensure that data formats of both succ and pred are same. - // We expect them to be same, so we can enforce this as assert. - // But assert can be too strict, so we enforce this as a check. - // If the check fails, then we do not merge two nodes. - // We also do same check for devices. - // if (data_format_pred != data_format_succ || T_pred != T_succ || + // Check if the data types and devices of both succ and pred are the same. + // Assert is not used, because it can be too strict. + // Don't need to check for data formats because it is not available in Pad. if (T_pred != T_succ || pred->assigned_device_name() != succ->assigned_device_name() || pred->def().device() != succ->def().device()) { return Status(error::Code::INVALID_ARGUMENT, - "data_format or T attribute or devices of Conv2D and " + "T attribute or devices of Conv2D and " "Pad do not match. Will skip node merge optimization"); } @@ -4159,11 +4155,10 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, } } - // 2. Get inputs from both the nodes. ( ? ? Explanation of the following) - // Find the 2 inputs from the Pad and the Filter input from the Conv2D. - // Get operand 0, 1 of conv2D. - CHECK_EQ(pred->in_edges().size(), 2); // Pad must have 2 inputs. - // Get operand 1 of add_bias??? + // 2. Get inputs from both the nodes. + + // Pad must have 2 inputs: "input" and paddings. + CHECK_EQ(pred->in_edges().size(), 2); // Conv2D must have 2 inputs: pad output and Filter CHECK_EQ(succ->in_edges().size(), 2); @@ -4497,8 +4492,8 @@ MklLayoutRewritePass::CheckForNodeRewrite(const Node* n) const { return nullptr; } - // We make an exception for __MklDummyConv2DWithBias and - // __MklConv2DBackpropFilterWithBias, __MklDummyPadWithConv2D since their names + // We make an exception for __MklDummyConv2DWithBias, + // __MklConv2DBackpropFilterWithBias, and __MklDummyPadWithConv2D since their names // do not match Mkl node names. if (n->type_string() != csinfo_.conv2d_with_bias && n->type_string() != csinfo_.pad_with_conv2d && diff --git a/tensorflow/core/kernels/BUILD b/tensorflow/core/kernels/BUILD index f14542068d..b057b78ace 100644 --- a/tensorflow/core/kernels/BUILD +++ b/tensorflow/core/kernels/BUILD @@ -22,6 +22,7 @@ package_group( "//learning/brain/research/sparse_matrix/...", "//learning/faster_training/...", "//tensorflow/...", + "//third_party/car/...", ], ) @@ -783,7 +784,7 @@ tf_kernel_library( tf_kernel_library( name = "quantize_and_dequantize_op", prefix = "quantize_and_dequantize_op", - deps = ARRAY_DEPS, + deps = ARRAY_DEPS + [":cwise_op"], ) tf_kernel_library( @@ -1130,7 +1131,6 @@ tf_cuda_cc_test( ], ) - tf_cc_test( name = "decode_wav_op_test", size = "small", @@ -2855,6 +2855,8 @@ tf_kernel_library( srcs = [] + if_mkl([ "mkl_batch_matmul_op.cc", ]), + # *impl.h are excluded by default from the CPU build, add explicitly. + hdrs = ["batch_matmul_op_impl.h"], # Override EIGEN_STRONG_INLINE to inline when --define=override_eigen_strong_inline=true, # to avoid long compiling time. See https://github.com/tensorflow/tensorflow/issues/10521 copts = if_override_eigen_strong_inline(["/DEIGEN_STRONG_INLINE=inline"]), @@ -3791,7 +3793,7 @@ tf_kernel_library( "spacetodepth_op.h", "spacetodepth_op_gpu.cu.cc", ], - visibility = ["//visibility:private"], + visibility = [":friends"], deps = [ "//tensorflow/core:framework", "//tensorflow/core:lib", @@ -4888,6 +4890,7 @@ filegroup( "fill_functor.cc", "fill_functor.h", "function_ops.cc", + "function_ops.h", "gather_functor.h", "gather_nd_op.cc", "gather_nd_op.h", @@ -5379,6 +5382,18 @@ cc_library( alwayslink = 1, ) +cc_library( + name = "android_whole_file_read_ops", + srcs = if_android(["whole_file_read_ops.cc"]), + copts = tf_copts(), + linkopts = ["-ldl"], + visibility = ["//visibility:public"], + deps = [ + "//tensorflow/core:android_tensorflow_lib_lite", + ], + alwayslink = 1, +) + # Quantization-specific OpKernels tf_kernel_library( @@ -6126,7 +6141,6 @@ tf_mkl_kernel_library( ] + if_mkl(["@mkl_dnn"]), ) - tf_mkl_kernel_library( name = "mkl_tfconv_op", prefix = "mkl_tfconv", diff --git a/tensorflow/core/kernels/mkl_conv_ops.cc b/tensorflow/core/kernels/mkl_conv_ops.cc index d4ec831cf2..b5ae312fa5 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.cc +++ b/tensorflow/core/kernels/mkl_conv_ops.cc @@ -979,9 +979,11 @@ class MklConv2DOp : public OpKernel { // if the paddings_tf is [ [0, 0] [1,2] [3,4] [0,0] ] // paddings = {0, 0, 1, 2, 3, 4, 0, 0} ; flat method is row major // then, values are: top = 1, bottom =2, left=3, right=4 - // For NCHW format, + // For NCHW format: // paddings[0], paddings[1], paddings[2], paddings[3] should be zero // similar explanation as NHWC format will apply. + int64 pad_top, pad_left; + int64 pad_bottom, pad_right; string data_format = ToString(data_format_); if(data_format == "NHWC"){ pad_top = paddings[2]; @@ -1004,8 +1006,6 @@ class MklConv2DOp : public OpKernel { private: std::vector strides_; std::vector dilations_; - int64 pad_top, pad_left; - int64 pad_bottom, pad_right; Padding padding_; TensorFormat data_format_; const int kInputIndex_Src = 0, kInputIndex_Filter = 1, kInputIndex_Bias = 2; @@ -1139,7 +1139,7 @@ class MklConv2DOp : public OpKernel { MklDummyOp); TF_CALL_float(REGISTER_MKL_CPU); -#endif +#endif // INTEL_MKL_ML } // namespace tensorflow #endif // INTEL_MKL diff --git a/tensorflow/core/kernels/mkl_conv_ops.h b/tensorflow/core/kernels/mkl_conv_ops.h index c6487a4512..aae4d767a2 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.h +++ b/tensorflow/core/kernels/mkl_conv_ops.h @@ -281,6 +281,14 @@ class MklDnnConvUtil { out_rows = out_rows + (pad_top + pad_bottom ) / stride_rows; out_cols = out_cols + (pad_left + pad_right ) / stride_cols; } + // Handle padding. MKL-DNN uses asymetric padding. + // But, if padEnabled, i.e., pad and conv op are fused, + // then, *pad_l and *pad_r are already set from pad op. + // In that case they need not set here. + else { + *pad_l = {static_cast(pad_top), static_cast(pad_left)}; + *pad_r = {static_cast(pad_bottom), static_cast(pad_right)}; + } // Tensorflow output is in data_format order. (NHWC or NCHW) TensorShape out_shape = ShapeFromFormat(data_format_, out_batch, out_rows, out_cols, out_depth); @@ -297,10 +305,6 @@ class MklDnnConvUtil { // Now handle padding. MKL-DNN uses asymetric padding. // But, if padEnabled, i.e., pad and conv op are fused, // then, *pad_l and *pad_r are already set from pad op - if(!padEnabled) { - *pad_l = {static_cast(pad_top), static_cast(pad_left)}; - *pad_r = {static_cast(pad_bottom), static_cast(pad_right)}; - } } // Calculate output and pad size of forward Convolution operator. diff --git a/tensorflow/core/kernels/mkl_fused_ops_test.cc b/tensorflow/core/kernels/mkl_fused_ops_test.cc index 216e8d0206..e408886861 100644 --- a/tensorflow/core/kernels/mkl_fused_ops_test.cc +++ b/tensorflow/core/kernels/mkl_fused_ops_test.cc @@ -1,4 +1,4 @@ -/* Copyright 2015 The TensorFlow Authors. All Rights Reserved. +/* 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. @@ -33,11 +33,11 @@ limitations under the License. namespace tensorflow { -// Helper class for converting MKL tesnors to TF tensor and comparing to +// Helper class for converting MKL tesnors to TF tensors and comparing to // expected values -const uint8 dummy_tensor[] = {0, 0, 0, 0, 0, 0, 0, 0}; -const TensorShape dummy_shape({8}); +static const uint8 dummy_tensor[] = {0, 0, 0, 0, 0, 0, 0, 0}; +static const TensorShape dummy_shape({8}); class ConvMklToTF : public OpsTestBase { public: -- GitLab From 819afabbeda709a94894c894515b62c85d236d50 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Mon, 13 Aug 2018 10:44:36 -0700 Subject: [PATCH 063/596] modifying the ifdef INTEL_MKL_ML to INTEL_MKL_ML_ONLY --- tensorflow/core/kernels/mkl_fused_ops_test.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/kernels/mkl_fused_ops_test.cc b/tensorflow/core/kernels/mkl_fused_ops_test.cc index e408886861..900325ac91 100644 --- a/tensorflow/core/kernels/mkl_fused_ops_test.cc +++ b/tensorflow/core/kernels/mkl_fused_ops_test.cc @@ -13,7 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. ==============================================================================*/ #ifdef INTEL_MKL -#ifndef INTEL_MKL_ML // We don't support fusion in MKL ML +#ifndef INTEL_MKL_ML_ONLY // We don't support fusion in MKL ML #include "tensorflow/cc/ops/const_op.h" #include "tensorflow/cc/ops/image_ops.h" #include "tensorflow/cc/ops/nn_ops.h" @@ -160,5 +160,5 @@ TEST_F(FusedPadConvOpTest, PaddingConvTestNchw) { Run(DT_FLOAT, image, filter, padding, expected, "NCHW"); } } // namespace tensorflow -#endif // INTEL_MKL_ML +#endif // INTEL_MKL_ML_ONLY #endif // INTEL_MKL -- GitLab From 6b292c27c7ad09a89c8b75c2505e6472b533a4e1 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Mon, 13 Aug 2018 15:03:18 -0700 Subject: [PATCH 064/596] formatted as Clang format for Google code compliance, replaced directive INTEL_MKL_ML by INTEL_MKL_ML_ONLY, and merged with master --- tensorflow/core/graph/mkl_layout_pass.cc | 69 +++++++++---------- tensorflow/core/graph/mkl_layout_pass_test.cc | 19 ++--- tensorflow/core/kernels/mkl_conv_ops.cc | 47 +++++++------ tensorflow/core/kernels/mkl_conv_ops.h | 42 +++++------ 4 files changed, 86 insertions(+), 91 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index 65b999b193..84e8ea8f70 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -2551,10 +2551,10 @@ class MklLayoutRewritePass : public GraphOptimizationPass { minfo_.push_back({csinfo_.conv2d_grad_filter, csinfo_.bias_add_grad, csinfo_.conv2d_grad_filter_with_bias, GetConv2DBackpropFilterOrBiasAddGrad}); - minfo_.push_back({csinfo_.pad, csinfo_.conv2d, - csinfo_.pad_with_conv2d, GetPadOrConv2D}); - //TODO : Need to check if pad is with zero or not - // if is zero then replace, if not then do not replace + minfo_.push_back( + {csinfo_.pad, csinfo_.conv2d, csinfo_.pad_with_conv2d, GetPadOrConv2D}); + // Merge Pad and Conv2d, only if the pad op is "Pad" + // Doesn't merge if pad op is "PadV2" or "MirrorPad" } // Standard interface to run pass @@ -2792,42 +2792,39 @@ class MklLayoutRewritePass : public GraphOptimizationPass { CHECK_NOTNULL(m); Node* n = nullptr; - const Node* conv_node; + const Node* conv_node; if (m->type_string() == csinfo_.pad) { // If m is Pad, then Conv2D is the output of Pad. for (const Edge* e : m->out_edges()) { - if (!e->IsControlEdge() && - e->dst()->type_string() == csinfo_.conv2d) { + if (!e->IsControlEdge() && e->dst()->type_string() == csinfo_.conv2d) { n = e->dst(); - conv_node = n; + conv_node = n; break; } } } else { CHECK_EQ(m->type_string(), csinfo_.conv2d); - // If m is conv2D, Go over all input edges + // If m is conv2D, Go over all input edges // and search for Pad Node. for (const Edge* e : m->in_edges()) { - if (!e->IsControlEdge() && - e->src()->type_string() == csinfo_.pad) { + if (!e->IsControlEdge() && e->src()->type_string() == csinfo_.pad) { n = e->src(); - conv_node = m; + conv_node = m; break; } } } - // Check if only VALID type of padding is used - // or not. + // Check if only VALID type of padding is used + // or not. if (n != nullptr) { string padding; TF_CHECK_OK(GetNodeAttr(conv_node->def(), "padding", &padding)); - if (padding != "VALID") - // Then do not merge. - // Only VALID type of padding in conv op can be + if (padding != "VALID") + // Then do not merge. + // Only VALID type of padding in conv op can be // merged with Pad op. n = nullptr; - } - else { + } else { VLOG(1) << "MklLayoutRewritePass: Could not find matching " << "Pad and Conv2D node for merging. Input node: " << m->DebugString(); @@ -3155,7 +3152,8 @@ class MklLayoutRewritePass : public GraphOptimizationPass { static void CopyAttrsFusedBatchNorm(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsLRN(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsPadWithConv2D(const Node* orig_node, NodeBuilder* nb); - static void CopyAttrsFromPadAndConv2D(const Node* orig_node1, const Node* orig_node2, + static void CopyAttrsFromPadAndConv2D(const Node* orig_node1, + const Node* orig_node2, NodeBuilder* nb); static void CopyAttrsPooling(const Node* orig_node, NodeBuilder* nb); static void CopyAttrsReshape(const Node* orig_node, NodeBuilder* nb); @@ -3356,7 +3354,7 @@ int MklLayoutRewritePass::SetUpContiguousInputs( // 2nd input (slot 1) of _MklConv2D and _MklConv2DWithBias. for (const Edge* e : filter_node->out_edges()) { if ((e->dst()->type_string() == csinfo_.mkl_conv2d || - // add check for mkl_pad_with_conv2d + // add check for mkl_pad_with_conv2d e->dst()->type_string() == csinfo_.mkl_pad_with_conv2d || e->dst()->type_string() == csinfo_.mkl_conv2d_with_bias) && e->dst_input() == kConv2DFilterInputSlotIdx @@ -3669,7 +3667,7 @@ void MklLayoutRewritePass::CopyAttrsConv2D(const Node* orig_node, // Used in rinfo when replacing __MklDummyPadWithConv2D by _MklPadWithConv2D void MklLayoutRewritePass::CopyAttrsPadWithConv2D(const Node* orig_node, - NodeBuilder* nb) { + NodeBuilder* nb) { DataType Tpaddings; DataType T; string data_format; @@ -3700,7 +3698,8 @@ void MklLayoutRewritePass::CopyAttrsPadWithConv2D(const Node* orig_node, // Used with MergePadWithConv2D void MklLayoutRewritePass::CopyAttrsFromPadAndConv2D(const Node* orig_node1, - const Node* orig_node2, NodeBuilder* nb) { + const Node* orig_node2, + NodeBuilder* nb) { DataType Tpaddings; DataType T; string data_format; @@ -4095,12 +4094,12 @@ Status MklLayoutRewritePass::MergeConv2DWithBiasAdd(std::unique_ptr* g, } Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, - Node* m, Node* n) { + Node* m, Node* n) { CHECK(((m->type_string() == csinfo_.pad && - n->type_string() == csinfo_.conv2d)) || - ((n->type_string() == csinfo_.pad && - m->type_string() == csinfo_.conv2d))); - + n->type_string() == csinfo_.conv2d)) || + ((n->type_string() == csinfo_.pad && + m->type_string() == csinfo_.conv2d))); + // Conv2D is successor node, and Pad predecessor node. Node* pred = m->type_string() == csinfo_.pad ? m : n; Node* succ = m->type_string() == csinfo_.pad ? n : m; @@ -4158,7 +4157,7 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, // 2. Get inputs from both the nodes. // Pad must have 2 inputs: "input" and paddings. - CHECK_EQ(pred->in_edges().size(), 2); + CHECK_EQ(pred->in_edges().size(), 2); // Conv2D must have 2 inputs: pad output and Filter CHECK_EQ(succ->in_edges().size(), 2); @@ -4174,8 +4173,8 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, nb.Input(pred_in[1].first, pred_in[1].second); // In2 (paddings) of Pad // Copy attributes from Pad and conv2D to PadWithConv2D. - CopyAttrsFromPadAndConv2D(const_cast(succ), const_cast(pred), - &nb); + CopyAttrsFromPadAndConv2D(const_cast(succ), + const_cast(pred), &nb); // Copy the device assigned to old node to new node. nb.Device(succ->def().device()); @@ -4186,7 +4185,7 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, CHECK_NOTNULL(new_node); // Incoming data edges from 'pred' node and 'succ' node to new 'new_node' - // node are already copied in BuildNode. + // node are already copied in BuildNode. // We handle control edges now. for (const Edge* e : pred->in_edges()) { if (e->IsControlEdge()) { @@ -4493,10 +4492,10 @@ MklLayoutRewritePass::CheckForNodeRewrite(const Node* n) const { } // We make an exception for __MklDummyConv2DWithBias, - // __MklConv2DBackpropFilterWithBias, and __MklDummyPadWithConv2D since their names - // do not match Mkl node names. + // __MklConv2DBackpropFilterWithBias, and __MklDummyPadWithConv2D since their + // names do not match Mkl node names. if (n->type_string() != csinfo_.conv2d_with_bias && - n->type_string() != csinfo_.pad_with_conv2d && + n->type_string() != csinfo_.pad_with_conv2d && n->type_string() != csinfo_.conv2d_grad_filter_with_bias && !mkl_op_registry::IsMklOp(mkl_op_registry::GetMklOpName(n->type_string()), T)) { diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index 2925b1bde0..d1b39ceeca 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -2013,11 +2013,11 @@ TEST_F(MklLayoutPassTest, Basic) { } // Test set 0: Pad + Conv2D; padding is VALID -// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// A = input(image), B = input(paddings), C= Pad = input of conv2D, // D=input(filter), E = Conv2D, Z = Zeta // C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) // After layout pass -// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); InitGraph( @@ -2049,10 +2049,10 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { } // Test set 0: Pad + Conv2D; padding is SAME -// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// A = input(image), B = input(paddings), C= Pad = input of conv2D, // D=input(filter), E = Conv2D, Z = Zeta // C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) -// After layout pass - No merging +// After layout pass - No merging TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Negative) { CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); InitGraph( @@ -2075,11 +2075,12 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Negative) { "node { name: 'Z' op: 'Zeta'" " attr {key: 'T' value { type: DT_FLOAT } }" " input: ['E', 'Y']}"); - EXPECT_EQ(DoMklLayoutOptimizationPass(), - "A(Input);B(Int32Input);C(Pad);D(Input);DMT/_0(Const);DMT/_1(Const);" - "E(_MklConv2D);Y(Input);Z(Zeta)|A->C;B->C:1;C->E;" - "C:control->DMT/_0:control;C:control->DMT/_1:control;" - "D->E:1;DMT/_0->E:2;DMT/_1->E:3;E->Z;Y->Z:1"); + EXPECT_EQ( + DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);C(Pad);D(Input);DMT/_0(Const);DMT/_1(Const);" + "E(_MklConv2D);Y(Input);Z(Zeta)|A->C;B->C:1;C->E;" + "C:control->DMT/_0:control;C:control->DMT/_1:control;" + "D->E:1;DMT/_0->E:2;DMT/_1->E:3;E->Z;Y->Z:1"); } // Test set 1: Conv2D + AddBias diff --git a/tensorflow/core/kernels/mkl_conv_ops.cc b/tensorflow/core/kernels/mkl_conv_ops.cc index 8c8be197f9..7ee9f66810 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.cc +++ b/tensorflow/core/kernels/mkl_conv_ops.cc @@ -963,39 +963,38 @@ class MklConv2DOp : public OpKernel { errors::Aborted("Operation received an exception:", error_msg)); } } - - void PadWithConvFusion(OpKernelContext* context, memory::dims &padding_left, - memory::dims &padding_right){ + + void PadWithConvFusion(OpKernelContext* context, memory::dims& padding_left, + memory::dims& padding_right) { const Tensor& paddings_tf = MklGetInput(context, 2); OP_REQUIRES(context, paddings_tf.dims() == 2, errors::InvalidArgument("paddings must be 2-dimensional: ", paddings_tf.shape().DebugString())); Tpadding* paddings = nullptr; // To get individual pad, need to flatten the tensor - paddings = static_cast(const_cast - (paddings_tf.flat().data())); + paddings = static_cast( + const_cast(paddings_tf.flat().data())); // For NHWC format: - // paddings[0], paddings[1], paddings[6], paddings[7] should be zero + // paddings[0], paddings[1], paddings[6], paddings[7] should be zero // if the paddings_tf is [ [0, 0] [1,2] [3,4] [0,0] ] // paddings = {0, 0, 1, 2, 3, 4, 0, 0} ; flat method is row major // then, values are: top = 1, bottom =2, left=3, right=4 - // For NCHW format: - // paddings[0], paddings[1], paddings[2], paddings[3] should be zero + // For NCHW format: + // paddings[0], paddings[1], paddings[2], paddings[3] should be zero // similar explanation as NHWC format will apply. - int64 pad_top, pad_left; - int64 pad_bottom, pad_right; + int64 pad_top, pad_left; + int64 pad_bottom, pad_right; string data_format = ToString(data_format_); - if(data_format == "NHWC"){ - pad_top = paddings[2]; - pad_bottom = paddings[3]; - pad_left = paddings[4]; - pad_right = paddings[5]; - } - else if (data_format == "NCHW"){ - pad_top = paddings[4]; - pad_bottom = paddings[5]; - pad_left = paddings[6]; - pad_right = paddings[7]; + if (data_format == "NHWC") { + pad_top = paddings[2]; + pad_bottom = paddings[3]; + pad_left = paddings[4]; + pad_right = paddings[5]; + } else if (data_format == "NCHW") { + pad_top = paddings[4]; + pad_bottom = paddings[5]; + pad_left = paddings[6]; + pad_right = paddings[7]; } // Create padding arrays for MKL DNN convolutions. // MKL-DNN uses asymetric padding. @@ -1124,13 +1123,13 @@ class MklConv2DOp : public OpKernel { .TypeConstraint("T") \ .TypeConstraint("Tpaddings") \ .Label(mkl_op_registry::kMklOpLabel), \ - MklConv2DOp); \ + MklConv2DOp); \ REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ .Device(DEVICE_CPU) \ .TypeConstraint("T") \ .TypeConstraint("Tpaddings") \ .Label(mkl_op_registry::kMklOpLabel), \ - MklConv2DOp); \ + MklConv2DOp); \ REGISTER_KERNEL_BUILDER(Name("__MklDummyPadWithConv2D") \ .Device(DEVICE_CPU) \ .TypeConstraint("T") \ @@ -1139,7 +1138,7 @@ class MklConv2DOp : public OpKernel { MklDummyOp); TF_CALL_float(REGISTER_MKL_CPU); -#endif // INTEL_MKL_ML +#endif // INTEL_MKL_ML_ONLY } // namespace tensorflow #endif // INTEL_MKL diff --git a/tensorflow/core/kernels/mkl_conv_ops.h b/tensorflow/core/kernels/mkl_conv_ops.h index 3955bd919d..cd24ae02c4 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.h +++ b/tensorflow/core/kernels/mkl_conv_ops.h @@ -230,9 +230,8 @@ class MklDnnConvUtil { virtual inline void GetOutputAndPadSizeInMklOrder( const TensorShape& input_shape, const TensorShape& filter_shape, const memory::dims& strides, const memory::dims& dilations, - memory::dims* output_dims_tf_order, - memory::dims* output_dims_mkl_order, memory::dims* pad_l, - memory::dims* pad_r, bool padEnabled=false) { + memory::dims* output_dims_tf_order, memory::dims* output_dims_mkl_order, + memory::dims* pad_l, memory::dims* pad_r, bool padEnabled = false) { CHECK_NOTNULL(output_dims_tf_order); CHECK_NOTNULL(output_dims_mkl_order); CHECK_NOTNULL(pad_l); @@ -269,17 +268,17 @@ class MklDnnConvUtil { dilation_cols, stride_cols, padding_, &out_cols, &pad_left, &pad_right)); // If padEnabled, i.e., pad and conv op are fused, then - // all pads are already passed from pad op through - // *pad_l and *pad_r - if(padEnabled) { - pad_top = static_cast((*pad_l)[0]); - pad_left = static_cast((*pad_l)[1]); - pad_bottom = static_cast((*pad_r)[0]); - pad_right = static_cast((*pad_r)[1]); - // update the out_rows and out_cols based on all - // sides of the pads coming from pad op. - out_rows = out_rows + (pad_top + pad_bottom ) / stride_rows; - out_cols = out_cols + (pad_left + pad_right ) / stride_cols; + // all pads are already passed from pad op through + // *pad_l and *pad_r + if (padEnabled) { + pad_top = static_cast((*pad_l)[0]); + pad_left = static_cast((*pad_l)[1]); + pad_bottom = static_cast((*pad_r)[0]); + pad_right = static_cast((*pad_r)[1]); + // update the out_rows and out_cols based on all + // sides of the pads coming from pad op. + out_rows = out_rows + (pad_top + pad_bottom) / stride_rows; + out_cols = out_cols + (pad_left + pad_right) / stride_cols; } // Handle padding. MKL-DNN uses asymetric padding. // But, if padEnabled, i.e., pad and conv op are fused, @@ -342,10 +341,9 @@ class MklDnnConvUtil { inline void GetConvFwdSizesInMklOrder( const TensorShape& input_shape, const TensorShape& filter_shape, memory::dims* input_dims, memory::dims* filter_dims, - memory::dims* strides, memory::dims *dilations, - memory::dims* output_dims_tf_order, - memory::dims* output_dims_mkl_order, memory::dims* pad_l, - memory::dims* pad_r, bool padEnabled=false) { + memory::dims* strides, memory::dims* dilations, + memory::dims* output_dims_tf_order, memory::dims* output_dims_mkl_order, + memory::dims* pad_l, memory::dims* pad_r, bool padEnabled = false) { CHECK_NOTNULL(input_dims); CHECK_NOTNULL(filter_dims); CHECK_NOTNULL(strides); @@ -361,15 +359,13 @@ class MklDnnConvUtil { if (!context_->status().ok()) return; GetStridesInMklOrder(strides); GetDilationsInMklOrder(dilations); - GetOutputAndPadSizeInMklOrder(input_shape, filter_shape, - *strides, *dilations, - output_dims_tf_order, output_dims_mkl_order, - pad_l, pad_r, padEnabled); + GetOutputAndPadSizeInMklOrder( + input_shape, filter_shape, *strides, *dilations, output_dims_tf_order, + output_dims_mkl_order, pad_l, pad_r, padEnabled); if (!context_->status().ok()) return; } }; - ///////////////////////////////////////////////////////////////////// /// Common class that implements Conv2DBackpropFilter and Input ///////////////////////////////////////////////////////////////////// -- GitLab From 53f2aefe86f8b50addd4b67eb20eb91135b1fac7 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Tue, 14 Aug 2018 13:23:16 -0700 Subject: [PATCH 065/596] fixed, so that now not allowing duplicate control edges, alos cleaning up the comments --- tensorflow/core/graph/mkl_layout_pass.cc | 15 ++++++--------- tensorflow/core/kernels/mkl_conv_ops.h | 4 ---- 2 files changed, 6 insertions(+), 13 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index 84e8ea8f70..9157080330 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -4189,16 +4189,14 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, // We handle control edges now. for (const Edge* e : pred->in_edges()) { if (e->IsControlEdge()) { - // Allow duplicate while adding control edge as it would fail (return - // NULL) if we try to add duplicate edge. - CHECK_NOTNULL((*g)->AddControlEdge(e->src(), new_node, true)); + //Don't allow duplicate edge + (*g)->AddControlEdge(e->src(), new_node, false); } } for (const Edge* e : succ->in_edges()) { if (e->IsControlEdge()) { - // Allow duplicate while adding control edge as it would fail (return - // NULL) if we try to add duplicate edge. - CHECK_NOTNULL((*g)->AddControlEdge(e->src(), new_node, true)); + //Don't allow duplicate edge + (*g)->AddControlEdge(e->src(), new_node, false); } } @@ -4206,9 +4204,8 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, // First, we will fix outgoing control edges from 'pred' node. for (const Edge* e : pred->out_edges()) { if (e->IsControlEdge()) { - // Allow duplicate while adding control edge as it would fail (return - // NULL) if we try to add duplicate edge. - CHECK_NOTNULL((*g)->AddControlEdge(new_node, e->dst(), true)); + //Don't allow duplicate edge + (*g)->AddControlEdge(new_node, e->dst(), false); } } diff --git a/tensorflow/core/kernels/mkl_conv_ops.h b/tensorflow/core/kernels/mkl_conv_ops.h index cd24ae02c4..ebaf1a9947 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.h +++ b/tensorflow/core/kernels/mkl_conv_ops.h @@ -300,10 +300,6 @@ class MklDnnConvUtil { mkldnn_sizes[MklDnnDims::Dim_H] = static_cast(out_rows); mkldnn_sizes[MklDnnDims::Dim_W] = static_cast(out_cols); *output_dims_mkl_order = mkldnn_sizes; - - // Now handle padding. MKL-DNN uses asymetric padding. - // But, if padEnabled, i.e., pad and conv op are fused, - // then, *pad_l and *pad_r are already set from pad op } // Calculate output and pad size of forward Convolution operator. -- GitLab From f6c9e054a042bf0f518a740380f3f96a28e8c5be Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Fri, 17 Aug 2018 12:45:06 -0700 Subject: [PATCH 066/596] not allowing duplicate edges, and, add two unit tests in mkl_layout_pass_test to test if common input and common output of pad an conv2D work correctly for pad+conv2D fusion --- tensorflow/core/graph/mkl_layout_pass.cc | 2 +- tensorflow/core/graph/mkl_layout_pass_test.cc | 212 ++++++++++++------ 2 files changed, 143 insertions(+), 71 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index 9157080330..6d99e57417 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -4214,7 +4214,7 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, if (e->IsControlEdge()) { // Allow duplicate while adding control edge as it would fail (return // NULL) if we try to add duplicate edge. - CHECK_NOTNULL((*g)->AddControlEdge(new_node, e->dst(), true)); + (*g)->AddControlEdge(new_node, e->dst(), false); } else { // Conv2D has only 1 output (at slot 0) and merged node also has only 1 // output (at slot 0). diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index d1b39ceeca..248520a7f4 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -1994,6 +1994,10 @@ REGISTER_OP("_MklInput2") .Output("o: uint8") .Output("o1: uint8") .SetIsStateful(); +REGISTER_OP("Output2") + .Input("i: float") + .Input("i1: float") + .SetIsStateful(); ///////////////////////////////////////////////////////////////////// // Unit tests related to node merge optiimization @@ -2012,76 +2016,6 @@ TEST_F(MklLayoutPassTest, Basic) { "A->C;A->D;B->C:1;B->D:1"); } -// Test set 0: Pad + Conv2D; padding is VALID -// A = input(image), B = input(paddings), C= Pad = input of conv2D, -// D=input(filter), E = Conv2D, Z = Zeta -// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) -// After layout pass -// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) -TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { - CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); - InitGraph( - "node { name: 'A' op: 'Input'}" - "node { name: 'B' op: 'Int32Input'}" - "node { name: 'C' op: 'Pad'" - " attr { key: 'T' value { type: DT_FLOAT } }" - " attr { key: 'Tpaddings' value { type: DT_INT32 } }" - " input: ['A', 'B']}" - "node { name: 'D' op: 'Input'}" - "node { name: 'E' op: 'Conv2D'" - " attr { key: 'T' value { type: DT_FLOAT } }" - " attr { key: 'data_format' value { s: 'NHWC' } }" - " attr { key: 'use_cudnn_on_gpu' value { b: false } }" - " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" - " attr { key: 'padding' value { s: 'VALID' } }" - " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" - " input: ['C', 'D'] }" - "node { name: 'Y' op: 'Input'}" - "node { name: 'Z' op: 'Zeta'" - " attr {key: 'T' value { type: DT_FLOAT } }" - " input: ['E', 'Y']}"); - EXPECT_EQ(DoMklLayoutOptimizationPass(), - "A(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" - "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" - "A:control->DMT/_0:control;A:control->DMT/_1:control;" - "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" - "DMT/_2->E:5;E->Z;Y->Z:1"); -} - -// Test set 0: Pad + Conv2D; padding is SAME -// A = input(image), B = input(paddings), C= Pad = input of conv2D, -// D=input(filter), E = Conv2D, Z = Zeta -// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) -// After layout pass - No merging -TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Negative) { - CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); - InitGraph( - "node { name: 'A' op: 'Input'}" - "node { name: 'B' op: 'Int32Input'}" - "node { name: 'C' op: 'Pad'" - " attr { key: 'T' value { type: DT_FLOAT } }" - " attr { key: 'Tpaddings' value { type: DT_INT32 } }" - " input: ['A', 'B']}" - "node { name: 'D' op: 'Input'}" - "node { name: 'E' op: 'Conv2D'" - " attr { key: 'T' value { type: DT_FLOAT } }" - " attr { key: 'data_format' value { s: 'NHWC' } }" - " attr { key: 'use_cudnn_on_gpu' value { b: false } }" - " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" - " attr { key: 'padding' value { s: 'SAME' } }" - " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" - " input: ['C', 'D'] }" - "node { name: 'Y' op: 'Input'}" - "node { name: 'Z' op: 'Zeta'" - " attr {key: 'T' value { type: DT_FLOAT } }" - " input: ['E', 'Y']}"); - EXPECT_EQ( - DoMklLayoutOptimizationPass(), - "A(Input);B(Int32Input);C(Pad);D(Input);DMT/_0(Const);DMT/_1(Const);" - "E(_MklConv2D);Y(Input);Z(Zeta)|A->C;B->C:1;C->E;" - "C:control->DMT/_0:control;C:control->DMT/_1:control;" - "D->E:1;DMT/_0->E:2;DMT/_1->E:3;E->Z;Y->Z:1"); -} // Test set 1: Conv2D + AddBias @@ -2389,6 +2323,144 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_ConvBpropInput_FilterFwd) { "E:3->G:4;F->G;F:control->DMT/_3:control;G->Z;X->Y:1;X->Z:1"); } +// Test set 3: Pad + Conv2D fusion +// padding is VALID type +// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// D=input(filter), E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) +// After layout pass +// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'D' op: 'Input'}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'VALID' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'D'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" + "A:control->DMT/_0:control;A:control->DMT/_1:control;" + "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;Y->Z:1"); +} +// Pad + Conv2D fusion with padding is VALID, +// Input node pointing to both Pad and Conv2D +// A = input(image), B = input(paddings), C= Pad +// E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,A); Z=Zeta(E,Y) +// After layout pass +// _MklPadWithConv2D(A, A, B, DMT/_0, DMT/_1, DMT/_2) +TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Common_Input) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'VALID' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'A'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;A->E:1;" + "A:control->DMT/_0:control;A:control->DMT/_1:control;" + "A:control->DMT/_2:control;B->E:2;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;Y->Z:1"); +} +// Pad + Conv2D with padding is VALID, +// Input node pointing to both Pad and Conv2D +// Output of both Pad and Conv2D feeds one node (Z as Output2) +// A = input(as image), B = input(as paddings), C= Pad +// E = Conv2D, Z = Output2 +// C=Pad(A,B); E=Conv2D(C,A); Z=Output(C,E) +// After layout pass - No merging, since Pad and Conv2D both +// feed to the same node (Z) +TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Common_InOutput) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'VALID' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'A'] }" + "node { name: 'Z' op: 'Output2'" + " input: ['C', 'E']}"); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);C(Pad);DMT/_0(Const);DMT/_1(Const);" + "E(_MklConv2D);Z(Output2)|A->C;A->E:1;B->C:1;C->E;C->Z;" + "C:control->DMT/_0:control;C:control->DMT/_1:control;" + "DMT/_0->E:2;DMT/_1->E:3;E->Z:1"); +} +// Pad + Conv2D; padding is SAME +// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// D=input(filter), E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) +// After layout pass - No merging +TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Negative) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'D' op: 'Input'}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'SAME' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'D'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + EXPECT_EQ( + DoMklLayoutOptimizationPass(), + "A(Input);B(Int32Input);C(Pad);D(Input);DMT/_0(Const);DMT/_1(Const);" + "E(_MklConv2D);Y(Input);Z(Zeta)|A->C;B->C:1;C->E;" + "C:control->DMT/_0:control;C:control->DMT/_1:control;" + "D->E:1;DMT/_0->E:2;DMT/_1->E:3;E->Z;Y->Z:1"); +} ///////////////////////////////////////////////////////////////////// // Unit tests related to rewriting node to Mkl node ///////////////////////////////////////////////////////////////////// -- GitLab From f8ec0f101bac066faa2e917ac714ca9eea310eac Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Thu, 13 Sep 2018 22:40:49 -0700 Subject: [PATCH 067/596] adding checks that pad fusion works only Conv2D --- tensorflow/core/kernels/mkl_conv_ops.cc | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/mkl_conv_ops.cc b/tensorflow/core/kernels/mkl_conv_ops.cc index 54670c8521..4b54ce1d52 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.cc +++ b/tensorflow/core/kernels/mkl_conv_ops.cc @@ -900,7 +900,10 @@ class MklConvOp : public OpKernel { bool isConv2D = (strides_.size() == 4); // TODO(Intel-tf) Add check to make sure padEnabled is true only for 2D - + if(!isConv2D){ + OP_REQUIRES(context, padEnabled, + errors::InvalidArgument("Pad+Conv fusion only works for 2D")); + } // Create memory for user data. // Describe how the inputs and outputs of Convolution look like. Also // specify buffers containing actual input and output data. -- GitLab From 0e87ed82815053b4f1c038975382d72282fdf97f Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Sun, 16 Sep 2018 11:20:26 -0700 Subject: [PATCH 068/596] Adding two unit tests for pad+conv2d fusion. They test if the two merging ops get control edge from a common op, then the merged node will have only one control edge. --- tensorflow/core/graph/mkl_layout_pass.cc | 21 +++- tensorflow/core/graph/mkl_layout_pass_test.cc | 111 ++++++++++++++++++ 2 files changed, 128 insertions(+), 4 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index ef8a2b0838..d3a4112ee9 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -4192,10 +4192,23 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, // 2. Get inputs from both the nodes. - // Pad must have 2 inputs: "input" and paddings. - CHECK_EQ(pred->in_edges().size(), 2); - // Conv2D must have 2 inputs: pad output and Filter - CHECK_EQ(succ->in_edges().size(), 2); + // Pad must have 2 data inputs: "input" and paddings. + int PadDataInputEdges = 0; + for (const Edge* e : pred->in_edges()) { + if (!e->IsControlEdge()) { + PadDataInputEdges++; + } + } + CHECK_EQ(PadDataInputEdges, 2); + + // Conv2D must have 2 data inputs: pad output and Filter + int ConvDataInputEdges = 0; + for (const Edge* e : succ->in_edges()) { + if (!e->IsControlEdge()) { + ConvDataInputEdges++; + } + } + CHECK_EQ(ConvDataInputEdges, 2); // We will use the node name of Conv2D as the name of new node // Build new node. We use same name as original node, but change the op diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index 248520a7f4..e9e234010c 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -1928,6 +1928,13 @@ static void InitGraph(const string& s, Graph* graph, class MklLayoutPassTest : public ::testing::Test { public: MklLayoutPassTest() : graph_(OpRegistry::Global()) {} + // Return Node* from the Node Name + Node* FindNode(const string& name) { + for (Node* node : graph_.nodes()) { + if (node->name() == name) return node; + } + LOG(FATAL) << name; + } void InitGraph(const string& s, const string& device = kCPUDevice) { ::tensorflow::InitGraph(s, &graph_, device); @@ -1998,6 +2005,9 @@ REGISTER_OP("Output2") .Input("i: float") .Input("i1: float") .SetIsStateful(); +REGISTER_OP("Output") + .Input("i: float") + .SetIsStateful(); ///////////////////////////////////////////////////////////////////// // Unit tests related to node merge optiimization @@ -2359,6 +2369,107 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" "DMT/_2->E:5;E->Z;Y->Z:1"); } +// Test if input control edges do not duplicate after merge. +// If both the merging ops have input control edge from a common op +// then, the merged op will have only one control edge from that +// common op. +// padding is VALID type +// A = input(image), A1 = input, B = input(paddings), +// C= Pad = input of conv2D, +// D=input(filter), E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) +// C:control->A1:control +// E:control->A1:control +// After layout pass +// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +// A1:control->E:control +TEST_F(MklLayoutPassTest, Output_ControlEdge_PadWithConv2D_Positive) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A1' op: 'Input'}" + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'D' op: 'Input'}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'VALID' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'D'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + Node* a1 = FindNode("A1"); + Node* c = FindNode("C"); + Node* e = FindNode("E"); + const Edge* edge = graph_.AddControlEdge(a1, c); + const Edge* edge_1 = graph_.AddControlEdge(a1, e); + ASSERT_TRUE(edge != nullptr); + ASSERT_TRUE(edge_1 != nullptr); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" + "A1:control->E:control;A:control->DMT/_0:control;A:control->DMT/_1:control;" + "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;Y->Z:1"); +} +// Test if output control edges does not duplicate after merge. +// If both the merging ops have output control edge to a common op, +// then after merge, the merged op will have only one control edge +// to that commom op. +// padding is VALID type +// A = input(image), B = input(paddings), C= Pad = input of conv2D, +// D=input(filter), E = Conv2D, Z = Zeta +// C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) +// C:control->A1:control +// E:control->A1:control +// After layout pass +// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +// E:control->A1:control (only one control edge) +TEST_F(MklLayoutPassTest, ControlEdge_PadWithConv2D_Positive) { + CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); + InitGraph( + "node { name: 'A1' op: 'Input'}" + "node { name: 'A' op: 'Input'}" + "node { name: 'B' op: 'Int32Input'}" + "node { name: 'C' op: 'Pad'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'Tpaddings' value { type: DT_INT32 } }" + " input: ['A', 'B']}" + "node { name: 'D' op: 'Input'}" + "node { name: 'E' op: 'Conv2D'" + " attr { key: 'T' value { type: DT_FLOAT } }" + " attr { key: 'data_format' value { s: 'NHWC' } }" + " attr { key: 'use_cudnn_on_gpu' value { b: false } }" + " attr { key: 'strides' value { list: {i: 1, i:1, i:1, i:1} } }" + " attr { key: 'padding' value { s: 'VALID' } }" + " attr { key: 'dilations' value { list: {i: 1, i:1, i:1, i:1} } }" + " input: ['C', 'D'] }" + "node { name: 'Y' op: 'Input'}" + "node { name: 'Z' op: 'Zeta'" + " attr {key: 'T' value { type: DT_FLOAT } }" + " input: ['E', 'Y']}"); + Node* a1 = FindNode("A1"); + Node* c = FindNode("C"); + Node* e = FindNode("E"); + const Edge* edge = graph_.AddControlEdge(c, a1); + const Edge* edge_1 = graph_.AddControlEdge(e, a1); + ASSERT_TRUE(edge != nullptr); + ASSERT_TRUE(edge_1 != nullptr); + EXPECT_EQ(DoMklLayoutOptimizationPass(), + "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" + "A:control->DMT/_0:control;A:control->DMT/_1:control;" + "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;E:control->A1:control;Y->Z:1"); +} // Pad + Conv2D fusion with padding is VALID, // Input node pointing to both Pad and Conv2D // A = input(image), B = input(paddings), C= Pad -- GitLab From 4e140eed6b4f6722b94cf85432d4519b8c5ce0bf Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Mon, 17 Sep 2018 10:05:40 -0700 Subject: [PATCH 069/596] changing the name of the unit tests --- tensorflow/core/graph/mkl_layout_pass_test.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index e9e234010c..9ad45a2cfd 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -1928,7 +1928,7 @@ static void InitGraph(const string& s, Graph* graph, class MklLayoutPassTest : public ::testing::Test { public: MklLayoutPassTest() : graph_(OpRegistry::Global()) {} - // Return Node* from the Node Name + // Ashraf added Node* FindNode(const string& name) { for (Node* node : graph_.nodes()) { if (node->name() == name) return node; @@ -2383,7 +2383,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { // After layout pass // _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) // A1:control->E:control -TEST_F(MklLayoutPassTest, Output_ControlEdge_PadWithConv2D_Positive) { +TEST_F(MklLayoutPassTest, Input_ControlEdge_PadWithConv2D_Positive) { CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); InitGraph( "node { name: 'A1' op: 'Input'}" -- GitLab From 2a8f7bcc59bc4e36ea88f4187028b4461f5f1072 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Tue, 18 Sep 2018 11:28:21 -0700 Subject: [PATCH 070/596] minor change in the two unit tests --- tensorflow/core/graph/mkl_layout_pass_test.cc | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index 9ad45a2cfd..60a7f138c8 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -2378,11 +2378,11 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { // C= Pad = input of conv2D, // D=input(filter), E = Conv2D, Z = Zeta // C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) -// C:control->A1:control -// E:control->A1:control -// After layout pass -// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +// A1:control->C:control // A1:control->E:control +// After layout pass: +// _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) +// A1:control->E:control (only one control edge) TEST_F(MklLayoutPassTest, Input_ControlEdge_PadWithConv2D_Positive) { CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); InitGraph( @@ -2411,8 +2411,8 @@ TEST_F(MklLayoutPassTest, Input_ControlEdge_PadWithConv2D_Positive) { Node* e = FindNode("E"); const Edge* edge = graph_.AddControlEdge(a1, c); const Edge* edge_1 = graph_.AddControlEdge(a1, e); - ASSERT_TRUE(edge != nullptr); - ASSERT_TRUE(edge_1 != nullptr); + ASSERT_NE(edge, nullptr); + ASSERT_NE(edge_1, nullptr); EXPECT_EQ(DoMklLayoutOptimizationPass(), "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" @@ -2430,10 +2430,10 @@ TEST_F(MklLayoutPassTest, Input_ControlEdge_PadWithConv2D_Positive) { // C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) // C:control->A1:control // E:control->A1:control -// After layout pass +// After layout pass: // _MklPadWithConv2D(A, D, B, DMT/_0, DMT/_1, DMT/_2) // E:control->A1:control (only one control edge) -TEST_F(MklLayoutPassTest, ControlEdge_PadWithConv2D_Positive) { +TEST_F(MklLayoutPassTest, Output_ControlEdge_PadWithConv2D_Positive) { CHECK_EQ(kTensorOrdering, MklTfTensorOrdering::TENSORS_CONTIGUOUS); InitGraph( "node { name: 'A1' op: 'Input'}" @@ -2461,8 +2461,8 @@ TEST_F(MklLayoutPassTest, ControlEdge_PadWithConv2D_Positive) { Node* e = FindNode("E"); const Edge* edge = graph_.AddControlEdge(c, a1); const Edge* edge_1 = graph_.AddControlEdge(e, a1); - ASSERT_TRUE(edge != nullptr); - ASSERT_TRUE(edge_1 != nullptr); + ASSERT_NE(edge, nullptr); + ASSERT_NE(edge_1, nullptr); EXPECT_EQ(DoMklLayoutOptimizationPass(), "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" -- GitLab From 777e6a4e194e4cc141feb6b250702c0e4946ca2d Mon Sep 17 00:00:00 2001 From: wangsiyu Date: Mon, 1 Oct 2018 13:51:36 +0800 Subject: [PATCH 071/596] Make colocations be compatible with DistributionStrategy in SyncReplicasOptimizer --- tensorflow/python/training/sync_replicas_optimizer.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/tensorflow/python/training/sync_replicas_optimizer.py b/tensorflow/python/training/sync_replicas_optimizer.py index 7afaa92699..99d2563fc6 100644 --- a/tensorflow/python/training/sync_replicas_optimizer.py +++ b/tensorflow/python/training/sync_replicas_optimizer.py @@ -27,6 +27,7 @@ from tensorflow.python.ops import state_ops from tensorflow.python.ops import variable_scope from tensorflow.python.ops import variables from tensorflow.python.platform import tf_logging as logging +from tensorflow.python.training import distribution_strategy_context from tensorflow.python.training import optimizer from tensorflow.python.training import queue_runner from tensorflow.python.training import session_manager @@ -245,7 +246,9 @@ class SyncReplicasOptimizer(optimizer.Optimizer): # local_anchor op will be placed on this worker task by default. local_anchor = control_flow_ops.no_op() # Colocating local_step variable prevents it being placed on the PS. - with ops.colocate_with(local_anchor): + distribution_strategy = ( + distribution_strategy_context.get_distribution_strategy()) + with distribution_strategy.colocate_vars_with(local_anchor): self._local_step = variable_scope.variable( initial_value=0, trainable=False, -- GitLab From f0cd69e047eb8f29538a27a41c17da20a1c59f2b Mon Sep 17 00:00:00 2001 From: himkt Date: Fri, 12 Oct 2018 13:17:58 +0900 Subject: [PATCH 072/596] Fix indentation in CRF1d --- tensorflow/contrib/crf/python/ops/crf.py | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/tensorflow/contrib/crf/python/ops/crf.py b/tensorflow/contrib/crf/python/ops/crf.py index 43bb43129b..7653b48e9c 100644 --- a/tensorflow/contrib/crf/python/ops/crf.py +++ b/tensorflow/contrib/crf/python/ops/crf.py @@ -38,12 +38,12 @@ tf_unary_scores, tf_sequence_lengths, tf_transition_params, _ = session.run( [unary_scores, sequence_lengths, transition_params, train_op]) for tf_unary_scores_, tf_sequence_length_ in zip(tf_unary_scores, tf_sequence_lengths): -# Remove padding. -tf_unary_scores_ = tf_unary_scores_[:tf_sequence_length_] + # Remove padding. + tf_unary_scores_ = tf_unary_scores_[:tf_sequence_length_] -# Compute the highest score and its tag sequence. -tf_viterbi_sequence, tf_viterbi_score = tf.contrib.crf.viterbi_decode( - tf_unary_scores_, tf_transition_params) + # Compute the highest score and its tag sequence. + tf_viterbi_sequence, tf_viterbi_score = tf.contrib.crf.viterbi_decode( + tf_unary_scores_, tf_transition_params) """ from __future__ import absolute_import -- GitLab From 650172a574504223ec2bdb328ed7c985389313d7 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 13 Oct 2018 20:30:26 +0000 Subject: [PATCH 073/596] Update test case for complex support of squared difference Signed-off-by: Yong Tang --- tensorflow/python/ops/math_ops_test.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/ops/math_ops_test.py b/tensorflow/python/ops/math_ops_test.py index 533a00e737..62645230ee 100644 --- a/tensorflow/python/ops/math_ops_test.py +++ b/tensorflow/python/ops/math_ops_test.py @@ -217,7 +217,7 @@ class SquaredDifferenceTest(test_util.TensorFlowTestCase): for dtype in [np.complex64, np.complex128]: x = np.array([[1+3j, 2+2j, 3+1j], [4-1j, 5-2j, 6-3j]], dtype=dtype) y = np.array([-3+1j, -2+2j, -1+3j], dtype=dtype) - z = (x - y) * (x - y) + z = np.conj(x - y) * (x - y) with test_util.device(use_gpu=False): z_tf = self.evaluate(math_ops.squared_difference(x, y)) self.assertAllClose(z, z_tf) -- GitLab From 3a06e557619ebaa5437d1506af058b858806e9c7 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 13 Oct 2018 20:30:47 +0000 Subject: [PATCH 074/596] Update squared difference implementation for complex types. Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_ops.h | 32 +++++++++++++++-------------- 1 file changed, 17 insertions(+), 15 deletions(-) diff --git a/tensorflow/core/kernels/cwise_ops.h b/tensorflow/core/kernels/cwise_ops.h index 06918075a4..5afb97dc52 100644 --- a/tensorflow/core/kernels/cwise_ops.h +++ b/tensorflow/core/kernels/cwise_ops.h @@ -296,27 +296,31 @@ struct less_equal : std::binary_function { } }; -// Functor that enables composition of multiple Eigen functors. -template -struct scalar_compose_op { +// Functor that enables squared difference functor. +template +struct scalar_squared_difference_op { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Scalar operator()(const Scalar& a, const Scalar& b) const { - return UnaryFunctor()(BinaryFunctor()(a, b)); + const Scalar v = scalar_difference_op()(a, b); + return scalar_product_op()(v, scalar_conjugate_op()(v)); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { - return UnaryFunctor().packetOp(BinaryFunctor().packetOp(a, b)); + const Packet v = scalar_difference_op().packetOp(a, b); + return scalar_product_op().packetOp(v, scalar_conjugate_op().packetOp(v)); } }; -template -struct functor_traits> { +template +struct functor_traits> { enum { - Cost = functor_traits::Cost + - functor_traits::Cost, - PacketAccess = functor_traits::PacketAccess && - functor_traits::PacketAccess + Cost = functor_traits>::Cost + + functor_traits>::Cost + + functor_traits>::Cost, + PacketAccess = functor_traits>::PacketAccess && + functor_traits>::PacketAccess && + functor_traits>::PacketAccess }; }; @@ -709,7 +713,7 @@ struct rint : base> {}; // pow(x, y) = x ^ y // maximum(x, y) = x > y ? x : y // minimum(x, y) = x < y ? x : y -// squared_difference(x, y) = (x - y) * (x - y) +// squared_difference(x, y) = conj(x - y) * (x - y) template struct add : base> { @@ -812,9 +816,7 @@ struct atan2 : base> {}; template struct squared_difference - : base, - Eigen::internal::scalar_difference_op>> {}; + : base> {}; template struct less : base, bool> {}; -- GitLab From 82642d91dbe6fbba87e6a582e396ca91df1f6440 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 13 Oct 2018 20:43:59 +0000 Subject: [PATCH 075/596] Fix `Experimental clang-format Check` error Signed-off-by: Yong Tang --- tensorflow/core/kernels/cwise_ops.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/cwise_ops.h b/tensorflow/core/kernels/cwise_ops.h index 5afb97dc52..2682a25868 100644 --- a/tensorflow/core/kernels/cwise_ops.h +++ b/tensorflow/core/kernels/cwise_ops.h @@ -308,7 +308,8 @@ struct scalar_squared_difference_op { EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE const Packet packetOp(const Packet& a, const Packet& b) const { const Packet v = scalar_difference_op().packetOp(a, b); - return scalar_product_op().packetOp(v, scalar_conjugate_op().packetOp(v)); + return scalar_product_op().packetOp( + v, scalar_conjugate_op().packetOp(v)); } }; -- GitLab From 02987bd1e3459ba7ecfe689bb182ce8e692b70c4 Mon Sep 17 00:00:00 2001 From: "Yuan (Terry) Tang" Date: Tue, 16 Oct 2018 09:26:55 -0400 Subject: [PATCH 076/596] Add bullet points so Reduction values are clearer --- tensorflow/python/ops/losses/losses_impl.py | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/tensorflow/python/ops/losses/losses_impl.py b/tensorflow/python/ops/losses/losses_impl.py index 8a8a81ab5c..03cdc69ae5 100644 --- a/tensorflow/python/ops/losses/losses_impl.py +++ b/tensorflow/python/ops/losses/losses_impl.py @@ -39,13 +39,14 @@ class Reduction(object): """Types of loss reduction. Contains the following values: - `NONE`: Un-reduced weighted losses with the same shape as input. - `SUM`: Scalar sum of weighted losses. - `MEAN`: Scalar `SUM` divided by sum of weights. - `SUM_OVER_BATCH_SIZE`: Scalar `SUM` divided by number of elements in losses. - `SUM_OVER_NONZERO_WEIGHTS`: Scalar `SUM` divided by number of non-zero + + * `NONE`: Un-reduced weighted losses with the same shape as input. + * `SUM`: Scalar sum of weighted losses. + * `MEAN`: Scalar `SUM` divided by sum of weights. + * `SUM_OVER_BATCH_SIZE`: Scalar `SUM` divided by number of elements in losses. + * `SUM_OVER_NONZERO_WEIGHTS`: Scalar `SUM` divided by number of non-zero weights. - `SUM_BY_NONZERO_WEIGHTS`: Same as `SUM_OVER_NONZERO_WEIGHTS`. + * `SUM_BY_NONZERO_WEIGHTS`: Same as `SUM_OVER_NONZERO_WEIGHTS`. """ NONE = "none" -- GitLab From 8d4805d8a3d2b2b48c69d9cd7f358996cb9a87a1 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Mon, 22 Oct 2018 11:09:58 -0700 Subject: [PATCH 077/596] cleaning up MKL ML code in LRN kernel --- tensorflow/core/kernels/mkl_lrn_op.cc | 680 +------------------------- tensorflow/core/ops/nn_ops.cc | 8 - 2 files changed, 2 insertions(+), 686 deletions(-) diff --git a/tensorflow/core/kernels/mkl_lrn_op.cc b/tensorflow/core/kernels/mkl_lrn_op.cc index 22ff4cd80f..4d46abb0a4 100644 --- a/tensorflow/core/kernels/mkl_lrn_op.cc +++ b/tensorflow/core/kernels/mkl_lrn_op.cc @@ -22,6 +22,7 @@ limitations under the License. #define EIGEN_USE_THREADS #include +#include "mkldnn.hpp" #include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" #include "tensorflow/core/framework/op_kernel.h" #include "tensorflow/core/framework/register_types.h" @@ -29,25 +30,18 @@ limitations under the License. #include "tensorflow/core/kernels/bounds_check.h" #include "tensorflow/core/kernels/ops_util.h" #include "tensorflow/core/lib/core/errors.h" +#include "tensorflow/core/util/mkl_util.h" #include "tensorflow/core/util/tensor_format.h" #if !defined(IS_MOBILE_PLATFORM) #include "tensorflow/core/util/work_sharder.h" #endif -#ifndef INTEL_MKL_ML_ONLY -#include "mkldnn.hpp" using mkldnn::lrn_across_channels; using mkldnn::lrn_backward; using mkldnn::lrn_forward; using mkldnn::prop_kind; using mkldnn::stream; -#else -#include "mkl_dnn.h" -#include "mkl_dnn_types.h" -#endif - -#include "tensorflow/core/util/mkl_util.h" namespace tensorflow { @@ -69,672 +63,6 @@ void GetBandMatrix(int depth, int depth_radius, } // namespace -#ifdef INTEL_MKL_ML_ONLY - -template -class MklLRNOp : public OpKernel { - public: - ~MklLRNOp() {} - - explicit MklLRNOp(OpKernelConstruction* context) : OpKernel(context) { - int64 depth_radius64; - OP_REQUIRES_OK(context, context->GetAttr("depth_radius", &depth_radius64)); - OP_REQUIRES( - context, - FastBoundsCheck(depth_radius64, std::numeric_limits::max()), - errors::InvalidArgument("depth_radius = ", depth_radius64, - " larger than int max")); - depth_radius_ = static_cast(depth_radius64); - - OP_REQUIRES_OK(context, context->GetAttr("bias", &bias_)); - OP_REQUIRES_OK(context, context->GetAttr("alpha", &alpha_)); - OP_REQUIRES_OK(context, context->GetAttr("beta", &beta_)); - workspace_enabled_ = false; - OP_REQUIRES_OK(context, - context->GetAttr("workspace_enabled", &workspace_enabled_)); - } - - void Compute(OpKernelContext* context) override { - MklLRNOpContext mkl_context; - - const Tensor& input = MklGetInput(context, 0); - GetMklShape(context, 0, &mkl_context.input_shape); - bool input_in_mkl_format = mkl_context.input_shape.IsMklTensor(); - - // Sanity checks - mkl_context.in_dims = input_in_mkl_format - ? mkl_context.input_shape.GetDimension() - : input.dims(); - OP_REQUIRES(context, mkl_context.in_dims == 4, - errors::InvalidArgument("input must be 4-dimensional")); - OP_REQUIRES( - context, - FastBoundsCheck(input.NumElements(), std::numeric_limits::max()), - errors::InvalidArgument("argument to LRN too large")); - - if (!input_in_mkl_format) { - mkl_context.MklDefaultToEigen(context, depth_radius_, bias_, alpha_, - beta_, input); - return; - } - - if (input_in_mkl_format) { - // MKL supports normalization over channel dimension only - if (mkl_context.input_shape.tf_dim_idx(mkl_context.in_dims - 1) == - MklDims::C) { - mkl_context.lt_input = - static_cast(mkl_context.input_shape.GetCurLayout()); - workspace_enabled_ = true; - } else { - Tensor converted_tensor = - ConvertMklToTF(context, input, mkl_context.input_shape); - mkl_context.MklDefaultToEigen(context, depth_radius_, bias_, alpha_, - beta_, converted_tensor); - return; - } - } - - int kernel_size = 2 * depth_radius_ + 1; - - CHECK_EQ(dnnLRNCreateForward_F32( - &mkl_context.lrn_fwd, NULL, mkl_context.lt_input, kernel_size, - static_cast(alpha_ * kernel_size), beta_, bias_), - E_SUCCESS); - - // Allocate output tensor and shape - Tensor* output = nullptr; - Tensor* workspace = nullptr; - - // Convert Inputs if needed - Tensor mkl_tmp_input_buf_tensor; - mkl_context.MklPrepareLRNInputs(context, &mkl_tmp_input_buf_tensor); - - // Allocate Layer Outputs - mkl_context.MklAllocateOutputs(context, &output, &workspace, - workspace_enabled_); - - Tensor mkl_tmp_workspace_buf_tensor; - mkl_context.MklPrepareLRNOutputs(context, output, workspace, - &mkl_tmp_workspace_buf_tensor, - workspace_enabled_); - - // Execute LRN. - CHECK_EQ(dnnExecute_F32(mkl_context.lrn_fwd, mkl_context.lrn_res), - E_SUCCESS); - - // Release MKL resources. - mkl_context.MklCleanup(); - } - - private: - typedef struct { - size_t in_dims; - size_t in_sizes[4]; - size_t in_strides[4]; - size_t out_sizes[4]; - size_t out_strides[4]; - MklShape input_shape; - dnnPrimitive_t lrn_fwd = nullptr; - dnnPrimitive_t convert_input = nullptr; - dnnLayout_t lt_input = nullptr; - dnnLayout_t lt_internal_input = nullptr; - dnnLayout_t lt_internal_workspace = nullptr; - dnnLayout_t lt_internal_output = nullptr; - void* lrn_res[dnnResourceNumber]; - - // Convert Inputs if needed - void MklPrepareLRNInputs(OpKernelContext* context, - Tensor* mkl_tmp_input_buf_tensor) { - const Tensor& input = MklGetInput(context, 0); - void* mkl_buf_input = - const_cast(static_cast(input.flat().data())); - - CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(<_internal_input, lrn_fwd, - dnnResourceSrc), - E_SUCCESS); - - void* mkl_buf_convert_input = nullptr; - bool mkl_convert_input = false; - mkl_convert_input = !dnnLayoutCompare_F32(lt_internal_input, lt_input); - - if (mkl_convert_input) { - CHECK_EQ(dnnConversionCreate_F32(&convert_input, lt_input, - lt_internal_input), - E_SUCCESS); - AllocTmpBuffer(context, mkl_tmp_input_buf_tensor, lt_internal_input, - &mkl_buf_convert_input); - CHECK_EQ(dnnConversionExecute_F32(convert_input, mkl_buf_input, - mkl_buf_convert_input), - E_SUCCESS); - dnnDelete_F32(convert_input); - } - - lrn_res[dnnResourceSrc] = - (mkl_convert_input) ? mkl_buf_convert_input : mkl_buf_input; - } - - // Allocate Layer Outputs - void MklAllocateOutputs(OpKernelContext* context, Tensor** output, - Tensor** workspace, bool workspace_enabled_) { - TensorShape mkl_output_tf_shape; /* First tensor */ - MklShape mkl_output_mkl_shape; /* Second tensor */ - - mkl_output_mkl_shape.SetMklTensor(true); - mkl_output_mkl_shape.SetMklLayout(lrn_fwd, dnnResourceDst); - mkl_output_mkl_shape.SetTfLayout(in_dims, input_shape.GetSizes(), - input_shape.GetStrides()); - mkl_output_mkl_shape.SetTfDimOrder(in_dims, - input_shape.GetTfToMklDimMap()); - mkl_output_tf_shape.AddDim( - dnnLayoutGetMemorySize_F32( - static_cast(mkl_output_mkl_shape.GetMklLayout())) / - sizeof(T)); - AllocateOutputSetMklShape(context, 0, output, - mkl_output_tf_shape /* First tensor */, - mkl_output_mkl_shape /* Second Tensor */); - - if (workspace_enabled_) { - TensorShape mkl_workspace_tf_shape; /* First tensor */ - MklShape mkl_workspace_mkl_shape; /* Second tensor */ - mkl_workspace_mkl_shape.SetMklTensor(false); - mkl_workspace_mkl_shape.SetMklLayout(lrn_fwd, dnnResourceWorkspace); - // Assumes workspace has same TF layout and TF dim order as input - mkl_workspace_mkl_shape.SetTfLayout(in_dims, input_shape.GetSizes(), - input_shape.GetStrides()); - mkl_workspace_mkl_shape.SetTfDimOrder(in_dims, - input_shape.GetTfToMklDimMap()); - mkl_workspace_tf_shape.AddDim( - dnnLayoutGetMemorySize_F32(static_cast( - mkl_workspace_mkl_shape.GetMklLayout())) / - sizeof(T)); - AllocateOutputSetMklShape(context, 1, workspace, - mkl_workspace_tf_shape /* First tensor */, - mkl_workspace_mkl_shape /* Second Tensor */); - } - } - - void MklPrepareLRNOutputs(OpKernelContext* context, Tensor* output, - Tensor* workspace, - Tensor* mkl_tmp_workspace_buf_tensor, - bool workspace_enabled_) { - CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(<_internal_workspace, lrn_fwd, - dnnResourceWorkspace), - E_SUCCESS); - - CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(<_internal_output, lrn_fwd, - dnnResourceDst), - E_SUCCESS); - - void* mkl_buf_output = - const_cast(static_cast(output->flat().data())); - lrn_res[dnnResourceDst] = mkl_buf_output; - - void* mkl_buf_workspace = nullptr; - if (workspace_enabled_) { - mkl_buf_workspace = const_cast( - static_cast(workspace->flat().data())); - } else { - AllocTmpBuffer(context, mkl_tmp_workspace_buf_tensor, - lt_internal_workspace, &mkl_buf_workspace); - } - lrn_res[dnnResourceWorkspace] = mkl_buf_workspace; - } - - // Fallback implementation - Taken from lrn_op.cc - // TODO(inteltf) Check if we can use EigenLRNOp directly instead of making a - // copy. - void MklDefaultToEigen(OpKernelContext* context, int depth_radius_, - float bias_, float alpha_, float beta_, - const Tensor& input) { - const int batch = static_cast(input.dim_size(0)); - const int rows = static_cast(input.dim_size(1)); - const int cols = static_cast(input.dim_size(2)); - const int depth = static_cast(input.dim_size(3)); - const int nodes = cols * rows; - - auto in_shaped = input.shaped({nodes * batch, depth}); - // Multiplying the input with the band matrix has the effect of reducing - // the - // correct patch along the depth. - Eigen::Tensor multiplier(depth, depth); - GetBandMatrix(depth, depth_radius_, &multiplier); - - Tensor *output, *workspace; - MklShape mkl_output_mkl_shape, mkl_workspace_mkl_shape; - mkl_output_mkl_shape.SetMklTensor(false); - mkl_output_mkl_shape.SetDimensions(4); - AllocateOutputSetMklShape(context, 0, &output, input.shape(), - mkl_output_mkl_shape); - - mkl_workspace_mkl_shape.SetMklTensor(false); - mkl_workspace_mkl_shape.SetDimensions(4); - AllocateOutputSetMklShape(context, 1, &workspace, input.shape(), - mkl_workspace_mkl_shape); - - auto out_shaped = output->shaped({nodes * batch, depth}); - Eigen::array dims = {{DimPair(1, 0)}}; - auto tmp = in_shaped.square().contract(multiplier, dims) * alpha_ + bias_; - if (beta_ == T(1)) { - out_shaped.device(context->eigen_cpu_device()) = - in_shaped * tmp.inverse(); - } else if (beta_ == T(0.5)) { - out_shaped.device(context->eigen_cpu_device()) = - in_shaped * tmp.rsqrt(); - } else { - out_shaped.device(context->eigen_cpu_device()) = - in_shaped * (tmp.log() * -beta_).exp(); - } - } - - // Release MKL resources. - void MklCleanup() { - dnnDelete_F32(lrn_fwd); - dnnLayoutDelete_F32(lt_internal_input); - dnnLayoutDelete_F32(lt_internal_workspace); - dnnLayoutDelete_F32(lt_internal_output); - } - } MklLRNOpContext; - - typedef typename Eigen::Tensor::DimensionPair DimPair; - - bool workspace_enabled_; - int depth_radius_; - float bias_; - float alpha_; - float beta_; -}; - -template -class MklLRNGradOp : public OpKernel { - public: - explicit MklLRNGradOp(OpKernelConstruction* context) : OpKernel(context) { - int64 depth_radius64; - OP_REQUIRES_OK(context, context->GetAttr("depth_radius", &depth_radius64)); - OP_REQUIRES( - context, - FastBoundsCheck(depth_radius64, std::numeric_limits::max()), - errors::InvalidArgument("depth_radius = ", depth_radius64, - " larger than int max")); - depth_radius_ = static_cast(depth_radius64); - OP_REQUIRES_OK(context, context->GetAttr("bias", &bias_)); - OP_REQUIRES_OK(context, context->GetAttr("alpha", &alpha_)); - OP_REQUIRES_OK(context, context->GetAttr("beta", &beta_)); - workspace_enabled_ = false; - OP_REQUIRES_OK(context, - context->GetAttr("workspace_enabled", &workspace_enabled_)); - } - - void Compute(OpKernelContext* context) override { - MklLRNGradOpContext mkl_context; - mkl_context.depth_radius_ = depth_radius_; - mkl_context.bias_ = bias_; - mkl_context.alpha_ = alpha_; - mkl_context.beta_ = beta_; - - const Tensor& in_grads = MklGetInput(context, 0); - const Tensor& in_image = MklGetInput(context, 1); - const Tensor& out_image = MklGetInput(context, 2); - - GetMklShape(context, 0, &mkl_context.ingrad_shape); - GetMklShape(context, 1, &mkl_context.inimage_shape); - GetMklShape(context, 2, &mkl_context.outimage_shape); - - bool ingrad_in_mkl_format = mkl_context.ingrad_shape.IsMklTensor(); - bool inimage_in_mkl_format = mkl_context.inimage_shape.IsMklTensor(); - bool outimage_in_mkl_format = mkl_context.outimage_shape.IsMklTensor(); - - mkl_context.in_dims = inimage_in_mkl_format - ? mkl_context.inimage_shape.GetDimension() - : in_image.dims(); - OP_REQUIRES(context, mkl_context.in_dims == 4, - errors::InvalidArgument("input images must be 4-dimensional")); - - if (!workspace_enabled_) { - mkl_context.MklDefaultToEigen(context); - return; - } - - if (ingrad_in_mkl_format || inimage_in_mkl_format) { - const MklShape* tmp_mkl_shape = (ingrad_in_mkl_format) - ? &mkl_context.ingrad_shape - : &mkl_context.inimage_shape; - if (tmp_mkl_shape->tf_dim_idx(mkl_context.in_dims - 1) != MklDims::C) { - // Fallback to eigen - mkl_context.MklDefaultToEigen(context); - return; - } else { // MKL supports normalization over channel dimension only - for (int i = 0; i < mkl_context.in_dims; i++) { - mkl_context.in_sizes[i] = mkl_context.out_sizes[i] = - tmp_mkl_shape->GetSizes()[i]; - mkl_context.in_strides[i] = mkl_context.out_strides[i] = - tmp_mkl_shape->GetStrides()[i]; - } - } - } else { - // Fallback to eigen - mkl_context.MklDefaultToEigen(context); - return; - } - - // Dimensions check for sanity purpose - if (ingrad_in_mkl_format) { - OP_REQUIRES( - context, mkl_context.ingrad_shape.GetDimension() == 4, - errors::InvalidArgument("input gradient must be 4-dimensional")); - } else { - OP_REQUIRES( - context, in_grads.dims() == 4, - errors::InvalidArgument("input gradient must be 4-dimensional")); - } - - if (outimage_in_mkl_format) { - OP_REQUIRES( - context, mkl_context.outimage_shape.GetDimension() == 4, - errors::InvalidArgument("Output image must be 4-dimensional")); - } else { - OP_REQUIRES( - context, out_image.dims() == 4, - errors::InvalidArgument("Output image must be 4-dimensional")); - } - - // Prepare mkl input layout - mkl_context.MklPrepareLRNInputsLayouts(context); - int ksize = 2 * depth_radius_ + 1; - - CHECK_EQ(dnnLRNCreateBackward_F32( - &mkl_context.lrn_bwd, NULL, mkl_context.lt_input, - mkl_context.lt_output, ksize, - static_cast(alpha_ * ksize), beta_, bias_), - E_SUCCESS); - - // Allocate output tensor and shape. - TensorShape mkl_output_tf_shape; /* First tensor */ - MklShape mkl_output_mkl_shape; /* Second tensor */ - mkl_output_mkl_shape.SetMklTensor(true); - CHECK_NE(mkl_context.lrn_bwd, nullptr); - mkl_output_mkl_shape.SetMklLayout(mkl_context.lrn_bwd, dnnResourceDiffSrc); - mkl_output_mkl_shape.SetTfLayout(mkl_context.in_dims, mkl_context.out_sizes, - mkl_context.out_strides); - if (ingrad_in_mkl_format) { - mkl_output_mkl_shape.SetTfDimOrder( - mkl_context.in_dims, mkl_context.ingrad_shape.GetTfToMklDimMap()); - } else { - mkl_output_mkl_shape.SetTfDimOrder( - mkl_context.in_dims, mkl_context.inimage_shape.GetTfToMklDimMap()); - } - mkl_output_tf_shape.AddDim( - dnnLayoutGetMemorySize_F32( - static_cast(mkl_output_mkl_shape.GetMklLayout())) / - sizeof(T)); - Tensor* output = nullptr; - AllocateOutputSetMklShape(context, 0, &output, mkl_output_tf_shape, - mkl_output_mkl_shape); - - // Get pointers to output data. - void* user_output = - const_cast(static_cast(output->flat().data())); - - Tensor mkl_tmp_input_buf_tensor, mkl_tmp_image_buf_tensor, - mkl_tmp_outimage_buf_tensor; - // Convert Inputs if needed - mkl_context.MklPrepareLRNGradInput(context, &mkl_tmp_input_buf_tensor, - &mkl_tmp_image_buf_tensor, - &mkl_tmp_outimage_buf_tensor); - - // We do not do any conversion for output. But we simply emit it - // in MKL format. - mkl_context.res_lrn_bwd[dnnResourceDiffSrc] = user_output; - // Execute LRN backward using dnnExecute - CHECK_EQ(dnnExecute_F32(mkl_context.lrn_bwd, mkl_context.res_lrn_bwd), - E_SUCCESS); - // Release MKL resources. - mkl_context.Mklcleanup(); - } - - private: - typedef struct { - int depth_radius_; - float bias_; - float alpha_; - float beta_; - size_t in_dims; - size_t in_sizes[4]; - size_t in_strides[4]; - size_t out_sizes[4]; - size_t out_strides[4]; - MklShape ingrad_shape, inimage_shape, outimage_shape; - dnnPrimitive_t lrn_bwd = nullptr; - dnnPrimitive_t convert_input = nullptr; - dnnLayout_t lt_input = nullptr; - dnnLayout_t lt_output = nullptr; - dnnLayout_t lt_bdw_input = nullptr; - dnnLayout_t lt_workspace = nullptr; - dnnLayout_t lt_internal_input = nullptr; - void* res_lrn_bwd[dnnResourceNumber]; - - // prepare mkl input - void MklPrepareLRNInputsLayouts(OpKernelContext* context) { - bool ingrad_in_mkl_format = ingrad_shape.IsMklTensor(); - bool inimage_in_mkl_format = inimage_shape.IsMklTensor(); - if (!ingrad_in_mkl_format) { - CHECK_EQ(dnnLayoutCreate_F32(<_input, in_dims, in_sizes, in_strides), - E_SUCCESS); - } else { - lt_input = static_cast(ingrad_shape.GetCurLayout()); - } - - if (!inimage_in_mkl_format) { - CHECK_EQ( - dnnLayoutCreate_F32(<_output, in_dims, out_sizes, out_strides), - E_SUCCESS); - } else { - lt_output = static_cast(inimage_shape.GetCurLayout()); - } - } - - // convert input if needed - void MklPrepareLRNGradInput(OpKernelContext* context, - Tensor* mkl_tmp_input_buf_tensor, - Tensor* mkl_tmp_image_buf_tensor, - Tensor* mkl_tmp_outimage_buf_tensor) { - const Tensor& in_grads = MklGetInput(context, 0); - const Tensor& in_image = MklGetInput(context, 1); - const Tensor& workspace = MklGetInput( - context, - 3); /*Worskpsace is enabled, get the buffer to the workspace */ - - void* user_input = const_cast( - static_cast(in_grads.flat().data())); - void* user_fwd_input = const_cast( - static_cast(in_image.flat().data())); - void* workspace_buffer = const_cast( - static_cast(workspace.flat().data())); - - CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(<_workspace, lrn_bwd, - dnnResourceWorkspace), - E_SUCCESS); - CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(<_bdw_input, lrn_bwd, - dnnResourceDiffDst), - E_SUCCESS); - CHECK_EQ(dnnLayoutCreateFromPrimitive_F32(<_internal_input, lrn_bwd, - dnnResourceSrc), - E_SUCCESS); - - bool ingrad_in_mkl_format = ingrad_shape.IsMklTensor(); - if (ingrad_in_mkl_format) { - if (!dnnLayoutCompare_F32(lt_bdw_input, lt_input)) { - AllocTmpBuffer(context, mkl_tmp_input_buf_tensor, lt_bdw_input, - &res_lrn_bwd[dnnResourceDiffDst]); - ingrad_shape.GetConvertedFlatData(lt_bdw_input, user_input, - res_lrn_bwd[dnnResourceDiffDst]); - } else { - res_lrn_bwd[dnnResourceDiffDst] = user_input; - } - } else { - if (!dnnLayoutCompare_F32(lt_bdw_input, lt_input)) { - CHECK_EQ( - dnnConversionCreate_F32(&convert_input, lt_input, lt_bdw_input), - E_SUCCESS); - - AllocTmpBuffer(context, mkl_tmp_input_buf_tensor, lt_bdw_input, - &res_lrn_bwd[dnnResourceDiffDst]); - CHECK_EQ(dnnConversionExecute_F32(convert_input, user_input, - res_lrn_bwd[dnnResourceDiffDst]), - E_SUCCESS); - dnnDelete_F32(convert_input); - } else { - res_lrn_bwd[dnnResourceDiffDst] = user_input; - } - } - - bool inimage_in_mkl_format = inimage_shape.IsMklTensor(); - if (inimage_in_mkl_format) { - if (!dnnLayoutCompare_F32( - lt_internal_input, - static_cast(inimage_shape.GetCurLayout()))) { - AllocTmpBuffer(context, mkl_tmp_image_buf_tensor, lt_internal_input, - &res_lrn_bwd[dnnResourceSrc]); - ingrad_shape.GetConvertedFlatData(lt_internal_input, user_fwd_input, - res_lrn_bwd[dnnResourceSrc]); - } else { - res_lrn_bwd[dnnResourceSrc] = user_fwd_input; - } - } else { - if (!dnnLayoutCompare_F32( - lt_internal_input, - static_cast(inimage_shape.GetCurLayout()))) { - CHECK_EQ(dnnConversionCreate_F32( - &convert_input, - static_cast(inimage_shape.GetCurLayout()), - lt_internal_input), - E_SUCCESS); - - AllocTmpBuffer(context, mkl_tmp_image_buf_tensor, lt_internal_input, - &res_lrn_bwd[dnnResourceSrc]); - CHECK_EQ(dnnConversionExecute_F32(convert_input, user_fwd_input, - res_lrn_bwd[dnnResourceSrc]), - E_SUCCESS); - dnnDelete_F32(convert_input); - } else { - res_lrn_bwd[dnnResourceSrc] = user_fwd_input; - } - } - - res_lrn_bwd[dnnResourceWorkspace] = workspace_buffer; - } - - // Fallback implementation - Taken from lrn_op.cc - // TODO(intelft) Check if we can use EigenLRNOp directly instead of making a - // copy. - void MklDefaultToEigen(OpKernelContext* context) { - Tensor in_grads; - Tensor in_image; - Tensor out_image; - - GetMklShape(context, 0, &ingrad_shape); - GetMklShape(context, 1, &inimage_shape); - GetMklShape(context, 2, &outimage_shape); - - if (ingrad_shape.IsMklTensor()) { - in_grads = - ConvertMklToTF(context, MklGetInput(context, 0), ingrad_shape); - } else { - in_grads = MklGetInput(context, 0); - } - - if (inimage_shape.IsMklTensor()) { - in_image = - ConvertMklToTF(context, MklGetInput(context, 1), inimage_shape); - } else { - in_image = MklGetInput(context, 1); - } - - if (outimage_shape.IsMklTensor()) { - out_image = - ConvertMklToTF(context, MklGetInput(context, 2), outimage_shape); - } else { - out_image = MklGetInput(context, 2); - } - - const int64 batch = static_cast(in_grads.dim_size(0)); - const int64 rows = static_cast(in_grads.dim_size(1)); - const int64 cols = static_cast(in_grads.dim_size(2)); - const int64 depth = static_cast(in_grads.dim_size(3)); - const auto nodes = cols * rows; - - auto grads_shaped = in_grads.shaped({nodes * batch, depth}); - - auto in_shaped = in_image.shaped({nodes * batch, depth}); - auto activations = out_image.shaped({nodes * batch, depth}); - - Tensor* output; - MklShape mkl_output_mkl_shape; - mkl_output_mkl_shape.SetMklTensor(false); - mkl_output_mkl_shape.SetDimensions(4); - AllocateOutputSetMklShape(context, 0, &output, in_grads.shape(), - mkl_output_mkl_shape); - - auto out_shaped = output->shaped({nodes * batch, depth}); - out_shaped.setZero(); - auto shard = [this, activations, in_shaped, grads_shaped, out_shaped, - depth](int64 begin, int64 end) { - for (int64 i = begin; i < end; ++i) { - for (int64 j = 0; j < depth; ++j) { - int64 depth_begin = std::max(0, j - depth_radius_); - int64 depth_end = std::min(depth, j + depth_radius_ + 1); - - T norm(0); - for (int64 k = depth_begin; k < depth_end; ++k) { - norm += in_shaped(i, k) * in_shaped(i, k); - } - norm = alpha_ * norm + bias_; - DCHECK_GT(norm, T(1e-6)); - for (int64 k = depth_begin; k < depth_end; ++k) { - T dyi = T(-2) * alpha_ * beta_ * in_shaped(i, k) * - activations(i, j) / norm; - if (k == j) { - dyi += Eigen::numext::pow(norm, -beta_); - } - dyi *= grads_shaped(i, j); - const_cast::Tensor&>(out_shaped)(i, k) += - dyi; - } - } - } - }; - auto worker_threads = - *(context->device()->tensorflow_cpu_worker_threads()); - Shard(worker_threads.num_threads, worker_threads.workers, nodes * batch, - depth * depth, shard); - } - - // release mkl resources - void Mklcleanup() { - bool ingrad_in_mkl_format = ingrad_shape.IsMklTensor(); - bool inimage_in_mkl_format = inimage_shape.IsMklTensor(); - if (!ingrad_in_mkl_format) { - CHECK_EQ(dnnLayoutDelete_F32(lt_input), E_SUCCESS); - } - - if (!inimage_in_mkl_format) { - CHECK_EQ(dnnLayoutDelete_F32(lt_output), E_SUCCESS); - } - dnnDelete_F32(lrn_bwd); - dnnLayoutDelete_F32(lt_bdw_input); - dnnLayoutDelete_F32(lt_workspace); - } - } MklLRNGradOpContext; - - typedef typename Eigen::Tensor::DimensionPair DimPair; - bool workspace_enabled_; - int depth_radius_; - float bias_; - float alpha_; - float beta_; -}; - -#else - template class MklLRNOp : public OpKernel { public: @@ -847,7 +175,6 @@ class MklLRNOp : public OpKernel { MklDnnData* src_dnn_data, MklDnnData* dst_dnn_data, MklDnnData* wksp_dnn_data = nullptr) { - // Check for input reorder src_dnn_data->CheckReorderToOpMem(lrn_fwd_desc.src_primitive_desc()); @@ -1160,7 +487,6 @@ class MklLRNGradOp : public OpKernel { MklDnnData* output_diff_src, const memory::primitive_desc& target_diff_dst_pd, const MklDnnData* workspace_dnn_data = nullptr) { - // Check for input reordering on the diff dst input input_gradient_diff_dst->CheckReorderToOpMem( lrn_bkwd_desc.diff_dst_primitive_desc()); @@ -1345,8 +671,6 @@ class MklLRNGradOp : public OpKernel { float beta_; }; -#endif // INTEL_MKL_ML_ONLY - #define REGISTER_MKL_LRN_CPU(T) \ REGISTER_KERNEL_BUILDER(Name("_MklLRN") \ .Device(DEVICE_CPU) \ diff --git a/tensorflow/core/ops/nn_ops.cc b/tensorflow/core/ops/nn_ops.cc index 38fe45936a..bee0add89c 100644 --- a/tensorflow/core/ops/nn_ops.cc +++ b/tensorflow/core/ops/nn_ops.cc @@ -2172,11 +2172,7 @@ REGISTER_OP("_MklLRN") .Input("input: T") .Input("mkl_input: uint8") .Output("output: T") -#ifdef INTEL_MKL_ML_ONLY - .Output("workspace: T") -#else .Output("workspace: uint8") -#endif .Output("mkl_output: uint8") .Output("mkl_workspace: uint8") .Attr("depth_radius: int = 5") @@ -2200,11 +2196,7 @@ REGISTER_OP("_MklLRNGrad") .Input("input_grads: T") .Input("input_image: T") .Input("output_image: T") -#ifdef INTEL_MKL_ML_ONLY - .Input("workspace: T") -#else .Input("workspace: uint8") -#endif .Input("mkl_input_grads: uint8") .Input("mkl_input_image: uint8") .Input("mkl_output_image: uint8") -- GitLab From c0814e3861c0b88caebc24d4ef1ce5e61a213f2e Mon Sep 17 00:00:00 2001 From: Rholais Lii Date: Tue, 23 Oct 2018 13:44:40 +0800 Subject: [PATCH 078/596] Fix comments to match usage Fix comments of `sparse_softmax_cross_entropy_with_logits` to match usage. --- tensorflow/python/ops/nn_ops.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/ops/nn_ops.py b/tensorflow/python/ops/nn_ops.py index e31d162285..2477271a22 100644 --- a/tensorflow/python/ops/nn_ops.py +++ b/tensorflow/python/ops/nn_ops.py @@ -1980,8 +1980,9 @@ def sparse_softmax_cross_entropy_with_logits( on `logits` internally for efficiency. Do not call this op with the output of `softmax`, as it will produce incorrect results. - A common use case is to have logits and labels of shape - `[batch_size, num_classes]`, but higher dimensions are supported, in which + A common use case is to have logits of shape + `[batch_size, num_classes]` and have labels of shape + `[batch_size]`, but higher dimensions are supported, in which case the `dim`-th dimension is assumed to be of size `num_classes`. `logits` must have the dtype of `float16`, `float32`, or `float64`, and `labels` must have the dtype of `int32` or `int64`. -- GitLab From 67e2c47e2ff1222e141c51ac2794aa9b9207a573 Mon Sep 17 00:00:00 2001 From: Matt Conley Date: Wed, 24 Oct 2018 13:30:35 -0700 Subject: [PATCH 079/596] Disable denormal test on ARM until the architecture is supported. --- tensorflow/python/kernel_tests/denormal_test.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/kernel_tests/denormal_test.py b/tensorflow/python/kernel_tests/denormal_test.py index 71a528c4aa..9f65a18c14 100644 --- a/tensorflow/python/kernel_tests/denormal_test.py +++ b/tensorflow/python/kernel_tests/denormal_test.py @@ -35,8 +35,8 @@ class DenormalTest(test.TestCase): self.assertEqual(tiny, tiny / 16 * 16) def _flushDenormalsTest(self, use_gpu, dtypes): - if platform.machine() == "ppc64le" or platform.machine() == "s390x": - # Disabled denormal_test on power/s390x platform + if platform.machine() == "ppc64le" or platform.machine() == "s390x" or platform.machine() == "aarch64": + # Disabled denormal_test on power/s390x/aarch64 platform # Check relevant discussion - https://github.com/tensorflow/tensorflow/issues/11902 return with self.cached_session(use_gpu=use_gpu): -- GitLab From 684ce69f568e89e34883d91b65c3983c333e53d9 Mon Sep 17 00:00:00 2001 From: Karl Lessard Date: Thu, 25 Oct 2018 22:18:22 -0400 Subject: [PATCH 080/596] Prevent memory leak by storing strings instead of StringPiece in vector --- tensorflow/core/common_runtime/eager/attr_builder.cc | 2 +- tensorflow/core/common_runtime/eager/attr_builder.h | 10 +++++----- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/tensorflow/core/common_runtime/eager/attr_builder.cc b/tensorflow/core/common_runtime/eager/attr_builder.cc index 5c8369de87..e9f2188656 100644 --- a/tensorflow/core/common_runtime/eager/attr_builder.cc +++ b/tensorflow/core/common_runtime/eager/attr_builder.cc @@ -99,7 +99,7 @@ Status AttrTypeMapForOp(const char* op_name, const AttrTypeMap** out) { #define DEFINE_SET_ATTR(value_type, value_field) \ template <> \ AttrBuilder& AttrBuilder::Set(StringPiece attr_name, value_type&& value) { \ - value_field.push_back(std::make_pair(attr_name, value)); \ + value_field.push_back(std::make_pair(string(attr_name), value)); \ return *this; \ } diff --git a/tensorflow/core/common_runtime/eager/attr_builder.h b/tensorflow/core/common_runtime/eager/attr_builder.h index c114ea4ba0..a48ba686d2 100644 --- a/tensorflow/core/common_runtime/eager/attr_builder.h +++ b/tensorflow/core/common_runtime/eager/attr_builder.h @@ -96,7 +96,7 @@ class AttrBuilder { template AttrBuilder& Set(StringPiece attr_name, T&& value) { MayBeInitializeNodeDef(); - SetInAttrValueMap(node_def_->mutable_attr(), attr_name, value); + SetInAttrValueMap(node_def_->mutable_attr(), string(attr_name), value); return *this; } @@ -107,7 +107,7 @@ class AttrBuilder { private: template - using AttrVec = tensorflow::gtl::InlinedVector, 2>; + using AttrVec = tensorflow::gtl::InlinedVector, 2>; void MayBeInitializeNodeDef(); // Fill `m` with the attr-value pairs set via AttrBuilder::Set() so far, as @@ -119,7 +119,7 @@ class AttrBuilder { void FillAttrValueMap(AttrValueMap* m, bool include_those_in_node_def) const; template - void SetInAttrValueMap(AttrValueMap* m, StringPiece attr_name, + void SetInAttrValueMap(AttrValueMap* m, const string& attr_name, T&& value) const { DCHECK(!node_def_finalized_) << "Calling SetInAttrValueMap after BuildNodeDef."; @@ -128,12 +128,12 @@ class AttrBuilder { AttrValue attr_value; if (found == nullptr) { SetAttrValue(value, &attr_value); - m->insert(AttrValueMap::value_type(string(attr_name), attr_value)); + m->insert(AttrValueMap::value_type(attr_name, attr_value)); } else { // TODO(ashankar): Do what is done in // NodeDefBuilder::CheckInconsistency(attr_name, *found, attr_value); SetAttrValue(std::forward(value), &attr_value); - (*m)[string(attr_name)] = attr_value; + (*m)[attr_name] = attr_value; } } -- GitLab From 59617ccaca8c5980f5418a0b612b040ac8d1afba Mon Sep 17 00:00:00 2001 From: Ouwen Huang Date: Tue, 30 Oct 2018 05:37:22 +0000 Subject: [PATCH 081/596] Added note on weight decay for tf.contrib.opt optimizers. --- .../python/training/weight_decay_optimizers.py | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py b/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py index 200b0d2008..1e8351b70f 100644 --- a/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py +++ b/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py @@ -59,6 +59,23 @@ class DecoupledWeightDecayExtension(object): Note that this extension decays weights BEFORE applying the update based on the gradient, i.e. this extension only has the desired behaviour for optimizers which do not depend on the value of'var' in the update step! + + Note: when applying a decay to the learning rate, be sure to manually apply + the decay to the `weight_decay` as well. For example: + + ```python + decay = tf.train.piecewise_constant(tf.train.get_global_step(), + [10000, 15000], [1e-1, 1e-2, 1e-3]) + lr = 1*decay + wd = 1e-4*decay + + # ... + + optimizer = tf.contrib.opt.MomentumWOptimizer(learning_rate=lr, + weight_decay=wd, + momentum=0.9, + use_nesterov=True) + ``` """ def __init__(self, weight_decay, **kwargs): -- GitLab From b0d15134f110ce380a0e769c4f415d41fbea2677 Mon Sep 17 00:00:00 2001 From: joaak <29533036+joaak@users.noreply.github.com> Date: Mon, 5 Nov 2018 15:51:37 -0500 Subject: [PATCH 082/596] update tokenizer code to remove bug --- .../image_captioning_with_attention.ipynb | 2301 +++++++++-------- 1 file changed, 1155 insertions(+), 1146 deletions(-) diff --git a/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb b/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb index 3acecd283c..09ea021c44 100644 --- a/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb +++ b/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb @@ -1,1184 +1,1193 @@ { - "nbformat": 4, - "nbformat_minor": 0, - "metadata": { + "cells": [ + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "K2s1A9eLRPEj" + }, + "source": [ + "##### Copyright 2018 The TensorFlow Authors.\n", + "\n", + "Licensed under the Apache License, Version 2.0 (the \"License\").\n" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "Cffg2i257iMS" + }, + "source": [ + "# Image Captioning with Attention\n", + "\n", + "
\n", + "\n", + " Run in Google Colab \n", + "\n", + "View source on GitHub
" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "QASbY_HGo4Lq" + }, + "source": [ + "Image captioning is the task of generating a caption for an image. Given an image like this:\n", + "\n", + "![Man Surfing](https://tensorflow.org/images/surf.jpg) \n", + "\n", + "[Image Source](https://commons.wikimedia.org/wiki/Surfing#/media/File:Surfing_in_Hawaii.jpg), License: Public Domain\n", + "\n", + "Our goal is to generate a caption, such as \"a surfer riding on a wave\". Here, we'll use an attention-based model. This enables us to see which parts of the image the model focuses on as it generates a caption.\n", + "\n", + "![Prediction](https://tensorflow.org/images/imcap_prediction.png)\n", + "\n", + "This model architecture below is similar to [Show, Attend and Tell: Neural Image Caption Generation with Visual Attention](https://arxiv.org/abs/1502.03044). \n", + "\n", + "The code uses [tf.keras](https://www.tensorflow.org/programmers_guide/keras) and [eager execution](https://www.tensorflow.org/programmers_guide/eager), which you can learn more about in the linked guides.\n", + "\n", + "This notebook is an end-to-end example. If you run it, it will download the [MS-COCO](http://cocodataset.org/#home) dataset, preprocess and cache a subset of the images using Inception V3, train an encoder-decoder model, and use it to generate captions on new images.\n", + "\n", + "The code requires TensorFlow version >=1.9. If you're running this in [Colab]()\n", + "\n", + "In this example, we're training on a relatively small amount of data as an example. On a single P100 GPU, this example will take about ~2 hours to train. We train on the first 30,000 captions (corresponding to about ~20,000 images depending on shuffling, as there are multiple captions per image in the dataset)\n" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { "colab": { - "name": "image_captioning_with_attention.ipynb", - "version": "0.3.2", - "views": {}, - "default_view": {}, - "provenance": [ - { - "file_id": "1HI8OK2sMjcx9CTWVn0122QAHOuXaOaMg", - "timestamp": 1530222436922 - } - ], - "private_outputs": true, - "collapsed_sections": [], - "toc_visible": true + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - "kernelspec": { - "display_name": "Python 3", - "language": "python", - "name": "python3" - }, - "accelerator": "GPU" + "colab_type": "code", + "id": "U8l4RJ0XRPEm" + }, + "outputs": [], + "source": [ + "# Import TensorFlow and enable eager execution\n", + "# This code requires TensorFlow version >=1.9\n", + "import tensorflow as tf\n", + "tf.enable_eager_execution()\n", + "\n", + "# We'll generate plots of attention in order to see which parts of an image\n", + "# our model focuses on during captioning\n", + "import matplotlib.pyplot as plt\n", + "\n", + "# Scikit-learn includes many helpful utilities\n", + "from sklearn.model_selection import train_test_split\n", + "from sklearn.utils import shuffle\n", + "\n", + "import re\n", + "import numpy as np\n", + "import os\n", + "import time\n", + "import json\n", + "from glob import glob\n", + "from PIL import Image\n", + "import pickle" + ] }, - "cells": [ - { - "metadata": { - "id": "K2s1A9eLRPEj", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "##### Copyright 2018 The TensorFlow Authors.\n", - "\n", - "Licensed under the Apache License, Version 2.0 (the \"License\").\n" - ] - }, - { - "metadata": { - "id": "Cffg2i257iMS", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "# Image Captioning with Attention\n", - "\n", - "
\n", - "\n", - " Run in Google Colab \n", - "\n", - "View source on GitHub
" - ] - }, - { - "metadata": { - "id": "QASbY_HGo4Lq", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "Image captioning is the task of generating a caption for an image. Given an image like this:\n", - "\n", - "![Man Surfing](https://tensorflow.org/images/surf.jpg) \n", - "\n", - "[Image Source](https://commons.wikimedia.org/wiki/Surfing#/media/File:Surfing_in_Hawaii.jpg), License: Public Domain\n", - "\n", - "Our goal is to generate a caption, such as \"a surfer riding on a wave\". Here, we'll use an attention-based model. This enables us to see which parts of the image the model focuses on as it generates a caption.\n", - "\n", - "![Prediction](https://tensorflow.org/images/imcap_prediction.png)\n", - "\n", - "This model architecture below is similar to [Show, Attend and Tell: Neural Image Caption Generation with Visual Attention](https://arxiv.org/abs/1502.03044). \n", - "\n", - "The code uses [tf.keras](https://www.tensorflow.org/programmers_guide/keras) and [eager execution](https://www.tensorflow.org/programmers_guide/eager), which you can learn more about in the linked guides.\n", - "\n", - "This notebook is an end-to-end example. If you run it, it will download the [MS-COCO](http://cocodataset.org/#home) dataset, preprocess and cache a subset of the images using Inception V3, train an encoder-decoder model, and use it to generate captions on new images.\n", - "\n", - "The code requires TensorFlow version >=1.9. If you're running this in [Colab]()\n", - "\n", - "In this example, we're training on a relatively small amount of data as an example. On a single P100 GPU, this example will take about ~2 hours to train. We train on the first 30,000 captions (corresponding to about ~20,000 images depending on shuffling, as there are multiple captions per image in the dataset)\n" - ] - }, - { - "metadata": { - "id": "U8l4RJ0XRPEm", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# Import TensorFlow and enable eager execution\n", - "# This code requires TensorFlow version >=1.9\n", - "import tensorflow as tf\n", - "tf.enable_eager_execution()\n", - "\n", - "# We'll generate plots of attention in order to see which parts of an image\n", - "# our model focuses on during captioning\n", - "import matplotlib.pyplot as plt\n", - "\n", - "# Scikit-learn includes many helpful utilities\n", - "from sklearn.model_selection import train_test_split\n", - "from sklearn.utils import shuffle\n", - "\n", - "import re\n", - "import numpy as np\n", - "import os\n", - "import time\n", - "import json\n", - "from glob import glob\n", - "from PIL import Image\n", - "import pickle" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "b6qbGw8MRPE5", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Download and prepare the MS-COCO dataset\n", - "\n", - "We will use the [MS-COCO dataset](http://cocodataset.org/#home) to train our model. This dataset contains >82,000 images, each of which has been annotated with at least 5 different captions. The code below will download and extract the dataset automatically. \n", - "\n", - "**Caution: large download ahead**. We'll use the training set, it's a 13GB file." - ] - }, - { - "metadata": { - "id": "krQuPYTtRPE7", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "annotation_zip = tf.keras.utils.get_file('captions.zip', \n", - " cache_subdir=os.path.abspath('.'),\n", - " origin = 'http://images.cocodataset.org/annotations/annotations_trainval2014.zip',\n", - " extract = True)\n", - "annotation_file = os.path.dirname(annotation_zip)+'/annotations/captions_train2014.json'\n", - "\n", - "name_of_zip = 'train2014.zip'\n", - "if not os.path.exists(os.path.abspath('.') + '/' + name_of_zip):\n", - " image_zip = tf.keras.utils.get_file(name_of_zip, \n", - " cache_subdir=os.path.abspath('.'),\n", - " origin = 'http://images.cocodataset.org/zips/train2014.zip',\n", - " extract = True)\n", - " PATH = os.path.dirname(image_zip)+'/train2014/'\n", - "else:\n", - " PATH = os.path.abspath('.')+'/train2014/'" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "aANEzb5WwSzg", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Optionally, limit the size of the training set for faster training\n", - "For this example, we'll select a subset of 30,000 captions and use these and the corresponding images to train our model. As always, captioning quality will improve if you choose to use more data." - ] - }, - { - "metadata": { - "id": "4G3b8x8_RPFD", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# read the json file\n", - "with open(annotation_file, 'r') as f:\n", - " annotations = json.load(f)\n", - "\n", - "# storing the captions and the image name in vectors\n", - "all_captions = []\n", - "all_img_name_vector = []\n", - "\n", - "for annot in annotations['annotations']:\n", - " caption = ' ' + annot['caption'] + ' '\n", - " image_id = annot['image_id']\n", - " full_coco_image_path = PATH + 'COCO_train2014_' + '%012d.jpg' % (image_id)\n", - " \n", - " all_img_name_vector.append(full_coco_image_path)\n", - " all_captions.append(caption)\n", - "\n", - "# shuffling the captions and image_names together\n", - "# setting a random state\n", - "train_captions, img_name_vector = shuffle(all_captions,\n", - " all_img_name_vector,\n", - " random_state=1)\n", - "\n", - "# selecting the first 30000 captions from the shuffled set\n", - "num_examples = 30000\n", - "train_captions = train_captions[:num_examples]\n", - "img_name_vector = img_name_vector[:num_examples]" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "mPBMgK34RPFL", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "len(train_captions), len(all_captions)" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "8cSW4u-ORPFQ", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Preprocess the images using InceptionV3\n", - "Next, we will use InceptionV3 (pretrained on Imagenet) to classify each image. We will extract features from the last convolutional layer. \n", - "\n", - "First, we will need to convert the images into the format inceptionV3 expects by:\n", - "* Resizing the image to (299, 299)\n", - "* Using the [preprocess_input](https://www.tensorflow.org/api_docs/python/tf/keras/applications/inception_v3/preprocess_input) method to place the pixels in the range of -1 to 1 (to match the format of the images used to train InceptionV3)." - ] - }, - { - "metadata": { - "id": "zXR0217aRPFR", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "def load_image(image_path):\n", - " img = tf.read_file(image_path)\n", - " img = tf.image.decode_jpeg(img, channels=3)\n", - " img = tf.image.resize_images(img, (299, 299))\n", - " img = tf.keras.applications.inception_v3.preprocess_input(img)\n", - " return img, image_path" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "MDvIu4sXRPFV", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Initialize InceptionV3 and load the pretrained Imagenet weights\n", - "\n", - "To do so, we'll create a tf.keras model where the output layer is the last convolutional layer in the InceptionV3 architecture. \n", - "* Each image is forwarded through the network and the vector that we get at the end is stored in a dictionary (image_name --> feature_vector). \n", - "* We use the last convolutional layer because we are using attention in this example. The shape of the output of this layer is ```8x8x2048```. \n", - "* We avoid doing this during training so it does not become a bottleneck. \n", - "* After all the images are passed through the network, we pickle the dictionary and save it to disk." - ] - }, - { - "metadata": { - "id": "RD3vW4SsRPFW", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "image_model = tf.keras.applications.InceptionV3(include_top=False, \n", - " weights='imagenet')\n", - "new_input = image_model.input\n", - "hidden_layer = image_model.layers[-1].output\n", - "\n", - "image_features_extract_model = tf.keras.Model(new_input, hidden_layer)" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "rERqlR3WRPGO", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Caching the features extracted from InceptionV3\n", - "\n", - "We will pre-process each image with InceptionV3 and cache the output to disk. Caching the output in RAM would be faster but memory intensive, requiring 8 \\* 8 \\* 2048 floats per image. At the time of writing, this would exceed the memory limitations of Colab (although these may change, an instance appears to have about 12GB of memory currently). \n", - "\n", - "Performance could be improved with a more sophisticated caching strategy (e.g., by sharding the images to reduce random access disk I/O) at the cost of more code.\n", - "\n", - "This will take about 10 minutes to run in Colab with a GPU. If you'd like to see a progress bar, you could: install [tqdm](https://github.com/tqdm/tqdm) (```!pip install tqdm```), then change this line: \n", - "\n", - "```for img, path in image_dataset:``` \n", - "\n", - "to:\n", - "\n", - "```for img, path in tqdm(image_dataset):```." - ] - }, - { - "metadata": { - "id": "Dx_fvbVgRPGQ", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# getting the unique images\n", - "encode_train = sorted(set(img_name_vector))\n", - "\n", - "# feel free to change the batch_size according to your system configuration\n", - "image_dataset = tf.data.Dataset.from_tensor_slices(\n", - " encode_train).map(load_image).batch(16)\n", - "\n", - "for img, path in image_dataset:\n", - " batch_features = image_features_extract_model(img)\n", - " batch_features = tf.reshape(batch_features, \n", - " (batch_features.shape[0], -1, batch_features.shape[3]))\n", - "\n", - " for bf, p in zip(batch_features, path):\n", - " path_of_feature = p.numpy().decode(\"utf-8\")\n", - " np.save(path_of_feature, bf.numpy())" - ], - "execution_count": 0, - "outputs": [] - }, - { - "metadata": { - "id": "nyqH3zFwRPFi", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Preprocess and tokenize the captions\n", - "\n", - "* First, we'll tokenize the captions (e.g., by splitting on spaces). This will give us a vocabulary of all the unique words in the data (e.g., \"surfing\", \"football\", etc).\n", - "* Next, we'll limit the vocabulary size to the top 5,000 words to save memory. We'll replace all other words with the token \"UNK\" (for unknown).\n", - "* Finally, we create a word --> index mapping and vice-versa.\n", - "* We will then pad all sequences to the be same length as the longest one. " - ] - }, - { - "metadata": { - "id": "HZfK8RhQRPFj", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# This will find the maximum length of any caption in our dataset\n", - "def calc_max_length(tensor):\n", - " return max(len(t) for t in tensor)" - ], - "execution_count": 0, - "outputs": [] + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "b6qbGw8MRPE5" + }, + "source": [ + "## Download and prepare the MS-COCO dataset\n", + "\n", + "We will use the [MS-COCO dataset](http://cocodataset.org/#home) to train our model. This dataset contains >82,000 images, each of which has been annotated with at least 5 different captions. The code below will download and extract the dataset automatically. \n", + "\n", + "**Caution: large download ahead**. We'll use the training set, it's a 13GB file." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "oJGE34aiRPFo", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# The steps above is a general process of dealing with text processing\n", - "\n", - "# choosing the top 5000 words from the vocabulary\n", - "top_k = 5000\n", - "tokenizer = tf.keras.preprocessing.text.Tokenizer(num_words=top_k, \n", - " oov_token=\"\", \n", - " filters='!\"#$%&()*+.,-/:;=?@[\\]^_`{|}~ ')\n", - "tokenizer.fit_on_texts(train_captions)\n", - "train_seqs = tokenizer.texts_to_sequences(train_captions)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "krQuPYTtRPE7" + }, + "outputs": [], + "source": [ + "annotation_zip = tf.keras.utils.get_file('captions.zip', \n", + " cache_subdir=os.path.abspath('.'),\n", + " origin = 'http://images.cocodataset.org/annotations/annotations_trainval2014.zip',\n", + " extract = True)\n", + "annotation_file = os.path.dirname(annotation_zip)+'/annotations/captions_train2014.json'\n", + "\n", + "name_of_zip = 'train2014.zip'\n", + "if not os.path.exists(os.path.abspath('.') + '/' + name_of_zip):\n", + " image_zip = tf.keras.utils.get_file(name_of_zip, \n", + " cache_subdir=os.path.abspath('.'),\n", + " origin = 'http://images.cocodataset.org/zips/train2014.zip',\n", + " extract = True)\n", + " PATH = os.path.dirname(image_zip)+'/train2014/'\n", + "else:\n", + " PATH = os.path.abspath('.')+'/train2014/'" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "aANEzb5WwSzg" + }, + "source": [ + "## Optionally, limit the size of the training set for faster training\n", + "For this example, we'll select a subset of 30,000 captions and use these and the corresponding images to train our model. As always, captioning quality will improve if you choose to use more data." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "8Q44tNQVRPFt", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "tokenizer.word_index = {key:value for key, value in tokenizer.word_index.items() if value <= top_k}\n", - "# putting token in the word2idx dictionary\n", - "tokenizer.word_index[tokenizer.oov_token] = top_k + 1\n", - "tokenizer.word_index[''] = 0" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "4G3b8x8_RPFD" + }, + "outputs": [], + "source": [ + "# read the json file\n", + "with open(annotation_file, 'r') as f:\n", + " annotations = json.load(f)\n", + "\n", + "# storing the captions and the image name in vectors\n", + "all_captions = []\n", + "all_img_name_vector = []\n", + "\n", + "for annot in annotations['annotations']:\n", + " caption = ' ' + annot['caption'] + ' '\n", + " image_id = annot['image_id']\n", + " full_coco_image_path = PATH + 'COCO_train2014_' + '%012d.jpg' % (image_id)\n", + " \n", + " all_img_name_vector.append(full_coco_image_path)\n", + " all_captions.append(caption)\n", + "\n", + "# shuffling the captions and image_names together\n", + "# setting a random state\n", + "train_captions, img_name_vector = shuffle(all_captions,\n", + " all_img_name_vector,\n", + " random_state=1)\n", + "\n", + "# selecting the first 30000 captions from the shuffled set\n", + "num_examples = 30000\n", + "train_captions = train_captions[:num_examples]\n", + "img_name_vector = img_name_vector[:num_examples]" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "0fpJb5ojRPFv", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# creating the tokenized vectors\n", - "train_seqs = tokenizer.texts_to_sequences(train_captions)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "mPBMgK34RPFL" + }, + "outputs": [], + "source": [ + "len(train_captions), len(all_captions)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "8cSW4u-ORPFQ" + }, + "source": [ + "## Preprocess the images using InceptionV3\n", + "Next, we will use InceptionV3 (pretrained on Imagenet) to classify each image. We will extract features from the last convolutional layer. \n", + "\n", + "First, we will need to convert the images into the format inceptionV3 expects by:\n", + "* Resizing the image to (299, 299)\n", + "* Using the [preprocess_input](https://www.tensorflow.org/api_docs/python/tf/keras/applications/inception_v3/preprocess_input) method to place the pixels in the range of -1 to 1 (to match the format of the images used to train InceptionV3)." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "olQArbgbRPF1", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# creating a reverse mapping (index -> word)\n", - "index_word = {value:key for key, value in tokenizer.word_index.items()}" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "zXR0217aRPFR" + }, + "outputs": [], + "source": [ + "def load_image(image_path):\n", + " img = tf.read_file(image_path)\n", + " img = tf.image.decode_jpeg(img, channels=3)\n", + " img = tf.image.resize_images(img, (299, 299))\n", + " img = tf.keras.applications.inception_v3.preprocess_input(img)\n", + " return img, image_path" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "MDvIu4sXRPFV" + }, + "source": [ + "## Initialize InceptionV3 and load the pretrained Imagenet weights\n", + "\n", + "To do so, we'll create a tf.keras model where the output layer is the last convolutional layer in the InceptionV3 architecture. \n", + "* Each image is forwarded through the network and the vector that we get at the end is stored in a dictionary (image_name --> feature_vector). \n", + "* We use the last convolutional layer because we are using attention in this example. The shape of the output of this layer is ```8x8x2048```. \n", + "* We avoid doing this during training so it does not become a bottleneck. \n", + "* After all the images are passed through the network, we pickle the dictionary and save it to disk." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "AidglIZVRPF4", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# padding each vector to the max_length of the captions\n", - "# if the max_length parameter is not provided, pad_sequences calculates that automatically\n", - "cap_vector = tf.keras.preprocessing.sequence.pad_sequences(train_seqs, padding='post')" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "RD3vW4SsRPFW" + }, + "outputs": [], + "source": [ + "image_model = tf.keras.applications.InceptionV3(include_top=False, \n", + " weights='imagenet')\n", + "new_input = image_model.input\n", + "hidden_layer = image_model.layers[-1].output\n", + "\n", + "image_features_extract_model = tf.keras.Model(new_input, hidden_layer)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "rERqlR3WRPGO" + }, + "source": [ + "## Caching the features extracted from InceptionV3\n", + "\n", + "We will pre-process each image with InceptionV3 and cache the output to disk. Caching the output in RAM would be faster but memory intensive, requiring 8 \\* 8 \\* 2048 floats per image. At the time of writing, this would exceed the memory limitations of Colab (although these may change, an instance appears to have about 12GB of memory currently). \n", + "\n", + "Performance could be improved with a more sophisticated caching strategy (e.g., by sharding the images to reduce random access disk I/O) at the cost of more code.\n", + "\n", + "This will take about 10 minutes to run in Colab with a GPU. If you'd like to see a progress bar, you could: install [tqdm](https://github.com/tqdm/tqdm) (```!pip install tqdm```), then change this line: \n", + "\n", + "```for img, path in image_dataset:``` \n", + "\n", + "to:\n", + "\n", + "```for img, path in tqdm(image_dataset):```." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "gL0wkttkRPGA", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# calculating the max_length \n", - "# used to store the attention weights\n", - "max_length = calc_max_length(train_seqs)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "Dx_fvbVgRPGQ" + }, + "outputs": [], + "source": [ + "# getting the unique images\n", + "encode_train = sorted(set(img_name_vector))\n", + "\n", + "# feel free to change the batch_size according to your system configuration\n", + "image_dataset = tf.data.Dataset.from_tensor_slices(\n", + " encode_train).map(load_image).batch(16)\n", + "\n", + "for img, path in image_dataset:\n", + " batch_features = image_features_extract_model(img)\n", + " batch_features = tf.reshape(batch_features, \n", + " (batch_features.shape[0], -1, batch_features.shape[3]))\n", + "\n", + " for bf, p in zip(batch_features, path):\n", + " path_of_feature = p.numpy().decode(\"utf-8\")\n", + " np.save(path_of_feature, bf.numpy())" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "nyqH3zFwRPFi" + }, + "source": [ + "## Preprocess and tokenize the captions\n", + "\n", + "* First, we'll tokenize the captions (e.g., by splitting on spaces). This will give us a vocabulary of all the unique words in the data (e.g., \"surfing\", \"football\", etc).\n", + "* Next, we'll limit the vocabulary size to the top 5,000 words to save memory. We'll replace all other words with the token \"UNK\" (for unknown).\n", + "* Finally, we create a word --> index mapping and vice-versa.\n", + "* We will then pad all sequences to the be same length as the longest one. " + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "M3CD75nDpvTI", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Split the data into training and testing" - ] + "colab_type": "code", + "id": "HZfK8RhQRPFj" + }, + "outputs": [], + "source": [ + "# This will find the maximum length of any caption in our dataset\n", + "def calc_max_length(tensor):\n", + " return max(len(t) for t in tensor)" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "iS7DDMszRPGF", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# Create training and validation sets using 80-20 split\n", - "img_name_train, img_name_val, cap_train, cap_val = train_test_split(img_name_vector, \n", - " cap_vector, \n", - " test_size=0.2, \n", - " random_state=0)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "oJGE34aiRPFo" + }, + "outputs": [], + "source": [ + "# The steps above is a general process of dealing with text processing\n", + "\n", + "# choosing the top 5000 words from the vocabulary\n", + "top_k = 5000\n", + "tokenizer = tf.keras.preprocessing.text.Tokenizer(num_words=top_k, \n", + " oov_token=\"\", \n", + " filters='!\"#$%&()*+.,-/:;=?@[\\]^_`{|}~ ')\n", + "tokenizer.fit_on_texts(train_captions)\n", + "train_seqs = tokenizer.texts_to_sequences(train_captions)" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "XmViPkRFRPGH", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "len(img_name_train), len(cap_train), len(img_name_val), len(cap_val)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "8Q44tNQVRPFt" + }, + "outputs": [], + "source": [ + "tokenizer.word_index[''] = 0" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "uEWM9xrYcg45", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Our images and captions are ready! Next, let's create a tf.data dataset to use for training our model.\n", - "\n" - ] + "colab_type": "code", + "id": "0fpJb5ojRPFv" + }, + "outputs": [], + "source": [ + "# creating the tokenized vectors\n", + "train_seqs = tokenizer.texts_to_sequences(train_captions)" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "Q3TnZ1ToRPGV", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# feel free to change these parameters according to your system's configuration\n", - "\n", - "BATCH_SIZE = 64\n", - "BUFFER_SIZE = 1000\n", - "embedding_dim = 256\n", - "units = 512\n", - "vocab_size = len(tokenizer.word_index)\n", - "# shape of the vector extracted from InceptionV3 is (64, 2048)\n", - "# these two variables represent that\n", - "features_shape = 2048\n", - "attention_features_shape = 64" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "olQArbgbRPF1" + }, + "outputs": [], + "source": [ + "# creating a reverse mapping (index -> word)\n", + "index_word = {value:key for key, value in tokenizer.word_index.items()}" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "SmZS2N0bXG3T", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# loading the numpy files \n", - "def map_func(img_name, cap):\n", - " img_tensor = np.load(img_name.decode('utf-8')+'.npy')\n", - " return img_tensor, cap" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "AidglIZVRPF4" + }, + "outputs": [], + "source": [ + "# padding each vector to the max_length of the captions\n", + "# if the max_length parameter is not provided, pad_sequences calculates that automatically\n", + "cap_vector = tf.keras.preprocessing.sequence.pad_sequences(train_seqs, padding='post')" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "FDF_Nm3tRPGZ", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "dataset = tf.data.Dataset.from_tensor_slices((img_name_train, cap_train))\n", - "\n", - "# using map to load the numpy files in parallel\n", - "# NOTE: Be sure to set num_parallel_calls to the number of CPU cores you have\n", - "# https://www.tensorflow.org/api_docs/python/tf/py_func\n", - "dataset = dataset.map(lambda item1, item2: tf.py_func(\n", - " map_func, [item1, item2], [tf.float32, tf.int32]), num_parallel_calls=8)\n", - "\n", - "# shuffling and batching\n", - "dataset = dataset.shuffle(BUFFER_SIZE)\n", - "# https://www.tensorflow.org/api_docs/python/tf/contrib/data/batch_and_drop_remainder\n", - "dataset = dataset.batch(BATCH_SIZE)\n", - "dataset = dataset.prefetch(1)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "gL0wkttkRPGA" + }, + "outputs": [], + "source": [ + "# calculating the max_length \n", + "# used to store the attention weights\n", + "max_length = calc_max_length(train_seqs)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "M3CD75nDpvTI" + }, + "source": [ + "## Split the data into training and testing" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "nrvoDphgRPGd", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Model\n", - "\n", - "Fun fact, the decoder below is identical to the one in the example for [Neural Machine Translation with Attention]( https://github.com/tensorflow/tensorflow/blob/master/tensorflow/contrib/eager/python/examples/nmt_with_attention/nmt_with_attention.ipynb).\n", - "\n", - "The model architecture is inspired by the [Show, Attend and Tell](https://arxiv.org/pdf/1502.03044.pdf) paper.\n", - "\n", - "* In this example, we extract the features from the lower convolutional layer of InceptionV3 giving us a vector of shape (8, 8, 2048). \n", - "* We squash that to a shape of (64, 2048).\n", - "* This vector is then passed through the CNN Encoder(which consists of a single Fully connected layer).\n", - "* The RNN(here GRU) attends over the image to predict the next word." - ] + "colab_type": "code", + "id": "iS7DDMszRPGF" + }, + "outputs": [], + "source": [ + "# Create training and validation sets using 80-20 split\n", + "img_name_train, img_name_val, cap_train, cap_val = train_test_split(img_name_vector, \n", + " cap_vector, \n", + " test_size=0.2, \n", + " random_state=0)" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "AAppCGLKRPGd", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "def gru(units):\n", - " # If you have a GPU, we recommend using the CuDNNGRU layer (it provides a \n", - " # significant speedup).\n", - " if tf.test.is_gpu_available():\n", - " return tf.keras.layers.CuDNNGRU(units, \n", - " return_sequences=True, \n", - " return_state=True, \n", - " recurrent_initializer='glorot_uniform')\n", - " else:\n", - " return tf.keras.layers.GRU(units, \n", - " return_sequences=True, \n", - " return_state=True, \n", - " recurrent_activation='sigmoid', \n", - " recurrent_initializer='glorot_uniform')" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "XmViPkRFRPGH" + }, + "outputs": [], + "source": [ + "len(img_name_train), len(cap_train), len(img_name_val), len(cap_val)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "uEWM9xrYcg45" + }, + "source": [ + "## Our images and captions are ready! Next, let's create a tf.data dataset to use for training our model.\n", + "\n" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "ja2LFTMSdeV3", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "class BahdanauAttention(tf.keras.Model):\n", - " def __init__(self, units):\n", - " super(BahdanauAttention, self).__init__()\n", - " self.W1 = tf.keras.layers.Dense(units)\n", - " self.W2 = tf.keras.layers.Dense(units)\n", - " self.V = tf.keras.layers.Dense(1)\n", - " \n", - " def call(self, features, hidden):\n", - " # features(CNN_encoder output) shape == (batch_size, 64, embedding_dim)\n", - " \n", - " # hidden shape == (batch_size, hidden_size)\n", - " # hidden_with_time_axis shape == (batch_size, 1, hidden_size)\n", - " hidden_with_time_axis = tf.expand_dims(hidden, 1)\n", - " \n", - " # score shape == (batch_size, 64, hidden_size)\n", - " score = tf.nn.tanh(self.W1(features) + self.W2(hidden_with_time_axis))\n", - " \n", - " # attention_weights shape == (batch_size, 64, 1)\n", - " # we get 1 at the last axis because we are applying score to self.V\n", - " attention_weights = tf.nn.softmax(self.V(score), axis=1)\n", - " \n", - " # context_vector shape after sum == (batch_size, hidden_size)\n", - " context_vector = attention_weights * features\n", - " context_vector = tf.reduce_sum(context_vector, axis=1)\n", - " \n", - " return context_vector, attention_weights" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "Q3TnZ1ToRPGV" + }, + "outputs": [], + "source": [ + "# feel free to change these parameters according to your system's configuration\n", + "\n", + "BATCH_SIZE = 64\n", + "BUFFER_SIZE = 1000\n", + "embedding_dim = 256\n", + "units = 512\n", + "vocab_size = len(tokenizer.word_index)\n", + "# shape of the vector extracted from InceptionV3 is (64, 2048)\n", + "# these two variables represent that\n", + "features_shape = 2048\n", + "attention_features_shape = 64" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "AZ7R1RxHRPGf", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "class CNN_Encoder(tf.keras.Model):\n", - " # Since we have already extracted the features and dumped it using pickle\n", - " # This encoder passes those features through a Fully connected layer\n", - " def __init__(self, embedding_dim):\n", - " super(CNN_Encoder, self).__init__()\n", - " # shape after fc == (batch_size, 64, embedding_dim)\n", - " self.fc = tf.keras.layers.Dense(embedding_dim)\n", - " \n", - " def call(self, x):\n", - " x = self.fc(x)\n", - " x = tf.nn.relu(x)\n", - " return x" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "SmZS2N0bXG3T" + }, + "outputs": [], + "source": [ + "# loading the numpy files \n", + "def map_func(img_name, cap):\n", + " img_tensor = np.load(img_name.decode('utf-8')+'.npy')\n", + " return img_tensor, cap" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "V9UbGQmERPGi", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "class RNN_Decoder(tf.keras.Model):\n", - " def __init__(self, embedding_dim, units, vocab_size):\n", - " super(RNN_Decoder, self).__init__()\n", - " self.units = units\n", - "\n", - " self.embedding = tf.keras.layers.Embedding(vocab_size, embedding_dim)\n", - " self.gru = gru(self.units)\n", - " self.fc1 = tf.keras.layers.Dense(self.units)\n", - " self.fc2 = tf.keras.layers.Dense(vocab_size)\n", - " \n", - " self.attention = BahdanauAttention(self.units)\n", - " \n", - " def call(self, x, features, hidden):\n", - " # defining attention as a separate model\n", - " context_vector, attention_weights = self.attention(features, hidden)\n", - " \n", - " # x shape after passing through embedding == (batch_size, 1, embedding_dim)\n", - " x = self.embedding(x)\n", - " \n", - " # x shape after concatenation == (batch_size, 1, embedding_dim + hidden_size)\n", - " x = tf.concat([tf.expand_dims(context_vector, 1), x], axis=-1)\n", - " \n", - " # passing the concatenated vector to the GRU\n", - " output, state = self.gru(x)\n", - " \n", - " # shape == (batch_size, max_length, hidden_size)\n", - " x = self.fc1(output)\n", - " \n", - " # x shape == (batch_size * max_length, hidden_size)\n", - " x = tf.reshape(x, (-1, x.shape[2]))\n", - " \n", - " # output shape == (batch_size * max_length, vocab)\n", - " x = self.fc2(x)\n", - "\n", - " return x, state, attention_weights\n", - "\n", - " def reset_state(self, batch_size):\n", - " return tf.zeros((batch_size, self.units))" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "FDF_Nm3tRPGZ" + }, + "outputs": [], + "source": [ + "dataset = tf.data.Dataset.from_tensor_slices((img_name_train, cap_train))\n", + "\n", + "# using map to load the numpy files in parallel\n", + "# NOTE: Be sure to set num_parallel_calls to the number of CPU cores you have\n", + "# https://www.tensorflow.org/api_docs/python/tf/py_func\n", + "dataset = dataset.map(lambda item1, item2: tf.py_func(\n", + " map_func, [item1, item2], [tf.float32, tf.int32]), num_parallel_calls=8)\n", + "\n", + "# shuffling and batching\n", + "dataset = dataset.shuffle(BUFFER_SIZE)\n", + "# https://www.tensorflow.org/api_docs/python/tf/contrib/data/batch_and_drop_remainder\n", + "dataset = dataset.batch(BATCH_SIZE)\n", + "dataset = dataset.prefetch(1)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "nrvoDphgRPGd" + }, + "source": [ + "## Model\n", + "\n", + "Fun fact, the decoder below is identical to the one in the example for [Neural Machine Translation with Attention]( https://github.com/tensorflow/tensorflow/blob/master/tensorflow/contrib/eager/python/examples/nmt_with_attention/nmt_with_attention.ipynb).\n", + "\n", + "The model architecture is inspired by the [Show, Attend and Tell](https://arxiv.org/pdf/1502.03044.pdf) paper.\n", + "\n", + "* In this example, we extract the features from the lower convolutional layer of InceptionV3 giving us a vector of shape (8, 8, 2048). \n", + "* We squash that to a shape of (64, 2048).\n", + "* This vector is then passed through the CNN Encoder(which consists of a single Fully connected layer).\n", + "* The RNN(here GRU) attends over the image to predict the next word." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "Qs_Sr03wRPGk", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "encoder = CNN_Encoder(embedding_dim)\n", - "decoder = RNN_Decoder(embedding_dim, units, vocab_size)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "AAppCGLKRPGd" + }, + "outputs": [], + "source": [ + "def gru(units):\n", + " # If you have a GPU, we recommend using the CuDNNGRU layer (it provides a \n", + " # significant speedup).\n", + " if tf.test.is_gpu_available():\n", + " return tf.keras.layers.CuDNNGRU(units, \n", + " return_sequences=True, \n", + " return_state=True, \n", + " recurrent_initializer='glorot_uniform')\n", + " else:\n", + " return tf.keras.layers.GRU(units, \n", + " return_sequences=True, \n", + " return_state=True, \n", + " recurrent_activation='sigmoid', \n", + " recurrent_initializer='glorot_uniform')" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "-bYN7xA0RPGl", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "optimizer = tf.train.AdamOptimizer()\n", - "\n", - "# We are masking the loss calculated for padding\n", - "def loss_function(real, pred):\n", - " mask = 1 - np.equal(real, 0)\n", - " loss_ = tf.nn.sparse_softmax_cross_entropy_with_logits(labels=real, logits=pred) * mask\n", - " return tf.reduce_mean(loss_)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "ja2LFTMSdeV3" + }, + "outputs": [], + "source": [ + "class BahdanauAttention(tf.keras.Model):\n", + " def __init__(self, units):\n", + " super(BahdanauAttention, self).__init__()\n", + " self.W1 = tf.keras.layers.Dense(units)\n", + " self.W2 = tf.keras.layers.Dense(units)\n", + " self.V = tf.keras.layers.Dense(1)\n", + " \n", + " def call(self, features, hidden):\n", + " # features(CNN_encoder output) shape == (batch_size, 64, embedding_dim)\n", + " \n", + " # hidden shape == (batch_size, hidden_size)\n", + " # hidden_with_time_axis shape == (batch_size, 1, hidden_size)\n", + " hidden_with_time_axis = tf.expand_dims(hidden, 1)\n", + " \n", + " # score shape == (batch_size, 64, hidden_size)\n", + " score = tf.nn.tanh(self.W1(features) + self.W2(hidden_with_time_axis))\n", + " \n", + " # attention_weights shape == (batch_size, 64, 1)\n", + " # we get 1 at the last axis because we are applying score to self.V\n", + " attention_weights = tf.nn.softmax(self.V(score), axis=1)\n", + " \n", + " # context_vector shape after sum == (batch_size, hidden_size)\n", + " context_vector = attention_weights * features\n", + " context_vector = tf.reduce_sum(context_vector, axis=1)\n", + " \n", + " return context_vector, attention_weights" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "PHod7t72RPGn", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Training\n", - "\n", - "* We extract the features stored in the respective `.npy` files and then pass those features through the encoder.\n", - "* The encoder output, hidden state(initialized to 0) and the decoder input (which is the start token) is passed to the decoder.\n", - "* The decoder returns the predictions and the decoder hidden state.\n", - "* The decoder hidden state is then passed back into the model and the predictions are used to calculate the loss.\n", - "* Use teacher forcing to decide the next input to the decoder.\n", - "* Teacher forcing is the technique where the target word is passed as the next input to the decoder.\n", - "* The final step is to calculate the gradients and apply it to the optimizer and backpropagate.\n" - ] + "colab_type": "code", + "id": "AZ7R1RxHRPGf" + }, + "outputs": [], + "source": [ + "class CNN_Encoder(tf.keras.Model):\n", + " # Since we have already extracted the features and dumped it using pickle\n", + " # This encoder passes those features through a Fully connected layer\n", + " def __init__(self, embedding_dim):\n", + " super(CNN_Encoder, self).__init__()\n", + " # shape after fc == (batch_size, 64, embedding_dim)\n", + " self.fc = tf.keras.layers.Dense(embedding_dim)\n", + " \n", + " def call(self, x):\n", + " x = self.fc(x)\n", + " x = tf.nn.relu(x)\n", + " return x" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "Vt4WZ5mhJE-E", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# adding this in a separate cell because if you run the training cell \n", - "# many times, the loss_plot array will be reset\n", - "loss_plot = []" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "V9UbGQmERPGi" + }, + "outputs": [], + "source": [ + "class RNN_Decoder(tf.keras.Model):\n", + " def __init__(self, embedding_dim, units, vocab_size):\n", + " super(RNN_Decoder, self).__init__()\n", + " self.units = units\n", + "\n", + " self.embedding = tf.keras.layers.Embedding(vocab_size, embedding_dim)\n", + " self.gru = gru(self.units)\n", + " self.fc1 = tf.keras.layers.Dense(self.units)\n", + " self.fc2 = tf.keras.layers.Dense(vocab_size)\n", + " \n", + " self.attention = BahdanauAttention(self.units)\n", + " \n", + " def call(self, x, features, hidden):\n", + " # defining attention as a separate model\n", + " context_vector, attention_weights = self.attention(features, hidden)\n", + " \n", + " # x shape after passing through embedding == (batch_size, 1, embedding_dim)\n", + " x = self.embedding(x)\n", + " \n", + " # x shape after concatenation == (batch_size, 1, embedding_dim + hidden_size)\n", + " x = tf.concat([tf.expand_dims(context_vector, 1), x], axis=-1)\n", + " \n", + " # passing the concatenated vector to the GRU\n", + " output, state = self.gru(x)\n", + " \n", + " # shape == (batch_size, max_length, hidden_size)\n", + " x = self.fc1(output)\n", + " \n", + " # x shape == (batch_size * max_length, hidden_size)\n", + " x = tf.reshape(x, (-1, x.shape[2]))\n", + " \n", + " # output shape == (batch_size * max_length, vocab)\n", + " x = self.fc2(x)\n", + "\n", + " return x, state, attention_weights\n", + "\n", + " def reset_state(self, batch_size):\n", + " return tf.zeros((batch_size, self.units))" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "UlA4VIQpRPGo", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "EPOCHS = 20\n", - "\n", - "for epoch in range(EPOCHS):\n", - " start = time.time()\n", - " total_loss = 0\n", - " \n", - " for (batch, (img_tensor, target)) in enumerate(dataset):\n", - " loss = 0\n", - " \n", - " # initializing the hidden state for each batch\n", - " # because the captions are not related from image to image\n", - " hidden = decoder.reset_state(batch_size=target.shape[0])\n", - "\n", - " dec_input = tf.expand_dims([tokenizer.word_index['']] * BATCH_SIZE, 1)\n", - " \n", - " with tf.GradientTape() as tape:\n", - " features = encoder(img_tensor)\n", - " \n", - " for i in range(1, target.shape[1]):\n", - " # passing the features through the decoder\n", - " predictions, hidden, _ = decoder(dec_input, features, hidden)\n", - "\n", - " loss += loss_function(target[:, i], predictions)\n", - " \n", - " # using teacher forcing\n", - " dec_input = tf.expand_dims(target[:, i], 1)\n", - " \n", - " total_loss += (loss / int(target.shape[1]))\n", - " \n", - " variables = encoder.variables + decoder.variables\n", - " \n", - " gradients = tape.gradient(loss, variables) \n", - " \n", - " optimizer.apply_gradients(zip(gradients, variables), tf.train.get_or_create_global_step())\n", - " \n", - " if batch % 100 == 0:\n", - " print ('Epoch {} Batch {} Loss {:.4f}'.format(epoch + 1, \n", - " batch, \n", - " loss.numpy() / int(target.shape[1])))\n", - " # storing the epoch end loss value to plot later\n", - " loss_plot.append(total_loss / len(cap_vector))\n", - " \n", - " print ('Epoch {} Loss {:.6f}'.format(epoch + 1, \n", - " total_loss/len(cap_vector)))\n", - " print ('Time taken for 1 epoch {} sec\\n'.format(time.time() - start))" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "Qs_Sr03wRPGk" + }, + "outputs": [], + "source": [ + "encoder = CNN_Encoder(embedding_dim)\n", + "decoder = RNN_Decoder(embedding_dim, units, vocab_size)" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "1Wm83G-ZBPcC", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "plt.plot(loss_plot)\n", - "plt.xlabel('Epochs')\n", - "plt.ylabel('Loss')\n", - "plt.title('Loss Plot')\n", - "plt.show()" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "-bYN7xA0RPGl" + }, + "outputs": [], + "source": [ + "optimizer = tf.train.AdamOptimizer()\n", + "\n", + "# We are masking the loss calculated for padding\n", + "def loss_function(real, pred):\n", + " mask = 1 - np.equal(real, 0)\n", + " loss_ = tf.nn.sparse_softmax_cross_entropy_with_logits(labels=real, logits=pred) * mask\n", + " return tf.reduce_mean(loss_)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "PHod7t72RPGn" + }, + "source": [ + "## Training\n", + "\n", + "* We extract the features stored in the respective `.npy` files and then pass those features through the encoder.\n", + "* The encoder output, hidden state(initialized to 0) and the decoder input (which is the start token) is passed to the decoder.\n", + "* The decoder returns the predictions and the decoder hidden state.\n", + "* The decoder hidden state is then passed back into the model and the predictions are used to calculate the loss.\n", + "* Use teacher forcing to decide the next input to the decoder.\n", + "* Teacher forcing is the technique where the target word is passed as the next input to the decoder.\n", + "* The final step is to calculate the gradients and apply it to the optimizer and backpropagate.\n" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "xGvOcLQKghXN", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Caption!\n", - "\n", - "* The evaluate function is similar to the training loop, except we don't use teacher forcing here. The input to the decoder at each time step is its previous predictions along with the hidden state and the encoder output.\n", - "* Stop predicting when the model predicts the end token.\n", - "* And store the attention weights for every time step." - ] + "colab_type": "code", + "id": "Vt4WZ5mhJE-E" + }, + "outputs": [], + "source": [ + "# adding this in a separate cell because if you run the training cell \n", + "# many times, the loss_plot array will be reset\n", + "loss_plot = []" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "RCWpDtyNRPGs", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "def evaluate(image):\n", - " attention_plot = np.zeros((max_length, attention_features_shape))\n", - "\n", - " hidden = decoder.reset_state(batch_size=1)\n", - "\n", - " temp_input = tf.expand_dims(load_image(image)[0], 0)\n", - " img_tensor_val = image_features_extract_model(temp_input)\n", - " img_tensor_val = tf.reshape(img_tensor_val, (img_tensor_val.shape[0], -1, img_tensor_val.shape[3]))\n", - "\n", - " features = encoder(img_tensor_val)\n", - "\n", - " dec_input = tf.expand_dims([tokenizer.word_index['']], 0)\n", - " result = []\n", - "\n", - " for i in range(max_length):\n", - " predictions, hidden, attention_weights = decoder(dec_input, features, hidden)\n", - "\n", - " attention_plot[i] = tf.reshape(attention_weights, (-1, )).numpy()\n", - "\n", - " predicted_id = tf.argmax(predictions[0]).numpy()\n", - " result.append(index_word[predicted_id])\n", - "\n", - " if index_word[predicted_id] == '':\n", - " return result, attention_plot\n", - "\n", - " dec_input = tf.expand_dims([predicted_id], 0)\n", - "\n", - " attention_plot = attention_plot[:len(result), :]\n", - " return result, attention_plot" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "UlA4VIQpRPGo" + }, + "outputs": [], + "source": [ + "EPOCHS = 20\n", + "\n", + "for epoch in range(EPOCHS):\n", + " start = time.time()\n", + " total_loss = 0\n", + " \n", + " for (batch, (img_tensor, target)) in enumerate(dataset):\n", + " loss = 0\n", + " \n", + " # initializing the hidden state for each batch\n", + " # because the captions are not related from image to image\n", + " hidden = decoder.reset_state(batch_size=target.shape[0])\n", + "\n", + " dec_input = tf.expand_dims([tokenizer.word_index['']] * BATCH_SIZE, 1)\n", + " \n", + " with tf.GradientTape() as tape:\n", + " features = encoder(img_tensor)\n", + " \n", + " for i in range(1, target.shape[1]):\n", + " # passing the features through the decoder\n", + " predictions, hidden, _ = decoder(dec_input, features, hidden)\n", + "\n", + " loss += loss_function(target[:, i], predictions)\n", + " \n", + " # using teacher forcing\n", + " dec_input = tf.expand_dims(target[:, i], 1)\n", + " \n", + " total_loss += (loss / int(target.shape[1]))\n", + " \n", + " variables = encoder.variables + decoder.variables\n", + " \n", + " gradients = tape.gradient(loss, variables) \n", + " \n", + " optimizer.apply_gradients(zip(gradients, variables), tf.train.get_or_create_global_step())\n", + " \n", + " if batch % 100 == 0:\n", + " print ('Epoch {} Batch {} Loss {:.4f}'.format(epoch + 1, \n", + " batch, \n", + " loss.numpy() / int(target.shape[1])))\n", + " # storing the epoch end loss value to plot later\n", + " loss_plot.append(total_loss / len(cap_vector))\n", + " \n", + " print ('Epoch {} Loss {:.6f}'.format(epoch + 1, \n", + " total_loss/len(cap_vector)))\n", + " print ('Time taken for 1 epoch {} sec\\n'.format(time.time() - start))" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "fD_y7PD6RPGt", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "def plot_attention(image, result, attention_plot):\n", - " temp_image = np.array(Image.open(image))\n", - "\n", - " fig = plt.figure(figsize=(10, 10))\n", - " \n", - " len_result = len(result)\n", - " for l in range(len_result):\n", - " temp_att = np.resize(attention_plot[l], (8, 8))\n", - " ax = fig.add_subplot(len_result//2, len_result//2, l+1)\n", - " ax.set_title(result[l])\n", - " img = ax.imshow(temp_image)\n", - " ax.imshow(temp_att, cmap='gray', alpha=0.6, extent=img.get_extent())\n", - "\n", - " plt.tight_layout()\n", - " plt.show()" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "1Wm83G-ZBPcC" + }, + "outputs": [], + "source": [ + "plt.plot(loss_plot)\n", + "plt.xlabel('Epochs')\n", + "plt.ylabel('Loss')\n", + "plt.title('Loss Plot')\n", + "plt.show()" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "xGvOcLQKghXN" + }, + "source": [ + "## Caption!\n", + "\n", + "* The evaluate function is similar to the training loop, except we don't use teacher forcing here. The input to the decoder at each time step is its previous predictions along with the hidden state and the encoder output.\n", + "* Stop predicting when the model predicts the end token.\n", + "* And store the attention weights for every time step." + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "io7ws3ReRPGv", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "# captions on the validation set\n", - "rid = np.random.randint(0, len(img_name_val))\n", - "image = img_name_val[rid]\n", - "real_caption = ' '.join([index_word[i] for i in cap_val[rid] if i not in [0]])\n", - "result, attention_plot = evaluate(image)\n", - "\n", - "print ('Real Caption:', real_caption)\n", - "print ('Prediction Caption:', ' '.join(result))\n", - "plot_attention(image, result, attention_plot)\n", - "# opening the image\n", - "Image.open(img_name_val[rid])" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "RCWpDtyNRPGs" + }, + "outputs": [], + "source": [ + "def evaluate(image):\n", + " attention_plot = np.zeros((max_length, attention_features_shape))\n", + "\n", + " hidden = decoder.reset_state(batch_size=1)\n", + "\n", + " temp_input = tf.expand_dims(load_image(image)[0], 0)\n", + " img_tensor_val = image_features_extract_model(temp_input)\n", + " img_tensor_val = tf.reshape(img_tensor_val, (img_tensor_val.shape[0], -1, img_tensor_val.shape[3]))\n", + "\n", + " features = encoder(img_tensor_val)\n", + "\n", + " dec_input = tf.expand_dims([tokenizer.word_index['']], 0)\n", + " result = []\n", + "\n", + " for i in range(max_length):\n", + " predictions, hidden, attention_weights = decoder(dec_input, features, hidden)\n", + "\n", + " attention_plot[i] = tf.reshape(attention_weights, (-1, )).numpy()\n", + "\n", + " predicted_id = tf.argmax(predictions[0]).numpy()\n", + " result.append(index_word[predicted_id])\n", + "\n", + " if index_word[predicted_id] == '':\n", + " return result, attention_plot\n", + "\n", + " dec_input = tf.expand_dims([predicted_id], 0)\n", + "\n", + " attention_plot = attention_plot[:len(result), :]\n", + " return result, attention_plot" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "Rprk3HEvZuxb", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "## Try it on your own images\n", - "For fun, below we've provided a method you can use to caption your own images with the model we've just trained. Keep in mind, it was trained on a relatively small amount of data, and your images may be different from the training data (so be prepared for weird results!)\n" - ] + "colab_type": "code", + "id": "fD_y7PD6RPGt" + }, + "outputs": [], + "source": [ + "def plot_attention(image, result, attention_plot):\n", + " temp_image = np.array(Image.open(image))\n", + "\n", + " fig = plt.figure(figsize=(10, 10))\n", + " \n", + " len_result = len(result)\n", + " for l in range(len_result):\n", + " temp_att = np.resize(attention_plot[l], (8, 8))\n", + " ax = fig.add_subplot(len_result//2, len_result//2, l+1)\n", + " ax.set_title(result[l])\n", + " img = ax.imshow(temp_image)\n", + " ax.imshow(temp_att, cmap='gray', alpha=0.6, extent=img.get_extent())\n", + "\n", + " plt.tight_layout()\n", + " plt.show()" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, - { - "metadata": { - "id": "9Psd1quzaAWg", - "colab_type": "code", - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - } - }, - "cell_type": "code", - "source": [ - "image_url = 'https://tensorflow.org/images/surf.jpg'\n", - "image_extension = image_url[-4:]\n", - "image_path = tf.keras.utils.get_file('image'+image_extension, \n", - " origin=image_url)\n", - "\n", - "result, attention_plot = evaluate(image_path)\n", - "print ('Prediction Caption:', ' '.join(result))\n", - "plot_attention(image_path, result, attention_plot)\n", - "# opening the image\n", - "Image.open(image_path)" - ], - "execution_count": 0, - "outputs": [] + "colab_type": "code", + "id": "io7ws3ReRPGv" + }, + "outputs": [], + "source": [ + "# captions on the validation set\n", + "rid = np.random.randint(0, len(img_name_val))\n", + "image = img_name_val[rid]\n", + "real_caption = ' '.join([index_word[i] for i in cap_val[rid] if i not in [0]])\n", + "result, attention_plot = evaluate(image)\n", + "\n", + "print ('Real Caption:', real_caption)\n", + "print ('Prediction Caption:', ' '.join(result))\n", + "plot_attention(image, result, attention_plot)\n", + "# opening the image\n", + "Image.open(img_name_val[rid])" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "Rprk3HEvZuxb" + }, + "source": [ + "## Try it on your own images\n", + "For fun, below we've provided a method you can use to caption your own images with the model we've just trained. Keep in mind, it was trained on a relatively small amount of data, and your images may be different from the training data (so be prepared for weird results!)\n" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "metadata": { + "colab": { + "autoexec": { + "startup": false, + "wait_interval": 0 + } }, + "colab_type": "code", + "id": "9Psd1quzaAWg" + }, + "outputs": [], + "source": [ + "image_url = 'https://tensorflow.org/images/surf.jpg'\n", + "image_extension = image_url[-4:]\n", + "image_path = tf.keras.utils.get_file('image'+image_extension, \n", + " origin=image_url)\n", + "\n", + "result, attention_plot = evaluate(image_path)\n", + "print ('Prediction Caption:', ' '.join(result))\n", + "plot_attention(image_path, result, attention_plot)\n", + "# opening the image\n", + "Image.open(image_path)" + ] + }, + { + "cell_type": "markdown", + "metadata": { + "colab_type": "text", + "id": "VJZXyJco6uLO" + }, + "source": [ + "# Next steps\n", + "\n", + "Congrats! You've just trained an image captioning model with attention. Next, we recommend taking a look at this example [Neural Machine Translation with Attention]( https://github.com/tensorflow/tensorflow/blob/master/tensorflow/contrib/eager/python/examples/nmt_with_attention/nmt_with_attention.ipynb). It uses a similar architecture to translate between Spanish and English sentences. You can also experiment with training the code in this notebook on a different dataset." + ] + } + ], + "metadata": { + "accelerator": "GPU", + "colab": { + "collapsed_sections": [], + "default_view": {}, + "name": "image_captioning_with_attention.ipynb", + "private_outputs": true, + "provenance": [ { - "metadata": { - "id": "VJZXyJco6uLO", - "colab_type": "text" - }, - "cell_type": "markdown", - "source": [ - "# Next steps\n", - "\n", - "Congrats! You've just trained an image captioning model with attention. Next, we recommend taking a look at this example [Neural Machine Translation with Attention]( https://github.com/tensorflow/tensorflow/blob/master/tensorflow/contrib/eager/python/examples/nmt_with_attention/nmt_with_attention.ipynb). It uses a similar architecture to translate between Spanish and English sentences. You can also experiment with training the code in this notebook on a different dataset." - ] + "file_id": "1HI8OK2sMjcx9CTWVn0122QAHOuXaOaMg", + "timestamp": 1530222436922 } - ] + ], + "toc_visible": true, + "version": "0.3.2", + "views": {} + }, + "kernelspec": { + "display_name": "Python 3", + "language": "python", + "name": "python3" + }, + "language_info": { + "codemirror_mode": { + "name": "ipython", + "version": 3 + }, + "file_extension": ".py", + "mimetype": "text/x-python", + "name": "python", + "nbconvert_exporter": "python", + "pygments_lexer": "ipython3", + "version": "3.6.3" + } + }, + "nbformat": 4, + "nbformat_minor": 2 } -- GitLab From ea684a74937fdf18ade43e42b9b320118f70d3c1 Mon Sep 17 00:00:00 2001 From: joaak <29533036+joaak@users.noreply.github.com> Date: Tue, 6 Nov 2018 16:49:18 -0500 Subject: [PATCH 083/596] replace index_word with tokenizer.index_word --- .../image_captioning_with_attention.ipynb | 25 +++---------------- 1 file changed, 3 insertions(+), 22 deletions(-) diff --git a/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb b/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb index 09ea021c44..12c5eff2b4 100644 --- a/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb +++ b/tensorflow/contrib/eager/python/examples/generative_examples/image_captioning_with_attention.ipynb @@ -441,25 +441,6 @@ "train_seqs = tokenizer.texts_to_sequences(train_captions)" ] }, - { - "cell_type": "code", - "execution_count": null, - "metadata": { - "colab": { - "autoexec": { - "startup": false, - "wait_interval": 0 - } - }, - "colab_type": "code", - "id": "olQArbgbRPF1" - }, - "outputs": [], - "source": [ - "# creating a reverse mapping (index -> word)\n", - "index_word = {value:key for key, value in tokenizer.word_index.items()}" - ] - }, { "cell_type": "code", "execution_count": null, @@ -1031,9 +1012,9 @@ " attention_plot[i] = tf.reshape(attention_weights, (-1, )).numpy()\n", "\n", " predicted_id = tf.argmax(predictions[0]).numpy()\n", - " result.append(index_word[predicted_id])\n", + " result.append(tokenizer.index_word[predicted_id])\n", "\n", - " if index_word[predicted_id] == '':\n", + " if tokenizer.index_word[predicted_id] == '':\n", " return result, attention_plot\n", "\n", " dec_input = tf.expand_dims([predicted_id], 0)\n", @@ -1092,7 +1073,7 @@ "# captions on the validation set\n", "rid = np.random.randint(0, len(img_name_val))\n", "image = img_name_val[rid]\n", - "real_caption = ' '.join([index_word[i] for i in cap_val[rid] if i not in [0]])\n", + "real_caption = ' '.join([tokenizer.index_word[i] for i in cap_val[rid] if i not in [0]])\n", "result, attention_plot = evaluate(image)\n", "\n", "print ('Real Caption:', real_caption)\n", -- GitLab From 8e4ec9ae62135adbc523470af1546c178a7f97c5 Mon Sep 17 00:00:00 2001 From: frreiss Date: Tue, 6 Nov 2018 12:55:29 -0800 Subject: [PATCH 084/596] Add missing random seed field to OrderedEnqueuer Fix whitespace Simplify changes Simplify changeset --- tensorflow/python/keras/utils/data_utils.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/keras/utils/data_utils.py b/tensorflow/python/keras/utils/data_utils.py index 01a9d61a84..8e36d4dea7 100644 --- a/tensorflow/python/keras/utils/data_utils.py +++ b/tensorflow/python/keras/utils/data_utils.py @@ -598,7 +598,7 @@ class OrderedEnqueuer(SequenceEnqueuer): def pool_fn(seqs): return multiprocessing.Pool(workers, initializer=init_pool_generator, - initargs=(seqs, self.random_seed)) + initargs=(seqs, None)) return pool_fn def _wait_queue(self): -- GitLab From 3ea1267b9758fdc5582948805cdd852b09f21f6b Mon Sep 17 00:00:00 2001 From: dianlujitao Date: Wed, 7 Nov 2018 13:24:28 +0800 Subject: [PATCH 085/596] Install abseil headers to cmake shared library build * Since commit 5f004516 tensorflow::StringPiece is replaced by absl::string_view, so abseil headers should be installed to shared library build to fix compilation error for out-of-source build. * To cleanly copy abseil headers, disable in source build for abseil to avoid src tree been polluted by cmake generated files. * Meanwhile, remove _build suffix from abseil_cpp product name since it's confusing. --- .../contrib/cmake/external/abseil_cpp.cmake | 15 ++++++--------- tensorflow/contrib/cmake/tf_shared_lib.cmake | 4 ++++ 2 files changed, 10 insertions(+), 9 deletions(-) diff --git a/tensorflow/contrib/cmake/external/abseil_cpp.cmake b/tensorflow/contrib/cmake/external/abseil_cpp.cmake index 4546dbdecc..b0fee24448 100644 --- a/tensorflow/contrib/cmake/external/abseil_cpp.cmake +++ b/tensorflow/contrib/cmake/external/abseil_cpp.cmake @@ -31,17 +31,17 @@ if (systemlib_ABSEIL_CPP) message(STATUS " abseil_cpp includes: ${ABSEIL_CPP_INCLUDE_DIR}") message(STATUS " abseil_cpp libraries: ${ABSEIL_CPP_LIBRARIES}") - add_custom_target(abseil_cpp_build) - list(APPEND tensorflow_EXTERNAL_DEPENDENCIES abseil_cpp_build) + add_custom_target(abseil_cpp) + list(APPEND tensorflow_EXTERNAL_DEPENDENCIES abseil_cpp) else (systemlib_ABSEIL_CPP) include (ExternalProject) - set(abseil_cpp_INCLUDE_DIR ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp_build) + set(abseil_cpp_INCLUDE_DIR ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp) set(abseil_cpp_URL https://github.com/abseil/abseil-cpp/archive/e01d95528ea2137a4a27a88d1f57c6cb260aafed.tar.gz) set(abseil_cpp_HASH SHA256=84043ed402d2a2a6ba4cdddb7e85118b1158fd81fe4ac3a14adc343d054c1e2e) - set(abseil_cpp_BUILD ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp_build) + set(abseil_cpp_BUILD ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp-build) if(WIN32) if(${CMAKE_GENERATOR} MATCHES "Visual Studio.*") @@ -80,15 +80,12 @@ else (systemlib_ABSEIL_CPP) ${abseil_cpp_BUILD}/absl/types/libabsl_bad_optional_access.a) endif() - ExternalProject_Add(abseil_cpp_build + ExternalProject_Add(abseil_cpp PREFIX abseil_cpp URL ${abseil_cpp_URL} URL_HASH ${abseil_cpp_HASH} DOWNLOAD_DIR "${DOWNLOAD_LOCATION}" - BUILD_IN_SOURCE 1 BUILD_BYPRODUCTS ${abseil_cpp_STATIC_LIBRARIES} - BUILD_COMMAND ${CMAKE_COMMAND} --build . --config Release - COMMAND ${CMAKE_COMMAND} --build . --config Release INSTALL_COMMAND "" CMAKE_CACHE_ARGS -DCMAKE_POSITION_INDEPENDENT_CODE:BOOL=${tensorflow_ENABLE_POSITION_INDEPENDENT_CODE} @@ -99,6 +96,6 @@ else (systemlib_ABSEIL_CPP) include_directories(${abseil_cpp_INCLUDE_DIR}) list(APPEND tensorflow_EXTERNAL_LIBRARIES ${abseil_cpp_STATIC_LIBRARIES}) - list(APPEND tensorflow_EXTERNAL_DEPENDENCIES abseil_cpp_build) + list(APPEND tensorflow_EXTERNAL_DEPENDENCIES abseil_cpp) endif (systemlib_ABSEIL_CPP) diff --git a/tensorflow/contrib/cmake/tf_shared_lib.cmake b/tensorflow/contrib/cmake/tf_shared_lib.cmake index fdf522f1fd..c1bdc35fc6 100644 --- a/tensorflow/contrib/cmake/tf_shared_lib.cmake +++ b/tensorflow/contrib/cmake/tf_shared_lib.cmake @@ -145,6 +145,10 @@ install(DIRECTORY ${tensorflow_source_dir}/third_party/eigen3/ # unsupported Eigen directory install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/eigen/src/eigen/unsupported/Eigen/ DESTINATION include/unsupported/Eigen) +# absl directory +install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/abseil_cpp/src/abseil_cpp/absl/ + DESTINATION include/absl + FILES_MATCHING PATTERN "*.h") # mkl if (tensorflow_ENABLE_MKL_SUPPORT) install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/mkl/src/mkl/include/ -- GitLab From c806b163c4d52cf80daecf2d63e3b76c7dc696e6 Mon Sep 17 00:00:00 2001 From: Nutti Date: Sat, 10 Nov 2018 10:41:09 +0900 Subject: [PATCH 086/596] OptimizationPass::POST_REWRITE_FOR_EXEC after Grappler optimization in PartitionedCallOp --- tensorflow/core/kernels/partitioned_function_ops.cc | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/tensorflow/core/kernels/partitioned_function_ops.cc b/tensorflow/core/kernels/partitioned_function_ops.cc index 71e506e5e6..72310f33ae 100644 --- a/tensorflow/core/kernels/partitioned_function_ops.cc +++ b/tensorflow/core/kernels/partitioned_function_ops.cc @@ -166,12 +166,6 @@ class PartitionedCallOp : public AsyncOpKernel { OptimizationPassRegistry::Global()->RunGrouping( OptimizationPassRegistry::POST_PLACEMENT, optimization_options), done); - OP_REQUIRES_OK_ASYNC( - ctx, - OptimizationPassRegistry::Global()->RunGrouping( - OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, - optimization_options), - done); Device* cpu_device; OP_REQUIRES_OK_ASYNC( @@ -184,6 +178,13 @@ class PartitionedCallOp : public AsyncOpKernel { device_set, cpu_device, &graph), done); + OP_REQUIRES_OK_ASYNC( + ctx, + OptimizationPassRegistry::Global()->RunGrouping( + OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, + optimization_options), + done); + std::unordered_map> subgraphs; OP_REQUIRES_OK_ASYNC( ctx, PartitionHelper(device_set, std::move(graph), &subgraphs), -- GitLab From 7d96d6fbd3cdbe215c9dce78f8227ef273b5d37a Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 10 Nov 2018 20:28:31 +0000 Subject: [PATCH 087/596] Fix ValueError by image.transform in eager mode This fix tries to address the issue raised in 23654 where in eager mode tf.contrib.image.transform will throw out ``` ValueError: The truth value of an array with more than one element is ambiguous. Use a.any() or a.all() ``` This fix addresses the issue. This fix fixes 23654. Signed-off-by: Yong Tang --- tensorflow/contrib/image/python/ops/image_ops.py | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/image/python/ops/image_ops.py b/tensorflow/contrib/image/python/ops/image_ops.py index d4fb99a017..b25a6f7b57 100644 --- a/tensorflow/contrib/image/python/ops/image_ops.py +++ b/tensorflow/contrib/image/python/ops/image_ops.py @@ -17,6 +17,7 @@ from __future__ import absolute_import from __future__ import division from __future__ import print_function +from tensorflow.python.eager import context from tensorflow.contrib.image.ops import gen_image_ops from tensorflow.contrib.util import loader from tensorflow.python.framework import common_shapes @@ -271,8 +272,11 @@ def transform(images, raise TypeError("Images should have rank between 2 and 4.") if output_shape is None: - output_shape = tensor_util.constant_value( - array_ops.shape(images)[1:3]) or array_ops.shape(images)[1:3] + output_shape = array_ops.shape(images)[1:3] + if not context.executing_eagerly(): + output_shape_value = tensor_util.constant_value(output_shape) + if output_shape_value is not None: + output_shape = output_shape_value output_shape = ops.convert_to_tensor( output_shape, dtypes.int32, name="output_shape") -- GitLab From 38455b2e111fa1acef3dd7dd00517b3fc1f1c38f Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 10 Nov 2018 20:30:22 +0000 Subject: [PATCH 088/596] Add test case for image.transform in eager mode Signed-off-by: Yong Tang --- .../contrib/image/python/kernel_tests/image_ops_test.py | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py b/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py index 4997c31a7f..ebf8a8adb3 100644 --- a/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py +++ b/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py @@ -281,6 +281,14 @@ class ImageOpsTest(test_util.TensorFlowTestCase): value.eval(), np.array([[4, 4], [4, 4]]).astype(dtype.as_numpy_dtype())) + @test_util.run_in_graph_and_eager_modes + def test_transform_eager(self): + image = constant_op.constant([[1., 2.], [3., 4.]]) + value = image_ops.transform(image, [1] * 8) + with self.test_session(use_gpu=True): + self.assertAllEqual( + self.evaluate(value), np.array([[4, 4], [4, 4]])) + class BipartiteMatchTest(test_util.TensorFlowTestCase): -- GitLab From c50685de2d680d7e76e4b586e14138f33272a9cb Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 10 Nov 2018 20:33:28 +0000 Subject: [PATCH 089/596] Pylint fix Signed-off-by: Yong Tang --- tensorflow/contrib/image/python/kernel_tests/image_ops_test.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py b/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py index ebf8a8adb3..ba5cdfebf9 100644 --- a/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py +++ b/tensorflow/contrib/image/python/kernel_tests/image_ops_test.py @@ -286,8 +286,7 @@ class ImageOpsTest(test_util.TensorFlowTestCase): image = constant_op.constant([[1., 2.], [3., 4.]]) value = image_ops.transform(image, [1] * 8) with self.test_session(use_gpu=True): - self.assertAllEqual( - self.evaluate(value), np.array([[4, 4], [4, 4]])) + self.assertAllEqual(self.evaluate(value), np.array([[4, 4], [4, 4]])) class BipartiteMatchTest(test_util.TensorFlowTestCase): -- GitLab From 6a2f6f8f68ffc55aca43c85ffb0be1cee17d99c8 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Wed, 14 Nov 2018 08:34:30 +0800 Subject: [PATCH 090/596] abseil build in windows --- tensorflow/contrib/cmake/CMakeLists.txt | 2 +- tensorflow/contrib/cmake/external/abseil_cpp.cmake | 9 ++++----- tensorflow/contrib/cmake/tf_python.cmake | 2 -- 3 files changed, 5 insertions(+), 8 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 39fe8ff00e..d2988a4c7c 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -108,7 +108,7 @@ if (NOT WIN32) # Options for linking other libraries option(systemlib_ZLIB "Use the system installed library as shared objects instead of downloading ZLIB and statically linking to it: ZLIB" OFF) - option(systemlib_ABSEIL_CPP "Use the system installed library as shared objects instead of downloading ABSEIL_CPP and statically linking to it: ABSEIL_CPP" OFF) + option(systemlib_ABSEIL_CPP "Use the system installed library as shared objects instead of downloading ABSEIL_CPP and statically linking to it: ABSEIL_CPP" ON) option(systemlib_ALL "Turn on every possible systemlib_* options" OFF) if (systemlib_ALL) diff --git a/tensorflow/contrib/cmake/external/abseil_cpp.cmake b/tensorflow/contrib/cmake/external/abseil_cpp.cmake index 4546dbdecc..efa8d86d69 100644 --- a/tensorflow/contrib/cmake/external/abseil_cpp.cmake +++ b/tensorflow/contrib/cmake/external/abseil_cpp.cmake @@ -39,8 +39,8 @@ else (systemlib_ABSEIL_CPP) include (ExternalProject) set(abseil_cpp_INCLUDE_DIR ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp_build) - set(abseil_cpp_URL https://github.com/abseil/abseil-cpp/archive/e01d95528ea2137a4a27a88d1f57c6cb260aafed.tar.gz) - set(abseil_cpp_HASH SHA256=84043ed402d2a2a6ba4cdddb7e85118b1158fd81fe4ac3a14adc343d054c1e2e) + set(abseil_cpp_URL https://github.com/abseil/abseil-cpp.git) + set(abseil_cpp_TAG master) set(abseil_cpp_BUILD ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp_build) if(WIN32) @@ -82,8 +82,7 @@ else (systemlib_ABSEIL_CPP) ExternalProject_Add(abseil_cpp_build PREFIX abseil_cpp - URL ${abseil_cpp_URL} - URL_HASH ${abseil_cpp_HASH} + GIT_REPOSITORY ${abseil_cpp_URL} DOWNLOAD_DIR "${DOWNLOAD_LOCATION}" BUILD_IN_SOURCE 1 BUILD_BYPRODUCTS ${abseil_cpp_STATIC_LIBRARIES} @@ -101,4 +100,4 @@ else (systemlib_ABSEIL_CPP) list(APPEND tensorflow_EXTERNAL_DEPENDENCIES abseil_cpp_build) -endif (systemlib_ABSEIL_CPP) +endif (systemlib_ABSEIL_CPP) \ No newline at end of file diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index 7647fe565d..4a9732d757 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -765,8 +765,6 @@ string(REPLACE "# BEGIN GENERATED FILES" "" api_init_files_text ${api_init_files string(REPLACE "# END GENERATED FILES" "" api_init_files_text ${api_init_files_text}) string(REPLACE "," ";" api_init_files_list ${api_init_files_text}) -message(STATUS ${api_init_files_text}) - set(api_init_files "") foreach(api_init_file ${api_init_files_list}) string(STRIP "${api_init_file}" api_init_file) -- GitLab From ada03a97ab77bdd58cd96d3a3fdca490278165d4 Mon Sep 17 00:00:00 2001 From: Abhinav Upadhyay Date: Wed, 14 Nov 2018 18:34:43 +0530 Subject: [PATCH 091/596] Fix a TypeError We cannot concatenate string and FailedPreConditionError --- tensorflow/contrib/tpu/python/tpu/keras_support.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/tpu/python/tpu/keras_support.py b/tensorflow/contrib/tpu/python/tpu/keras_support.py index 08f58a5f5b..ac72e3f55d 100644 --- a/tensorflow/contrib/tpu/python/tpu/keras_support.py +++ b/tensorflow/contrib/tpu/python/tpu/keras_support.py @@ -132,7 +132,7 @@ def _tpu_session_context(): An error occurred connecting or initializing your TPU. The session has been reset. re-run keras_to_tpu_model to create a new session. -""" + e) +""" + str(e)) def setup_tpu_session(cluster_resolver): -- GitLab From a3a8dfb385be626748bc31a1e10806af5d35dfd5 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Wed, 14 Nov 2018 22:32:17 +0800 Subject: [PATCH 092/596] nccl is removed from contrib --- tensorflow/contrib/cmake/CMakeLists.txt | 10 +++++++++- tensorflow/contrib/cmake/external/abseil_cpp.cmake | 6 +++--- tensorflow/contrib/cmake/modules/FindAbseilCpp.cmake | 6 +++--- tensorflow/contrib/cmake/tf_python.cmake | 3 --- 4 files changed, 15 insertions(+), 10 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index d2988a4c7c..3da938a8d0 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -108,7 +108,7 @@ if (NOT WIN32) # Options for linking other libraries option(systemlib_ZLIB "Use the system installed library as shared objects instead of downloading ZLIB and statically linking to it: ZLIB" OFF) - option(systemlib_ABSEIL_CPP "Use the system installed library as shared objects instead of downloading ABSEIL_CPP and statically linking to it: ABSEIL_CPP" ON) + option(systemlib_ABSEIL_CPP "Use the system installed library as shared objects instead of downloading ABSEIL_CPP and statically linking to it: ABSEIL_CPP" OFF) option(systemlib_ALL "Turn on every possible systemlib_* options" OFF) if (systemlib_ALL) @@ -293,6 +293,14 @@ else (systemlib_ZLIB) ${zlib_STATIC_LIBRARIES}) endif (systemlib_ZLIB) +if (systemlib_ABSEIL_CPP) + set(tensorflow_EXTERNAL_LIBRARIES ${tensorflow_EXTERNAL_LIBRARIES} + ${abseil_cpp_LIBRARIES}) +else (systemlib_ABSEIL_CPP) + set(tensorflow_EXTERNAL_LIBRARIES ${tensorflow_EXTERNAL_LIBRARIES} + ${abseil_cpp_STATIC_LIBRARIES}) +endif (systemlib_ABSEIL_CPP) + set(tensorflow_EXTERNAL_DEPENDENCIES zlib_copy_headers_to_destination gif_copy_headers_to_destination diff --git a/tensorflow/contrib/cmake/external/abseil_cpp.cmake b/tensorflow/contrib/cmake/external/abseil_cpp.cmake index efa8d86d69..a97704650a 100644 --- a/tensorflow/contrib/cmake/external/abseil_cpp.cmake +++ b/tensorflow/contrib/cmake/external/abseil_cpp.cmake @@ -47,10 +47,10 @@ else (systemlib_ABSEIL_CPP) if(${CMAKE_GENERATOR} MATCHES "Visual Studio.*") set(abseil_cpp_STATIC_LIBRARIES ${abseil_cpp_BUILD}/absl/base/Release/absl_base.lib - ${abseil_cpp_BUILD}/absl/base/Release/absl_spinlock_wait.lib + ${abseil_cpp_BUILD}/absl/base/Release/absl_internal_spinlock_wait.lib ${abseil_cpp_BUILD}/absl/base/Release/absl_dynamic_annotations.lib - ${abseil_cpp_BUILD}/absl/base/Release/absl_malloc_internal.lib - ${abseil_cpp_BUILD}/absl/base/Release/absl_throw_delegate.lib + ${abseil_cpp_BUILD}/absl/base/Release/absl_internal_malloc_internal.lib + ${abseil_cpp_BUILD}/absl/base/Release/absl_internal_throw_delegate.lib ${abseil_cpp_BUILD}/absl/numeric/Release/absl_int128.lib ${abseil_cpp_BUILD}/absl/strings/Release/absl_strings.lib ${abseil_cpp_BUILD}/absl/strings/Release/str_format_internal.lib diff --git a/tensorflow/contrib/cmake/modules/FindAbseilCpp.cmake b/tensorflow/contrib/cmake/modules/FindAbseilCpp.cmake index d4f8bb1bec..944ae3997a 100644 --- a/tensorflow/contrib/cmake/modules/FindAbseilCpp.cmake +++ b/tensorflow/contrib/cmake/modules/FindAbseilCpp.cmake @@ -24,10 +24,10 @@ if(EXISTS "${ABSEIL_CPP_INCLUDE_DIR}" AND NOT "${ABSEIL_CPP_INCLUDE_DIR}" STREQU # search all libraries if no COMPONENTS was requested set(AbseilCpp_FIND_COMPONENTS "absl_algorithm;absl_any;absl_bad_any_cast" - "absl_bad_optional_access;absl_base absl_container;absl_debugging" + "absl_bad_optional_access;absl_base;absl_container;absl_debugging" "absl_dynamic_annotations;absl_examine_stack;absl_failure_signal_handler" - "absl_int128;absl_leak_check;absl_malloc_internal;absl_memory;absl_meta" - "absl_numeric;absl_optional;absl_span;absl_spinlock_wait;absl_stack_consumption" + "absl_int128;absl_leak_check;absl_internal_malloc_internal;absl_memory;absl_meta" + "absl_numeric;absl_optional;absl_span;absl_internal_spinlock_wait;absl_stack_consumption" "absl_stacktrace;absl_str_format;absl_strings;absl_symbolize;absl_synchronization" "absl_throw_delegate;absl_time;absl_utility;str_format_extension_internal" "str_format_internal;test_instance_tracker_lib") diff --git a/tensorflow/contrib/cmake/tf_python.cmake b/tensorflow/contrib/cmake/tf_python.cmake index 4a9732d757..8faccf8d55 100755 --- a/tensorflow/contrib/cmake/tf_python.cmake +++ b/tensorflow/contrib/cmake/tf_python.cmake @@ -398,11 +398,8 @@ GENERATE_PYTHON_OP_LIB("contrib_layers_sparse_feature_cross_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/layers/ops/gen_sparse_feature_cross_op.py) GENERATE_PYTHON_OP_LIB("contrib_memory_stats_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/memory_stats/ops/gen_memory_stats_ops.py) -GENERATE_PYTHON_OP_LIB("contrib_nccl_ops" - DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/nccl/ops/gen_nccl_ops.py) GENERATE_PYTHON_OP_LIB("contrib_periodic_resample_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/periodic_resample/python/ops/gen_periodic_resample_op.py) - GENERATE_PYTHON_OP_LIB("contrib_nearest_neighbor_ops" DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/tf_python/tensorflow/contrib/nearest_neighbor/ops/gen_nearest_neighbor_ops.py) GENERATE_PYTHON_OP_LIB("contrib_resampler_ops" -- GitLab From fbd2c76afc0d884161845e256c7191dc915ec8c3 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Wed, 14 Nov 2018 23:01:27 +0800 Subject: [PATCH 093/596] remove change on pip setup --- tensorflow/tools/pip_package/setup.py | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/tensorflow/tools/pip_package/setup.py b/tensorflow/tools/pip_package/setup.py index afcc7a8601..e15655ea44 100644 --- a/tensorflow/tools/pip_package/setup.py +++ b/tensorflow/tools/pip_package/setup.py @@ -45,10 +45,8 @@ DOCLINES = __doc__.split('\n') # This version string is semver compatible, but incompatible with pip. # For pip, we will remove all '-' characters from this string, and use the # result for pip. - _VERSION = '1.12.0-rc0' - REQUIRED_PACKAGES = [ 'absl-py >= 0.1.6', 'astor >= 0.6.0', @@ -60,8 +58,7 @@ REQUIRED_PACKAGES = [ 'protobuf >= 3.6.1', 'tensorboard >= 1.12.0, < 1.13.0', 'tensorflow_estimator >= 1.10.0', - 'termcolor >= 1.1.0', - 'absl-py >= 0.1.9' + 'termcolor >= 1.1.0' ] if sys.byteorder == 'little': -- GitLab From 7f82d6d20981307ce9c5e40c39d1f510d523e3bb Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Wed, 14 Nov 2018 23:09:31 +0800 Subject: [PATCH 094/596] remove change in non cmake folder --- tensorflow/core/util/cuda_device_functions.h | 14 -------------- tensorflow/core/util/cuda_launch_config.h | 4 ---- tensorflow/tools/pip_package/setup.py | 2 +- 3 files changed, 1 insertion(+), 19 deletions(-) diff --git a/tensorflow/core/util/cuda_device_functions.h b/tensorflow/core/util/cuda_device_functions.h index c05e410abc..2fc142a905 100644 --- a/tensorflow/core/util/cuda_device_functions.h +++ b/tensorflow/core/util/cuda_device_functions.h @@ -30,23 +30,9 @@ limitations under the License. #include #include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" -#ifdef _WIN32 -#include "cuda.h" -#include "device_functions.h" -#else #include "cuda/include/cuda.h" -#include "cuda/include/device_functions.h" -#endif #include "tensorflow/core/platform/types.h" -#if CUDA_VERSION >= 7050 -#ifdef _WIN32 -#include "cuda_fp16.h" -#else -#include "cuda/include/cuda_fp16.h" -#endif -#endif // CUDA_VERSION >= 7050 - namespace tensorflow { namespace detail { diff --git a/tensorflow/core/util/cuda_launch_config.h b/tensorflow/core/util/cuda_launch_config.h index af7ea94ff4..080d4067ce 100644 --- a/tensorflow/core/util/cuda_launch_config.h +++ b/tensorflow/core/util/cuda_launch_config.h @@ -21,11 +21,7 @@ limitations under the License. #include #include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" -#ifdef PLATFORM_GOOGLE #include "cuda/include/cuda.h" -#else -#include "cuda.h" -#endif #include "tensorflow/core/framework/op_kernel.h" #include "tensorflow/core/platform/logging.h" #include "tensorflow/core/platform/stream_executor.h" diff --git a/tensorflow/tools/pip_package/setup.py b/tensorflow/tools/pip_package/setup.py index e15655ea44..07475cc0c4 100644 --- a/tensorflow/tools/pip_package/setup.py +++ b/tensorflow/tools/pip_package/setup.py @@ -58,7 +58,7 @@ REQUIRED_PACKAGES = [ 'protobuf >= 3.6.1', 'tensorboard >= 1.12.0, < 1.13.0', 'tensorflow_estimator >= 1.10.0', - 'termcolor >= 1.1.0' + 'termcolor >= 1.1.0', ] if sys.byteorder == 'little': -- GitLab From 4d150cb81e103759d2f54f179ded573b8c7f07e0 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Wed, 14 Nov 2018 23:13:20 +0800 Subject: [PATCH 095/596] remove change in non cmake folder --- tensorflow/core/platform/default/logging.h | 6 ------ tensorflow/core/util/cuda_device_functions.h | 1 - 2 files changed, 7 deletions(-) diff --git a/tensorflow/core/platform/default/logging.h b/tensorflow/core/platform/default/logging.h index 483f095ee7..08a692fff7 100644 --- a/tensorflow/core/platform/default/logging.h +++ b/tensorflow/core/platform/default/logging.h @@ -187,12 +187,6 @@ string* MakeCheckOpString(const T1& v1, const T2& v2, const char* exprtext) { return comb.NewString(); } -// for MSVC build, the max and min function maybe defined in other macros -#ifdef _WIN32 -#undef max -#undef min -#endif - // Helper functions for CHECK_OP macro. // The (int, int) specialization works around the issue that the compiler // will not instantiate the template version of the function on values of diff --git a/tensorflow/core/util/cuda_device_functions.h b/tensorflow/core/util/cuda_device_functions.h index 2fc142a905..b91f8bb8ef 100644 --- a/tensorflow/core/util/cuda_device_functions.h +++ b/tensorflow/core/util/cuda_device_functions.h @@ -29,7 +29,6 @@ limitations under the License. #include #include #include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" - #include "cuda/include/cuda.h" #include "tensorflow/core/platform/types.h" -- GitLab From 6af7756a91da9a548224c4a836cce078edc77538 Mon Sep 17 00:00:00 2001 From: Jacky Ko Date: Thu, 15 Nov 2018 22:02:49 +0800 Subject: [PATCH 096/596] PR requested changes --- tensorflow/contrib/cmake/CMakeLists.txt | 2 +- .../contrib/cmake/external/abseil_cpp.cmake | 10 +++---- tensorflow/contrib/cmake/tf_c.cmake | 27 ------------------- 3 files changed, 5 insertions(+), 34 deletions(-) diff --git a/tensorflow/contrib/cmake/CMakeLists.txt b/tensorflow/contrib/cmake/CMakeLists.txt index 3da938a8d0..2ad9ae42a1 100644 --- a/tensorflow/contrib/cmake/CMakeLists.txt +++ b/tensorflow/contrib/cmake/CMakeLists.txt @@ -12,7 +12,7 @@ if(WIN32) endif() # Project -project(tensorflow VERSION 1.8.0 LANGUAGES C CXX) +project(tensorflow VERSION 1.12.0 LANGUAGES C CXX) # Set C++14 as standard for the whole project set(CMAKE_CXX_STANDARD 14) diff --git a/tensorflow/contrib/cmake/external/abseil_cpp.cmake b/tensorflow/contrib/cmake/external/abseil_cpp.cmake index a97704650a..8b76f37858 100644 --- a/tensorflow/contrib/cmake/external/abseil_cpp.cmake +++ b/tensorflow/contrib/cmake/external/abseil_cpp.cmake @@ -39,19 +39,16 @@ else (systemlib_ABSEIL_CPP) include (ExternalProject) set(abseil_cpp_INCLUDE_DIR ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp_build) - set(abseil_cpp_URL https://github.com/abseil/abseil-cpp.git) - set(abseil_cpp_TAG master) + set(abseil_cpp_URL https://github.com/abseil/abseil-cpp/archive/e01d95528ea2137a4a27a88d1f57c6cb260aafed.tar.gz) + set(abseil_cpp_HASH SHA256=84043ed402d2a2a6ba4cdddb7e85118b1158fd81fe4ac3a14adc343d054c1e2e) set(abseil_cpp_BUILD ${CMAKE_BINARY_DIR}/abseil_cpp/src/abseil_cpp_build) if(WIN32) if(${CMAKE_GENERATOR} MATCHES "Visual Studio.*") set(abseil_cpp_STATIC_LIBRARIES ${abseil_cpp_BUILD}/absl/base/Release/absl_base.lib - ${abseil_cpp_BUILD}/absl/base/Release/absl_internal_spinlock_wait.lib ${abseil_cpp_BUILD}/absl/base/Release/absl_dynamic_annotations.lib ${abseil_cpp_BUILD}/absl/base/Release/absl_internal_malloc_internal.lib - ${abseil_cpp_BUILD}/absl/base/Release/absl_internal_throw_delegate.lib - ${abseil_cpp_BUILD}/absl/numeric/Release/absl_int128.lib ${abseil_cpp_BUILD}/absl/strings/Release/absl_strings.lib ${abseil_cpp_BUILD}/absl/strings/Release/str_format_internal.lib ${abseil_cpp_BUILD}/absl/types/Release/absl_bad_optional_access.lib) @@ -82,7 +79,8 @@ else (systemlib_ABSEIL_CPP) ExternalProject_Add(abseil_cpp_build PREFIX abseil_cpp - GIT_REPOSITORY ${abseil_cpp_URL} + URL ${abseil_cpp_URL} + URL_HASH ${abseil_cpp_HASH} DOWNLOAD_DIR "${DOWNLOAD_LOCATION}" BUILD_IN_SOURCE 1 BUILD_BYPRODUCTS ${abseil_cpp_STATIC_LIBRARIES} diff --git a/tensorflow/contrib/cmake/tf_c.cmake b/tensorflow/contrib/cmake/tf_c.cmake index 7de56147d7..a04142bd24 100644 --- a/tensorflow/contrib/cmake/tf_c.cmake +++ b/tensorflow/contrib/cmake/tf_c.cmake @@ -13,33 +13,6 @@ # limitations under the License. # ============================================================================== -# 1. Resolve the installed version of Python (for Python.h and python). -# TODO(mrry): Parameterize the build script to enable Python 3 building. -if(NOT PYTHON_INCLUDE_DIR) - set(PYTHON_NOT_FOUND false) - exec_program("${PYTHON_EXECUTABLE}" - ARGS "-c \"import distutils.sysconfig; print(distutils.sysconfig.get_python_inc())\"" - OUTPUT_VARIABLE PYTHON_INCLUDE_DIR - RETURN_VALUE PYTHON_NOT_FOUND) - if(${PYTHON_NOT_FOUND}) - message(FATAL_ERROR - "Cannot get Python include directory. Is distutils installed?") - endif(${PYTHON_NOT_FOUND}) -endif(NOT PYTHON_INCLUDE_DIR) - -# 2. Resolve the installed version of NumPy (for numpy/arrayobject.h). -if(NOT NUMPY_INCLUDE_DIR) - set(NUMPY_NOT_FOUND false) - exec_program("${PYTHON_EXECUTABLE}" - ARGS "-c \"import numpy; print(numpy.get_include())\"" - OUTPUT_VARIABLE NUMPY_INCLUDE_DIR - RETURN_VALUE NUMPY_NOT_FOUND) - if(${NUMPY_NOT_FOUND}) - message(FATAL_ERROR - "Cannot get NumPy include directory: Is NumPy installed?") - endif(${NUMPY_NOT_FOUND}) -endif(NOT NUMPY_INCLUDE_DIR) - ######################################################## # tf_c_framework library ######################################################## -- GitLab From 71ba0ec86e7cde759006f17979f71e863602182c Mon Sep 17 00:00:00 2001 From: Ouwen Huang Date: Thu, 15 Nov 2018 14:28:48 -0500 Subject: [PATCH 097/596] Update weight_decay_optimizers.py --- .../opt/python/training/weight_decay_optimizers.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py b/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py index 1e8351b70f..8b8065c678 100644 --- a/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py +++ b/tensorflow/contrib/opt/python/training/weight_decay_optimizers.py @@ -64,10 +64,10 @@ class DecoupledWeightDecayExtension(object): the decay to the `weight_decay` as well. For example: ```python - decay = tf.train.piecewise_constant(tf.train.get_global_step(), - [10000, 15000], [1e-1, 1e-2, 1e-3]) - lr = 1*decay - wd = 1e-4*decay + schedule = tf.train.piecewise_constant(tf.train.get_global_step(), + [10000, 15000], [1e-0, 1e-1, 1e-2]) + lr = 1e-1 * schedule() + wd = lambda: 1e-4 * schedule() # ... -- GitLab From 489e181be77b83b2b631f48968aaf40897001838 Mon Sep 17 00:00:00 2001 From: Siju Date: Fri, 16 Nov 2018 11:56:50 +0530 Subject: [PATCH 098/596] Update graph_transformations.h --- .../lite/toco/graph_transformations/graph_transformations.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/lite/toco/graph_transformations/graph_transformations.h b/tensorflow/lite/toco/graph_transformations/graph_transformations.h index 73a90c8239..187b584b69 100644 --- a/tensorflow/lite/toco/graph_transformations/graph_transformations.h +++ b/tensorflow/lite/toco/graph_transformations/graph_transformations.h @@ -139,7 +139,7 @@ DECLARE_GRAPH_TRANSFORMATION(MakeInitialDequantizeOperator) DECLARE_GRAPH_TRANSFORMATION(MoveBinaryOperatorBeforeReshape) DECLARE_GRAPH_TRANSFORMATION(PropagateActivationFunctionIntoConstants) DECLARE_GRAPH_TRANSFORMATION(PropagateArrayDataTypes) -DECLARE_GRAPH_TRANSFORMATION(PropagateFakeQuantNumBits); +DECLARE_GRAPH_TRANSFORMATION(PropagateFakeQuantNumBits) DECLARE_GRAPH_TRANSFORMATION(PropagateFixedSizes) DECLARE_GRAPH_TRANSFORMATION(HardcodeMinMax) DECLARE_GRAPH_TRANSFORMATION(Quantize) -- GitLab From 27171a09e5812d3c8d237c69aa5d53250e7f1696 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sun, 18 Nov 2018 01:28:50 +0000 Subject: [PATCH 099/596] Fix deprecated div While running test I noticed the following warning: ``` WARNING:tensorflow:From /usr/local/lib/python2.7/dist-packages/tensorflow/python/ops/nn_ops.py:2744: div (from tensorflow.python.ops.math_ops) is deprecated and will be removed in a future version. Instructions for updating: Deprecated in favor of operator or tf.math.divide. ``` This fix fixes the deprecated warning. Signed-off-by: Yong Tang --- tensorflow/python/ops/nn_ops.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/ops/nn_ops.py b/tensorflow/python/ops/nn_ops.py index 21008fc392..223a37c87e 100644 --- a/tensorflow/python/ops/nn_ops.py +++ b/tensorflow/python/ops/nn_ops.py @@ -2741,7 +2741,7 @@ def dropout(x, keep_prob, noise_shape=None, seed=None, name=None): # pylint: di noise_shape, seed=seed, dtype=x.dtype) # 0. if [keep_prob, 1.0) and 1. if [1.0, 1.0 + keep_prob) binary_tensor = math_ops.floor(random_tensor) - ret = math_ops.div(x, keep_prob) * binary_tensor + ret = math_ops.divide(x, keep_prob) * binary_tensor if not context.executing_eagerly(): ret.set_shape(x.get_shape()) return ret -- GitLab From e320fba1e9349dee60ba1e06e1f6bbc08c2a85c1 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Thu, 22 Nov 2018 20:47:59 +0000 Subject: [PATCH 100/596] Update re2 library to 2018-10-01 This fix updates re2 library to the latest release of 2018-10-01 Signed-off-by: Yong Tang --- tensorflow/workspace.bzl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tensorflow/workspace.bzl b/tensorflow/workspace.bzl index 7ad094c507..065a695453 100755 --- a/tensorflow/workspace.bzl +++ b/tensorflow/workspace.bzl @@ -168,12 +168,12 @@ def tf_workspace(path_prefix = "", tf_repo_name = ""): tf_http_archive( name = "com_googlesource_code_re2", - sha256 = "803c7811146edeef8f91064de37c6f19136ff01a2a8cdb3230e940b2fd9f07fe", - strip_prefix = "re2-2018-07-01", + sha256 = "a31397714a353587413d307337d0b58f8a2e20e2b9d02f2e24e3463fa4eeda81", + strip_prefix = "re2-2018-10-01", system_build_file = clean_dep("//third_party/systemlibs:re2.BUILD"), urls = [ - "https://mirror.bazel.build/github.com/google/re2/archive/2018-07-01.tar.gz", - "https://github.com/google/re2/archive/2018-07-01.tar.gz", + "https://mirror.bazel.build/github.com/google/re2/archive/2018-10-01.tar.gz", + "https://github.com/google/re2/archive/2018-10-01.tar.gz", ], ) -- GitLab From 48809b87793882266f01b7b40bc9e4a6e0f18f57 Mon Sep 17 00:00:00 2001 From: AG Ramesh Date: Fri, 23 Nov 2018 11:44:19 -0800 Subject: [PATCH 101/596] Fixed merge errors and clang format issues --- tensorflow/core/graph/mkl_layout_pass.cc | 6 +- tensorflow/core/graph/mkl_layout_pass_test.cc | 47 +++-- tensorflow/core/kernels/mkl_conv_ops.cc | 180 +++++++++--------- tensorflow/core/kernels/mkl_conv_ops.h | 14 +- tensorflow/core/kernels/mkl_fused_ops_test.cc | 2 - tensorflow/core/ops/nn_ops.cc | 3 +- 6 files changed, 121 insertions(+), 131 deletions(-) diff --git a/tensorflow/core/graph/mkl_layout_pass.cc b/tensorflow/core/graph/mkl_layout_pass.cc index 8de0fc6083..de1a982b9d 100644 --- a/tensorflow/core/graph/mkl_layout_pass.cc +++ b/tensorflow/core/graph/mkl_layout_pass.cc @@ -2289,13 +2289,13 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, // We handle control edges now. for (const Edge* e : pred->in_edges()) { if (e->IsControlEdge()) { - //Don't allow duplicate edge + // Don't allow duplicate edge (*g)->AddControlEdge(e->src(), new_node, false); } } for (const Edge* e : succ->in_edges()) { if (e->IsControlEdge()) { - //Don't allow duplicate edge + // Don't allow duplicate edge (*g)->AddControlEdge(e->src(), new_node, false); } } @@ -2304,7 +2304,7 @@ Status MklLayoutRewritePass::MergePadWithConv2D(std::unique_ptr* g, // First, we will fix outgoing control edges from 'pred' node. for (const Edge* e : pred->out_edges()) { if (e->IsControlEdge()) { - //Don't allow duplicate edge + // Don't allow duplicate edge (*g)->AddControlEdge(new_node, e->dst(), false); } } diff --git a/tensorflow/core/graph/mkl_layout_pass_test.cc b/tensorflow/core/graph/mkl_layout_pass_test.cc index 0c8d7f7dbb..fa059f1194 100644 --- a/tensorflow/core/graph/mkl_layout_pass_test.cc +++ b/tensorflow/core/graph/mkl_layout_pass_test.cc @@ -138,13 +138,8 @@ REGISTER_OP("_MklInput2") .Output("o: uint8") .Output("o1: uint8") .SetIsStateful(); -REGISTER_OP("Output2") - .Input("i: float") - .Input("i1: float") - .SetIsStateful(); -REGISTER_OP("Output") - .Input("i: float") - .SetIsStateful(); +REGISTER_OP("Output2").Input("i: float").Input("i1: float").SetIsStateful(); +REGISTER_OP("Output").Input("i: float").SetIsStateful(); ///////////////////////////////////////////////////////////////////// // Unit tests related to node merge optiimization @@ -163,7 +158,6 @@ TEST_F(MklLayoutPassTest, Basic) { "A->C;A->D;B->C:1;B->D:1"); } - // Test set 1: Conv2D + AddBias // C=Conv2D(A,B); E=BiasAdd(C,D); Z=Zeta(E,Y) @@ -470,7 +464,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_Conv2DWithBias_ConvBpropInput_FilterFwd) { "E:3->G:4;F->G;F:control->DMT/_3:control;G->Z;X->Y:1;X->Z:1"); } -// Test set 3: Pad + Conv2D fusion +// Test set 3: Pad + Conv2D fusion // padding is VALID type // A = input(image), B = input(paddings), C= Pad = input of conv2D, // D=input(filter), E = Conv2D, Z = Zeta @@ -508,10 +502,10 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Positive) { } // Test if input control edges do not duplicate after merge. // If both the merging ops have input control edge from a common op -// then, the merged op will have only one control edge from that +// then, the merged op will have only one control edge from that // common op. // padding is VALID type -// A = input(image), A1 = input, B = input(paddings), +// A = input(image), A1 = input, B = input(paddings), // C= Pad = input of conv2D, // D=input(filter), E = Conv2D, Z = Zeta // C=Pad(A,B); E=Conv2D(C,D); Z=Zeta(E,Y) @@ -550,12 +544,14 @@ TEST_F(MklLayoutPassTest, Input_ControlEdge_PadWithConv2D_Positive) { const Edge* edge_1 = graph_.AddControlEdge(a1, e); ASSERT_NE(edge, nullptr); ASSERT_NE(edge_1, nullptr); - EXPECT_EQ(DoMklLayoutOptimizationPass(), - "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" - "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" - "A1:control->E:control;A:control->DMT/_0:control;A:control->DMT/_1:control;" - "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" - "DMT/_2->E:5;E->Z;Y->Z:1"); + EXPECT_EQ( + DoMklLayoutOptimizationPass(), + "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" + "A1:control->E:control;A:control->DMT/_0:control;A:control->DMT/" + "_1:control;" + "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;Y->Z:1"); } // Test if output control edges does not duplicate after merge. // If both the merging ops have output control edge to a common op, @@ -600,16 +596,17 @@ TEST_F(MklLayoutPassTest, Output_ControlEdge_PadWithConv2D_Positive) { const Edge* edge_1 = graph_.AddControlEdge(e, a1); ASSERT_NE(edge, nullptr); ASSERT_NE(edge_1, nullptr); - EXPECT_EQ(DoMklLayoutOptimizationPass(), - "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" - "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" - "A:control->DMT/_0:control;A:control->DMT/_1:control;" - "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" - "DMT/_2->E:5;E->Z;E:control->A1:control;Y->Z:1"); + EXPECT_EQ( + DoMklLayoutOptimizationPass(), + "A(Input);A1(Input);B(Int32Input);D(Input);DMT/_0(Const);DMT/_1(Const);" + "DMT/_2(Const);E(_MklPadWithConv2D);Y(Input);Z(Zeta)|A->E;" + "A:control->DMT/_0:control;A:control->DMT/_1:control;" + "A:control->DMT/_2:control;B->E:2;D->E:1;DMT/_0->E:3;DMT/_1->E:4;" + "DMT/_2->E:5;E->Z;E:control->A1:control;Y->Z:1"); } // Pad + Conv2D fusion with padding is VALID, // Input node pointing to both Pad and Conv2D -// A = input(image), B = input(paddings), C= Pad +// A = input(image), B = input(paddings), C= Pad // E = Conv2D, Z = Zeta // C=Pad(A,B); E=Conv2D(C,A); Z=Zeta(E,Y) // After layout pass @@ -645,7 +642,7 @@ TEST_F(MklLayoutPassTest, NodeMerge_PadWithConv2D_Common_Input) { // Pad + Conv2D with padding is VALID, // Input node pointing to both Pad and Conv2D // Output of both Pad and Conv2D feeds one node (Z as Output2) -// A = input(as image), B = input(as paddings), C= Pad +// A = input(as image), B = input(as paddings), C= Pad // E = Conv2D, Z = Output2 // C=Pad(A,B); E=Conv2D(C,A); Z=Output(C,E) // After layout pass - No merging, since Pad and Conv2D both diff --git a/tensorflow/core/kernels/mkl_conv_ops.cc b/tensorflow/core/kernels/mkl_conv_ops.cc index cfc36d1495..9193d00592 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.cc +++ b/tensorflow/core/kernels/mkl_conv_ops.cc @@ -465,19 +465,18 @@ class MklConvOp : public OpKernel { filter.shape().DebugString())); for (int i = 0; i < 3; i++) { - OP_REQUIRES( - context, - FastBoundsCheck(filter.dim_size(i), std::numeric_limits::max()), - errors::InvalidArgument("filter too large")); + OP_REQUIRES(context, FastBoundsCheck(filter.dim_size(i), + std::numeric_limits::max()), + errors::InvalidArgument("filter too large")); } const int64 input_depth = input_in_mkl_format ? GetMklTensorDim(mkl_context.input_shape, 'C') : GetTensorDim(input, data_format_, 'C'); - OP_REQUIRES(context, input_depth == filter.dim_size(2), - errors::InvalidArgument( - "input and filter must have the same depth: ", input_depth, - " vs ", filter.dim_size(2))); + OP_REQUIRES( + context, input_depth == filter.dim_size(2), + errors::InvalidArgument("input and filter must have the same depth: ", + input_depth, " vs ", filter.dim_size(2))); // The last dimension for filter is out_depth. const int out_depth = static_cast(filter.dim_size(3)); @@ -486,10 +485,9 @@ class MklConvOp : public OpKernel { const int64 input_rows_raw = input_in_mkl_format ? GetMklTensorDim(mkl_context.input_shape, 'H') : GetTensorDim(input, data_format_, 'H'); - OP_REQUIRES( - context, - FastBoundsCheck(input_rows_raw, std::numeric_limits::max()), - errors::InvalidArgument("Input rows too large")); + OP_REQUIRES(context, FastBoundsCheck(input_rows_raw, + std::numeric_limits::max()), + errors::InvalidArgument("Input rows too large")); const int input_rows = static_cast(input_rows_raw); const int filter_rows = static_cast(filter.dim_size(0)); @@ -498,10 +496,9 @@ class MklConvOp : public OpKernel { const int64 input_cols_raw = input_in_mkl_format ? GetMklTensorDim(mkl_context.input_shape, 'W') : GetTensorDim(input, data_format_, 'W'); - OP_REQUIRES( - context, - FastBoundsCheck(input_cols_raw, std::numeric_limits::max()), - errors::InvalidArgument("Input cols too large")); + OP_REQUIRES(context, FastBoundsCheck(input_cols_raw, + std::numeric_limits::max()), + errors::InvalidArgument("Input cols too large")); const int input_cols = static_cast(input_cols_raw); const int filter_cols = static_cast(filter.dim_size(1)); @@ -509,10 +506,9 @@ class MklConvOp : public OpKernel { const int64 input_batch_raw = input_in_mkl_format ? GetMklTensorDim(mkl_context.input_shape, 'N') : GetTensorDim(input, data_format_, 'N'); - OP_REQUIRES( - context, - FastBoundsCheck(input_batch_raw, std::numeric_limits::max()), - errors::InvalidArgument("batch is too large")); + OP_REQUIRES(context, FastBoundsCheck(input_batch_raw, + std::numeric_limits::max()), + errors::InvalidArgument("batch is too large")); const int batch = static_cast(input_batch_raw); // For now we take the stride from the second and third dimensions only (we @@ -850,8 +846,8 @@ REGISTER_KERNEL_BUILDER(Name("_MklConv2DWithBias") // Base class for convolution forward operations template + typename Toutput, typename Ttemp_output, typename Tpadding, + bool biasEnabled, bool padEnabled> class MklConvOp : public OpKernel { public: ~MklConvOp() {} @@ -894,17 +890,15 @@ class MklConvOp : public OpKernel { OP_REQUIRES(context, dilations_.size() == 5, errors::InvalidArgument("Dilation rates field must " "specify 5 dimensions")); - OP_REQUIRES(context, - (GetTensorDim(dilations_, data_format_, 'N') == 1 && - GetTensorDim(dilations_, data_format_, 'C') == 1), + OP_REQUIRES(context, (GetTensorDim(dilations_, data_format_, 'N') == 1 && + GetTensorDim(dilations_, data_format_, 'C') == 1), errors::InvalidArgument( "Current implementation does not yet support " "dilations rates in the batch and depth dimensions.")); OP_REQUIRES( - context, - (GetTensorDim(dilations_, data_format_, '0') > 0 && - GetTensorDim(dilations_, data_format_, '1') > 0 && - GetTensorDim(dilations_, data_format_, '2') > 0), + context, (GetTensorDim(dilations_, data_format_, '0') > 0 && + GetTensorDim(dilations_, data_format_, '1') > 0 && + GetTensorDim(dilations_, data_format_, '2') > 0), errors::InvalidArgument("Dilated rates should be larger than 0.")); } } @@ -940,9 +934,9 @@ class MklConvOp : public OpKernel { auto src_tf_shape = GetTfShape(context, kInputIndex_Src); auto filter_tf_shape = GetTfShape(context, kInputIndex_Filter); conv_utl.GetConvFwdSizesInMklOrder( - src_tf_shape, filter_tf_shape, &src_dims, &filter_dims, - &strides, &dilations, &dst_dims_tf_order, &dst_dims_mkl_order, - &padding_left, &padding_right, padEnabled); + src_tf_shape, filter_tf_shape, &src_dims, &filter_dims, &strides, + &dilations, &dst_dims_tf_order, &dst_dims_mkl_order, &padding_left, + &padding_right, padEnabled); if (!context->status().ok()) return; // Check for corner case - if there is nothing to compute, return. @@ -974,9 +968,10 @@ class MklConvOp : public OpKernel { bool isConv2D = (strides_.size() == 4); // TODO(Intel-tf) Add check to make sure padEnabled is true only for 2D - if(!isConv2D){ - OP_REQUIRES(context, padEnabled, - errors::InvalidArgument("Pad+Conv fusion only works for 2D")); + if (!isConv2D) { + OP_REQUIRES( + context, !padEnabled, + errors::InvalidArgument("Pad+Conv fusion only works for 2D")); } // Create memory for user data. // Describe how the inputs and outputs of Convolution look like. Also @@ -1211,7 +1206,6 @@ class MklConvOp : public OpKernel { const int kInputIndex_Pad = 2; const int kOutputIndex_Dst = 0, kOutputIndex_Filter = 1; const int kDilationH = 0, kDilationW = 1; - // Allocate filter output tensor. void AllocateFilterOutputTensor( @@ -1282,7 +1276,7 @@ template class MklQuantizedConv2DOp : public MklConvOp { + int32, biasEnabled, false> { public: virtual ~MklQuantizedConv2DOp() { if (this->input_bias_ != nullptr) { @@ -1297,13 +1291,13 @@ class MklQuantizedConv2DOp } explicit MklQuantizedConv2DOp(OpKernelConstruction* context) - : MklConvOp(context) {} + : MklConvOp(context) {} void Compute(OpKernelContext* context) override { // Compute int32 output tensor - MklConvOp::Compute(context); + MklConvOp::Compute(context); // Compute additional outputs: min/max scalars. int bias_index_offset; @@ -1349,8 +1343,8 @@ class MklQuantizedConv2DOp protected: void ExtendConvFwdParams(OpKernelContext* context, MklConvFwdParams& params) override { - MklConvOp::ExtendConvFwdParams(context, params); + MklConvOp::ExtendConvFwdParams(context, params); // When the output type is quint8, the output data id requantized // into quint8. A post_op "output_scale" is added to do the conversion. @@ -1561,11 +1555,11 @@ class MklQuantizedConv2DSumReluOp } } // TODO(mdfaijul): Add cleaner code for non-mkl tensor - MklConvOp::AllocateOutputTensor(context, conv_prim_desc, - output_dims_mkl_order, - output_tf_format, - output_tensor); + MklConvOp::AllocateOutputTensor(context, conv_prim_desc, + output_dims_mkl_order, + output_tf_format, + output_tensor); const Tensor& summand = MklGetInput(context, summand_idx); if (summand.dtype() != DT_FLOAT) TF_CHECK_OK(Status(error::Code::FAILED_PRECONDITION, @@ -1583,8 +1577,8 @@ class MklQuantizedConv2DSumReluOp const float max_filter = context->input(5 + bias_index_offset).flat()(0); - reorder_sum_scale = 255.0 * 127.0 / - (std::max(std::abs(max_input), std::abs(min_input)) * + reorder_sum_scale = + 255.0 * 127.0 / (std::max(std::abs(max_input), std::abs(min_input)) * std::max(std::abs(max_filter), std::abs(min_filter))); std::vector scales; scales.push_back(reorder_sum_scale); @@ -1833,52 +1827,56 @@ REGISTER_KERNEL_BUILDER( MklQuantizedConv2DSumReluOp); #endif // INTEL_MKL_ML - // Register 2D operations -#define REGISTER_MKL_CPU_2D(T) \ - REGISTER_KERNEL_BUILDER(Name("_MklConv2D") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .Label(mkl_op_registry::kMklOpLabel), \ - MklConvOp); \ - REGISTER_KERNEL_BUILDER(Name("_MklConv2DWithBias") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .Label(mkl_op_registry::kMklOpLabel), \ - MklConvOp); \ - REGISTER_KERNEL_BUILDER(Name("__MklDummyConv2DWithBias") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .Label(mkl_op_registry::kMklOpLabel), \ - MklDummyOp); \ - REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .TypeConstraint("Tpaddings") \ - .Label(mkl_op_registry::kMklOpLabel), \ - MklConvOp); \ - REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .TypeConstraint("Tpaddings") \ - .Label(mkl_op_registry::kMklOpLabel), \ - MklConvOp); \ - REGISTER_KERNEL_BUILDER(Name("__MklDummyPadWithConv2D") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .TypeConstraint("Tpaddings") \ - .Label(mkl_op_registry::kMklOpLabel), \ +#define REGISTER_MKL_CPU_2D(T) \ + REGISTER_KERNEL_BUILDER(Name("_MklConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConvOp); \ + REGISTER_KERNEL_BUILDER(Name("_MklConv2DWithBias") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConvOp); \ + REGISTER_KERNEL_BUILDER(Name("__MklDummyConv2DWithBias") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklDummyOp); \ + REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .TypeConstraint("Tpaddings") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConvOp); \ + REGISTER_KERNEL_BUILDER(Name("_MklPadWithConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .TypeConstraint("Tpaddings") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConvOp); \ + REGISTER_KERNEL_BUILDER(Name("__MklDummyPadWithConv2D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .TypeConstraint("Tpaddings") \ + .Label(mkl_op_registry::kMklOpLabel), \ MklDummyOp); -TF_CALL_float(REGISTER_MKL_CPU); +TF_CALL_float(REGISTER_MKL_CPU_2D); // Register 3D operations -#define REGISTER_MKL_CPU_3D(T) \ - REGISTER_KERNEL_BUILDER(Name("_MklConv3D") \ - .Device(DEVICE_CPU) \ - .TypeConstraint("T") \ - .Label(mkl_op_registry::kMklOpLabel), \ - MklConvOp); +#define REGISTER_MKL_CPU_3D(T) \ + REGISTER_KERNEL_BUILDER( \ + Name("_MklConv3D") \ + .Device(DEVICE_CPU) \ + .TypeConstraint("T") \ + .Label(mkl_op_registry::kMklOpLabel), \ + MklConvOp); TF_CALL_float(REGISTER_MKL_CPU_3D); } // namespace tensorflow diff --git a/tensorflow/core/kernels/mkl_conv_ops.h b/tensorflow/core/kernels/mkl_conv_ops.h index 8c71c20879..963826a73a 100644 --- a/tensorflow/core/kernels/mkl_conv_ops.h +++ b/tensorflow/core/kernels/mkl_conv_ops.h @@ -17,8 +17,8 @@ limitations under the License. #define TENSORFLOW_CORE_KERNELS_MKL_CONV_OPS_H_ #include -#include #include +#include #include "mkldnn.hpp" #include "tensorflow/core/framework/numeric_op.h" @@ -85,7 +85,7 @@ class MklDnnConvUtil { } // Calculate Convolution dilations - virtual inline void GetDilationsInMklOrder(memory::dims *dilations) { + virtual inline void GetDilationsInMklOrder(memory::dims* dilations) { // For now we take the dilation from the second and third dimensions only // (we do not support dilation on the batch or depth dimension). CHECK_NOTNULL(dilations); @@ -195,9 +195,8 @@ class MklDnnConvUtil { filter_shape.DebugString())); for (int i = 0; i < ((strides_.size() == 4) ? 3 : 5); i++) { - OP_REQUIRES(context_, - FastBoundsCheck(filter_shape.dim_size(i), - std::numeric_limits::max()), + OP_REQUIRES(context_, FastBoundsCheck(filter_shape.dim_size(i), + std::numeric_limits::max()), errors::InvalidArgument("filter too large")); } @@ -463,8 +462,8 @@ class MklDnnConvUtil { input_tf_shape.DebugString())); } - GetOutputAndPadSizeInMklOrder(input_tf_shape, filter_tf_shape, - strides, dilations, output_dims_tf_order, + GetOutputAndPadSizeInMklOrder(input_tf_shape, filter_tf_shape, strides, + dilations, output_dims_tf_order, output_dims_mkl_order, pad_l, pad_r); } @@ -556,7 +555,6 @@ class MklConvBackpropCommonOp : public OpKernel { TensorFormat data_format_; // NCHW or NHWC }; - ///////////////////////////////////////////////////////////////////// /// Dummy Mkl op that is just used for operators that are intermediate /// output of node fusion in the graph diff --git a/tensorflow/core/kernels/mkl_fused_ops_test.cc b/tensorflow/core/kernels/mkl_fused_ops_test.cc index 900325ac91..991fb08093 100644 --- a/tensorflow/core/kernels/mkl_fused_ops_test.cc +++ b/tensorflow/core/kernels/mkl_fused_ops_test.cc @@ -13,7 +13,6 @@ See the License for the specific language governing permissions and limitations under the License. ==============================================================================*/ #ifdef INTEL_MKL -#ifndef INTEL_MKL_ML_ONLY // We don't support fusion in MKL ML #include "tensorflow/cc/ops/const_op.h" #include "tensorflow/cc/ops/image_ops.h" #include "tensorflow/cc/ops/nn_ops.h" @@ -160,5 +159,4 @@ TEST_F(FusedPadConvOpTest, PaddingConvTestNchw) { Run(DT_FLOAT, image, filter, padding, expected, "NCHW"); } } // namespace tensorflow -#endif // INTEL_MKL_ML_ONLY #endif // INTEL_MKL diff --git a/tensorflow/core/ops/nn_ops.cc b/tensorflow/core/ops/nn_ops.cc index 8afbe0333a..0b99542c5c 100644 --- a/tensorflow/core/ops/nn_ops.cc +++ b/tensorflow/core/ops/nn_ops.cc @@ -1649,7 +1649,7 @@ REGISTER_OP("_MklPadWithConv2D") .Attr(GetConvnetDataFormatAttrString()) .Attr("dilations: list(int) = [1, 1, 1, 1]") .Attr("Tpaddings: {int32, int64} = DT_INT32") - .SetShapeFn(shape_inference::Conv2DShape) + .SetShapeFn(shape_inference::Conv2DShape) .Doc(R"doc( MKL version of Pad and Conv2D operator. Uses MKL DNN APIs to perform Pad and 2D convolution to the output of convolution. @@ -2159,7 +2159,6 @@ NOTE Do not invoke this operator directly in Python. Graph rewrite pass is expected to invoke these operators. )doc"); - REGISTER_OP("_MklAvgPool3DGrad") .Input("orig_input_shape: int32") .Input("grad: T") -- GitLab From 768b36822e0d5b988a697e0c9e3b65302b051630 Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Mon, 26 Nov 2018 15:45:08 -0800 Subject: [PATCH 102/596] fixing the buildifer error of tensorflow/core/kernels/BUILD --- tensorflow/core/kernels/BUILD | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/tensorflow/core/kernels/BUILD b/tensorflow/core/kernels/BUILD index 1759a7f790..efa571b23b 100644 --- a/tensorflow/core/kernels/BUILD +++ b/tensorflow/core/kernels/BUILD @@ -29,26 +29,26 @@ package_group( load( "//tensorflow:tensorflow.bzl", + "cc_header_only_library", "if_android", + "if_not_windows", + "tf_cc_binary", "tf_cc_test", "tf_cc_test_mkl", "tf_cc_tests", - "tf_cc_binary", "tf_copts", "tf_cuda_library", - "tf_opts_nortti_if_android", "tf_kernel_library", "tf_mkl_kernel_library", - "cc_header_only_library", - "if_not_windows", + "tf_opts_nortti_if_android", ) load("@local_config_sycl//sycl:build_defs.bzl", "if_sycl") load("//tensorflow:tensorflow.bzl", "tf_cuda_cc_test") load("//tensorflow:tensorflow.bzl", "tf_cuda_cc_tests") load( "//tensorflow/core:platform/default/build_config.bzl", - "tf_proto_library", "tf_kernel_tests_linkstatic", + "tf_proto_library", ) load( "//tensorflow/core:platform/default/build_config_root.bzl", @@ -6714,10 +6714,10 @@ tf_cc_test_mkl( srcs = ["mkl_fused_ops_test.cc"], linkstatic = 1, deps = [ - ":mkl_conv_op", - ":mkl_tfconv_op", ":conv_ops", ":image", + ":mkl_conv_op", + ":mkl_tfconv_op", ":ops_testutil", ":ops_util", "//tensorflow/cc:cc_ops", @@ -6730,8 +6730,9 @@ tf_cc_test_mkl( "//tensorflow/core:test", "//tensorflow/core:test_main", "//tensorflow/core:testlib", - ] + ], ) + tf_mkl_kernel_library( name = "mkl_transpose_op", srcs = [ -- GitLab From 453335db8382418dc84593d044f63a995adc025a Mon Sep 17 00:00:00 2001 From: mbhuiyan Date: Mon, 26 Nov 2018 16:57:34 -0800 Subject: [PATCH 103/596] applying proper coding style --- tensorflow/core/kernels/mkl_lrn_op.cc | 42 ++++++++++++--------------- 1 file changed, 19 insertions(+), 23 deletions(-) diff --git a/tensorflow/core/kernels/mkl_lrn_op.cc b/tensorflow/core/kernels/mkl_lrn_op.cc index 4d46abb0a4..407ce5d653 100644 --- a/tensorflow/core/kernels/mkl_lrn_op.cc +++ b/tensorflow/core/kernels/mkl_lrn_op.cc @@ -23,7 +23,6 @@ limitations under the License. #define EIGEN_USE_THREADS #include #include "mkldnn.hpp" -#include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" #include "tensorflow/core/framework/op_kernel.h" #include "tensorflow/core/framework/register_types.h" #include "tensorflow/core/framework/tensor.h" @@ -32,6 +31,7 @@ limitations under the License. #include "tensorflow/core/lib/core/errors.h" #include "tensorflow/core/util/mkl_util.h" #include "tensorflow/core/util/tensor_format.h" +#include "third_party/eigen3/unsupported/Eigen/CXX11/Tensor" #if !defined(IS_MOBILE_PLATFORM) #include "tensorflow/core/util/work_sharder.h" @@ -71,11 +71,10 @@ class MklLRNOp : public OpKernel { explicit MklLRNOp(OpKernelConstruction* context) : OpKernel(context) { int64 depth_radius64; OP_REQUIRES_OK(context, context->GetAttr("depth_radius", &depth_radius64)); - OP_REQUIRES( - context, - FastBoundsCheck(depth_radius64, std::numeric_limits::max()), - errors::InvalidArgument("depth_radius = ", depth_radius64, - " larger than int max")); + OP_REQUIRES(context, FastBoundsCheck(depth_radius64, + std::numeric_limits::max()), + errors::InvalidArgument("depth_radius = ", depth_radius64, + " larger than int max")); depth_radius_ = static_cast(depth_radius64); OP_REQUIRES_OK(context, context->GetAttr("bias", &bias_)); @@ -161,9 +160,9 @@ class MklLRNOp : public OpKernel { PrepareAndExecuteNet(lrn_prim_desc, &src_dnn_data, &dst_dnn_data, &workspace_dnn_data); } catch (mkldnn::error& e) { - string error_msg = "Status: " + std::to_string(e.status) + - ", message: " + string(e.message) + ", in file " + - string(__FILE__) + ":" + std::to_string(__LINE__); + string error_msg = "Status: " + std::to_string(e.status) + ", message: " + + string(e.message) + ", in file " + string(__FILE__) + + ":" + std::to_string(__LINE__); OP_REQUIRES_OK( context, errors::Aborted("Operation received an exception:", error_msg)); @@ -292,16 +291,14 @@ class MklLRNOp : public OpKernel { if (src_dnn_shape.IsMklTensor()) { OP_REQUIRES(context, src_dnn_shape.GetDimension() == 4, errors::InvalidArgument("input must be 4-dimensional")); - OP_REQUIRES(context, - FastBoundsCheck(src_tensor.NumElements(), - std::numeric_limits::max()), + OP_REQUIRES(context, FastBoundsCheck(src_tensor.NumElements(), + std::numeric_limits::max()), errors::InvalidArgument("argument to LRN too large")); } else { OP_REQUIRES(context, src_tensor.dims() == 4, errors::InvalidArgument("input must be 4-dimensional")); - OP_REQUIRES(context, - FastBoundsCheck(src_tensor.NumElements(), - std::numeric_limits::max()), + OP_REQUIRES(context, FastBoundsCheck(src_tensor.NumElements(), + std::numeric_limits::max()), errors::InvalidArgument("argument to LRN too large")); } } @@ -321,11 +318,10 @@ class MklLRNGradOp : public OpKernel { explicit MklLRNGradOp(OpKernelConstruction* context) : OpKernel(context) { int64 depth_radius64; OP_REQUIRES_OK(context, context->GetAttr("depth_radius", &depth_radius64)); - OP_REQUIRES( - context, - FastBoundsCheck(depth_radius64, std::numeric_limits::max()), - errors::InvalidArgument("depth_radius = ", depth_radius64, - " larger than int max")); + OP_REQUIRES(context, FastBoundsCheck(depth_radius64, + std::numeric_limits::max()), + errors::InvalidArgument("depth_radius = ", depth_radius64, + " larger than int max")); depth_radius_ = static_cast(depth_radius64); OP_REQUIRES_OK(context, context->GetAttr("bias", &bias_)); OP_REQUIRES_OK(context, context->GetAttr("alpha", &alpha_)); @@ -432,9 +428,9 @@ class MklLRNGradOp : public OpKernel { memory::primitive_desc(target_diff_dst_md, cpu_engine), &workspace_dnn_data); } catch (mkldnn::error& e) { - string error_msg = "Status: " + std::to_string(e.status) + - ", message: " + string(e.message) + ", in file " + - string(__FILE__) + ":" + std::to_string(__LINE__); + string error_msg = "Status: " + std::to_string(e.status) + ", message: " + + string(e.message) + ", in file " + string(__FILE__) + + ":" + std::to_string(__LINE__); OP_REQUIRES_OK( context, errors::Aborted("Operation received an exception:", error_msg)); -- GitLab From b607192183e06ef18f79e7c6dd8812c043fc1ab3 Mon Sep 17 00:00:00 2001 From: Serge Panev Date: Tue, 27 Nov 2018 10:57:57 +0100 Subject: [PATCH 104/596] Fix comparison between signed and unsigned integer expressions in tensor_format.h --- tensorflow/core/framework/dataset.h | 2 +- tensorflow/core/util/tensor_format.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/framework/dataset.h b/tensorflow/core/framework/dataset.h index 9b11449b30..e0041492ee 100644 --- a/tensorflow/core/framework/dataset.h +++ b/tensorflow/core/framework/dataset.h @@ -162,7 +162,7 @@ class GraphDefBuilderWrapper { const std::vector>& attrs, Node** output) { std::vector> enumerated_inputs(inputs.size()); - for (int i = 0; i < inputs.size(); i++) { + for (int i = 0; static_cast(i) < inputs.size(); i++) { enumerated_inputs[i] = std::make_pair(i, inputs[i]); } return AddDataset(dataset, enumerated_inputs, {}, attrs, output); diff --git a/tensorflow/core/util/tensor_format.h b/tensorflow/core/util/tensor_format.h index b0c349dd90..68674cb282 100644 --- a/tensorflow/core/util/tensor_format.h +++ b/tensorflow/core/util/tensor_format.h @@ -498,7 +498,7 @@ inline TensorShape ShapeFromFormat(TensorFormat format, int64 N, dim_sizes[GetTensorBatchDimIndex(dims, format)] = N; for (int dim = 0; static_cast(dim) < spatial.size(); dim++) { auto dim_size = spatial[dim]; - if (format == FORMAT_NHWC_VECT_W && dim == spatial.size() - 1) { + if (format == FORMAT_NHWC_VECT_W && static_cast(dim) == spatial.size() - 1) { CHECK_EQ(0, dim_size % 4) << "FORMAT_NHWC_VECT_W requires W to be a multiple of 4, but W=" << dim_size; -- GitLab From 6e436c4f058175c62dca5aa37b9b95c15c251ea9 Mon Sep 17 00:00:00 2001 From: manhyuk Date: Tue, 27 Nov 2018 19:08:47 +0900 Subject: [PATCH 105/596] fix typo --- tensorflow/compiler/xla/service/algebraic_simplifier_test.cc | 2 +- tensorflow/compiler/xla/shape_util.cc | 2 +- tensorflow/python/ops/ragged/ragged_tensor.py | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/tensorflow/compiler/xla/service/algebraic_simplifier_test.cc b/tensorflow/compiler/xla/service/algebraic_simplifier_test.cc index 24c35464ad..054d518671 100644 --- a/tensorflow/compiler/xla/service/algebraic_simplifier_test.cc +++ b/tensorflow/compiler/xla/service/algebraic_simplifier_test.cc @@ -1617,7 +1617,7 @@ TEST_F(AlgebraicSimplifierTest, ReshapeOfTransposeOfRngToRng) { (AlgebraicSimplifierOptions(bitcasting_callback()))); EXPECT_TRUE(simplifier.Run(m.get()).ValueOrDie()); - // Verify that that reshape(transpose(rng)) is replace by a single rng of the + // Verify that reshape(transpose(rng)) is replace by a single rng of the // same shape as the reshape. EXPECT_THAT(computation->root_instruction(), op::Rng()); EXPECT_TRUE(ShapeUtil::Equal(computation->root_instruction()->shape(), diff --git a/tensorflow/compiler/xla/shape_util.cc b/tensorflow/compiler/xla/shape_util.cc index b05ec209cc..c185b19687 100644 --- a/tensorflow/compiler/xla/shape_util.cc +++ b/tensorflow/compiler/xla/shape_util.cc @@ -1165,7 +1165,7 @@ Status ForEachMutableSubshapeHelper( // Let the argument `permutation` be P. This is a permutation over `shape`'s // dimensions, so our return value will be a shape with dims P.I = P. Our // goal is to construct a layout permutation L* that we can apply to P such - // that that the physical dimension ordering of the returned shape is the same + // that the physical dimension ordering of the returned shape is the same // as that of the original shape, namely L'. // // Our returned shape has dims P and layout L*, so its in-memory layout is diff --git a/tensorflow/python/ops/ragged/ragged_tensor.py b/tensorflow/python/ops/ragged/ragged_tensor.py index abb27fc3c0..1d4bbd592d 100644 --- a/tensorflow/python/ops/ragged/ragged_tensor.py +++ b/tensorflow/python/ops/ragged/ragged_tensor.py @@ -64,7 +64,7 @@ class RaggedTensor(object): a 3-D `RaggedTensor` that stores the fixed-size word embedding for each word in a sentence, for each sentence in a batch, could be written as `[num_sentences, (num_words), embedding_size]`. The parentheses around - `(num_words)` indicate that that dimension is ragged, and that the length + `(num_words)` indicate that dimension is ragged, and that the length of each element list in that dimension may vary for each item. ### Component Tensors -- GitLab From 4b3a92f1bf905a58daa878999a9e5891d9b9e931 Mon Sep 17 00:00:00 2001 From: Serge Panev Date: Tue, 27 Nov 2018 22:18:58 +0100 Subject: [PATCH 106/596] Simpler loop --- tensorflow/core/framework/dataset.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/core/framework/dataset.h b/tensorflow/core/framework/dataset.h index e0041492ee..927cf22c68 100644 --- a/tensorflow/core/framework/dataset.h +++ b/tensorflow/core/framework/dataset.h @@ -162,7 +162,7 @@ class GraphDefBuilderWrapper { const std::vector>& attrs, Node** output) { std::vector> enumerated_inputs(inputs.size()); - for (int i = 0; static_cast(i) < inputs.size(); i++) { + for (size_t i = 0; i < inputs.size(); i++) { enumerated_inputs[i] = std::make_pair(i, inputs[i]); } return AddDataset(dataset, enumerated_inputs, {}, attrs, output); -- GitLab From 59ebe545b3385c4c36d3b1602671d109e44ea38c Mon Sep 17 00:00:00 2001 From: Fei Hu Date: Fri, 16 Nov 2018 16:20:05 -0800 Subject: [PATCH 107/596] Add CancellationManager to LoopCondOp --- tensorflow/core/kernels/control_flow_ops.cc | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/tensorflow/core/kernels/control_flow_ops.cc b/tensorflow/core/kernels/control_flow_ops.cc index 1587eb5114..21aabd9295 100644 --- a/tensorflow/core/kernels/control_flow_ops.cc +++ b/tensorflow/core/kernels/control_flow_ops.cc @@ -600,6 +600,15 @@ LoopCondOp::LoopCondOp(OpKernelConstruction* context) : OpKernel(context) {} LoopCondOp::~LoopCondOp() = default; void LoopCondOp::Compute(OpKernelContext* context) { + CancellationManager* cm = context->cancellation_manager(); + bool already_cancelled = cm->IsCancelled(); + + if (already_cancelled) { + Tensor continue_running(false); + context->set_output(0, continue_running); + return; + } + context->set_output(0, context->input(0)); } -- GitLab From c0b128c45396560f26769a293525c60f76850a3f Mon Sep 17 00:00:00 2001 From: Fei Hu Date: Fri, 16 Nov 2018 16:20:52 -0800 Subject: [PATCH 108/596] Add a python test for while_loop timeout --- .../python/kernel_tests/control_flow_ops_py_test.py | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/tensorflow/python/kernel_tests/control_flow_ops_py_test.py b/tensorflow/python/kernel_tests/control_flow_ops_py_test.py index 37654abd18..595d4ff37b 100644 --- a/tensorflow/python/kernel_tests/control_flow_ops_py_test.py +++ b/tensorflow/python/kernel_tests/control_flow_ops_py_test.py @@ -1988,6 +1988,16 @@ class ControlFlowTest(test.TestCase): for i in xrange(10): self.assertEqual([i], q.dequeue().eval()) + def testWhileTimeOut(self): + run_options = config_pb2.RunOptions(timeout_in_ms=1) + with self.cached_session() as sess: + n = constant_op.constant(0) + c = lambda x: True + b = lambda x: math_ops.add(x, 1) + r = control_flow_ops.while_loop(c, b, [n]) + with self.assertRaises(errors_impl.DeadlineExceededError): + sess.run(r, options=run_options) + @test_util.disable_control_flow_v2("b/117119329 (stack)") def testWhileStack_1(self): with self.cached_session(): -- GitLab From fc5392dedec126f988788a597edb55021fb07b60 Mon Sep 17 00:00:00 2001 From: Fei Hu Date: Tue, 20 Nov 2018 22:25:51 -0500 Subject: [PATCH 109/596] Raise an error when loop execution is cancelled --- tensorflow/core/kernels/control_flow_ops.cc | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/tensorflow/core/kernels/control_flow_ops.cc b/tensorflow/core/kernels/control_flow_ops.cc index 21aabd9295..61547adb73 100644 --- a/tensorflow/core/kernels/control_flow_ops.cc +++ b/tensorflow/core/kernels/control_flow_ops.cc @@ -602,12 +602,8 @@ LoopCondOp::~LoopCondOp() = default; void LoopCondOp::Compute(OpKernelContext* context) { CancellationManager* cm = context->cancellation_manager(); bool already_cancelled = cm->IsCancelled(); - - if (already_cancelled) { - Tensor continue_running(false); - context->set_output(0, continue_running); - return; - } + OP_REQUIRES(context, !already_cancelled, + errors::Cancelled("Loop execution was cancelled.")); context->set_output(0, context->input(0)); } -- GitLab From 5bb0553a1b3bd580bd1502ab6b339ca3a1f0b5df Mon Sep 17 00:00:00 2001 From: Fei Hu Date: Tue, 27 Nov 2018 22:05:43 -0800 Subject: [PATCH 110/596] Handle the case that CancellationManager is null in eager mode --- tensorflow/core/kernels/control_flow_ops.cc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/tensorflow/core/kernels/control_flow_ops.cc b/tensorflow/core/kernels/control_flow_ops.cc index 61547adb73..b3bdff2575 100644 --- a/tensorflow/core/kernels/control_flow_ops.cc +++ b/tensorflow/core/kernels/control_flow_ops.cc @@ -601,9 +601,11 @@ LoopCondOp::~LoopCondOp() = default; void LoopCondOp::Compute(OpKernelContext* context) { CancellationManager* cm = context->cancellation_manager(); - bool already_cancelled = cm->IsCancelled(); - OP_REQUIRES(context, !already_cancelled, - errors::Cancelled("Loop execution was cancelled.")); + if (cm != nullptr) { + bool already_cancelled = cm->IsCancelled(); + OP_REQUIRES(context, !already_cancelled, + errors::Cancelled("Loop execution was cancelled.")); + } context->set_output(0, context->input(0)); } -- GitLab From f3db4e300b782afbff93223777ed1b5cd3f1f7cd Mon Sep 17 00:00:00 2001 From: Yuxin Wu Date: Wed, 28 Nov 2018 21:56:27 -0800 Subject: [PATCH 111/596] Fix deprecated use of `sparse_to_dense`. Calling `sparse_to_dense` gives a deprecation warning that asks users to use `sparse.to_dense`: https://github.com/tensorflow/tensorflow/blob/71f40f044450736cd6acd29e92ffbfc0e571ee14/tensorflow/python/ops/sparse_ops.py#L952-L955 However, `sparse.to_dense` calls `sparse_to_dense`, which again produces the deprecation warning. --- tensorflow/python/ops/sparse_ops.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/ops/sparse_ops.py b/tensorflow/python/ops/sparse_ops.py index 245080cb26..914616ff6d 100644 --- a/tensorflow/python/ops/sparse_ops.py +++ b/tensorflow/python/ops/sparse_ops.py @@ -1422,7 +1422,7 @@ def sparse_tensor_to_dense(sp_input, """ sp_input = _convert_to_sparse_tensor(sp_input) - return sparse_to_dense( + return gen_sparse_ops.sparse_to_dense( sp_input.indices, sp_input.dense_shape, sp_input.values, -- GitLab From 0c7a31a168aae02e323c97bc6b81f2d3f19cbb2a Mon Sep 17 00:00:00 2001 From: Nutti Date: Thu, 29 Nov 2018 22:32:54 +0900 Subject: [PATCH 112/596] Fix: clang-format error --- tensorflow/core/kernels/partitioned_function_ops.cc | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/tensorflow/core/kernels/partitioned_function_ops.cc b/tensorflow/core/kernels/partitioned_function_ops.cc index 72310f33ae..8aac78f56f 100644 --- a/tensorflow/core/kernels/partitioned_function_ops.cc +++ b/tensorflow/core/kernels/partitioned_function_ops.cc @@ -179,10 +179,9 @@ class PartitionedCallOp : public AsyncOpKernel { done); OP_REQUIRES_OK_ASYNC( - ctx, - OptimizationPassRegistry::Global()->RunGrouping( - OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, - optimization_options), + ctx, OptimizationPassRegistry::Global()->RunGrouping( + OptimizationPassRegistry::POST_REWRITE_FOR_EXEC, + optimization_options), done); std::unordered_map> subgraphs; -- GitLab From d00013e72cd3a1a4805395eb8e66748dcf387295 Mon Sep 17 00:00:00 2001 From: Wen yun Date: Wed, 31 Oct 2018 15:22:14 +0800 Subject: [PATCH 113/596] fix the case when input value are MirroredVariable for assign_moving_average --- .../distribute/python/moving_averages_test.py | 20 +++++++++++++++++++ .../python/distribute/mirrored_strategy.py | 3 +++ 2 files changed, 23 insertions(+) diff --git a/tensorflow/contrib/distribute/python/moving_averages_test.py b/tensorflow/contrib/distribute/python/moving_averages_test.py index c492d8bafc..da3353b2d5 100644 --- a/tensorflow/contrib/distribute/python/moving_averages_test.py +++ b/tensorflow/contrib/distribute/python/moving_averages_test.py @@ -139,6 +139,26 @@ class AssignMovingAveragesTest(test.TestCase, parameterized.TestCase): (2.0 * 0.25 + 0.0) / (1.0 * 0.25 + 1.0)], var.eval()) + @combinations.generate(all_combinations) + def testAssignVariable(self, distribution): + def replica_fn(): + var = variables.Variable([10.0, 11.0]) + # Here we expect to check the case when input value are variable. + val = variables.Variable([1., 2.]) + decay = 0.25 + assign = moving_averages.assign_moving_average( + var, val, decay, zero_debias=False) + return var, assign + + with distribution.scope(), self.cached_session() as sess: + var, assign = distribution.call_for_each_replica(replica_fn) + variables.global_variables_initializer().run() + self.assertAllClose([10.0, 11.0], var.eval()) + sess.run(distribution.unwrap(assign)) + self.assertAllClose( + [10 * 0.25 + 1. * (1 - 0.25), + 11 * 0.25 + 2. * (1 - 0.25)], + var.eval()) if __name__ == "__main__": test.main() diff --git a/tensorflow/python/distribute/mirrored_strategy.py b/tensorflow/python/distribute/mirrored_strategy.py index d6d40df5ce..3cd5cf09c0 100644 --- a/tensorflow/python/distribute/mirrored_strategy.py +++ b/tensorflow/python/distribute/mirrored_strategy.py @@ -598,6 +598,9 @@ class MirroredExtended(distribute_lib.DistributionStrategyExtended): return self._cross_device_ops def _reduce_to(self, reduce_op, value, destinations): + if (isinstance(value, values.Mirrored) and + reduce_op == reduce_util.ReduceOp.MEAN): + return value assert not isinstance(value, values.Mirrored) if not isinstance(value, values.DistributedValues): # This function handles reducing values that are not PerReplica or -- GitLab From 7d931d7b85f65b6145643fdc638aabbce779ab21 Mon Sep 17 00:00:00 2001 From: Neargye Date: Mon, 23 Apr 2018 00:12:31 +0500 Subject: [PATCH 114/596] change toolchain to clang --- tensorflow/contrib/android/cmake/build.gradle | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tensorflow/contrib/android/cmake/build.gradle b/tensorflow/contrib/android/cmake/build.gradle index 17a57b99fd..9e7fd317f3 100644 --- a/tensorflow/contrib/android/cmake/build.gradle +++ b/tensorflow/contrib/android/cmake/build.gradle @@ -22,8 +22,8 @@ android { } externalNativeBuild { cmake { - arguments '-DANDROID_TOOLCHAIN=gcc', - '-DANDROID_STL=gnustl_static' + arguments '-DANDROID_TOOLCHAIN=clang', + '-DANDROID_STL=c++_static' } } } -- GitLab From f56a6058a44edb2dc6172a0723fb92fd63e7a36f Mon Sep 17 00:00:00 2001 From: neargye Date: Fri, 30 Nov 2018 16:57:33 +0500 Subject: [PATCH 115/596] clean-up --- tensorflow/contrib/android/cmake/build.gradle | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/contrib/android/cmake/build.gradle b/tensorflow/contrib/android/cmake/build.gradle index 9e7fd317f3..ddec08894f 100644 --- a/tensorflow/contrib/android/cmake/build.gradle +++ b/tensorflow/contrib/android/cmake/build.gradle @@ -70,7 +70,7 @@ if (ndkDir == null || ndkDir == "") { ndkDir = System.getenv('ANDROID_NDK_HOME') } -if(! Os.isFamily(Os.FAMILY_WINDOWS)) { +if (!Os.isFamily(Os.FAMILY_WINDOWS)) { // This script is for non-Windows OS. For Windows OS, MANUALLY build // (or copy the built) libs/headers to the // ${TENSORFLOW_ROOT_DIR}/tensorflow/contrib/makefile/gen -- GitLab From 29c1506fecd37c853f4c5d53e0a199c04f213436 Mon Sep 17 00:00:00 2001 From: Fei Hu Date: Fri, 30 Nov 2018 16:10:51 -0800 Subject: [PATCH 116/596] Change the variable names to be more readable --- .../data/kernel_tests/list_files_test.py | 33 ++++++++++--------- 1 file changed, 17 insertions(+), 16 deletions(-) diff --git a/tensorflow/python/data/kernel_tests/list_files_test.py b/tensorflow/python/data/kernel_tests/list_files_test.py index 789f1ab6de..d4e0cdeb97 100644 --- a/tensorflow/python/data/kernel_tests/list_files_test.py +++ b/tensorflow/python/data/kernel_tests/list_files_test.py @@ -82,27 +82,27 @@ class ListFilesTest(test_base.DatasetTestBase): dataset = dataset_ops.Dataset.list_files( path.join(self.tmp_dir, '*'), shuffle=True, seed=37) - full_filenames = [compat.as_bytes(path.join(self.tmp_dir, filename)) + expected_filenames = [compat.as_bytes(path.join(self.tmp_dir, filename)) for filename in filenames] - all_produced_filenames = [] + all_actual_filenames = [] for _ in range(3): - produced_filenames = [] + actual_filenames = [] next_element = self.getNext(dataset, requires_initialization=True) try: while True: - produced_filenames.append(self.evaluate(next_element())) + actual_filenames.append(self.evaluate(next_element())) except errors.OutOfRangeError: pass - all_produced_filenames.append(produced_filenames) + all_actual_filenames.append(actual_filenames) # Each run should produce the same set of filenames, which may be - # different from the order of `full_filenames`. - self.assertItemsEqual(full_filenames, all_produced_filenames[0]) + # different from the order of `expected_filenames`. + self.assertItemsEqual(expected_filenames, all_actual_filenames[0]) # However, the different runs should produce filenames in the same order # as each other. - self.assertEqual(all_produced_filenames[0], all_produced_filenames[1]) - self.assertEqual(all_produced_filenames[0], all_produced_filenames[2]) + self.assertEqual(all_actual_filenames[0], all_actual_filenames[1]) + self.assertEqual(all_actual_filenames[0], all_actual_filenames[2]) # TODO(b/117581999): eager mode assertion fail wrapped, debug. def tesSkipEagerEmptyDirectoryInitializer(self): @@ -169,16 +169,17 @@ class ListFilesTest(test_base.DatasetTestBase): path.join(self.tmp_dir, '*'), shuffle=False).repeat(2) next_element = self.getNext(dataset) - full_filenames = [] - produced_filenames = [] + expected_filenames = [] + actual_filenames = [] for filename in filenames * 2: - full_filenames.append(compat.as_bytes(path.join(self.tmp_dir, filename))) - produced_filenames.append(compat.as_bytes(self.evaluate(next_element()))) + expected_filenames.append( + compat.as_bytes(path.join(self.tmp_dir, filename))) + actual_filenames.append(compat.as_bytes(self.evaluate(next_element()))) with self.assertRaises(errors.OutOfRangeError): self.evaluate(next_element()) - self.assertItemsEqual(full_filenames, produced_filenames) - self.assertEqual(produced_filenames[:len(filenames)], - produced_filenames[len(filenames):]) + self.assertItemsEqual(expected_filenames, actual_filenames) + self.assertEqual(actual_filenames[:len(filenames)], + actual_filenames[len(filenames):]) def testMultiplePatternsAsList(self): filenames = ['a.txt', 'b.py', 'c.py', 'd.pyc'] -- GitLab From cb57c4dd8d9cadcb0c25c0940d7cb4f08d5aa530 Mon Sep 17 00:00:00 2001 From: Fei Hu Date: Fri, 30 Nov 2018 16:32:34 -0800 Subject: [PATCH 117/596] Fix the coding style --- tensorflow/python/data/kernel_tests/list_files_test.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/data/kernel_tests/list_files_test.py b/tensorflow/python/data/kernel_tests/list_files_test.py index d4e0cdeb97..093d6cb3ca 100644 --- a/tensorflow/python/data/kernel_tests/list_files_test.py +++ b/tensorflow/python/data/kernel_tests/list_files_test.py @@ -83,7 +83,7 @@ class ListFilesTest(test_base.DatasetTestBase): path.join(self.tmp_dir, '*'), shuffle=True, seed=37) expected_filenames = [compat.as_bytes(path.join(self.tmp_dir, filename)) - for filename in filenames] + for filename in filenames] all_actual_filenames = [] for _ in range(3): -- GitLab From fe9a9dbb8d0ef118b15beb4724f256190fd04d13 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Fri, 30 Nov 2018 16:30:48 -0600 Subject: [PATCH 118/596] Split convolution invocation into preparation and actual invocation - split DoConvolve into: PrepareForConvolution DoConvolve - split DoConvolveBackwardData into: PrepareForConvolutionBackwardData DoConvolveBackwardData - split DoConvolveBackwardFilter into: PrepareForConvolutionBackwardFilter DoConvolveBackwardFilter PrepareForConvolutionXXX would allocate scratch memory. DoConolveXXX would invoke actual convolution algorithms. Implement forward convoution, backward input convolution, backward filter convolution on CUDA path. --- tensorflow/stream_executor/cuda/cuda_dnn.cc | 459 +++++++++++++++----- tensorflow/stream_executor/cuda/cuda_dnn.h | 196 ++++++++- tensorflow/stream_executor/dnn.h | 180 ++++++-- tensorflow/stream_executor/stream.cc | 172 ++++++-- 4 files changed, 827 insertions(+), 180 deletions(-) diff --git a/tensorflow/stream_executor/cuda/cuda_dnn.cc b/tensorflow/stream_executor/cuda/cuda_dnn.cc index 1f2e2f48bb..387afefc21 100644 --- a/tensorflow/stream_executor/cuda/cuda_dnn.cc +++ b/tensorflow/stream_executor/cuda/cuda_dnn.cc @@ -2380,7 +2380,7 @@ bool ShouldIncludeWinogradNonfusedAlgo( } // namespace template -port::Status CudnnSupport::DoConvolveImpl( +port::Status CudnnSupport::PrepareForConvolutionImpl( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::FilterDescriptor& filter_descriptor, @@ -2389,6 +2389,34 @@ port::Status CudnnSupport::DoConvolveImpl( const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, dnn::DataType accumulator_type, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + cudnnDataType_t cudnn_type = GetCudnnDataType(); + CudnnTensorDescriptor input_nd(input_descriptor, cudnn_type); + CudnnTensorDescriptor output_nd(output_descriptor, cudnn_type); + CudnnFilterDescriptor filter(filter_descriptor, cudnn_type); + CudnnConvolutionDescriptor conv(convolution_descriptor, + ToCudnnDataType(accumulator_type)); + + auto cudnn = cudnn_->GetHandle(parent_, stream); + + SE_ASSIGN_OR_RETURN(*algorithm_desc, + GetCudnnConvolutionForwardAlgorithm( + stream, cudnn, algorithm_config, input_nd, filter, + conv, output_nd, scratch_allocator, scratch_memory)); + + return port::Status::OK(); +} + +template +port::Status CudnnSupport::DoConvolveImpl( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, + dnn::DataType accumulator_type, const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) { cudnnDataType_t cudnn_type = GetCudnnDataType(); CudnnTensorDescriptor input_nd(input_descriptor, cudnn_type); @@ -2411,12 +2439,6 @@ port::Status CudnnSupport::DoConvolveImpl( const bool is_profiling = output_profile_result != nullptr; - DeviceMemory scratch; - SE_ASSIGN_OR_RETURN(dnn::AlgorithmDesc algo_desc, - GetCudnnConvolutionForwardAlgorithm( - stream, cudnn, algorithm_config, input_nd, filter, - conv, output_nd, scratch_allocator, &scratch)); - std::unique_ptr timer; if (is_profiling) { timer.reset(new CUDATimer(parent_)); // NOLINT @@ -2432,7 +2454,7 @@ port::Status CudnnSupport::DoConvolveImpl( // memory. See nvbugs/2138754, b/80018418. if (CUDNN_VERSION < 7300) { SE_RETURN_IF_ERROR([&] { - if (algo_desc.algo_id() != CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING) { + if (algorithm_desc.algo_id() != CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING) { return port::Status::OK(); } if (input_descriptor.ndims() < 3) { @@ -2457,7 +2479,8 @@ port::Status CudnnSupport::DoConvolveImpl( }()); } - if (algo_desc.algo_id() == CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED && + if (algorithm_desc.algo_id() == + CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED && !ShouldIncludeWinogradNonfusedAlgo(input_descriptor, output_descriptor)) { return port::Status(port::error::FAILED_PRECONDITION, "This configuration has potential integer overflow in " @@ -2469,18 +2492,19 @@ port::Status CudnnSupport::DoConvolveImpl( /*alpha=*/alpha, /*srcDesc=*/input_nd.handle(), /*srcData=*/input_data.opaque(), /*filterDesc=*/filter.handle(), /*filterData=*/filter_data.opaque(), /*convDesc=*/conv.handle(), - /*algo=*/ToConvForwardAlgo(algo_desc), /*workSpace=*/scratch.opaque(), - /*workSpaceSizeInBytes=*/scratch.size(), /*beta=*/beta, + /*algo=*/ToConvForwardAlgo(algorithm_desc), + /*workSpace=*/scratch_memory->opaque(), + /*workSpaceSizeInBytes=*/scratch_memory->size(), /*beta=*/beta, /*yDesc=*/output_nd.handle(), /*y=*/output_data->opaque())); if (is_profiling) { if (!timer->Stop(AsCUDAStream(stream))) { return port::Status(port::error::INTERNAL, "Failed to stop timer"); } - output_profile_result->set_algorithm(algo_desc); + output_profile_result->set_algorithm(algorithm_desc); output_profile_result->set_elapsed_time_in_ms( timer->GetElapsedMilliseconds()); - output_profile_result->set_scratch_size(scratch.size()); + output_profile_result->set_scratch_size(scratch_memory->size()); } return port::Status::OK(); @@ -2877,7 +2901,7 @@ port::Status CudnnSupport::DoBatchNormalizationBackwardImpl( return port::Status::OK(); } -bool CudnnSupport::DoConvolve( +bool CudnnSupport::PrepareForConvolution( Stream* stream, const dnn::BatchDescriptor& batch_descriptor, const DeviceMemory& input_data, const dnn::FilterDescriptor& filter_descriptor, @@ -2886,12 +2910,70 @@ bool CudnnSupport::DoConvolve( const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + return IsStatusOk(PrepareForConvolutionImpl( + stream, batch_descriptor, input_data, filter_descriptor, + filter_data, convolution_descriptor, output_descriptor, + output_data, dnn::DataType::kFloat, scratch_allocator, + algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); +} + +bool CudnnSupport::PrepareForConvolution( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + return IsStatusOk(PrepareForConvolutionImpl( + stream, batch_descriptor, input_data, filter_descriptor, + filter_data, convolution_descriptor, output_descriptor, + output_data, dnn::DataType::kDouble, scratch_allocator, + algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); +} + +bool CudnnSupport::PrepareForConvolution( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + dnn::DataType acc_type = + CudnnEnvVar::IsEnabled() + ? dnn::DataType::kFloat + : dnn::DataType::kHalf; + return IsStatusOk( + PrepareForConvolutionImpl( + stream, batch_descriptor, input_data, filter_descriptor, filter_data, + convolution_descriptor, output_descriptor, output_data, acc_type, + scratch_allocator, algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); +} + +bool CudnnSupport::DoConvolve( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) { return IsStatusOk( DoConvolveImpl(stream, batch_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, - output_data, dnn::DataType::kFloat, scratch_allocator, - algorithm_config, output_profile_result), + output_data, dnn::DataType::kFloat, algorithm_desc, + scratch_memory, output_profile_result), /*report_error=*/!output_profile_result); } @@ -2902,14 +2984,14 @@ bool CudnnSupport::DoConvolve( const DeviceMemory& filter_data, const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, - DeviceMemory* output_data, ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + DeviceMemory* output_data, const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) { return IsStatusOk( DoConvolveImpl(stream, batch_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, - output_data, dnn::DataType::kDouble, scratch_allocator, - algorithm_config, output_profile_result), + output_data, dnn::DataType::kDouble, algorithm_desc, + scratch_memory, output_profile_result), /*report_error=*/!output_profile_result); } @@ -2920,8 +3002,9 @@ bool CudnnSupport::DoConvolve( const DeviceMemory& filter_data, const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, - DeviceMemory* output_data, ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + DeviceMemory* output_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) { dnn::DataType acc_type = CudnnEnvVar::IsEnabled() @@ -2930,7 +3013,7 @@ bool CudnnSupport::DoConvolve( return IsStatusOk( DoConvolveImpl(stream, batch_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, - output_data, acc_type, scratch_allocator, algorithm_config, + output_data, acc_type, algorithm_desc, scratch_memory, output_profile_result), /*report_error=*/!output_profile_result); } @@ -3066,7 +3149,7 @@ bool CudnnSupport::DoTransformTensor(Stream* stream, } template -port::Status CudnnSupport::DoConvolveBackwardDataImpl( +port::Status CudnnSupport::PrepareForConvolutionBackwardDataImpl( Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, const dnn::BatchDescriptor& output_descriptor, @@ -3076,6 +3159,36 @@ port::Status CudnnSupport::DoConvolveBackwardDataImpl( DeviceMemory* backward_input_data, dnn::DataType accumulator_type, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + cudnnDataType_t cudnn_type = GetCudnnDataType(); + auto cudnn = cudnn_->GetHandle(parent_, stream); + + CudnnTensorDescriptor out_back_nd(output_descriptor, cudnn_type); + CudnnTensorDescriptor in_back_nd(input_descriptor, cudnn_type); + CudnnFilterDescriptor filter(filter_descriptor, cudnn_type); + CudnnConvolutionDescriptor conv(convolution_descriptor, + ToCudnnDataType(accumulator_type)); + + SE_ASSIGN_OR_RETURN( + *algorithm_desc, + GetCudnnConvolutionBackwardDataAlgorithm( + stream, cudnn, algorithm_config, in_back_nd, filter, conv, + out_back_nd, scratch_allocator, scratch_memory)); + + return port::Status::OK(); +} + +template +port::Status CudnnSupport::DoConvolveBackwardDataImpl( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, dnn::DataType accumulator_type, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) { cudnnDataType_t cudnn_type = GetCudnnDataType(); // Alpha is the scaling factor for input. @@ -3099,12 +3212,6 @@ port::Status CudnnSupport::DoConvolveBackwardDataImpl( const bool is_profiling = output_profile_result != nullptr; - DeviceMemory scratch; - SE_ASSIGN_OR_RETURN(dnn::AlgorithmDesc algo_desc, - GetCudnnConvolutionBackwardDataAlgorithm( - stream, cudnn, algorithm_config, in_back_nd, filter, - conv, out_back_nd, scratch_allocator, &scratch)); - std::unique_ptr timer; if (is_profiling) { timer.reset(new CUDATimer(parent_)); // NOLINT @@ -3116,7 +3223,8 @@ port::Status CudnnSupport::DoConvolveBackwardDataImpl( } } - if (algo_desc.algo_id() == CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED && + if (algorithm_desc.algo_id() == + CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED && !ShouldIncludeWinogradNonfusedAlgo(input_descriptor, output_descriptor)) { return port::Status(port::error::FAILED_PRECONDITION, "This configuration has potential integer overflow in " @@ -3126,44 +3234,44 @@ port::Status CudnnSupport::DoConvolveBackwardDataImpl( // Cudnn 7.1.4 has a bug if the workspace of the following convolution is not // zero-initialized, nvbugs/2254619. if (CUDNN_VERSION >= 7000 && CUDNN_VERSION < 7300 && - algo_desc.algo_id() == CUDNN_CONVOLUTION_BWD_DATA_ALGO_1 && - cudnn_type == CUDNN_DATA_HALF && algo_desc.tensor_ops_enabled() && + algorithm_desc.algo_id() == CUDNN_CONVOLUTION_BWD_DATA_ALGO_1 && + cudnn_type == CUDNN_DATA_HALF && algorithm_desc.tensor_ops_enabled() && input_descriptor.layout() == dnn::DataLayout::kBatchYXDepth && filter_descriptor.layout() == dnn::FilterLayout::kOutputInputYX && output_descriptor.layout() == dnn::DataLayout::kBatchDepthYX && (convolution_descriptor.vertical_filter_stride() > 1 || convolution_descriptor.horizontal_filter_stride() > 1)) { - stream->ThenMemZero(&scratch, scratch.size()); + stream->ThenMemZero(scratch_memory, scratch_memory->size()); } - RETURN_IF_CUDNN_ERROR( - cudnnConvolutionBackwardData(cudnn.handle(), - /*alpha=*/alpha, - /*wDesc=*/filter.handle(), - /*w=*/filter_data.opaque(), - /*dyDesc=*/out_back_nd.handle(), - /*dy=*/backward_output_data.opaque(), - /*convDesc=*/conv.handle(), - /*algo=*/ToConvBackwardDataAlgo(algo_desc), - /*workSpace=*/scratch.opaque(), - /*workSpaceSizeInBytes=*/scratch.size(), - /*beta=*/beta, - /*dxDesc=*/in_back_nd.handle(), - /*dx=*/backward_input_data->opaque())); + RETURN_IF_CUDNN_ERROR(cudnnConvolutionBackwardData( + cudnn.handle(), + /*alpha=*/alpha, + /*wDesc=*/filter.handle(), + /*w=*/filter_data.opaque(), + /*dyDesc=*/out_back_nd.handle(), + /*dy=*/backward_output_data.opaque(), + /*convDesc=*/conv.handle(), + /*algo=*/ToConvBackwardDataAlgo(algorithm_desc), + /*workSpace=*/scratch_memory->opaque(), + /*workSpaceSizeInBytes=*/scratch_memory->size(), + /*beta=*/beta, + /*dxDesc=*/in_back_nd.handle(), + /*dx=*/backward_input_data->opaque())); if (is_profiling) { if (!timer->Stop(AsCUDAStream(stream))) { return port::Status(port::error::INTERNAL, "Failed to stop timer"); } - output_profile_result->set_algorithm(algo_desc); + output_profile_result->set_algorithm(algorithm_desc); output_profile_result->set_elapsed_time_in_ms( timer->GetElapsedMilliseconds()); - output_profile_result->set_scratch_size(scratch.size()); + output_profile_result->set_scratch_size(scratch_memory->size()); } return port::Status::OK(); } -bool CudnnSupport::DoConvolveBackwardData( +bool CudnnSupport::PrepareForConvolutionBackwardData( Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, const dnn::BatchDescriptor& output_descriptor, @@ -3173,17 +3281,17 @@ bool CudnnSupport::DoConvolveBackwardData( DeviceMemory* backward_input_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) { + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { return IsStatusOk( - DoConvolveBackwardDataImpl( + PrepareForConvolutionBackwardDataImpl( stream, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, dnn::DataType::kDouble, scratch_allocator, - algorithm_config, output_profile_result), - /*report_error=*/!output_profile_result); + algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); } -bool CudnnSupport::DoConvolveBackwardData( +bool CudnnSupport::PrepareForConvolutionBackwardData( Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, const dnn::BatchDescriptor& output_descriptor, @@ -3193,17 +3301,17 @@ bool CudnnSupport::DoConvolveBackwardData( DeviceMemory* backward_input_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) { + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { return IsStatusOk( - DoConvolveBackwardDataImpl( + PrepareForConvolutionBackwardDataImpl( stream, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, dnn::DataType::kFloat, scratch_allocator, - algorithm_config, output_profile_result), - /*report_error=*/!output_profile_result); + algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); } -bool CudnnSupport::DoConvolveBackwardData( +bool CudnnSupport::PrepareForConvolutionBackwardData( Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, const dnn::BatchDescriptor& output_descriptor, @@ -3213,22 +3321,86 @@ bool CudnnSupport::DoConvolveBackwardData( DeviceMemory* backward_input_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) { + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { dnn::DataType acc_type = CudnnEnvVar::IsEnabled() ? dnn::DataType::kFloat : dnn::DataType::kHalf; return IsStatusOk( - DoConvolveBackwardDataImpl( + PrepareForConvolutionBackwardDataImpl( stream, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, acc_type, scratch_allocator, algorithm_config, - output_profile_result), + algorithm_desc, scratch_memory), + /*report_error=*/true); +} + +bool CudnnSupport::DoConvolveBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) { + return IsStatusOk( + DoConvolveBackwardDataImpl( + stream, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, dnn::DataType::kDouble, algorithm_desc, + scratch_memory, output_profile_result), + /*report_error=*/!output_profile_result); +} + +bool CudnnSupport::DoConvolveBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) { + return IsStatusOk( + DoConvolveBackwardDataImpl( + stream, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, dnn::DataType::kFloat, algorithm_desc, + scratch_memory, output_profile_result), + /*report_error=*/!output_profile_result); +} + +bool CudnnSupport::DoConvolveBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) { + dnn::DataType acc_type = + CudnnEnvVar::IsEnabled() + ? dnn::DataType::kFloat + : dnn::DataType::kHalf; + return IsStatusOk( + DoConvolveBackwardDataImpl(stream, filter_descriptor, filter_data, + output_descriptor, backward_output_data, + convolution_descriptor, input_descriptor, + backward_input_data, acc_type, algorithm_desc, + scratch_memory, output_profile_result), /*report_error=*/!output_profile_result); } template -port::Status CudnnSupport::DoConvolveBackwardFilterImpl( +port::Status CudnnSupport::PrepareForConvolutionBackwardFilterImpl( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -3238,6 +3410,36 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( DeviceMemory* backward_filter_data, dnn::DataType accumulator_type, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + cudnnDataType_t cudnn_type = GetCudnnDataType(); + auto cudnn = cudnn_->GetHandle(parent_, stream); + + CudnnTensorDescriptor out_back_nd(output_descriptor, cudnn_type); + CudnnTensorDescriptor input_nd(input_descriptor, cudnn_type); + CudnnFilterDescriptor filter(filter_descriptor, cudnn_type); + CudnnConvolutionDescriptor conv(convolution_descriptor, + ToCudnnDataType(accumulator_type)); + + SE_ASSIGN_OR_RETURN( + *algorithm_desc, + GetCudnnConvolutionBackwardFilterAlgorithm( + stream, cudnn, algorithm_config, input_nd, filter, conv, out_back_nd, + scratch_allocator, scratch_memory)); + + return port::Status::OK(); +} + +template +port::Status CudnnSupport::DoConvolveBackwardFilterImpl( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, dnn::DataType accumulator_type, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) { cudnnDataType_t cudnn_type = GetCudnnDataType(); // Alpha is the scaling factor for input. @@ -3261,12 +3463,6 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( const bool is_profiling = output_profile_result != nullptr; - DeviceMemory scratch; - SE_ASSIGN_OR_RETURN(dnn::AlgorithmDesc algo_desc, - GetCudnnConvolutionBackwardFilterAlgorithm( - stream, cudnn, algorithm_config, input_nd, filter, - conv, out_back_nd, scratch_allocator, &scratch)); - std::unique_ptr timer; if (is_profiling) { timer.reset(new CUDATimer(parent_)); // NOLINT @@ -3282,7 +3478,8 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( // results. See nvbugs/2072856 if (CUDNN_VERSION < 7300) { SE_RETURN_IF_ERROR([&] { - if (algo_desc.algo_id() != CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT_TILING) { + if (algorithm_desc.algo_id() != + CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT_TILING) { return port::Status::OK(); } if (output_descriptor.height() > 1 && output_descriptor.width() > 1) { @@ -3308,7 +3505,8 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( }()); } - if (algo_desc.algo_id() == CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED && + if (algorithm_desc.algo_id() == + CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED && !ShouldIncludeWinogradNonfusedAlgo(input_descriptor, output_descriptor)) { return port::Status(port::error::FAILED_PRECONDITION, "This configuration has potential integer overflow in " @@ -3324,7 +3522,7 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( // // See nvbugs/2379553. if (CUDNN_VERSION >= 7100 && CUDNN_VERSION < 7300 && - algo_desc.algo_id() == CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1 && + algorithm_desc.algo_id() == CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1 && cudnn_type == CUDNN_DATA_HALF && input_descriptor.layout() == dnn::DataLayout::kBatchYXDepth && filter_descriptor.layout() == dnn::FilterLayout::kOutputYXInput && @@ -3342,9 +3540,9 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( /*diffDesc=*/out_back_nd.handle(), /*diffData=*/backward_output_data.opaque(), /*convDesc=*/conv.handle(), - /*algo=*/ToConvBackwardFilterAlgo(algo_desc), - /*workSpace=*/scratch.opaque(), - /*workSpaceSizeInBytes=*/scratch.size(), + /*algo=*/ToConvBackwardFilterAlgo(algorithm_desc), + /*workSpace=*/scratch_memory->opaque(), + /*workSpaceSizeInBytes=*/scratch_memory->size(), /*beta=*/beta, /*gradDesc=*/filter.handle(), /*dw=*/backward_filter_data->opaque())); @@ -3352,16 +3550,16 @@ port::Status CudnnSupport::DoConvolveBackwardFilterImpl( if (!timer->Stop(AsCUDAStream(stream))) { return port::Status(port::error::INTERNAL, "Failed to stop timer"); } - output_profile_result->set_algorithm(algo_desc); + output_profile_result->set_algorithm(algorithm_desc); output_profile_result->set_elapsed_time_in_ms( timer->GetElapsedMilliseconds()); - output_profile_result->set_scratch_size(scratch.size()); + output_profile_result->set_scratch_size(scratch_memory->size()); } return port::Status::OK(); } -bool CudnnSupport::DoConvolveBackwardFilter( +bool CudnnSupport::PrepareForConvolutionBackwardFilter( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -3371,18 +3569,17 @@ bool CudnnSupport::DoConvolveBackwardFilter( DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) { + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { return IsStatusOk( - DoConvolveBackwardFilterImpl( + PrepareForConvolutionBackwardFilterImpl( stream, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, - backward_filter_data, dnn::DataType::kDouble, - - scratch_allocator, algorithm_config, output_profile_result), - /*report_error=*/!output_profile_result); + backward_filter_data, dnn::DataType::kDouble, scratch_allocator, + algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); } -bool CudnnSupport::DoConvolveBackwardFilter( +bool CudnnSupport::PrepareForConvolutionBackwardFilter( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -3392,18 +3589,17 @@ bool CudnnSupport::DoConvolveBackwardFilter( DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) { - return IsStatusOk(DoConvolveBackwardFilterImpl( - stream, input_descriptor, input_data, output_descriptor, - backward_output_data, convolution_descriptor, - filter_descriptor, backward_filter_data, - - dnn::DataType::kFloat, scratch_allocator, - algorithm_config, output_profile_result), - /*report_error=*/!output_profile_result); + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { + return IsStatusOk( + PrepareForConvolutionBackwardFilterImpl( + stream, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, dnn::DataType::kFloat, scratch_allocator, + algorithm_config, algorithm_desc, scratch_memory), + /*report_error=*/true); } -bool CudnnSupport::DoConvolveBackwardFilter( +bool CudnnSupport::PrepareForConvolutionBackwardFilter( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -3413,20 +3609,83 @@ bool CudnnSupport::DoConvolveBackwardFilter( DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) { + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory) { dnn::DataType acc_type = CudnnEnvVar::IsEnabled() ? dnn::DataType::kFloat : dnn::DataType::kHalf; return IsStatusOk( - DoConvolveBackwardFilterImpl( + PrepareForConvolutionBackwardFilterImpl( stream, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, backward_filter_data, acc_type, scratch_allocator, algorithm_config, - output_profile_result), + algorithm_desc, scratch_memory), + /*report_error=*/true); +} + +bool CudnnSupport::DoConvolveBackwardFilter( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) { + return IsStatusOk( + DoConvolveBackwardFilterImpl( + stream, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, dnn::DataType::kDouble, algorithm_desc, + scratch_memory, output_profile_result), /*report_error=*/!output_profile_result); } +bool CudnnSupport::DoConvolveBackwardFilter( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) { + return IsStatusOk( + DoConvolveBackwardFilterImpl( + stream, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, dnn::DataType::kFloat, algorithm_desc, + scratch_memory, output_profile_result), + /*report_error=*/!output_profile_result); +} + +bool CudnnSupport::DoConvolveBackwardFilter( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) { + dnn::DataType acc_type = + CudnnEnvVar::IsEnabled() + ? dnn::DataType::kFloat + : dnn::DataType::kHalf; + return IsStatusOk(DoConvolveBackwardFilterImpl( + stream, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, + filter_descriptor, backward_filter_data, acc_type, + algorithm_desc, scratch_memory, output_profile_result), + /*report_error=*/!output_profile_result); +} + template port::Status CudnnSupport::DoConvolveBackwardBiasImpl( Stream* stream, const dnn::BatchDescriptor& input_descriptor, diff --git a/tensorflow/stream_executor/cuda/cuda_dnn.h b/tensorflow/stream_executor/cuda/cuda_dnn.h index 0641be140d..0b8f4d035c 100644 --- a/tensorflow/stream_executor/cuda/cuda_dnn.h +++ b/tensorflow/stream_executor/cuda/cuda_dnn.h @@ -252,6 +252,43 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* scale_backprop, DeviceMemory* offset_backprop) override; + bool PrepareForConvolution( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + + bool PrepareForConvolution( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + + bool PrepareForConvolution( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + bool DoConvolve(Stream* stream, const dnn::BatchDescriptor& batch_descriptor, const DeviceMemory& input_data, const dnn::FilterDescriptor& filter_descriptor, @@ -259,8 +296,8 @@ class CudnnSupport : public dnn::DnnSupport { const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; bool DoConvolve(Stream* stream, const dnn::BatchDescriptor& batch_descriptor, @@ -270,8 +307,8 @@ class CudnnSupport : public dnn::DnnSupport { const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; bool DoConvolve(Stream* stream, const dnn::BatchDescriptor& batch_descriptor, @@ -281,8 +318,8 @@ class CudnnSupport : public dnn::DnnSupport { const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; bool DoFusedConvolve( @@ -384,7 +421,20 @@ class CudnnSupport : public dnn::DnnSupport { return false; } - bool DoConvolveBackwardData( + bool PrepareForConvolutionBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + + bool PrepareForConvolutionBackwardData( Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, const dnn::BatchDescriptor& output_descriptor, @@ -394,6 +444,32 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* backward_input_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + + bool PrepareForConvolutionBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + + bool DoConvolveBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; bool DoConvolveBackwardData( @@ -404,8 +480,8 @@ class CudnnSupport : public dnn::DnnSupport { const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& input_descriptor, DeviceMemory* backward_input_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; bool DoConvolveBackwardData( @@ -416,11 +492,11 @@ class CudnnSupport : public dnn::DnnSupport { const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& input_descriptor, DeviceMemory* backward_input_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; - bool DoConvolveBackwardFilter( + bool PrepareForConvolutionBackwardFilter( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -430,9 +506,10 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) override; + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; - bool DoConvolveBackwardFilter( + bool PrepareForConvolutionBackwardFilter( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -442,9 +519,10 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - dnn::ProfileResult* output_profile_result) override; + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; - bool DoConvolveBackwardFilter( + bool PrepareForConvolutionBackwardFilter( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, const dnn::BatchDescriptor& output_descriptor, @@ -454,6 +532,43 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) override; + + bool DoConvolveBackwardFilter( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) override; + + bool DoConvolveBackwardFilter( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + dnn::ProfileResult* output_profile_result) override; + + bool DoConvolveBackwardFilter( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) override; bool DoConvolveBackwardBias( @@ -663,8 +778,8 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* offset_backprop); template - port::Status DoConvolveImpl( - Stream* stream, const dnn::BatchDescriptor& input_descriptor, + port::Status PrepareForConvolutionImpl( + Stream* stream, const dnn::BatchDescriptor& batch_descriptor, const DeviceMemory& input_data, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, @@ -673,6 +788,19 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* output_data, dnn::DataType accumulator_type, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory); + + template + port::Status DoConvolveImpl( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory* output_data, dnn::DataType accumulator_type, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result); template @@ -693,7 +821,7 @@ class CudnnSupport : public dnn::DnnSupport { dnn::ProfileResult* output_profile_result); template - port::Status DoConvolveBackwardDataImpl( + port::Status PrepareForConvolutionBackwardDataImpl( Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, const dnn::BatchDescriptor& output_descriptor, @@ -703,19 +831,45 @@ class CudnnSupport : public dnn::DnnSupport { DeviceMemory* backward_input_data, dnn::DataType accumulator_type, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory); + + template + port::Status DoConvolveBackwardDataImpl( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, dnn::DataType accumulator_type, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result); template - port::Status DoConvolveBackwardFilterImpl( + port::Status PrepareForConvolutionBackwardFilterImpl( Stream* stream, const dnn::BatchDescriptor& input_descriptor, const DeviceMemory& input_data, - const dnn::BatchDescriptor& output_descriptor, + const dnn::BatchDescriptor& output_descriptor_in, DeviceMemory backward_output_data, const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::FilterDescriptor& filter_descriptor, DeviceMemory* backward_filter_data, dnn::DataType accumulator_type, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, DeviceMemory* scratch_memory); + + template + port::Status DoConvolveBackwardFilterImpl( + Stream* stream, const dnn::BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, dnn::DataType accumulator_type, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result); template diff --git a/tensorflow/stream_executor/dnn.h b/tensorflow/stream_executor/dnn.h index c044a356ef..322232d263 100644 --- a/tensorflow/stream_executor/dnn.h +++ b/tensorflow/stream_executor/dnn.h @@ -730,6 +730,7 @@ class PoolingDescriptor { class AlgorithmDesc { public: typedef int64 Index; + AlgorithmDesc() : AlgorithmDesc(0, false) {} AlgorithmDesc(Index a, bool use_tensor_ops) { proto_.set_algo_id(a); proto_.set_math_type(use_tensor_ops ? AlgorithmProto::TENSOR_OP_MATH @@ -1175,6 +1176,43 @@ class DnnSupport { return false; } + virtual bool PrepareForConvolution( + Stream* stream, const BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const ConvolutionDescriptor& convolution_descriptor, + const BatchDescriptor& output_descriptor, + DeviceMemory* output_data, ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + + virtual bool PrepareForConvolution( + Stream* stream, const BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const ConvolutionDescriptor& convolution_descriptor, + const BatchDescriptor& output_descriptor, + DeviceMemory* output_data, ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + + virtual bool PrepareForConvolution( + Stream* stream, const BatchDescriptor& batch_descriptor, + const DeviceMemory& input_data, + const FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const ConvolutionDescriptor& convolution_descriptor, + const BatchDescriptor& output_descriptor, + DeviceMemory* output_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + // Enqueues a single-precision convolution operation onto the stream. // // Arguments (all borrowed): @@ -1188,10 +1226,10 @@ class DnnSupport { // output_descriptor: dimensions of the output layer. // output_data: un-owned device memory region in which to place the // convolution result. - // scratch_allocator: un-owned, may-be-null object that may allocate scratch - // space in order to speed up the convolution operation. - // algorithm_config: specifies which algorithm should be used for the + // algorithm_desc: specifies which algorithm should be used for the // operation. + // scratch: un-owned device memory for scratch space in order to speed up + // the convolution operation. // output_profile_result: the output profile result for this call. The // profiling is only enabled when this is not nullptr. // @@ -1216,8 +1254,9 @@ class DnnSupport { const DeviceMemory& filter_data, const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, - DeviceMemory* output_data, ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + DeviceMemory* output_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; // Enqueues a double-precision convolution operation onto the stream. @@ -1229,8 +1268,9 @@ class DnnSupport { const DeviceMemory& filter_data, const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, - DeviceMemory* output_data, ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + DeviceMemory* output_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, dnn::ProfileResult* output_profile_result) = 0; // Enqueues a half-precision convolution operation onto the stream. @@ -1243,8 +1283,8 @@ class DnnSupport { const dnn::ConvolutionDescriptor& convolution_descriptor, const dnn::BatchDescriptor& output_descriptor, DeviceMemory* output_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; // Return a list of algorithms supported by the forward convolution pass. @@ -1300,6 +1340,45 @@ class DnnSupport { const BatchDescriptor& output_descriptor, DeviceMemory* output_data) = 0; + virtual bool PrepareForConvolutionBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + + virtual bool PrepareForConvolutionBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + + virtual bool PrepareForConvolutionBackwardData( + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, + const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + // Enqueues a single-precision backward convolution (for data) operation onto // the stream. // @@ -1319,15 +1398,15 @@ class DnnSupport { // scratch_allocator: un-owned, may-be-null object that may allocate scratch // space in order to speed up the convolution operation. virtual bool DoConvolveBackwardData( - Stream* stream, const FilterDescriptor& filter_descriptor, + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, - const BatchDescriptor& output_descriptor, + const dnn::BatchDescriptor& output_descriptor, DeviceMemory backward_output_data, - const ConvolutionDescriptor& convolution_descriptor, - const BatchDescriptor& input_descriptor, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, DeviceMemory* backward_input_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; // Return a list of algorithms supported by the backward convolution pass for @@ -1337,28 +1416,67 @@ class DnnSupport { std::vector* out_algorithms); virtual bool DoConvolveBackwardData( - Stream* stream, const FilterDescriptor& filter_descriptor, + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, - const BatchDescriptor& output_descriptor, + const dnn::BatchDescriptor& output_descriptor, DeviceMemory backward_output_data, - const ConvolutionDescriptor& convolution_descriptor, - const BatchDescriptor& input_descriptor, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, DeviceMemory* backward_input_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; virtual bool DoConvolveBackwardData( - Stream* stream, const FilterDescriptor& filter_descriptor, + Stream* stream, const dnn::FilterDescriptor& filter_descriptor, const DeviceMemory& filter_data, + const dnn::BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const dnn::ConvolutionDescriptor& convolution_descriptor, + const dnn::BatchDescriptor& input_descriptor, + DeviceMemory* backward_input_data, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, + ProfileResult* output_profile_result) = 0; + + virtual bool PrepareForConvolutionBackwardFilter( + Stream* stream, const BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const ConvolutionDescriptor& convolution_descriptor, + const FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + + virtual bool PrepareForConvolutionBackwardFilter( + Stream* stream, const BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, + const BatchDescriptor& output_descriptor, + DeviceMemory backward_output_data, + const ConvolutionDescriptor& convolution_descriptor, + const FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, + ScratchAllocator* scratch_allocator, + const dnn::AlgorithmConfig& algorithm_config, + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; + + virtual bool PrepareForConvolutionBackwardFilter( + Stream* stream, const BatchDescriptor& input_descriptor, + const DeviceMemory& input_data, const BatchDescriptor& output_descriptor, DeviceMemory backward_output_data, const ConvolutionDescriptor& convolution_descriptor, - const BatchDescriptor& input_descriptor, - DeviceMemory* backward_input_data, + const FilterDescriptor& filter_descriptor, + DeviceMemory* backward_filter_data, ScratchAllocator* scratch_allocator, const dnn::AlgorithmConfig& algorithm_config, - ProfileResult* output_profile_result) = 0; + dnn::AlgorithmDesc* algorithm_desc, + DeviceMemory* scratch_memory) = 0; // Enqueues a single-precision backward convolution (for filter) operation // onto the stream. @@ -1387,8 +1505,8 @@ class DnnSupport { const ConvolutionDescriptor& convolution_descriptor, const FilterDescriptor& filter_descriptor, DeviceMemory* backward_filter_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; // Return a list of algorithms supported by the backward convolution pass for @@ -1405,8 +1523,8 @@ class DnnSupport { const ConvolutionDescriptor& convolution_descriptor, const FilterDescriptor& filter_descriptor, DeviceMemory* backward_filter_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; virtual bool DoConvolveBackwardFilter( @@ -1417,8 +1535,8 @@ class DnnSupport { const ConvolutionDescriptor& convolution_descriptor, const FilterDescriptor& filter_descriptor, DeviceMemory* backward_filter_data, - ScratchAllocator* scratch_allocator, - const dnn::AlgorithmConfig& algorithm_config, + const dnn::AlgorithmDesc& algorithm_desc, + DeviceMemory* scratch_memory, ProfileResult* output_profile_result) = 0; // Enqueues a single-precision backward convolution (for bias) operation onto diff --git a/tensorflow/stream_executor/stream.cc b/tensorflow/stream_executor/stream.cc index 3edc66cde8..4503127bee 100644 --- a/tensorflow/stream_executor/stream.cc +++ b/tensorflow/stream_executor/stream.cc @@ -549,11 +549,16 @@ Stream &Stream::ThenConvolveWithScratch( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - CheckError(dnn->DoConvolve( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + CheckError(dnn->PrepareForConvolution( this, input_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, output, scratch_allocator, - dnn::AlgorithmConfig(), - /*output_profile_result=*/nullptr)); + dnn::AlgorithmConfig(), &algorithm_desc, &scratch_memory)); + CheckError(dnn->DoConvolve( + this, input_descriptor, input_data, filter_descriptor, filter_data, + convolution_descriptor, output_descriptor, output, algorithm_desc, + &scratch_memory, nullptr)); } else { SetErrorAndLogNoDnnSupport(); } @@ -576,11 +581,16 @@ Stream &Stream::ThenConvolveWithScratch( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - CheckError(dnn->DoConvolve( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + CheckError(dnn->PrepareForConvolution( this, input_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, output, scratch_allocator, - dnn::AlgorithmConfig(), - /*output_profile_result=*/nullptr)); + dnn::AlgorithmConfig(), &algorithm_desc, &scratch_memory)); + CheckError(dnn->DoConvolve( + this, input_descriptor, input_data, filter_descriptor, filter_data, + convolution_descriptor, output_descriptor, output, algorithm_desc, + &scratch_memory, nullptr)); } else { SetErrorAndLogNoDnnSupport(); } @@ -758,10 +768,18 @@ Stream &Stream::ThenConvolveWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolve( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolution( this, input_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, output, scratch_allocator, - algorithm_config, output_profile_result); + algorithm_config, &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolve( + this, input_descriptor, input_data, filter_descriptor, filter_data, + convolution_descriptor, output_descriptor, output, algorithm_desc, + &scratch_memory, output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -789,10 +807,18 @@ Stream &Stream::ThenConvolveWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolve( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolution( this, input_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, output, scratch_allocator, - algorithm_config, output_profile_result); + algorithm_config, &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolve( + this, input_descriptor, input_data, filter_descriptor, filter_data, + convolution_descriptor, output_descriptor, output, algorithm_desc, + &scratch_memory, output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -820,10 +846,18 @@ Stream &Stream::ThenConvolveWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolve( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolution( this, input_descriptor, input_data, filter_descriptor, filter_data, convolution_descriptor, output_descriptor, output, scratch_allocator, - algorithm_config, output_profile_result); + algorithm_config, &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolve( + this, input_descriptor, input_data, filter_descriptor, filter_data, + convolution_descriptor, output_descriptor, output, algorithm_desc, + &scratch_memory, output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -969,10 +1003,17 @@ Stream &Stream::ThenConvolveBackwardDataWithScratch( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - CheckError(dnn->DoConvolveBackwardData( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + CheckError(dnn->PrepareForConvolutionBackwardData( this, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, scratch_allocator, dnn::AlgorithmConfig(), + &algorithm_desc, &scratch_memory)); + CheckError(dnn->DoConvolveBackwardData( + this, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, algorithm_desc, &scratch_memory, /*output_profile_result=*/nullptr)); } else { SetErrorAndLogNoDnnSupport(); @@ -999,11 +1040,20 @@ Stream &Stream::ThenConvolveBackwardDataWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolveBackwardData( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolutionBackwardData( this, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, scratch_allocator, algorithm_config, - output_profile_result); + &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolveBackwardData( + this, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, algorithm_desc, &scratch_memory, + output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -1032,11 +1082,20 @@ Stream &Stream::ThenConvolveBackwardDataWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolveBackwardData( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolutionBackwardData( this, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, scratch_allocator, algorithm_config, - output_profile_result); + &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolveBackwardData( + this, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, algorithm_desc, &scratch_memory, + output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -1065,11 +1124,20 @@ Stream &Stream::ThenConvolveBackwardDataWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolveBackwardData( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolutionBackwardData( this, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, scratch_allocator, algorithm_config, - output_profile_result); + &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolveBackwardData( + this, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, algorithm_desc, &scratch_memory, + output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -1096,10 +1164,17 @@ Stream &Stream::ThenConvolveBackwardDataWithScratch( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - CheckError(dnn->DoConvolveBackwardData( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + CheckError(dnn->PrepareForConvolutionBackwardData( this, filter_descriptor, filter_data, output_descriptor, backward_output_data, convolution_descriptor, input_descriptor, backward_input_data, scratch_allocator, dnn::AlgorithmConfig(), + &algorithm_desc, &scratch_memory)); + CheckError(dnn->DoConvolveBackwardData( + this, filter_descriptor, filter_data, output_descriptor, + backward_output_data, convolution_descriptor, input_descriptor, + backward_input_data, algorithm_desc, &scratch_memory, /*output_profile_result=*/nullptr)); } else { SetErrorAndLogNoDnnSupport(); @@ -1138,10 +1213,17 @@ Stream &Stream::ThenConvolveBackwardFilterWithScratch( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - CheckError(dnn->DoConvolveBackwardFilter( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + CheckError(dnn->PrepareForConvolutionBackwardFilter( this, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, backward_filter_data, scratch_allocator, dnn::AlgorithmConfig(), + &algorithm_desc, &scratch_memory)); + CheckError(dnn->DoConvolveBackwardFilter( + this, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, algorithm_desc, &scratch_memory, /*output_profile_result=*/nullptr)); } else { SetErrorAndLogNoDnnSupport(); @@ -1168,11 +1250,20 @@ Stream &Stream::ThenConvolveBackwardFilterWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolveBackwardFilter( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolutionBackwardFilter( this, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, backward_filter_data, scratch_allocator, algorithm_config, - output_profile_result); + &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolveBackwardFilter( + this, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, algorithm_desc, &scratch_memory, + output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -1201,11 +1292,20 @@ Stream &Stream::ThenConvolveBackwardFilterWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolveBackwardFilter( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolutionBackwardFilter( this, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, backward_filter_data, scratch_allocator, algorithm_config, - output_profile_result); + &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolveBackwardFilter( + this, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, algorithm_desc, &scratch_memory, + output_profile_result); + } if (!status && !output_profile_result) { SetError(); } @@ -1232,10 +1332,17 @@ Stream &Stream::ThenConvolveBackwardFilterWithScratch( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - CheckError(dnn->DoConvolveBackwardFilter( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + CheckError(dnn->PrepareForConvolutionBackwardFilter( this, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, backward_filter_data, scratch_allocator, dnn::AlgorithmConfig(), + &algorithm_desc, &scratch_memory)); + CheckError(dnn->DoConvolveBackwardFilter( + this, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, algorithm_desc, &scratch_memory, /*output_profile_result=*/nullptr)); } else { SetErrorAndLogNoDnnSupport(); @@ -1262,11 +1369,20 @@ Stream &Stream::ThenConvolveBackwardFilterWithAlgorithm( if (ok()) { if (dnn::DnnSupport *dnn = parent_->AsDnn()) { - auto status = dnn->DoConvolveBackwardFilter( + DeviceMemory scratch_memory; + dnn::AlgorithmDesc algorithm_desc; + auto status = dnn->PrepareForConvolutionBackwardFilter( this, input_descriptor, input_data, output_descriptor, backward_output_data, convolution_descriptor, filter_descriptor, backward_filter_data, scratch_allocator, algorithm_config, - output_profile_result); + &algorithm_desc, &scratch_memory); + if (status) { + status = dnn->DoConvolveBackwardFilter( + this, input_descriptor, input_data, output_descriptor, + backward_output_data, convolution_descriptor, filter_descriptor, + backward_filter_data, algorithm_desc, &scratch_memory, + output_profile_result); + } if (!status && !output_profile_result) { SetError(); } -- GitLab From 0a8e110e8e4558fd99244316051ce0ff729119e9 Mon Sep 17 00:00:00 2001 From: Ruoxin Sang Date: Sat, 1 Dec 2018 00:52:39 -0800 Subject: [PATCH 119/596] Use Keras v2 optimizer for TPUStrategy in keras_test.py. PiperOrigin-RevId: 223625692 --- tensorflow/contrib/distribute/python/keras_test.py | 10 +--------- 1 file changed, 1 insertion(+), 9 deletions(-) diff --git a/tensorflow/contrib/distribute/python/keras_test.py b/tensorflow/contrib/distribute/python/keras_test.py index 1d00281974..435f10358f 100644 --- a/tensorflow/contrib/distribute/python/keras_test.py +++ b/tensorflow/contrib/distribute/python/keras_test.py @@ -1267,17 +1267,9 @@ class TestDistributionStrategyCorrectness(test.TestCase, # We have initialized the model to the same weight for the distribution # and non-distribution run. model.set_weights(initial_weights) - # TODO(b/120245072): Also use gradient_descent_keras.SGD for - # TPUStrategy. - # pylint: disable=line-too-long - if with_distribution and with_distribution.__class__.__name__ == 'TPUStrategy': - # pylint: enable=line-too-long - optimizer = gradient_descent.GradientDescentOptimizer(0.5) - else: - optimizer = gradient_descent_keras.SGD(0.5) model.compile( loss=keras.losses.mean_squared_error, - optimizer=optimizer, + optimizer=gradient_descent_keras.SGD(0.5), distribute=with_distribution) training_inputs, eval_inputs, predict_inputs = ( -- GitLab From 6866480d9f97d14c6fc5e2488e396181be3b43c1 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Sat, 1 Dec 2018 01:02:22 -0800 Subject: [PATCH 120/596] compat: Update forward compatibility horizon to 2018-12-01 PiperOrigin-RevId: 223626306 --- tensorflow/python/compat/compat.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/compat/compat.py b/tensorflow/python/compat/compat.py index 156680e081..a7e8bef670 100644 --- a/tensorflow/python/compat/compat.py +++ b/tensorflow/python/compat/compat.py @@ -32,7 +32,7 @@ from tensorflow.python.ops import variable_scope from tensorflow.python.util import tf_contextlib from tensorflow.python.util.tf_export import tf_export -_FORWARD_COMPATIBILITY_HORIZON = datetime.date(2018, 11, 30) +_FORWARD_COMPATIBILITY_HORIZON = datetime.date(2018, 12, 1) @tf_export("compat.forward_compatible") -- GitLab From d15b49732f9bec02571f20b2c8c3a0b62d9888fb Mon Sep 17 00:00:00 2001 From: Andrew Selle Date: Sat, 1 Dec 2018 01:38:43 -0800 Subject: [PATCH 121/596] Fix defaulting of axis in TF 2.0 softmax/entropy code. PiperOrigin-RevId: 223628497 --- tensorflow/python/ops/nn_ops.py | 4 +++- tensorflow/tools/api/golden/v1/tensorflow.nn.pbtxt | 2 +- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/ops/nn_ops.py b/tensorflow/python/ops/nn_ops.py index cec532ea86..225904854f 100644 --- a/tensorflow/python/ops/nn_ops.py +++ b/tensorflow/python/ops/nn_ops.py @@ -2365,7 +2365,7 @@ def softmax_cross_entropy_with_logits_v2(labels, logits, axis=-1, name=None): @tf_export(v1=["nn.softmax_cross_entropy_with_logits_v2"]) @deprecated_args(None, "dim is deprecated, use axis instead", "dim") def softmax_cross_entropy_with_logits_v2_helper( - labels, logits, axis=-1, name=None, dim=None): + labels, logits, axis=None, name=None, dim=None): """Computes softmax cross entropy between `logits` and `labels`. Measures the probability error in discrete classification tasks in which the @@ -2419,6 +2419,8 @@ def softmax_cross_entropy_with_logits_v2_helper( # results. axis = deprecated_argument_lookup("axis", axis, "dim", dim) del dim + if axis is None: + axis = -1 with ops.name_scope(name, "softmax_cross_entropy_with_logits", [logits, labels]) as name: diff --git a/tensorflow/tools/api/golden/v1/tensorflow.nn.pbtxt b/tensorflow/tools/api/golden/v1/tensorflow.nn.pbtxt index 3697912ec2..40e20f8c91 100644 --- a/tensorflow/tools/api/golden/v1/tensorflow.nn.pbtxt +++ b/tensorflow/tools/api/golden/v1/tensorflow.nn.pbtxt @@ -326,7 +326,7 @@ tf_module { } member_method { name: "softmax_cross_entropy_with_logits_v2" - argspec: "args=[\'labels\', \'logits\', \'axis\', \'name\', \'dim\'], varargs=None, keywords=None, defaults=[\'-1\', \'None\', \'None\'], " + argspec: "args=[\'labels\', \'logits\', \'axis\', \'name\', \'dim\'], varargs=None, keywords=None, defaults=[\'None\', \'None\', \'None\'], " } member_method { name: "softplus" -- GitLab From 35c19c850a2c0cfb13cda56a493fa9163936d66c Mon Sep 17 00:00:00 2001 From: Asim Shankar Date: Sat, 1 Dec 2018 07:59:21 -0800 Subject: [PATCH 122/596] Java tests: Fix nightly builds. For reasons not quite clear to me, the nightly builds (which are currently using bazel 0.16.1) see failing Java tests with errors like: java.lang.NoSuchMethodError: java.nio.DoubleBuffer.flip()Ljava/nio/DoubleBuffer; when calling "flip()" or "clear()" on a java.nio.DoubleBuffer object. This is symptomatic of cross compiling in Java (e.g. javac -source=1.9 -target=1.8) and this change works around that problem. This does not happen with bazel 0.17.0 on the same test machine. (Bringing me to the "reasons not quite clear to me" mentioned above). PiperOrigin-RevId: 223643679 --- .../test/java/org/tensorflow/TensorTest.java | 46 ++++++++++++++----- 1 file changed, 34 insertions(+), 12 deletions(-) diff --git a/tensorflow/java/src/test/java/org/tensorflow/TensorTest.java b/tensorflow/java/src/test/java/org/tensorflow/TensorTest.java index 1bd00a763d..3229cce277 100644 --- a/tensorflow/java/src/test/java/org/tensorflow/TensorTest.java +++ b/tensorflow/java/src/test/java/org/tensorflow/TensorTest.java @@ -21,6 +21,7 @@ import static org.junit.Assert.assertEquals; import static org.junit.Assert.assertTrue; import static org.junit.Assert.fail; +import java.nio.Buffer; import java.nio.ByteBuffer; import java.nio.ByteOrder; import java.nio.DoubleBuffer; @@ -100,7 +101,7 @@ public class TensorTest { : ByteOrder.LITTLE_ENDIAN) .asDoubleBuffer() .put(doubles); - buf.flip(); + flipBuffer(buf); try (Tensor t = Tensor.create(new long[] {doubles.length}, buf)) { double[] actual = new double[doubles.length]; assertArrayEquals(doubles, t.copyTo(actual), EPSILON); @@ -179,30 +180,30 @@ public class TensorTest { { ByteBuffer bbuf = ByteBuffer.allocate(1024).order(ByteOrder.nativeOrder()); - bbuf.clear(); // FLOAT + clearBuffer(bbuf); // FLOAT tfloats.writeTo(bbuf); assertEquals(tfloats.numBytes(), bbuf.position()); - bbuf.flip(); + flipBuffer(bbuf); assertEquals(floats[0], bbuf.asFloatBuffer().get(0), EPSILON); - bbuf.clear(); // DOUBLE + clearBuffer(bbuf); // DOUBLE tdoubles.writeTo(bbuf); assertEquals(tdoubles.numBytes(), bbuf.position()); - bbuf.flip(); + flipBuffer(bbuf); assertEquals(doubles[0], bbuf.asDoubleBuffer().get(0), EPSILON); - bbuf.clear(); // INT32 + clearBuffer(bbuf); // INT32 tints.writeTo(bbuf); assertEquals(tints.numBytes(), bbuf.position()); - bbuf.flip(); + flipBuffer(bbuf); assertEquals(ints[0], bbuf.asIntBuffer().get(0)); - bbuf.clear(); // INT64 + clearBuffer(bbuf); // INT64 tlongs.writeTo(bbuf); assertEquals(tlongs.numBytes(), bbuf.position()); - bbuf.flip(); + flipBuffer(bbuf); assertEquals(longs[0], bbuf.asLongBuffer().get(0)); - bbuf.clear(); // BOOL + clearBuffer(bbuf); // BOOL tbools.writeTo(bbuf); assertEquals(tbools.numBytes(), bbuf.position()); - bbuf.flip(); + flipBuffer(bbuf); assertEquals(bools[0], bbuf.get(0) != 0); } @@ -254,7 +255,7 @@ public class TensorTest { : ByteOrder.LITTLE_ENDIAN) .asDoubleBuffer(); tdoubles.writeTo(foreignBuf); - foreignBuf.flip(); + flipBuffer(foreignBuf); double[] actual = new double[foreignBuf.remaining()]; foreignBuf.get(actual); assertArrayEquals(doubles, actual, EPSILON); @@ -547,4 +548,25 @@ public class TensorTest { // expected. } } + + // Workaround for cross compiliation + // (e.g., javac -source 1.9 -target 1.8). + // + // In Java 8 and prior, subclasses of java.nio.Buffer (e.g., java.nio.DoubleBuffer) inherited the + // "flip()" and "clear()" methods from java.nio.Buffer resulting in the signature: + // Buffer flip(); + // In Java 9 these subclasses had their own methods like: + // DoubleBuffer flip(); + // As a result, compiling for 1.9 source for a target of JDK 1.8 would result in errors at runtime + // like: + // + // java.lang.NoSuchMethodError: java.nio.DoubleBuffer.flip()Ljava/nio/DoubleBuffer + private static void flipBuffer(Buffer buf) { + buf.flip(); + } + + // See comment for flipBuffer() + private static void clearBuffer(Buffer buf) { + buf.clear(); + } } -- GitLab From 0e672efe02940c803f7201d515c9ed0f36135e90 Mon Sep 17 00:00:00 2001 From: John Lin Date: Sun, 2 Dec 2018 00:04:22 +0800 Subject: [PATCH 123/596] Remove redundant word "the" --- .../tools/dockerfiles/dockerfiles/cpu-devel-jupyter.Dockerfile | 2 +- tensorflow/tools/dockerfiles/dockerfiles/cpu-devel.Dockerfile | 2 +- tensorflow/tools/dockerfiles/dockerfiles/cpu-jupyter.Dockerfile | 2 +- tensorflow/tools/dockerfiles/dockerfiles/cpu.Dockerfile | 2 +- .../tools/dockerfiles/dockerfiles/gpu-devel-jupyter.Dockerfile | 2 +- tensorflow/tools/dockerfiles/dockerfiles/gpu-devel.Dockerfile | 2 +- tensorflow/tools/dockerfiles/dockerfiles/gpu-jupyter.Dockerfile | 2 +- tensorflow/tools/dockerfiles/dockerfiles/gpu.Dockerfile | 2 +- tensorflow/tools/dockerfiles/spec.yml | 2 +- tensorflow/tools/docs/parser.py | 2 +- 10 files changed, 10 insertions(+), 10 deletions(-) diff --git a/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel-jupyter.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel-jupyter.Dockerfile index 9577920e13..14ddf08199 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel-jupyter.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel-jupyter.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel.Dockerfile index 2f4a3d6beb..16973b47af 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/cpu-devel.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/cpu-jupyter.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/cpu-jupyter.Dockerfile index b663cca8e2..d8fabadec2 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/cpu-jupyter.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/cpu-jupyter.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/cpu.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/cpu.Dockerfile index 323f89155b..857b5e2047 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/cpu.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/cpu.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel-jupyter.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel-jupyter.Dockerfile index 17c1d5b904..9ecaec38c2 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel-jupyter.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel-jupyter.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel.Dockerfile index bcac1f7015..c79bc3cf4c 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/gpu-devel.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/gpu-jupyter.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/gpu-jupyter.Dockerfile index cb0eb5a3d3..acfe4d8607 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/gpu-jupyter.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/gpu-jupyter.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/dockerfiles/gpu.Dockerfile b/tensorflow/tools/dockerfiles/dockerfiles/gpu.Dockerfile index e8e6ceafe2..f36a21eaf0 100644 --- a/tensorflow/tools/dockerfiles/dockerfiles/gpu.Dockerfile +++ b/tensorflow/tools/dockerfiles/dockerfiles/gpu.Dockerfile @@ -16,7 +16,7 @@ # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented -# throughout. Please refer to the the TensorFlow dockerfiles documentation +# throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. ARG UBUNTU_VERSION=16.04 diff --git a/tensorflow/tools/dockerfiles/spec.yml b/tensorflow/tools/dockerfiles/spec.yml index 5049e8dcfb..4826ddd8e2 100644 --- a/tensorflow/tools/dockerfiles/spec.yml +++ b/tensorflow/tools/dockerfiles/spec.yml @@ -17,7 +17,7 @@ header: | # THIS IS A GENERATED DOCKERFILE. # # This file was assembled from multiple pieces, whose use is documented - # throughout. Please refer to the the TensorFlow dockerfiles documentation + # throughout. Please refer to the TensorFlow dockerfiles documentation # for more information. # A combinatorial explosion of Docker images and Dockerfiles. diff --git a/tensorflow/tools/docs/parser.py b/tensorflow/tools/docs/parser.py index 83b4bf8128..6dc18ee8dc 100644 --- a/tensorflow/tools/docs/parser.py +++ b/tensorflow/tools/docs/parser.py @@ -39,7 +39,7 @@ def is_free_function(py_object, full_name, index): """Check if input is a free function (and not a class- or static method). Args: - py_object: The the object in question. + py_object: The object in question. full_name: The full name of the object, like `tf.module.symbol`. index: The {full_name:py_object} dictionary for the public API. -- GitLab From d2c1774b55da71d0998435663ffab88cb58afc08 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Sat, 1 Dec 2018 09:41:49 -0800 Subject: [PATCH 124/596] Add a file containing all files that will not be used outside the OSS repository. PiperOrigin-RevId: 223647591 --- tensorflow/opensource_only.files | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) create mode 100644 tensorflow/opensource_only.files diff --git a/tensorflow/opensource_only.files b/tensorflow/opensource_only.files new file mode 100644 index 0000000000..8ff1645b98 --- /dev/null +++ b/tensorflow/opensource_only.files @@ -0,0 +1,17 @@ +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/py3/BUILD +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/gcc-nvcc-cuda9.0/BUILD +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/gcc-nvcc/BUILD +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/cuda10.0-cudnn7/cuda/build_defs.bzl +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/cuda10.0-cudnn7/cuda/BUILD +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/gcc-nvcc-cuda10.0/BUILD +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/cuda9.0-cudnn7/cuda/build_defs.bzl +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/cuda9.0-cudnn7/cuda/BUILD +tensorflow/third_party/toolchains/preconfig/ubuntu14.04/nccl2/BUILD +tensorflow/third_party/toolchains/preconfig/generate/workspace.bzl +tensorflow/third_party/toolchains/preconfig/generate/containers.bzl +tensorflow/third_party/toolchains/preconfig/generate/generate.bzl +tensorflow/third_party/toolchains/preconfig/generate/BUILD +tensorflow/third_party/toolchains/preconfig/win_1803/bazel_018/BUILD +tensorflow/third_party/toolchains/preconfig/win_1803/bazel_018/dummy_toolchain.bzl +tensorflow/third_party/toolchains/preconfig/win_1803/py36/BUILD +tensorflow/third_party/toolchains/preconfig/win_1803/BUILD \ No newline at end of file -- GitLab From f1edf205b290c7fdeefe0b73d27c50a3cab67dcb Mon Sep 17 00:00:00 2001 From: Derek Murray Date: Sat, 1 Dec 2018 10:10:16 -0800 Subject: [PATCH 125/596] [tf.data] Add internal method for testing whether a dataset has a `Dataset.repeat(None)`. This utility function can be used to generate a warning when an indefinitely repeated `Dataset` is passed to an API that expects a finite `Dataset`. PiperOrigin-RevId: 223648690 --- .../python/data/experimental/__init__.py | 2 + .../data/experimental/kernel_tests/BUILD | 12 +++ .../has_indefinite_repeat_test.py | 53 ++++++++++++ tensorflow/python/data/experimental/ops/BUILD | 11 +++ .../experimental/ops/has_indefinite_repeat.py | 81 +++++++++++++++++++ .../v1/tensorflow.data.experimental.pbtxt | 4 + .../v2/tensorflow.data.experimental.pbtxt | 4 + 7 files changed, 167 insertions(+) create mode 100644 tensorflow/python/data/experimental/kernel_tests/has_indefinite_repeat_test.py create mode 100644 tensorflow/python/data/experimental/ops/has_indefinite_repeat.py diff --git a/tensorflow/python/data/experimental/__init__.py b/tensorflow/python/data/experimental/__init__.py index 8a1048513a..a72dd379e1 100644 --- a/tensorflow/python/data/experimental/__init__.py +++ b/tensorflow/python/data/experimental/__init__.py @@ -35,6 +35,7 @@ See [Importing Data](https://tensorflow.org/guide/datasets) for an overview. @@TFRecordWriter @@ThreadingOptions +@@assume_finite @@bucket_by_sequence_length @@choose_from_datasets @@copy_to_device @@ -82,6 +83,7 @@ from tensorflow.python.data.experimental.ops.grouping import bucket_by_sequence_ from tensorflow.python.data.experimental.ops.grouping import group_by_reducer from tensorflow.python.data.experimental.ops.grouping import group_by_window from tensorflow.python.data.experimental.ops.grouping import Reducer +from tensorflow.python.data.experimental.ops.has_indefinite_repeat import assume_finite from tensorflow.python.data.experimental.ops.interleave_ops import choose_from_datasets from tensorflow.python.data.experimental.ops.interleave_ops import parallel_interleave from tensorflow.python.data.experimental.ops.interleave_ops import sample_from_datasets diff --git a/tensorflow/python/data/experimental/kernel_tests/BUILD b/tensorflow/python/data/experimental/kernel_tests/BUILD index d7ca5a70e4..f2b2e2d0c1 100644 --- a/tensorflow/python/data/experimental/kernel_tests/BUILD +++ b/tensorflow/python/data/experimental/kernel_tests/BUILD @@ -211,6 +211,18 @@ py_test( ], ) +py_test( + name = "has_indefinite_repeat_test", + srcs = ["has_indefinite_repeat_test.py"], + srcs_version = "PY2AND3", + deps = [ + "//tensorflow/python/data/experimental/ops:has_indefinite_repeat", + "//tensorflow/python/data/kernel_tests:test_base", + "//tensorflow/python/data/ops:dataset_ops", + "@absl_py//absl/testing:parameterized", + ], +) + py_test( name = "ignore_errors_test", srcs = ["ignore_errors_test.py"], diff --git a/tensorflow/python/data/experimental/kernel_tests/has_indefinite_repeat_test.py b/tensorflow/python/data/experimental/kernel_tests/has_indefinite_repeat_test.py new file mode 100644 index 0000000000..6b2569c317 --- /dev/null +++ b/tensorflow/python/data/experimental/kernel_tests/has_indefinite_repeat_test.py @@ -0,0 +1,53 @@ +# 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. +# ============================================================================== +"""Tests for `tf.data.experimental.has_indefinite_repeat()`.""" +from __future__ import absolute_import +from __future__ import division +from __future__ import print_function + +from absl.testing import parameterized + +from tensorflow.python.data.experimental.ops import has_indefinite_repeat +from tensorflow.python.data.kernel_tests import test_base +from tensorflow.python.data.ops import dataset_ops +from tensorflow.python.platform import test + + +class HasIndefiniteRepeat(test_base.DatasetTestBase, parameterized.TestCase): + """Tests for `tf.data.experimental.has_indefinite_repeat()`.""" + + @parameterized.named_parameters( + ("NoRepeat", dataset_ops.Dataset.range(10), False), + ("FiniteRepeat", dataset_ops.Dataset.range(10).repeat(2), False), + ("FiniteRepeatNotAtEnd", dataset_ops.Dataset.range(10).repeat(2).skip(1), + False), + ("InfiniteRepeat", dataset_ops.Dataset.range(10).repeat(), True), + ("InfiniteRepeatNotAtEnd", dataset_ops.Dataset.range(10).repeat().skip(1), + True), + ("InfiniteRepeatThenFiniteRepeat", + dataset_ops.Dataset.range(10).repeat().repeat(2), True), + ("ConcatenateFiniteAndInfinite", + dataset_ops.Dataset.range(10).repeat(2).concatenate( + dataset_ops.Dataset.range(10).repeat()), True), + ("AssumeFinite", dataset_ops.Dataset.range(10).repeat().apply( + has_indefinite_repeat.assume_finite()), False), + ) + def testHasIndefiniteRepeat(self, dataset, expected_result): + self.assertEqual( + has_indefinite_repeat.has_indefinite_repeat(dataset), expected_result) + + +if __name__ == "__main__": + test.main() diff --git a/tensorflow/python/data/experimental/ops/BUILD b/tensorflow/python/data/experimental/ops/BUILD index 50f5127833..c4a4781aad 100644 --- a/tensorflow/python/data/experimental/ops/BUILD +++ b/tensorflow/python/data/experimental/ops/BUILD @@ -28,6 +28,16 @@ py_library( ], ) +py_library( + name = "has_indefinite_repeat", + srcs = ["has_indefinite_repeat.py"], + srcs_version = "PY2AND3", + deps = [ + "//tensorflow/python:tensor_util", + "//tensorflow/python/data/ops:dataset_ops", + ], +) + py_library( name = "iterator_ops", srcs = [ @@ -428,6 +438,7 @@ py_library( ":filter_for_shard_ops", ":get_single_element", ":grouping", + ":has_indefinite_repeat", ":indexed_dataset_ops", ":interleave_ops", ":map_defun", diff --git a/tensorflow/python/data/experimental/ops/has_indefinite_repeat.py b/tensorflow/python/data/experimental/ops/has_indefinite_repeat.py new file mode 100644 index 0000000000..cf5e3f811d --- /dev/null +++ b/tensorflow/python/data/experimental/ops/has_indefinite_repeat.py @@ -0,0 +1,81 @@ +# 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. +# ============================================================================== +"""Finiteness analysis of `Dataset` objects.""" +from __future__ import absolute_import +from __future__ import division +from __future__ import print_function + +from tensorflow.python.data.ops import dataset_ops +from tensorflow.python.framework import tensor_util +from tensorflow.python.util.tf_export import tf_export + + +class _AssumeFiniteDataset(dataset_ops.UnaryUnchangedStructureDataset): + + def __init__(self, input_dataset): + super(_AssumeFiniteDataset, self).__init__(input_dataset) + self._input_dataset = input_dataset + + def _as_variant_tensor(self): + return self._input_dataset._as_variant_tensor() # pylint: disable=protected-access + + +@tf_export("data.experimental.assume_finite") +def assume_finite(): + """Assume that the input is finite, even if it contains `Dataset.repeat()`. + + Training libraries may analyze a `tf.data.Dataset` to determine if it is + finite or infinite (e.g. because it contains an indefinite + `tf.data.Dataset.repeat` transformation). Since that analysis may be + imprecise, this transformation allows the user to annotate a dataset + explicitly as being finite. + + Returns: + A `Dataset` transformation function, which can be passed to + `tf.data.Dataset.apply`. + """ + + def _apply_fn(dataset): + return _AssumeFiniteDataset(dataset) + + return _apply_fn + + +def has_indefinite_repeat(dataset): + """Returns `True` if `dataset` or any of its inputs is `Dataset.repeat()`. + + NOTE: For simplicity, this analysis does not attempt to analyze nested + datasets (e.g. in a function passed to `tf.data.Dataset.flat_map`). If the + analysis is incorrect, you can apply `tf.data.experimental.assume_finite()` + to the dataset to override it. + + Args: + dataset: A `tf.data.Dataset`. + + Returns: + `True` if `dataset` or any of its inputs is repeated indefinitely. + """ + # pylint: disable=protected-access + if isinstance(dataset, dataset_ops.DatasetV1Adapter): + return has_indefinite_repeat(dataset._dataset) + elif isinstance(dataset, dataset_ops.RepeatDataset): + count = tensor_util.constant_value(dataset._count) + return count == -1 or has_indefinite_repeat(dataset._inputs()[0]) + elif isinstance(dataset, _AssumeFiniteDataset): + return False + else: + return any( + has_indefinite_repeat(input_dataset) + for input_dataset in dataset._inputs()) diff --git a/tensorflow/tools/api/golden/v1/tensorflow.data.experimental.pbtxt b/tensorflow/tools/api/golden/v1/tensorflow.data.experimental.pbtxt index ad10b82283..27ea063b66 100644 --- a/tensorflow/tools/api/golden/v1/tensorflow.data.experimental.pbtxt +++ b/tensorflow/tools/api/golden/v1/tensorflow.data.experimental.pbtxt @@ -52,6 +52,10 @@ tf_module { name: "Counter" argspec: "args=[\'start\', \'step\', \'dtype\'], varargs=None, keywords=None, defaults=[\'0\', \'1\', \"\"], " } + member_method { + name: "assume_finite" + argspec: "args=[], varargs=None, keywords=None, defaults=None" + } member_method { name: "bucket_by_sequence_length" argspec: "args=[\'element_length_func\', \'bucket_boundaries\', \'bucket_batch_sizes\', \'padded_shapes\', \'padding_values\', \'pad_to_bucket_boundary\', \'no_padding\'], varargs=None, keywords=None, defaults=[\'None\', \'None\', \'False\', \'False\'], " diff --git a/tensorflow/tools/api/golden/v2/tensorflow.data.experimental.pbtxt b/tensorflow/tools/api/golden/v2/tensorflow.data.experimental.pbtxt index ad10b82283..27ea063b66 100644 --- a/tensorflow/tools/api/golden/v2/tensorflow.data.experimental.pbtxt +++ b/tensorflow/tools/api/golden/v2/tensorflow.data.experimental.pbtxt @@ -52,6 +52,10 @@ tf_module { name: "Counter" argspec: "args=[\'start\', \'step\', \'dtype\'], varargs=None, keywords=None, defaults=[\'0\', \'1\', \"\"], " } + member_method { + name: "assume_finite" + argspec: "args=[], varargs=None, keywords=None, defaults=None" + } member_method { name: "bucket_by_sequence_length" argspec: "args=[\'element_length_func\', \'bucket_boundaries\', \'bucket_batch_sizes\', \'padded_shapes\', \'padding_values\', \'pad_to_bucket_boundary\', \'no_padding\'], varargs=None, keywords=None, defaults=[\'None\', \'None\', \'False\', \'False\'], " -- GitLab From a4f06947c61ef6f69057f1207c2532a3551ff4f0 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Sat, 1 Dec 2018 11:54:43 -0800 Subject: [PATCH 126/596] Make grouped convolutions use additional spatial dimension instead of a reduce window. PiperOrigin-RevId: 223651863 --- .../convolution_feature_group_converter.cc | 154 ++++------- tensorflow/compiler/xla/tests/BUILD | 26 ++ .../xla/tests/grouped_convolution_test.cc | 245 ++++++++++++++++++ 3 files changed, 319 insertions(+), 106 deletions(-) create mode 100644 tensorflow/compiler/xla/tests/grouped_convolution_test.cc diff --git a/tensorflow/compiler/xla/service/convolution_feature_group_converter.cc b/tensorflow/compiler/xla/service/convolution_feature_group_converter.cc index 09c3f32860..07d6680a72 100644 --- a/tensorflow/compiler/xla/service/convolution_feature_group_converter.cc +++ b/tensorflow/compiler/xla/service/convolution_feature_group_converter.cc @@ -271,130 +271,72 @@ Status ConvolutionVisitor::HandleConvolution(HloInstruction* convolution) { convolution, std::move(new_convolution))); } else { int64 activation_input_feature_dim = dim_numbers.input_feature_dimension(); - auto activation = convolution->mutable_operand(0); int64 output_feature = filter->shape().dimensions(kernel_output_feature_dim); - int64 input_feature = - activation->shape().dimensions(activation_input_feature_dim); - // If group_count == output_feature, then we map those grouped convolutions - // onto depthwise convolution + reduce. E.g., we would turn + // onto depthwise convolution. This is done by adding an additional spatial + // dimension to the activations, kernel, and the output. + // E.g., we would turn // [2, 12]{B, IF} conv [3, 4]{IF, OF} into - // [2, 12]{B, IF} depth conv [1, 12]{IF, OF}, and then use a reduce window - // of {1, 3} on the generated [2, 12] output to produce the final result of - // [2, 4]. + // [3, 2, 4]{S, B, IF} depth conv [3, 1, 4]{S, IF, OF}, where S is the + // additional spatial dimension. The generated convolution output will be + // [1, 2, 4]{S, B, OF} and then reshape the output back to [2, 4] {B, OF}. + if (group_count == output_feature && !filter_expansion_) { - Shape reshaped_filter_shape = filter->shape(); + auto filter = convolution->mutable_operand(1); + auto activation = convolution->mutable_operand(0); - if (kernel_input_feature_dim < kernel_output_feature_dim) { - // Transpose IF and OF on the kernel. - std::vector filter_dims; - for (int64 i = 0; i < dim_numbers.kernel_spatial_dimensions().size(); - ++i) { - filter_dims.push_back(dim_numbers.kernel_spatial_dimensions(i)); - } - filter_dims.push_back(kernel_output_feature_dim); - filter_dims.push_back(kernel_input_feature_dim); - - Shape transposed_filter = filter->shape(); - auto& dimensions = *transposed_filter.mutable_dimensions(); - std::swap(dimensions[kernel_input_feature_dim], - dimensions[kernel_output_feature_dim]); - - filter = add(HloInstruction::CreateTranspose(transposed_filter, filter, - filter_dims)); - } else { - // For depthwise convolutions, we want the kernel input feature - // dimension to be smaller than the output feature dimension. If that's - // not the case, we swap the dimensions. + // Add spatial dimension to the activation, and reshape. + Shape reshaped_activation_shape = activation->shape(); + ShapeUtil::AppendMajorDimension(group_size, &reshaped_activation_shape); - auto& dimensions = *reshaped_filter_shape.mutable_dimensions(); - std::swap(dimensions[kernel_input_feature_dim], - dimensions[kernel_output_feature_dim]); + int64 new_spatial_dim = reshaped_activation_shape.dimensions().size() - 1; - dim_numbers.set_kernel_input_feature_dimension( - kernel_output_feature_dim); + reshaped_activation_shape.set_dimensions(activation_input_feature_dim, + group_count); + activation = add( + HloInstruction::CreateReshape(reshaped_activation_shape, activation)); - dim_numbers.set_kernel_output_feature_dimension( - kernel_input_feature_dim); - std::swap(kernel_output_feature_dim, kernel_input_feature_dim); - } + // Add spatial dimension to the filter, and reshape. + Shape reshaped_filter_shape = filter->shape(); + ShapeUtil::AppendMajorDimension(1, &reshaped_filter_shape); - reshaped_filter_shape.set_dimensions(kernel_input_feature_dim, 1); - reshaped_filter_shape.set_dimensions(kernel_output_feature_dim, - group_count * group_size); - auto reshaped_filter = + filter = add(HloInstruction::CreateReshape(reshaped_filter_shape, filter)); - Shape reshaped_convolution_shape = convolution->shape(); - reshaped_convolution_shape.set_dimensions( - dim_numbers.output_feature_dimension(), group_count * group_size); - auto new_convolution = add(HloInstruction::CreateConvolve( - reshaped_convolution_shape, convolution->mutable_operand(0), - reshaped_filter, /*feature_group_count=*/input_feature, - convolution->window(), dim_numbers, convolution->precision_config())); - - // Create the reduce window. - Window window; - for (int64 i = 0; i < new_convolution->shape().dimensions_size(); ++i) { - auto* dim = window.add_dimensions(); - dim->set_padding_low(0); - dim->set_padding_high(0); - dim->set_window_dilation(1); - dim->set_base_dilation(1); - if (i == dim_numbers.output_feature_dimension()) { - dim->set_stride(group_size); - dim->set_size(group_size); - } else { - dim->set_stride(1); - dim->set_size(1); - } - } + Shape new_output_shape = convolution->shape(); + ShapeUtil::AppendMajorDimension(1, &new_output_shape); + + // Edit convolution dimension numbers. Note that kernel_input_feature_dim + // now becomes a spatial dimension, and the newly added dimension of size + // 1 is the new kernel_input_feature_dim. + dim_numbers.add_input_spatial_dimensions(new_spatial_dim); + dim_numbers.add_kernel_spatial_dimensions(kernel_input_feature_dim); + dim_numbers.set_kernel_input_feature_dimension(new_spatial_dim); + dim_numbers.add_output_spatial_dimensions(new_spatial_dim); + + // Add window for the new spatial dimension. + Window new_window = convolution->window(); + auto* dim = new_window.add_dimensions(); + dim->set_window_dilation(1); + dim->set_base_dilation(1); + dim->set_stride(1); + dim->set_size(group_size); - auto reduce_window_shape = new_convolution->shape(); - reduce_window_shape.set_dimensions(dim_numbers.output_feature_dimension(), - group_count); - - auto zero_literal = LiteralUtil::CreateR0(0.0f); - TF_ASSIGN_OR_RETURN(zero_literal, zero_literal.Convert(F32)); - auto zero = add(HloInstruction::CreateConstant(std::move(zero_literal))); - - auto reduce_function = [&]() -> HloComputation* { - HloComputation::Builder b("add_computation"); - Shape shape = ShapeUtil::MakeShape(F32, {}); - auto lhs = - b.AddInstruction(HloInstruction::CreateParameter(0, shape, "lhs")); - auto rhs = - b.AddInstruction(HloInstruction::CreateParameter(1, shape, "rhs")); - auto scalar_op = b.AddInstruction( - HloInstruction::CreateBinary(shape, HloOpcode::kAdd, lhs, rhs)); - return computation_->parent()->AddEmbeddedComputation( - b.Build(scalar_op)); - }; - - // Ensure that data input to reduce window is of type F32. - if (primitive_util::BitWidth(new_convolution->shape().element_type()) < - primitive_util::BitWidth(F32)) { - Shape convert_shape = new_convolution->shape(); - convert_shape.set_element_type(F32); - new_convolution = add(HloInstruction::CreateBitcastConvert( - convert_shape, new_convolution)); - } - - auto reduce_window = add(HloInstruction::CreateReduceWindow( - reduce_window_shape, new_convolution, zero, window, - reduce_function())); + auto new_convolution = add(HloInstruction::CreateConvolve( + new_output_shape, activation, filter, group_count, new_window, + dim_numbers, convolution->precision_config())); - Shape convert_back_shape = reduce_window->shape(); - convert_back_shape.set_element_type(activation->shape().element_type()); + // Delete the extra spatial dimension, and reshape. + Shape reshaped_convolution_shape = + ShapeUtil::DeleteDimension(new_spatial_dim, new_convolution->shape()); + auto reshaped_convolution = HloInstruction::CreateReshape( + reshaped_convolution_shape, new_convolution); - // Convert reduced data back to the original data type. - auto reduce_window_converted = HloInstruction::CreateBitcastConvert( - convert_back_shape, reduce_window); TF_RETURN_IF_ERROR(computation_->ReplaceWithNewInstruction( - convolution, std::move(reduce_window_converted))); + convolution, std::move(reshaped_convolution))); } else { // The filter expansion mechanism adds zeroes in the kernel. diff --git a/tensorflow/compiler/xla/tests/BUILD b/tensorflow/compiler/xla/tests/BUILD index 2c18e2fd10..48bbc4e472 100644 --- a/tensorflow/compiler/xla/tests/BUILD +++ b/tensorflow/compiler/xla/tests/BUILD @@ -324,6 +324,32 @@ xla_test( ], ) +xla_test( + name = "grouped_convolution_test", + timeout = "long", + srcs = ["grouped_convolution_test.cc"], + blacklisted_backends = [ + # disabled because of a break b/119590850. + "gpu", + # disabled because it times out. + "cpu", + ], + shard_count = 50, + deps = [ + "//tensorflow/compiler/xla:execution_options_util", + "//tensorflow/compiler/xla:status_macros", + "//tensorflow/compiler/xla:test", + "//tensorflow/compiler/xla/client:xla_computation", + "//tensorflow/compiler/xla/service:bfloat16_normalization", + "//tensorflow/compiler/xla/service:despecializer", + "//tensorflow/compiler/xla/service:hlo_parser", + "//tensorflow/compiler/xla/tests:client_library_test_base", + "//tensorflow/compiler/xla/tests:hlo_test_base", + "//tensorflow/compiler/xla/tests:xla_internal_test_main", + "@com_google_absl//absl/types:optional", + ], +) + xla_test( name = "check_execution_arity_test", srcs = ["check_execution_arity_test.cc"], diff --git a/tensorflow/compiler/xla/tests/grouped_convolution_test.cc b/tensorflow/compiler/xla/tests/grouped_convolution_test.cc new file mode 100644 index 0000000000..8f7049910e --- /dev/null +++ b/tensorflow/compiler/xla/tests/grouped_convolution_test.cc @@ -0,0 +1,245 @@ +/* Copyright 2017 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 "absl/types/optional.h" +#include "tensorflow/compiler/xla/client/xla_computation.h" +#include "tensorflow/compiler/xla/execution_options_util.h" +#include "tensorflow/compiler/xla/service/bfloat16_normalization.h" +#include "tensorflow/compiler/xla/service/despecializer.h" +#include "tensorflow/compiler/xla/service/hlo_parser.h" +#include "tensorflow/compiler/xla/status_macros.h" +#include "tensorflow/compiler/xla/test.h" +#include "tensorflow/compiler/xla/tests/client_library_test_base.h" +#include "tensorflow/compiler/xla/tests/hlo_test_base.h" +#include "tensorflow/compiler/xla/tests/test_macros.h" + +namespace xla { +namespace { + +string GetFloatDataType(bool use_bfloat16) { + return use_bfloat16 ? "bf16" : "f32"; +} + +struct GroupedConvolution2DSpec { + int64 input_feature, output_feature, window, stride, pad, lhs_dilate; + int64 group_size, group_count; + std::vector activation_dims; + std::vector activation_layout; + std::vector kernel_dims; + std::vector kernel_layout; + std::vector output_dims; + std::vector output_layout; +}; + +class GroupedConvolution2DTest + : public HloTestBase, + public ::testing::WithParamInterface< + ::testing::tuple> {}; + +static std::vector GetConv2DTestCases() { + std::vector config_set; + // Add to this set if you want a new test configuration. + // Rule : the penultimate number must be divisible by the last number. + std::vector> config_options = {{8, 2, 2, 1, 1024, 128}, + {512, 3, 3, 144, 1024, 16}, + {256, 3, 3, 129, 512, 64}, + {64, 1, 2, 127, 32, 8}, + {256, 3, 3, 256, 1024, 4}}; + + for (auto option : config_options) { + int64 output_feature = option[0]; + int64 activation_size = option[1]; + int64 kernel_size = option[2]; + int64 batch = option[3]; + int64 input_feature = option[4]; + int64 group_size = option[5]; + + std::vector kernel_layout = {3, 2, 1, 0}; + GroupedConvolution2DSpec config; + config.group_size = group_size; + config.group_count = input_feature / group_size; + config.output_feature = output_feature; + config.window = kernel_size; + + config.activation_dims = {batch, activation_size, activation_size, + input_feature}; + config.activation_layout = {3, 0, 2, 1}; + + config.kernel_dims = {kernel_size, kernel_size, group_size, output_feature}; + config.kernel_layout = {3, 2, 1, 0}; + + if (activation_size == 1 && kernel_size == 2) { + // Test for outer dim. + config.output_dims = {batch, activation_size + kernel_size - 1, + activation_size + kernel_size, output_feature}; + } else if (output_feature == 256) { + // Restrict dilation-based tests only to one feature configuration. + config.stride = activation_size - 1; + config.pad = 0; + config.lhs_dilate = output_feature / 32; + config.output_dims = {batch, output_feature / 32, + activation_size - kernel_size + 1, output_feature}; + } else { + config.stride = config.pad = config.lhs_dilate = -1; + config.output_dims = {batch, activation_size - kernel_size + 1, + activation_size - kernel_size + 1, output_feature}; + } + + // Try this layout for all kernel shapes. + config.output_layout = {3, 0, 2, 1}; + config_set.push_back(config); + + // Try other layouts only for certain kernel shapes. + if (kernel_size % 2 == 0) { + config.activation_layout = {0, 3, 2, 1}; + config_set.push_back(config); + + config.output_layout = {0, 3, 2, 1}; + config_set.push_back(config); + + config.activation_layout = {3, 0, 2, 1}; + config_set.push_back(config); + } + } + + return config_set; +} + +string GroupedConvolution2DTestDataToString( + const ::testing::TestParamInfo< + ::testing::tuple>& data) { + const auto& spec = ::testing::get<0>(data.param); + const string data_type = GetFloatDataType(::testing::get<1>(data.param)); + string str = absl::StrCat( + "activation_dims_", absl::StrJoin(spec.activation_dims, "x"), + "_activation_layout_", absl::StrJoin(spec.activation_layout, "_"), + "_kernel_dims_", absl::StrJoin(spec.kernel_dims, "x"), "_kernel_layout_", + absl::StrJoin(spec.kernel_layout, "_"), "_output_dims_", + absl::StrJoin(spec.output_dims, "x"), "_output_layout_", + absl::StrJoin(spec.output_layout, "_"), data_type); + // -1 indicates non-existence. + if (spec.stride != -1) { + absl::StrAppend(&str, "_lhs_dilation_", spec.lhs_dilate, "x1"); + } + + // Test names are not allowed to contain the '-' character. + absl::c_replace(str, '-', 'n'); + return str; +} + +string BuildHloTextGroupedConvolution2D(const GroupedConvolution2DSpec& spec, + bool use_bfloat16) { + const string data_type = GetFloatDataType(use_bfloat16); + if (spec.activation_dims[1] == 1 && spec.kernel_dims[1] == 2) { + // Check for outer dim. + return absl::StrFormat( + R"( + HloModule TensorFlowDepthwiseConv + + ENTRY main { + activation = %s[%s]{%s} parameter(0) + kernel = %s[%s]{%s} parameter(1) + ROOT conv = %s[%s]{%s} convolution(%s[%s]{%s} activation, %s[%s]{%s} kernel), + window={size=%dx%d pad=1_1x%d_%d rhs_dilate=1x%d}, dim_labels=b01f_01io->b01f, + feature_group_count=%d + } + )", + data_type, absl::StrJoin(spec.activation_dims, ","), + absl::StrJoin(spec.activation_layout, ","), data_type, + absl::StrJoin(spec.kernel_dims, ","), + absl::StrJoin(spec.kernel_layout, ","), data_type, + absl::StrJoin(spec.output_dims, ","), + absl::StrJoin(spec.output_layout, ","), data_type, + absl::StrJoin(spec.activation_dims, ","), + absl::StrJoin(spec.activation_layout, ","), data_type, + absl::StrJoin(spec.kernel_dims, ","), + absl::StrJoin(spec.kernel_layout, ","), spec.window, spec.window, + spec.window, spec.window, spec.window, spec.group_count); + + } else if (spec.stride == -1) { + // Check for basic, non-dilated cases. + return absl::StrFormat( + R"( + HloModule TensorFlowDepthwiseConv + + ENTRY main { + activation = %s[%s]{%s} parameter(0) + kernel = %s[%s]{%s} parameter(1) + ROOT conv = %s[%s]{%s} convolution(%s[%s]{%s} activation, %s[%s]{%s} kernel), + window={size=%dx%d}, dim_labels=b01f_01io->b01f, + feature_group_count=%d + } + )", + data_type, absl::StrJoin(spec.activation_dims, ","), + absl::StrJoin(spec.activation_layout, ","), data_type, + absl::StrJoin(spec.kernel_dims, ","), + absl::StrJoin(spec.kernel_layout, ","), data_type, + absl::StrJoin(spec.output_dims, ","), + absl::StrJoin(spec.output_layout, ","), data_type, + absl::StrJoin(spec.activation_dims, ","), + absl::StrJoin(spec.activation_layout, ","), data_type, + absl::StrJoin(spec.kernel_dims, ","), + absl::StrJoin(spec.kernel_layout, ","), spec.window, spec.window, + spec.group_count); + } else { + // Check for base dilations. + return absl::StrFormat( + R"( + HloModule TensorFlowDepthwiseConv + + ENTRY main { + activation = %s[%s]{%s} parameter(0) + kernel = %s[%s]{%s} parameter(1) + ROOT conv = %s[%s]{%s} convolution(%s[%s]{%s} activation, %s[%s]{%s} kernel), + window={size=%dx%d stride=%dx1 pad=%d_%dx0_0 lhs_dilate=%dx1}, + dim_labels=b01f_01io->b01f, feature_group_count=%d + } + )", + data_type, absl::StrJoin(spec.activation_dims, ","), + absl::StrJoin(spec.activation_layout, ","), data_type, + absl::StrJoin(spec.kernel_dims, ","), + absl::StrJoin(spec.kernel_layout, ","), data_type, + absl::StrJoin(spec.output_dims, ","), + absl::StrJoin(spec.output_layout, ","), data_type, + absl::StrJoin(spec.activation_dims, ","), + absl::StrJoin(spec.activation_layout, ","), data_type, + absl::StrJoin(spec.kernel_dims, ","), + absl::StrJoin(spec.kernel_layout, ","), spec.window, spec.window, + spec.stride, 0, 0, spec.lhs_dilate, spec.group_count); + } +} + +XLA_TEST_P(GroupedConvolution2DTest, DoIt) { + const GroupedConvolution2DSpec& spec = ::testing::get<0>(GetParam()); + bool use_bfloat16 = ::testing::get<1>(GetParam()); + const string hlo_text = BuildHloTextGroupedConvolution2D(spec, use_bfloat16); + + EXPECT_TRUE(RunAndCompare(hlo_text, ErrorSpec{0.01, 0.01}, + [](HloModule* module) -> Status { + BFloat16MixedPrecisionRemoval remover; + TF_RETURN_IF_ERROR(remover.Run(module).status()); + Despecializer despecializer; + return despecializer.Run(module).status(); + })); +} + +INSTANTIATE_TEST_CASE_P( + GroupedConvolution2DTestWithRandomIndices, GroupedConvolution2DTest, + ::testing::Combine(::testing::ValuesIn(GetConv2DTestCases()), + ::testing::Bool()), + GroupedConvolution2DTestDataToString); + +} // namespace +} // namespace xla -- GitLab From 8584f21392772170b007ee3b3fbfed17fe19e32f Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 1 Dec 2018 20:58:09 +0000 Subject: [PATCH 127/596] Fix SparseDenseCwise's broadcasting issue This fix tries to address the issue raised in 24072. In `sparse_dense_cwise_mul/add` operations the broadcasting only support dense to sparse, though the validation was not captured. This fix fixes the validation in SparseDenseBinaryOpShared so that error could be thrown correctly. This fix fixes 24072. Signed-off-by: Yong Tang --- tensorflow/core/kernels/sparse_dense_binary_op_shared.cc | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc b/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc index ac48202ada..3a6b66302f 100644 --- a/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc +++ b/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc @@ -88,12 +88,11 @@ class SparseDenseBinaryOpShared : public OpKernel { const auto rhs_dims = BCast::FromShape(dense_t->shape()); BCast b(lhs_dims, rhs_dims, false); // false for keeping the same num dims. - // True iff (size(lhs) > size(rhs)), or (sizes equal, lhs cwise rhs). + // True iff (size(lhs) >= size(rhs)) and all dims in lhs is smaller or equal to dims in rhs (from right to left). auto VecGreaterEq = [](ArraySlice lhs, ArraySlice rhs) { - if (lhs.size() > rhs.size()) return true; if (lhs.size() < rhs.size()) return false; - for (size_t i = 0; i < lhs.size(); ++i) { - if (lhs[i] < rhs[i]) return false; + for (size_t i = 0; i < rhs.size(); ++i) { + if (lhs[lhs.size() - 1 - i] < rhs[rhs.size() - 1 - i]) return false; } return true; }; -- GitLab From a7cd4dbea9f276160ebadf82178e77ae5c8d557e Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 1 Dec 2018 21:02:04 +0000 Subject: [PATCH 128/596] Add test case for sparse_dense_cwise shape validation. Signed-off-by: Yong Tang --- tensorflow/python/kernel_tests/sparse_ops_test.py | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/tensorflow/python/kernel_tests/sparse_ops_test.py b/tensorflow/python/kernel_tests/sparse_ops_test.py index 75f65e6251..f58832a89e 100644 --- a/tensorflow/python/kernel_tests/sparse_ops_test.py +++ b/tensorflow/python/kernel_tests/sparse_ops_test.py @@ -22,6 +22,7 @@ import numpy as np from tensorflow.python.framework import constant_op from tensorflow.python.framework import dtypes +from tensorflow.python.framework import errors from tensorflow.python.framework import ops from tensorflow.python.framework import sparse_tensor from tensorflow.python.framework import test_util @@ -798,6 +799,17 @@ class SparseMathOpsTest(test_util.TensorFlowTestCase): result_tensor.values).eval() self.assertAllEqual(result_np, res_densified) + @test_util.run_deprecated_v1 + def testCwiseShapeValidation(self): + # Test case for GitHub 24072. + with self.session(use_gpu=False): + a = array_ops.ones([3, 4, 1], dtype=dtypes.int32) + b = sparse_tensor.SparseTensor([[0, 0, 1, 0], [0, 0, 3, 0]], [10, 20], [1, 1, 4, 2]) + c = a * b + with self.assertRaisesRegexp(errors.InvalidArgumentError, + "broadcasts dense to sparse only; got incompatible shapes"): + c.eval() + @test_util.run_deprecated_v1 def testCwiseDivAndMul(self): np.random.seed(1618) -- GitLab From 676c6ea316f4dda962d8b2e29855c040020533f5 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 1 Dec 2018 21:04:41 +0000 Subject: [PATCH 129/596] Pylint fix Signed-off-by: Yong Tang --- tensorflow/python/kernel_tests/sparse_ops_test.py | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/tensorflow/python/kernel_tests/sparse_ops_test.py b/tensorflow/python/kernel_tests/sparse_ops_test.py index f58832a89e..50b6239185 100644 --- a/tensorflow/python/kernel_tests/sparse_ops_test.py +++ b/tensorflow/python/kernel_tests/sparse_ops_test.py @@ -804,10 +804,12 @@ class SparseMathOpsTest(test_util.TensorFlowTestCase): # Test case for GitHub 24072. with self.session(use_gpu=False): a = array_ops.ones([3, 4, 1], dtype=dtypes.int32) - b = sparse_tensor.SparseTensor([[0, 0, 1, 0], [0, 0, 3, 0]], [10, 20], [1, 1, 4, 2]) + b = sparse_tensor.SparseTensor( + [[0, 0, 1, 0], [0, 0, 3, 0]], [10, 20], [1, 1, 4, 2]) c = a * b - with self.assertRaisesRegexp(errors.InvalidArgumentError, - "broadcasts dense to sparse only; got incompatible shapes"): + with self.assertRaisesRegexp( + errors.InvalidArgumentError, + "broadcasts dense to sparse only; got incompatible shapes"): c.eval() @test_util.run_deprecated_v1 -- GitLab From 28e034bd293db09a0ecc707f71f65fe6c5dd2943 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sat, 1 Dec 2018 21:28:19 +0000 Subject: [PATCH 130/596] Fix `Experimental clang-format Check` Signed-off-by: Yong Tang --- tensorflow/core/kernels/sparse_dense_binary_op_shared.cc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc b/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc index 3a6b66302f..d7460363fc 100644 --- a/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc +++ b/tensorflow/core/kernels/sparse_dense_binary_op_shared.cc @@ -88,7 +88,8 @@ class SparseDenseBinaryOpShared : public OpKernel { const auto rhs_dims = BCast::FromShape(dense_t->shape()); BCast b(lhs_dims, rhs_dims, false); // false for keeping the same num dims. - // True iff (size(lhs) >= size(rhs)) and all dims in lhs is smaller or equal to dims in rhs (from right to left). + // True iff (size(lhs) >= size(rhs)) and all dims in lhs is smaller or equal + // to dims in rhs (from right to left). auto VecGreaterEq = [](ArraySlice lhs, ArraySlice rhs) { if (lhs.size() < rhs.size()) return false; for (size_t i = 0; i < rhs.size(); ++i) { -- GitLab From cc8e28b72b52718db5f22830e1d529d8e077b537 Mon Sep 17 00:00:00 2001 From: Yong Tang Date: Sun, 2 Dec 2018 01:55:11 +0000 Subject: [PATCH 131/596] Fix broken link in lite api docs This fix fixes broken link in lite apis docs Signed-off-by: Yong Tang --- tensorflow/lite/g3doc/apis.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/lite/g3doc/apis.md b/tensorflow/lite/g3doc/apis.md index e9fa24bff1..60ce11d685 100644 --- a/tensorflow/lite/g3doc/apis.md +++ b/tensorflow/lite/g3doc/apis.md @@ -347,7 +347,7 @@ interpreter.runForMultipleInputsOutputs(inputs, map_of_indices_to_outputs); where each entry in `inputs` corresponds to an input tensor and `map_of_indices_to_outputs` maps indices of output tensors to the corresponding output data. In both cases the tensor indices should correspond to -the values given to the [TensorFlow Lite Optimized Converter](https://github.com/tensorflow/tensorflow/tree/master/tensorflow/lite/toco/g3doc/cmdline_examples.md) +the values given to the [TensorFlow Lite Optimized Converter](convert/cmdline_examples.md) when the model was created. Be aware that the order of tensors in `input` must match the order given to the `TensorFlow Lite Optimized Converter`. -- GitLab From c297ff16bc798fb78e06460d00bcef52853601fd Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Sun, 2 Dec 2018 01:02:28 -0800 Subject: [PATCH 132/596] compat: Update forward compatibility horizon to 2018-12-02 PiperOrigin-RevId: 223684627 --- tensorflow/python/compat/compat.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/compat/compat.py b/tensorflow/python/compat/compat.py index a7e8bef670..8ec90ee759 100644 --- a/tensorflow/python/compat/compat.py +++ b/tensorflow/python/compat/compat.py @@ -32,7 +32,7 @@ from tensorflow.python.ops import variable_scope from tensorflow.python.util import tf_contextlib from tensorflow.python.util.tf_export import tf_export -_FORWARD_COMPATIBILITY_HORIZON = datetime.date(2018, 12, 1) +_FORWARD_COMPATIBILITY_HORIZON = datetime.date(2018, 12, 2) @tf_export("compat.forward_compatible") -- GitLab From 3188c2a38b9327e58940e7040a4ba7c72c014718 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Sun, 2 Dec 2018 03:42:28 -0800 Subject: [PATCH 133/596] Fixing LD_LIBRARY_PATH env variable. PiperOrigin-RevId: 223692658 --- .../tools/ci_build/Dockerfile.rbe.cuda10.0-cudnn7-ubuntu14.04 | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/tools/ci_build/Dockerfile.rbe.cuda10.0-cudnn7-ubuntu14.04 b/tensorflow/tools/ci_build/Dockerfile.rbe.cuda10.0-cudnn7-ubuntu14.04 index 85b9d94313..03de89b717 100644 --- a/tensorflow/tools/ci_build/Dockerfile.rbe.cuda10.0-cudnn7-ubuntu14.04 +++ b/tensorflow/tools/ci_build/Dockerfile.rbe.cuda10.0-cudnn7-ubuntu14.04 @@ -31,7 +31,7 @@ ENV PATH /usr/local/cuda/bin:${PATH} # -Wl,-rpath-link=/usr/local/cuda/lib64/stubs to all binaries transitively # depending on libcuda. Optimally, builds targeting cuda would do that # internally. -ENV LIBRARY_PATH /usr/local/cuda/lib64/stubs +ENV LD_LIBRARY_PATH /usr/local/cuda/lib64/stubs LABEL com.nvidia.cudnn.version="${CUDNN_VERSION}" -- GitLab From b18da39f10f998a21bfcda43184ba3a80243dde4 Mon Sep 17 00:00:00 2001 From: Serge Panev Date: Sun, 2 Dec 2018 19:24:56 +0100 Subject: [PATCH 134/596] Fix with clang-format-3.6.0 --- tensorflow/core/util/tensor_format.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/util/tensor_format.h b/tensorflow/core/util/tensor_format.h index 68674cb282..a296fb447e 100644 --- a/tensorflow/core/util/tensor_format.h +++ b/tensorflow/core/util/tensor_format.h @@ -498,7 +498,8 @@ inline TensorShape ShapeFromFormat(TensorFormat format, int64 N, dim_sizes[GetTensorBatchDimIndex(dims, format)] = N; for (int dim = 0; static_cast(dim) < spatial.size(); dim++) { auto dim_size = spatial[dim]; - if (format == FORMAT_NHWC_VECT_W && static_cast(dim) == spatial.size() - 1) { + if (format == FORMAT_NHWC_VECT_W && + static_cast(dim) == spatial.size() - 1) { CHECK_EQ(0, dim_size % 4) << "FORMAT_NHWC_VECT_W requires W to be a multiple of 4, but W=" << dim_size; -- GitLab From bf70c447d68070e81c1dbc2ebf0b6b0d8ae2c597 Mon Sep 17 00:00:00 2001 From: Jiri Simsa Date: Sun, 2 Dec 2018 11:02:30 -0800 Subject: [PATCH 135/596] [tf.data] Consolidating experimental C++ kernels. PiperOrigin-RevId: 223710135 --- tensorflow/c/c_api_experimental.cc | 26 +- tensorflow/contrib/data/python/ops/sliding.py | 4 +- ...perimentalBytesProducedStatsDataset.pbtxt} | 3 +- ...i_def_ExperimentalDatasetToTFRecord.pbtxt} | 2 +- ...perimentalDenseToSparseBatchDataset.pbtxt} | 3 +- ...f_ExperimentalGroupByReducerDataset.pbtxt} | 2 +- ...ef_ExperimentalGroupByWindowDataset.pbtxt} | 3 +- ...def_ExperimentalLatencyStatsDataset.pbtxt} | 3 +- ..._def_ExperimentalMapAndBatchDataset.pbtxt} | 2 +- ...perimentalParallelInterleaveDataset.pbtxt} | 3 +- ...def_ExperimentalParseExampleDataset.pbtxt} | 3 +- ...xperimentalSetStatsAggregatorDataset.pbtxt | 4 + ...ef_ExperimentalSlidingWindowDataset.pbtxt} | 3 +- ...t => api_def_ExperimentalSqlDataset.pbtxt} | 3 +- ...ef_ExperimentalStatsAggregatorHandle.pbtxt | 5 + ..._ExperimentalStatsAggregatorSummary.pbtxt} | 3 +- ... api_def_ExperimentalUnbatchDataset.pbtxt} | 3 +- .../base_api/api_def_MapAndBatchDataset.pbtxt | 53 - .../api_def_SetStatsAggregatorDataset.pbtxt | 3 - .../api_def_StatsAggregatorHandle.pbtxt | 4 - .../api_def_BytesProducedStatsDataset.pbtxt | 4 - .../api_def_DenseToSparseBatchDataset.pbtxt | 4 - .../api_def_GroupByWindowDataset.pbtxt | 4 - .../api_def_LatencyStatsDataset.pbtxt | 4 - .../api_def_MapAndBatchDataset.pbtxt | 4 - .../api_def_ParallelInterleaveDataset.pbtxt | 4 - .../api_def_ParseExampleDataset.pbtxt | 4 - .../api_def_SetStatsAggregatorDataset.pbtxt | 4 - .../python_api/api_def_SlideDataset.pbtxt | 4 - .../python_api/api_def_SqlDataset.pbtxt | 4 - .../api_def_StatsAggregatorHandle.pbtxt | 4 - .../api_def_StatsAggregatorSummary.pbtxt | 4 - .../python_api/api_def_UnbatchDataset.pbtxt | 4 - .../optimizers/data/graph_test_utils.cc | 2 +- .../optimizers/data/latency_all_edges.cc | 2 +- .../optimizers/data/latency_all_edges_test.cc | 7 +- .../optimizers/data/make_numa_aware.cc | 2 +- .../optimizers/data/make_numa_aware_test.cc | 6 +- .../optimizers/data/map_and_batch_fusion.cc | 2 +- .../data/map_and_batch_fusion_test.cc | 21 +- tensorflow/core/kernels/data/BUILD | 156 -- .../core/kernels/data/experimental/BUILD | 222 ++- .../dense_to_sparse_batch_dataset_op.cc | 5 +- .../group_by_reducer_dataset_op.cc | 5 +- .../group_by_window_dataset_op.cc | 5 +- .../experimental/identity_indexed_dataset.cc | 163 -- .../data/experimental/indexed_dataset.h | 119 -- ...dexed_dataset.cc => indexed_dataset_op.cc} | 249 ++- .../map_and_batch_dataset_op.cc | 10 +- .../parallel_interleave_dataset_op.cc | 1085 +++++++++++ .../parse_example_dataset_op.cc | 5 +- .../set_stats_aggregator_dataset_op.cc} | 5 +- .../sliding_window_dataset_op.cc} | 11 +- .../kernels/data/{ => experimental}/sql/BUILD | 0 .../{ => experimental}/sql/driver_manager.cc | 4 +- .../{ => experimental}/sql/driver_manager.h | 8 +- .../{ => experimental}/sql/query_connection.h | 6 +- .../sql/sqlite_query_connection.cc | 4 +- .../sql/sqlite_query_connection.h | 8 +- .../sql_dataset_op.cc} | 7 +- .../stats_aggregator_ops.cc | 10 +- .../{ => experimental}/stats_dataset_ops.cc | 10 +- .../to_tf_record_op.cc} | 4 +- .../{ => experimental}/unbatch_dataset_op.cc | 2 +- .../data/parallel_interleave_dataset_op.cc | 1051 +---------- .../core/ops/compat/ops_history.v1.pbtxt | 1622 ++++++++--------- tensorflow/core/ops/dataset_ops.cc | 230 --- .../core/ops/experimental_dataset_ops.cc | 214 ++- .../python/data/experimental/ops/batching.py | 8 +- .../python/data/experimental/ops/grouping.py | 68 +- .../data/experimental/ops/parsing_ops.py | 4 +- .../python/data/experimental/ops/readers.py | 8 +- .../data/experimental/ops/stats_aggregator.py | 6 +- .../python/data/experimental/ops/stats_ops.py | 12 +- .../python/data/experimental/ops/writers.py | 4 +- tensorflow/python/data/ops/dataset_ops.py | 2 +- tensorflow/python/data/ops/readers.py | 3 +- 77 files changed, 2583 insertions(+), 2981 deletions(-) rename tensorflow/core/api_def/base_api/{api_def_BytesProducedStatsDataset.pbtxt => api_def_ExperimentalBytesProducedStatsDataset.pbtxt} (56%) rename tensorflow/core/api_def/base_api/{api_def_DatasetToTFRecord.pbtxt => api_def_ExperimentalDatasetToTFRecord.pbtxt} (91%) rename tensorflow/core/api_def/base_api/{api_def_DenseToSparseBatchDataset.pbtxt => api_def_ExperimentalDenseToSparseBatchDataset.pbtxt} (89%) rename tensorflow/core/api_def/base_api/{api_def_GroupByReducerDataset.pbtxt => api_def_ExperimentalGroupByReducerDataset.pbtxt} (97%) rename tensorflow/core/api_def/base_api/{api_def_GroupByWindowDataset.pbtxt => api_def_ExperimentalGroupByWindowDataset.pbtxt} (82%) rename tensorflow/core/api_def/base_api/{api_def_LatencyStatsDataset.pbtxt => api_def_ExperimentalLatencyStatsDataset.pbtxt} (58%) rename tensorflow/core/api_def/base_api/{api_def_MapAndBatchDatasetV2.pbtxt => api_def_ExperimentalMapAndBatchDataset.pbtxt} (96%) rename tensorflow/core/api_def/base_api/{api_def_ParallelInterleaveDataset.pbtxt => api_def_ExperimentalParallelInterleaveDataset.pbtxt} (90%) rename tensorflow/core/api_def/base_api/{api_def_ParseExampleDataset.pbtxt => api_def_ExperimentalParseExampleDataset.pbtxt} (96%) create mode 100644 tensorflow/core/api_def/base_api/api_def_ExperimentalSetStatsAggregatorDataset.pbtxt rename tensorflow/core/api_def/base_api/{api_def_SlideDataset.pbtxt => api_def_ExperimentalSlidingWindowDataset.pbtxt} (88%) rename tensorflow/core/api_def/base_api/{api_def_SqlDataset.pbtxt => api_def_ExperimentalSqlDataset.pbtxt} (87%) create mode 100644 tensorflow/core/api_def/base_api/api_def_ExperimentalStatsAggregatorHandle.pbtxt rename tensorflow/core/api_def/base_api/{api_def_StatsAggregatorSummary.pbtxt => api_def_ExperimentalStatsAggregatorSummary.pbtxt} (56%) rename tensorflow/core/api_def/base_api/{api_def_UnbatchDataset.pbtxt => api_def_ExperimentalUnbatchDataset.pbtxt} (57%) delete mode 100644 tensorflow/core/api_def/base_api/api_def_MapAndBatchDataset.pbtxt delete mode 100644 tensorflow/core/api_def/base_api/api_def_SetStatsAggregatorDataset.pbtxt delete mode 100644 tensorflow/core/api_def/base_api/api_def_StatsAggregatorHandle.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_BytesProducedStatsDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_DenseToSparseBatchDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_GroupByWindowDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_LatencyStatsDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_MapAndBatchDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_ParallelInterleaveDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_ParseExampleDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_SetStatsAggregatorDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_SlideDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_SqlDataset.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_StatsAggregatorHandle.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_StatsAggregatorSummary.pbtxt delete mode 100644 tensorflow/core/api_def/python_api/api_def_UnbatchDataset.pbtxt rename tensorflow/core/kernels/data/{ => experimental}/dense_to_sparse_batch_dataset_op.cc (98%) rename tensorflow/core/kernels/data/{ => experimental}/group_by_reducer_dataset_op.cc (99%) rename tensorflow/core/kernels/data/{ => experimental}/group_by_window_dataset_op.cc (99%) delete mode 100644 tensorflow/core/kernels/data/experimental/identity_indexed_dataset.cc delete mode 100644 tensorflow/core/kernels/data/experimental/indexed_dataset.h rename tensorflow/core/kernels/data/experimental/{indexed_dataset.cc => indexed_dataset_op.cc} (62%) rename tensorflow/core/kernels/data/{ => experimental}/map_and_batch_dataset_op.cc (99%) create mode 100644 tensorflow/core/kernels/data/experimental/parallel_interleave_dataset_op.cc rename tensorflow/core/kernels/data/{ => experimental}/parse_example_dataset_op.cc (99%) rename tensorflow/core/kernels/data/{stats_aggregator_dataset_op.cc => experimental/set_stats_aggregator_dataset_op.cc} (98%) rename tensorflow/core/kernels/data/{slide_dataset_op.cc => experimental/sliding_window_dataset_op.cc} (97%) rename tensorflow/core/kernels/data/{ => experimental}/sql/BUILD (100%) rename tensorflow/core/kernels/data/{ => experimental}/sql/driver_manager.cc (88%) rename tensorflow/core/kernels/data/{ => experimental}/sql/driver_manager.h (81%) rename tensorflow/core/kernels/data/{ => experimental}/sql/query_connection.h (92%) rename tensorflow/core/kernels/data/{ => experimental}/sql/sqlite_query_connection.cc (97%) rename tensorflow/core/kernels/data/{ => experimental}/sql/sqlite_query_connection.h (84%) rename tensorflow/core/kernels/data/{sql_dataset_ops.cc => experimental/sql_dataset_op.cc} (96%) rename tensorflow/core/kernels/data/{ => experimental}/stats_aggregator_ops.cc (95%) rename tensorflow/core/kernels/data/{ => experimental}/stats_dataset_ops.cc (97%) rename tensorflow/core/kernels/data/{writer_ops.cc => experimental/to_tf_record_op.cc} (97%) rename tensorflow/core/kernels/data/{ => experimental}/unbatch_dataset_op.cc (99%) diff --git a/tensorflow/c/c_api_experimental.cc b/tensorflow/c/c_api_experimental.cc index 3693cc8599..38e29aa74a 100644 --- a/tensorflow/c/c_api_experimental.cc +++ b/tensorflow/c/c_api_experimental.cc @@ -6530,7 +6530,7 @@ library { } } node_def { - name: "ParallelInterleaveDataset/cycle_length" + name: "ExperimentalParallelInterleaveDataset/cycle_length" op: "Const" attr { key: "dtype" @@ -6551,7 +6551,7 @@ library { } } node_def { - name: "ParallelInterleaveDataset/block_length" + name: "ExperimentalParallelInterleaveDataset/block_length" op: "Const" attr { key: "dtype" @@ -6572,7 +6572,7 @@ library { } } node_def { - name: "ParallelInterleaveDataset/sloppy" + name: "ExperimentalParallelInterleaveDataset/sloppy" op: "Const" attr { key: "dtype" @@ -6593,7 +6593,7 @@ library { } } node_def { - name: "ParallelInterleaveDataset/buffer_output_elements" + name: "ExperimentalParallelInterleaveDataset/buffer_output_elements" op: "Const" attr { key: "dtype" @@ -6614,7 +6614,7 @@ library { } } node_def { - name: "ParallelInterleaveDataset/prefetch_input_elements" + name: "ExperimentalParallelInterleaveDataset/prefetch_input_elements" op: "Const" attr { key: "dtype" @@ -6635,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 { @@ -6742,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" diff --git a/tensorflow/contrib/data/python/ops/sliding.py b/tensorflow/contrib/data/python/ops/sliding.py index bcc383587c..a3ab6b8d72 100644 --- a/tensorflow/contrib/data/python/ops/sliding.py +++ b/tensorflow/contrib/data/python/ops/sliding.py @@ -22,7 +22,7 @@ from tensorflow.python.data.util import nest from tensorflow.python.framework import dtypes from tensorflow.python.framework import ops from tensorflow.python.framework import tensor_shape -from tensorflow.python.ops import gen_dataset_ops +from tensorflow.python.ops import gen_experimental_dataset_ops as ged_ops from tensorflow.python.util import deprecation @@ -41,7 +41,7 @@ class _SlideDataset(dataset_ops.UnaryDataset): window_shift, dtype=dtypes.int64, name="window_shift") def _as_variant_tensor(self): - return gen_dataset_ops.slide_dataset( + return ged_ops.experimental_sliding_window_dataset( self._input_dataset._as_variant_tensor(), # pylint: disable=protected-access window_size=self._window_size, window_shift=self._window_shift, diff --git a/tensorflow/core/api_def/base_api/api_def_BytesProducedStatsDataset.pbtxt b/tensorflow/core/api_def/base_api/api_def_ExperimentalBytesProducedStatsDataset.pbtxt similarity index 56% rename from tensorflow/core/api_def/base_api/api_def_BytesProducedStatsDataset.pbtxt rename to tensorflow/core/api_def/base_api/api_def_ExperimentalBytesProducedStatsDataset.pbtxt index 73df11b2f7..dc296162ae 100644 --- a/tensorflow/core/api_def/base_api/api_def_BytesProducedStatsDataset.pbtxt +++ b/tensorflow/core/api_def/base_api/api_def_ExperimentalBytesProducedStatsDataset.pbtxt @@ -1,4 +1,5 @@ op { - graph_op_name: "BytesProducedStatsDataset" + graph_op_name: "ExperimentalBytesProducedStatsDataset" + visibility: HIDDEN summary: "Records the bytes size of each element of `input_dataset` in a StatsAggregator." } diff --git a/tensorflow/core/api_def/base_api/api_def_DatasetToTFRecord.pbtxt b/tensorflow/core/api_def/base_api/api_def_ExperimentalDatasetToTFRecord.pbtxt similarity index 91% rename from tensorflow/core/api_def/base_api/api_def_DatasetToTFRecord.pbtxt rename to tensorflow/core/api_def/base_api/api_def_ExperimentalDatasetToTFRecord.pbtxt index e1b8a9abdd..085d20d7bf 100644 --- a/tensorflow/core/api_def/base_api/api_def_DatasetToTFRecord.pbtxt +++ b/tensorflow/core/api_def/base_api/api_def_ExperimentalDatasetToTFRecord.pbtxt @@ -1,5 +1,5 @@ op { - graph_op_name: "DatasetToTFRecord" + graph_op_name: "ExperimentalDatasetToTFRecord" visibility: HIDDEN in_arg { name: "input_dataset" diff --git a/tensorflow/core/api_def/base_api/api_def_DenseToSparseBatchDataset.pbtxt b/tensorflow/core/api_def/base_api/api_def_ExperimentalDenseToSparseBatchDataset.pbtxt similarity index 89% rename from tensorflow/core/api_def/base_api/api_def_DenseToSparseBatchDataset.pbtxt rename to tensorflow/core/api_def/base_api/api_def_ExperimentalDenseToSparseBatchDataset.pbtxt index e275cfdd3d..8ebd6d88a8 100644 --- a/tensorflow/core/api_def/base_api/api_def_DenseToSparseBatchDataset.pbtxt +++ b/tensorflow/core/api_def/base_api/api_def_ExperimentalDenseToSparseBatchDataset.pbtxt @@ -1,5 +1,6 @@ op { - graph_op_name: "DenseToSparseBatchDataset" + graph_op_name: "ExperimentalDenseToSparseBatchDataset" + visibility: HIDDEN in_arg { name: "input_dataset" description: <