diff --git a/CMakeLists.txt b/CMakeLists.txt index dcff6b54ca..c7d743e193 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -13,7 +13,6 @@ # limitations under the License cmake_minimum_required(VERSION 3.0) - set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake") set(PROJ_ROOT ${CMAKE_CURRENT_SOURCE_DIR}) set(PROJ_BINARY_ROOT ${CMAKE_CURRENT_BINARY_DIR}) @@ -37,6 +36,8 @@ include(simd) ################################ Configurations ####################################### option(WITH_GPU "Compile PaddlePaddle with NVIDIA GPU" ${CUDA_FOUND}) option(WITH_AVX "Compile PaddlePaddle with AVX intrinsics" ${AVX_FOUND}) +option(WITH_MKLDNN "Compile PaddlePaddle with mkl-dnn support." OFF) +option(WITH_MKLML "Compile PaddlePaddle with mklml package." OFF) option(WITH_DSO "Compile PaddlePaddle with dynamic linked CUDA" ON) option(WITH_TESTING "Compile PaddlePaddle with unit testing" ON) option(WITH_SWIG_PY "Compile PaddlePaddle with inference api" ON) @@ -75,6 +76,10 @@ if(ANDROID) "Disable PYTHON when cross-compiling for Android" FORCE) set(WITH_RDMA OFF CACHE STRING "Disable RDMA when cross-compiling for Android" FORCE) + set(WITH_MKLDNN OFF CACHE STRING + "Disable MKLDNN when cross-compiling for Android" FORCE) + set(WITH_MKLML OFF CACHE STRING + "Disable MKLML package when cross-compiling for Android" FORCE) endif(ANDROID) set(THIRD_PARTY_PATH "${CMAKE_BINARY_DIR}/third_party" CACHE STRING @@ -88,6 +93,7 @@ endif() ######################################################################################## +include(external/mklml) # download mklml package include(external/zlib) # download, build, install zlib include(external/gflags) # download, build, install gflags include(external/glog) # download, build, install glog @@ -95,6 +101,7 @@ include(external/gtest) # download, build, install gtest include(external/protobuf) # download, build, install protobuf include(external/python) # download, build, install python include(external/openblas) # download, build, install openblas +include(external/mkldnn) # download, build, install mkldnn include(external/swig) # download, build, install swig include(external/warpctc) # download, build, install warpctc include(external/any) # download libn::any @@ -136,6 +143,10 @@ if(WITH_GPU) endif(NOT WITH_DSO) endif(WITH_GPU) +if(WITH_MKLDNN) + list(APPEND EXTERNAL_LIBS ${MKLDNN_LIBRARY} ${MKLDNN_IOMP_LIB}) +endif() + if(USE_NNPACK) include(external/nnpack) list(APPEND EXTERNAL_LIBS ${NNPACK_LIBS}) diff --git a/cmake/cblas.cmake b/cmake/cblas.cmake index 913f711aff..854066fd1d 100644 --- a/cmake/cblas.cmake +++ b/cmake/cblas.cmake @@ -15,23 +15,44 @@ set(CBLAS_FOUND OFF) -## Find MKL First. -set(INTEL_ROOT "/opt/intel" CACHE PATH "Folder contains intel libs") -set(MKL_ROOT ${INTEL_ROOT}/mkl CACHE PATH "Folder contains MKL") +## Find MKLML First. +if(WITH_MKLML AND MKLML_INC_DIR AND MKLML_LIB) + set(CBLAS_FOUND ON) + set(CBLAS_PROVIDER MKLML) + set(CBLAS_INC_DIR ${MKLML_INC_DIR}) + set(CBLAS_LIBRARIES ${MKLML_LIB}) + + add_definitions(-DPADDLE_USE_MKLML) + add_definitions(-DLAPACK_FOUND) + + message(STATUS "Found cblas and lapack in MKLML " + "(include: ${CBLAS_INC_DIR}, library: ${CBLAS_LIBRARIES})") + return() +endif() + +## Then find MKL. +set(INTEL_MKL_ROOT "/opt/intel/mkl" CACHE PATH "Folder contains intel mkl libs") +set(MKL_ROOT $ENV{MKL_ROOT} CACHE PATH "Folder contains env MKL") + +set(MKL_INCLUDE_SEARCH_PATHS + ${MKL_ROOT}/include + ${INTEL_MKL_ROOT}/include) +set(MKL_LIB_SEARCH_PATHS + ${MKL_ROOT}/lib + ${MKL_ROOT}/lib/intel64 + ${INTEL_MKL_ROOT}/lib + ${INTEL_MKL_ROOT}/lib/intel64) find_path(MKL_INC_DIR mkl.h PATHS - ${MKL_ROOT}/include) + ${MKL_INCLUDE_SEARCH_PATHS}) find_path(MKL_LAPACK_INC_DIR mkl_lapacke.h PATHS - ${MKL_ROOT}/include) + ${MKL_INCLUDE_SEARCH_PATHS}) find_library(MKL_CORE_LIB NAMES mkl_core PATHS - ${MKL_ROOT}/lib - ${MKL_ROOT}/lib/intel64) + ${MKL_LIB_SEARCH_PATHS}) find_library(MKL_SEQUENTIAL_LIB NAMES mkl_sequential PATHS - ${MKL_ROOT}/lib - ${MKL_ROOT}/lib/intel64) + ${MKL_LIB_SEARCH_PATHS}) find_library(MKL_INTEL_LP64 NAMES mkl_intel_lp64 PATHS - ${MKL_ROOT}/lib - ${MKL_ROOT}/lib/intel64) + ${MKL_LIB_SEARCH_PATHS}) if(MKL_LAPACK_INC_DIR AND MKL_INC_DIR AND MKL_CORE_LIB AND MKL_SEQUENTIAL_LIB AND MKL_INTEL_LP64) set(CBLAS_FOUND ON) diff --git a/cmake/configure.cmake b/cmake/configure.cmake index 7afab5d534..69220e03fe 100644 --- a/cmake/configure.cmake +++ b/cmake/configure.cmake @@ -67,6 +67,30 @@ else() include_directories(${CUDA_TOOLKIT_INCLUDE}) endif(NOT WITH_GPU) +if(WITH_MKLDNN) + add_definitions(-DPADDLE_USE_MKLDNN) + if (WITH_MKLML AND MKLDNN_IOMP_DIR) + message(STATUS "Enable Intel OpenMP at ${MKLDNN_IOMP_DIR}") + set(OPENMP_FLAGS "-fopenmp") + set(CMAKE_C_CREATE_SHARED_LIBRARY_FORBIDDEN_FLAGS ${OPENMP_FLAGS}) + set(CMAKE_CXX_CREATE_SHARED_LIBRARY_FORBIDDEN_FLAGS ${OPENMP_FLAGS}) + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${MKLDNN_IOMP_DIR} -liomp5 -Wl,--as-needed") + set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -L${MKLDNN_IOMP_DIR} -liomp5 -Wl,--as-needed") + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OPENMP_FLAGS}") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OPENMP_FLAGS}") + else() + find_package(OpenMP) + if(OPENMP_FOUND) + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OpenMP_C_FLAGS}") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") + else() + message(WARNING "Can not find OpenMP." + "Some performance features in MKLDNN may not be available") + endif() + endif() + +endif(WITH_MKLDNN) + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${SIMD_FLAG}") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${SIMD_FLAG}") diff --git a/cmake/external/gtest.cmake b/cmake/external/gtest.cmake index 77e06e983e..e3970073a1 100644 --- a/cmake/external/gtest.cmake +++ b/cmake/external/gtest.cmake @@ -34,9 +34,15 @@ IF(WITH_TESTING) "${GTEST_INSTALL_DIR}/lib/libgtest_main.a" CACHE FILEPATH "gtest main libraries." FORCE) ENDIF(WIN32) + IF(WITH_MKLML) + # wait for mklml downloading completed + SET(GTEST_DEPENDS ${MKLML_PROJECT}) + ENDIF() + ExternalProject_Add( extern_gtest ${EXTERNAL_PROJECT_LOG_ARGS} + DEPENDS ${GTEST_DEPENDS} GIT_REPOSITORY "https://github.com/google/googletest.git" GIT_TAG "release-1.8.0" PREFIX ${GTEST_SOURCES_DIR} diff --git a/cmake/external/mkldnn.cmake b/cmake/external/mkldnn.cmake new file mode 100644 index 0000000000..eff15de73f --- /dev/null +++ b/cmake/external/mkldnn.cmake @@ -0,0 +1,72 @@ +# Copyright (c) 2017 PaddlePaddle Authors. All Rights Reserve. +# +# 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. + +IF(NOT ${WITH_MKLDNN}) + return() +ENDIF(NOT ${WITH_MKLDNN}) + +INCLUDE(ExternalProject) + +SET(MKLDNN_PROJECT "extern_mkldnn") +SET(MKLDNN_SOURCES_DIR ${THIRD_PARTY_PATH}/mkldnn) +SET(MKLDNN_INSTALL_ROOT ${CMAKE_INSTALL_PREFIX}) +IF(NOT "$ENV{HOME}" STREQUAL "/root") + SET(MKLDNN_INSTALL_ROOT "$ENV{HOME}") +ENDIF() + +SET(MKLDNN_INSTALL_DIR "${MKLDNN_INSTALL_ROOT}/opt/paddle/third_party/mkldnn") +SET(MKLDNN_INCLUDE_DIR "${MKLDNN_INSTALL_DIR}/include" CACHE PATH "mkldnn include directory." FORCE) + +IF(WIN32) + MESSAGE(WARNING "It is not supported compiling with mkldnn in windows Paddle yet." + "Force WITH_MKLDNN=OFF") + SET(WITH_MKLDNN OFF) + return() +ELSE(WIN32) + SET(MKLDNN_LIBRARY "${MKLDNN_INSTALL_DIR}/lib/libmkldnn.so" CACHE FILEPATH "mkldnn library." FORCE) + MESSAGE(STATUS "Set ${MKLDNN_INSTALL_DIR}/lib to runtime path") + SET(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE) + #SET(CMAKE_MACOSX_RPATH 1) # hold for MacOS + SET(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}" "${MKLDNN_INSTALL_DIR}/lib") +ENDIF(WIN32) + +INCLUDE_DIRECTORIES(${MKLDNN_INCLUDE_DIR}) + +IF(${CBLAS_PROVIDER} STREQUAL "MKLML") + SET(MKLDNN_DEPENDS ${MKLML_PROJECT}) + SET(MKLDNN_MKLROOT ${MKLML_ROOT}) + SET(MKLDNN_IOMP_LIB ${MKLML_IOMP_LIB}) + SET(MKLDNN_IOMP_DIR ${MKLML_LIB_DIR}) +ENDIF() + +ExternalProject_Add( + ${MKLDNN_PROJECT} + ${EXTERNAL_PROJECT_LOG_ARGS} + DEPENDS ${MKLDNN_DEPENDS} + GIT_REPOSITORY "https://github.com/01org/mkl-dnn.git" + GIT_TAG "v0.9" + PREFIX ${MKLDNN_SOURCES_DIR} + CONFIGURE_COMMAND mkdir -p /build + BUILD_COMMAND cd /build + && cmake .. -DCMAKE_INSTALL_PREFIX=${MKLDNN_INSTALL_DIR} -DMKLROOT=${MKLDNN_MKLROOT} + && $(MAKE) + INSTALL_COMMAND cd /build && $(MAKE) install + UPDATE_COMMAND "" +) + +ADD_LIBRARY(mkldnn SHARED IMPORTED GLOBAL) +SET_PROPERTY(TARGET mkldnn PROPERTY IMPORTED_LOCATION ${MKLDNN_LIBRARY}) +ADD_DEPENDENCIES(mkldnn ${MKLDNN_PROJECT}) +MESSAGE(STATUS "Mkldnn library: ${MKLDNN_LIBRARY}") +LIST(APPEND external_project_dependencies mkldnn) diff --git a/cmake/external/mklml.cmake b/cmake/external/mklml.cmake new file mode 100644 index 0000000000..3f940756a4 --- /dev/null +++ b/cmake/external/mklml.cmake @@ -0,0 +1,64 @@ +# Copyright (c) 2017 PaddlePaddle Authors. All Rights Reserve. +# +# 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. + +IF(NOT ${WITH_MKLML}) + return() +ENDIF(NOT ${WITH_MKLML}) + +INCLUDE(ExternalProject) + +SET(MKLML_PROJECT "extern_mklml") +SET(MKLML_VER "mklml_lnx_2018.0.20170425") +SET(MKLML_URL "https://github.com/01org/mkl-dnn/releases/download/v0.9/${MKLML_VER}.tgz") +SET(MKLML_SOURCE_DIR "${THIRD_PARTY_PATH}/mklml") +SET(MKLML_DOWNLOAD_DIR "${MKLML_SOURCE_DIR}/src/${MKLML_PROJECT}") +SET(MKLML_DST_DIR "opt/paddle/third_party/mklml") +SET(MKLML_INSTALL_ROOT "${CMAKE_INSTALL_PREFIX}") +IF(NOT "$ENV{HOME}" STREQUAL "/root") + SET(MKLML_INSTALL_ROOT "$ENV{HOME}") +ENDIF() + +SET(MKLML_INSTALL_DIR ${MKLML_INSTALL_ROOT}/${MKLML_DST_DIR}) +SET(MKLML_ROOT ${MKLML_INSTALL_DIR}/${MKLML_VER}) +SET(MKLML_INC_DIR ${MKLML_ROOT}/include) +SET(MKLML_LIB_DIR ${MKLML_ROOT}/lib) +SET(MKLML_LIB ${MKLML_LIB_DIR}/libmklml_intel.so) +SET(MKLML_IOMP_LIB ${MKLML_LIB_DIR}/libiomp5.so) +SET(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}" "${MKLML_ROOT}/lib") + +INCLUDE_DIRECTORIES(${MKLML_INC_DIR}) + +SET(mklml_cmakefile ${MKLML_DOWNLOAD_DIR}/CMakeLists.txt) +FILE(WRITE ${mklml_cmakefile} "PROJECT(MKLML)\n" + "cmake_minimum_required(VERSION 3.0)\n" + "install(DIRECTORY ${MKLML_VER}\n" + " DESTINATION ${MKLML_DST_DIR})\n") + +ExternalProject_Add( + ${MKLML_PROJECT} + ${EXTERNAL_PROJECT_LOG_ARGS} + PREFIX ${MKLML_SOURCE_DIR} + DOWNLOAD_DIR ${MKLML_DOWNLOAD_DIR} + DOWNLOAD_COMMAND wget --no-check-certificate -O ${MKLML_DOWNLOAD_DIR}/${MKLML_VER}.tgz ${MKLML_URL} + && tar -xzf ${MKLML_DOWNLOAD_DIR}/${MKLML_VER}.tgz + DOWNLOAD_NO_PROGRESS 1 + UPDATE_COMMAND "" + CMAKE_ARGS -DCMAKE_INSTALL_PREFIX=${MKLML_INSTALL_ROOT} + CMAKE_CACHE_ARGS -DCMAKE_INSTALL_PREFIX:PATH=${MKLML_INSTALL_ROOT} +) + +ADD_LIBRARY(mklml SHARED IMPORTED GLOBAL) +SET_PROPERTY(TARGET mklml PROPERTY IMPORTED_LOCATION ${MKLML_LIB}) +ADD_DEPENDENCIES(mklml ${MKLML_PROJECT}) +LIST(APPEND external_project_dependencies mklml) diff --git a/cmake/flags.cmake b/cmake/flags.cmake index c31e62fc08..34fd348893 100644 --- a/cmake/flags.cmake +++ b/cmake/flags.cmake @@ -124,6 +124,7 @@ set(GPU_COMMON_FLAGS -Wno-error=literal-suffix -Wno-error=unused-local-typedefs -Wno-error=unused-function # Warnings in Numpy Header. + -Wno-error=array-bounds # Warnings in Eigen::array ) if (APPLE) diff --git a/cmake/generic.cmake b/cmake/generic.cmake index e42e75c12a..534be0abe2 100644 --- a/cmake/generic.cmake +++ b/cmake/generic.cmake @@ -290,8 +290,22 @@ function(go_library TARGET_NAME) set(${TARGET_NAME}_LIB_NAME "${CMAKE_STATIC_LIBRARY_PREFIX}${TARGET_NAME}${CMAKE_STATIC_LIBRARY_SUFFIX}" CACHE STRING "output library name for target ${TARGET_NAME}") endif() - # Add dummy code to support `make target_name` under Terminal Command set(dummyfile ${CMAKE_CURRENT_BINARY_DIR}/${TARGET_NAME}_dummy.c) + + # This custom command will always run since it depends on a not + # existing file. + add_custom_command( + OUTPUT dummy_rebulid_${TARGET_NAME} + COMMAND cmake -E touch ${dummyfile} + ) + # Create a custom target that depends on the custom command output + # file, so the custom command can be referenced as a dependency by + # `add_dependencies`. + add_custom_target(rebuild_${TARGET_NAME} + DEPENDS dummy_rebulid_${TARGET_NAME} + ) + + # Add dummy code to support `make target_name` under Terminal Command file(WRITE ${dummyfile} "const char * dummy = \"${dummyfile}\";") if (go_library_SHARED OR go_library_shared) add_library(${TARGET_NAME} SHARED ${dummyfile}) @@ -302,6 +316,12 @@ function(go_library TARGET_NAME) add_dependencies(${TARGET_NAME} ${go_library_DEPS}) endif(go_library_DEPS) + # The "source file" of the library is `${dummyfile}` which never + # change, so the target will never rebuild. Make the target depends + # on the custom command that touches the library "source file", so + # rebuild will always happen. + add_dependencies(${TARGET_NAME} rebuild_${TARGET_NAME}) + set(${TARGET_NAME}_LIB_PATH "${CMAKE_CURRENT_BINARY_DIR}/${${TARGET_NAME}_LIB_NAME}" CACHE STRING "output library path for target ${TARGET_NAME}") file(GLOB GO_SOURCE RELATIVE "${CMAKE_CURRENT_SOURCE_DIR}" "*.go") diff --git a/go/cmd/master/master.go b/go/cmd/master/master.go index 9eaf8c04ae..287da69491 100644 --- a/go/cmd/master/master.go +++ b/go/cmd/master/master.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package main import ( diff --git a/go/cmd/pserver/pserver.go b/go/cmd/pserver/pserver.go index 652d7ba315..20094fbab4 100644 --- a/go/cmd/pserver/pserver.go +++ b/go/cmd/pserver/pserver.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package main import ( diff --git a/go/connection/conn.go b/go/connection/conn.go index 977e8cc123..ffa8db689d 100644 --- a/go/connection/conn.go +++ b/go/connection/conn.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package connection import ( diff --git a/go/master/CMakeLists.txt b/go/master/CMakeLists.txt index 30531e6469..93efa4eaf7 100644 --- a/go/master/CMakeLists.txt +++ b/go/master/CMakeLists.txt @@ -1,3 +1,17 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# if(WITH_TESTING) go_test(master_test) endif() diff --git a/go/master/c/CMakeLists.txt b/go/master/c/CMakeLists.txt index d900850be0..082d9f3f59 100644 --- a/go/master/c/CMakeLists.txt +++ b/go/master/c/CMakeLists.txt @@ -1 +1,15 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# go_library(paddle_master SHARED DEPS paddle_go_optimizer) diff --git a/go/master/c/client.go b/go/master/c/client.go index 2cbe164c7b..9f5733075f 100644 --- a/go/master/c/client.go +++ b/go/master/c/client.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package main /* diff --git a/go/master/client.go b/go/master/client.go index 90b9947097..7f33090dc7 100644 --- a/go/master/client.go +++ b/go/master/client.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master import ( diff --git a/go/master/client_internal_test.go b/go/master/client_internal_test.go index 70dc09bf94..ee305e2c80 100644 --- a/go/master/client_internal_test.go +++ b/go/master/client_internal_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master import ( diff --git a/go/master/client_test.go b/go/master/client_test.go index bc92dc5ac9..a90062c753 100644 --- a/go/master/client_test.go +++ b/go/master/client_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master_test import ( diff --git a/go/master/etcd_client.go b/go/master/etcd_client.go index 69dc6a8268..607e726251 100644 --- a/go/master/etcd_client.go +++ b/go/master/etcd_client.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master import ( diff --git a/go/master/inmem_store.go b/go/master/inmem_store.go index 57e75dc4e0..ffd663f7f0 100644 --- a/go/master/inmem_store.go +++ b/go/master/inmem_store.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master import "sync" diff --git a/go/master/service.go b/go/master/service.go index 262735f421..2766720c28 100644 --- a/go/master/service.go +++ b/go/master/service.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master import ( diff --git a/go/master/service_internal_test.go b/go/master/service_internal_test.go index 9c0d1d0a39..69a882fc33 100644 --- a/go/master/service_internal_test.go +++ b/go/master/service_internal_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package master import "testing" diff --git a/go/pserver/CMakeLists.txt b/go/pserver/CMakeLists.txt index 6267040a6e..4fe0a8cb02 100644 --- a/go/pserver/CMakeLists.txt +++ b/go/pserver/CMakeLists.txt @@ -1,3 +1,17 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# if(WITH_TESTING) go_test(pserver_test DEPS paddle_go_optimizer) endif() diff --git a/go/pserver/client/CMakeLists.txt b/go/pserver/client/CMakeLists.txt index 0052bb460b..e295611060 100644 --- a/go/pserver/client/CMakeLists.txt +++ b/go/pserver/client/CMakeLists.txt @@ -1,3 +1,17 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# if(WITH_TESTING) go_test(pserver_client_test DEPS paddle_go_optimizer) endif() diff --git a/go/pserver/client/c/CMakeLists.txt b/go/pserver/client/c/CMakeLists.txt index c6333eab55..a932791c7c 100644 --- a/go/pserver/client/c/CMakeLists.txt +++ b/go/pserver/client/c/CMakeLists.txt @@ -1,3 +1,17 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# cc_library(paddle_go_optimizer DEPS paddle_optimizer paddle_proto glog gflags protobuf) target_link_libraries(paddle_go_optimizer stdc++ m) diff --git a/go/pserver/client/c/cclient.go b/go/pserver/client/c/cclient.go index 718b4304c8..24cd922ffe 100644 --- a/go/pserver/client/c/cclient.go +++ b/go/pserver/client/c/cclient.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package main /* diff --git a/go/pserver/client/c/test/CMakeLists.txt b/go/pserver/client/c/test/CMakeLists.txt index dce8645ce7..3724ccb60b 100644 --- a/go/pserver/client/c/test/CMakeLists.txt +++ b/go/pserver/client/c/test/CMakeLists.txt @@ -1,2 +1,16 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# cc_test(test_cclient SRCS test_cclient.c DEPS paddle_pserver_cclient paddle_go_optimizer) add_style_check_target(test_cclient test_cclient.c) diff --git a/go/pserver/client/c/test/test_cclient.c b/go/pserver/client/c/test/test_cclient.c index 8eababbe33..f9b9967434 100644 --- a/go/pserver/client/c/test/test_cclient.c +++ b/go/pserver/client/c/test/test_cclient.c @@ -1,3 +1,17 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + #include #include diff --git a/go/pserver/client/client.go b/go/pserver/client/client.go index b4a45e1c21..ddb749d629 100644 --- a/go/pserver/client/client.go +++ b/go/pserver/client/client.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package client import ( diff --git a/go/pserver/client/client_test.go b/go/pserver/client/client_test.go index 5c89882a29..b630d434dc 100644 --- a/go/pserver/client/client_test.go +++ b/go/pserver/client/client_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package client_test import ( diff --git a/go/pserver/client/etcd_client.go b/go/pserver/client/etcd_client.go index 953065b427..b6ff1fec8a 100644 --- a/go/pserver/client/etcd_client.go +++ b/go/pserver/client/etcd_client.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package client import ( @@ -66,10 +80,10 @@ func (p *EtcdClient) List() []Server { for { for i := 0; i < psDesired; i++ { ctx, cancel := context.WithTimeout(context.Background(), p.timeout) - cancel() psKey := pserver.PsPath + strconv.Itoa(i) log.Debugf("checking %s", psKey) resp, err := p.client.Get(ctx, psKey) + cancel() if err != nil { log.Infof("Get psKey= %s error, %v", psKey, err) time.Sleep(p.timeout) diff --git a/go/pserver/etcd_client.go b/go/pserver/etcd_client.go index e70e826975..98ff8ce827 100644 --- a/go/pserver/etcd_client.go +++ b/go/pserver/etcd_client.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package pserver import ( diff --git a/go/pserver/optimizer.go b/go/pserver/optimizer.go index 151a3f8033..709160d45d 100644 --- a/go/pserver/optimizer.go +++ b/go/pserver/optimizer.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package pserver // #cgo CFLAGS: -I ../../ diff --git a/go/pserver/optimizer_test.go b/go/pserver/optimizer_test.go index d19e9de92e..d001e6993e 100644 --- a/go/pserver/optimizer_test.go +++ b/go/pserver/optimizer_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package pserver import ( diff --git a/go/pserver/service.go b/go/pserver/service.go index c723959d6b..46738413f0 100644 --- a/go/pserver/service.go +++ b/go/pserver/service.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package pserver import ( diff --git a/go/pserver/service_test.go b/go/pserver/service_test.go index a191f689fe..988f3b5acb 100644 --- a/go/pserver/service_test.go +++ b/go/pserver/service_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package pserver_test import ( diff --git a/go/utils/networkhelper/CMakeLists.txt b/go/utils/networkhelper/CMakeLists.txt index db6cf211d8..9233264ff3 100644 --- a/go/utils/networkhelper/CMakeLists.txt +++ b/go/utils/networkhelper/CMakeLists.txt @@ -1,3 +1,17 @@ +# Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. +# +# 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. +# if(WITH_TESTING) go_test(network_helper_test) endif() diff --git a/go/utils/networkhelper/helper.go b/go/utils/networkhelper/helper.go index fbeaea8f5e..c3fc747bda 100644 --- a/go/utils/networkhelper/helper.go +++ b/go/utils/networkhelper/helper.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package networkhelper import ( diff --git a/go/utils/networkhelper/helper_test.go b/go/utils/networkhelper/helper_test.go index 4208f9e358..0bc02ad42a 100644 --- a/go/utils/networkhelper/helper_test.go +++ b/go/utils/networkhelper/helper_test.go @@ -1,3 +1,17 @@ +// Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +// 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. + package networkhelper import "testing" diff --git a/paddle/framework/CMakeLists.txt b/paddle/framework/CMakeLists.txt index eb34164623..760d84e51e 100644 --- a/paddle/framework/CMakeLists.txt +++ b/paddle/framework/CMakeLists.txt @@ -1,23 +1,25 @@ # ddim lib -cc_library(enforce SRCS enforce.cc DEPS glog) -cc_test(enforce_test SRCS enforce_test.cc DEPS enforce) cc_library(ddim SRCS ddim.cc DEPS eigen3) cc_test(ddim_test SRCS ddim_test.cc DEPS ddim) nv_test(dim_test SRCS dim_test.cu DEPS ddim) -cc_library(tensor SRCS tensor.cc DEPS ddim place enforce paddle_memory) + +cc_library(tensor SRCS tensor.cc DEPS ddim place paddle_memory) cc_test(tensor_test SRCS tensor_test.cc DEPS tensor) +cc_test(eigen_test SRCS eigen_test.cc DEPS tensor) + cc_test(variable_test SRCS variable_test.cc) cc_test(scope_test SRCS scope_test.cc) + proto_library(attr_type SRCS attr_type.proto) proto_library(op_proto SRCS op_proto.proto DEPS attr_type) -cc_test(op_proto_test SRCS op_proto_test.cc DEPS op_proto protobuf) proto_library(op_desc SRCS op_desc.proto DEPS attr_type) +cc_test(op_proto_test SRCS op_proto_test.cc DEPS op_proto protobuf) cc_test(op_desc_test SRCS op_desc_test.cc DEPS op_desc protobuf) cc_library(operator SRCS operator.cc DEPS op_desc device_context tensor) cc_test(operator_test SRCS operator_test.cc DEPS operator op_registry) -cc_library(op_registry SRCS op_registry.cc DEPS op_proto op_desc enforce) +cc_library(op_registry SRCS op_registry.cc DEPS op_proto op_desc) cc_test(op_registry_test SRCS op_registry_test.cc DEPS op_registry operator) py_proto_compile(framework_py_proto SRCS attr_type.proto op_proto.proto op_desc.proto) diff --git a/paddle/framework/attr_checker.h b/paddle/framework/attr_checker.h index f2d88f3cb0..ea5614a45f 100644 --- a/paddle/framework/attr_checker.h +++ b/paddle/framework/attr_checker.h @@ -6,7 +6,7 @@ #include #include #include -#include "paddle/framework/enforce.h" +#include "paddle/platform/enforce.h" namespace paddle { namespace framework { diff --git a/paddle/framework/ddim.cc b/paddle/framework/ddim.cc index d2ef85afe5..545c1dcc2a 100644 --- a/paddle/framework/ddim.cc +++ b/paddle/framework/ddim.cc @@ -13,7 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/framework/ddim.h" -#include "paddle/framework/enforce.h" +#include "paddle/platform/enforce.h" namespace paddle { namespace framework { diff --git a/paddle/framework/ddim.h b/paddle/framework/ddim.h index 070850375d..9fcc657edc 100644 --- a/paddle/framework/ddim.h +++ b/paddle/framework/ddim.h @@ -19,7 +19,7 @@ limitations under the License. */ #include #include #include "paddle/framework/dim.h" -#include "paddle/framework/enforce.h" +#include "paddle/platform/enforce.h" #include "unsupported/Eigen/CXX11/Tensor" namespace paddle { @@ -119,17 +119,6 @@ int arity(const DDim& ddim); std::ostream& operator<<(std::ostream&, const DDim&); -template -Eigen::DSizes ToEigenDSizes(const DDim& dims) { - int rank = arity(dims); - PADDLE_ENFORCE(rank == NDIMS, "DDim and NDIMS must be same"); - Eigen::DSizes dsizes; - for (int d = 0; d < rank; d++) { - dsizes[d] = dims[d]; - } - return dsizes; -} - } // namespace framework } // namespace paddle diff --git a/paddle/framework/eigen.h b/paddle/framework/eigen.h new file mode 100644 index 0000000000..5f3358c69b --- /dev/null +++ b/paddle/framework/eigen.h @@ -0,0 +1,84 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once + +#include "paddle/framework/tensor.h" +#include "unsupported/Eigen/CXX11/Tensor" + +namespace paddle { +namespace framework { + +// EigenDim converts paddle::platform::DDim into Eigen::DSizes. +template +struct EigenDim { + using Type = Eigen::DSizes; + + static Type From(const DDim& dims) { + PADDLE_ENFORCE(arity(dims) == D, "D must match arity(DDim)"); + Type ret; + for (int d = 0; d < arity(dims); d++) { + ret[d] = dims[d]; + } + return ret; + } +}; + +// Interpret paddle::platform::Tensor as EigenTensor and EigenConstTensor. +template +struct EigenTensor { + // TODO(qijun) Now, default type in unaligned, and we will make a benchmark on + // the speed of aligned and unaligned version in future. + using Type = Eigen::TensorMap>; + + using ConstType = + Eigen::TensorMap>; + + static Type From(Tensor& tensor, DDim dims) { + return Type(tensor.data(), EigenDim::From(dims)); + } + + static Type From(Tensor& tensor) { return From(tensor, tensor.dims_); } + + static ConstType From(const Tensor& tensor, DDim dims) { + return ConstType(tensor.data(), EigenDim::From(dims)); + } + + static ConstType From(const Tensor& tensor) { + return From(tensor, tensor.dims_); + } +}; + +template +struct EigenMatrix : public EigenTensor {}; + +template +struct EigenVector : public EigenTensor { + // Flatten reshapes a Tensor into an EigenVector. + static typename EigenVector::Type Flatten(Tensor& tensor) { + return EigenVector::From( + tensor, make_ddim({static_cast(product(tensor.dims_))})); + } + + static typename EigenVector::ConstType Flatten(const Tensor& tensor) { + return EigenVector::From( + tensor, make_ddim({static_cast(product(tensor.dims_))})); + } +}; + +} // namespace framework +} // namespace paddle diff --git a/paddle/framework/eigen_test.cc b/paddle/framework/eigen_test.cc new file mode 100644 index 0000000000..a9fa728e49 --- /dev/null +++ b/paddle/framework/eigen_test.cc @@ -0,0 +1,101 @@ +/* + Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + 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 "paddle/framework/eigen.h" +#include + +namespace paddle { +namespace framework { + +TEST(EigenDim, From) { + EigenDim<3>::Type ed = EigenDim<3>::From(make_ddim({1, 2, 3})); + ASSERT_EQ(1, ed[0]); + ASSERT_EQ(2, ed[1]); + ASSERT_EQ(3, ed[2]); +} + +TEST(Eigen, Tensor) { + Tensor t; + float* p = t.mutable_data(make_ddim({1, 2, 3}), platform::CPUPlace()); + for (int i = 0; i < 1 * 2 * 3; i++) { + p[i] = static_cast(i); + } + + EigenTensor::Type et = EigenTensor::From(t); + + ASSERT_EQ(1, et.dimension(0)); + ASSERT_EQ(2, et.dimension(1)); + ASSERT_EQ(3, et.dimension(2)); + + for (int i = 0; i < 1; i++) { + for (int j = 0; j < 2; j++) { + for (int k = 0; k < 3; k++) { + ASSERT_NEAR((i * 2 + j) * 3 + k, et(i, j, k), 1e-6f); + } + } + } +} + +TEST(Eigen, VectorFrom) { + Tensor t; + float* p = t.mutable_data(make_ddim({6}), platform::CPUPlace()); + for (int i = 0; i < 6; i++) { + p[i] = static_cast(i); + } + + EigenVector::Type ev = EigenVector::From(t); + + ASSERT_EQ(6, ev.dimension(0)); + + for (int i = 0; i < 6; i++) { + ASSERT_NEAR(i, ev(i), 1e-6f); + } +} + +TEST(Eigen, VectorFlatten) { + Tensor t; + float* p = t.mutable_data(make_ddim({1, 2, 3}), platform::CPUPlace()); + for (int i = 0; i < 1 * 2 * 3; i++) { + p[i] = static_cast(i); + } + + EigenVector::Type ev = EigenVector::Flatten(t); + + ASSERT_EQ(1 * 2 * 3, ev.dimension(0)); + + for (int i = 0; i < 1 * 2 * 3; i++) { + ASSERT_NEAR(i, ev(i), 1e-6f); + } +} + +TEST(Eigen, Matrix) { + Tensor t; + float* p = t.mutable_data(make_ddim({2, 3}), platform::CPUPlace()); + for (int i = 0; i < 2 * 3; i++) { + p[i] = static_cast(i); + } + + EigenMatrix::Type em = EigenMatrix::From(t); + + ASSERT_EQ(2, em.dimension(0)); + ASSERT_EQ(3, em.dimension(1)); + + for (int i = 0; i < 2; i++) { + for (int j = 0; j < 3; j++) { + ASSERT_NEAR(i * 3 + j, em(i, j), 1e-6f); + } + } +} + +} // namespace framework +} // namespace paddle diff --git a/paddle/framework/enforce.cc b/paddle/framework/enforce.cc deleted file mode 100644 index 644930ff98..0000000000 --- a/paddle/framework/enforce.cc +++ /dev/null @@ -1,15 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - - 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 "paddle/framework/enforce.h" diff --git a/paddle/framework/enforce.h b/paddle/framework/enforce.h deleted file mode 100644 index ffce8148e9..0000000000 --- a/paddle/framework/enforce.h +++ /dev/null @@ -1,75 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. -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. */ - -#pragma once -#include -#include -#include -#include - -namespace paddle { -namespace framework { - -/** - * @brief Enforce exception. Inherits std::exception - * - * All enforce condition not met, will throw an EnforceNotMet exception. - */ -class EnforceNotMet : public std::exception { - public: - EnforceNotMet(const std::string& msg, const char* file, int fileline) { - std::ostringstream sout; - sout << msg << " at [" << file << ":" << fileline << "];"; - all_msg_ = sout.str(); - } - - const char* what() const noexcept override { return all_msg_.c_str(); } - - private: - std::string all_msg_; -}; - -// From https://stackoverflow.com/questions/30130930/ -// __buildin_expect is in C++ 11 standard. Since the condition which enforced -// should be true in most situation, it will make the compiler generate faster -// code by adding `UNLIKELY` macro. -#define UNLIKELY(condition) __builtin_expect(static_cast(condition), 0) - -/** - * @brief Throw a EnforceNotMet exception, automatically filled __FILE__ & - * __LINE__ - * - * This macro take __VA_ARGS__, user can pass any type if that type can - * serialize to std::ostream - */ -#define PADDLE_THROW(...) \ - do { \ - throw ::paddle::framework::EnforceNotMet( \ - ::paddle::string::Sprintf(__VA_ARGS__), __FILE__, __LINE__); \ - } while (0) - -/** - * @brief Enforce a condition, otherwise throw an EnforceNotMet - */ -#ifdef NDEBUG -#define PADDLE_ENFORCE(condition, ...) \ - do { \ - if (UNLIKELY(!(condition))) { \ - PADDLE_THROW(__VA_ARGS__); \ - } \ - } while (0) -#else -#define PADDLE_ENFORCE(condition, ...) \ - CHECK(condition) << ::paddle::string::Sprintf(__VA_ARGS__); -#endif - -} // namespace framework -} // namespace paddle diff --git a/paddle/framework/net.cc b/paddle/framework/net.cc index 407a69fda6..139425b356 100644 --- a/paddle/framework/net.cc +++ b/paddle/framework/net.cc @@ -39,19 +39,22 @@ void PlainNet::CompleteAddOp(bool calc) { output_set.insert(opt); } } + inputs_.reserve(input_set.size()); std::copy(input_set.begin(), input_set.end(), std::back_inserter(inputs_)); + std::sort(inputs_.begin(), inputs_.end()); outputs_.reserve(output_set.size()); + std::copy(output_set.begin(), output_set.end(), std::back_inserter(outputs_)); + std::sort(outputs_.begin(), outputs_.end()); + std::vector tmp_index; tmp_index.reserve(temp_output.size()); - int idx = 0; - for (auto& opt : output_set) { - if (Contains(temp_output, opt)) { - tmp_index.push_back(idx); + int output_len = static_cast(outputs_.size()); + for (int i = 0; i < output_len; ++i) { + if (Contains(temp_output, outputs_[i])) { + tmp_index.push_back(i); } - outputs_.push_back(opt); - ++idx; } attrs_["temporary_index"] = tmp_index; @@ -59,7 +62,7 @@ void PlainNet::CompleteAddOp(bool calc) { std::string PlainNet::DebugString() const { std::ostringstream os; - os << this->type_ << ":" << std::endl; + os << OperatorBase::DebugString() << std::endl; for (auto& op : ops_) { std::istringstream is(op->DebugString()); for (std::string line; std::getline(is, line);) { diff --git a/paddle/framework/net.h b/paddle/framework/net.h index 19c5fa223b..b2c64a8675 100644 --- a/paddle/framework/net.h +++ b/paddle/framework/net.h @@ -39,7 +39,7 @@ namespace framework { */ class Net : public OperatorBase { public: - virtual void AddOp(const OperatorPtr& op) = 0; + virtual void AddOp(const std::shared_ptr& op) = 0; virtual void CompleteAddOp(bool calc) = 0; }; @@ -57,7 +57,7 @@ class PlainNet : public Net { * Infer all the operators' input and output variables' shapes, will be called * before every mini-batch */ - void InferShape(const ScopePtr& scope) const override { + void InferShape(const std::shared_ptr& scope) const override { for (auto& op : ops_) { op->InferShape(scope); } @@ -70,7 +70,7 @@ class PlainNet : public Net { * scope will be used instead. If no OpContext is provicded, default context * will be used. */ - void Run(const ScopePtr& scope, + void Run(const std::shared_ptr& scope, const platform::DeviceContext& dev_ctx) const override { for (auto& op : ops_) { op->Run(scope, dev_ctx); @@ -80,7 +80,7 @@ class PlainNet : public Net { /** * @brief Add an operator by ptr */ - void AddOp(const OperatorPtr& op) override { + void AddOp(const std::shared_ptr& op) override { PADDLE_ENFORCE(!add_op_done_, "Cannot AddOp when this network is sealed"); ops_.push_back(op); } @@ -89,7 +89,7 @@ class PlainNet : public Net { std::string DebugString() const override; - std::vector ops_; + std::vector> ops_; private: bool add_op_done_{false}; diff --git a/paddle/framework/net_op_test.cc b/paddle/framework/net_op_test.cc index f5e1c22400..c179042c81 100644 --- a/paddle/framework/net_op_test.cc +++ b/paddle/framework/net_op_test.cc @@ -10,10 +10,10 @@ static int run_cnt = 0; class TestOp : public pd::OperatorBase { public: - void InferShape(const paddle::framework::ScopePtr& scope) const override { + void InferShape(const std::shared_ptr& scope) const override { ++infer_shape_cnt; } - void Run(const paddle::framework::ScopePtr& scope, + void Run(const std::shared_ptr& scope, const paddle::platform::DeviceContext& dev_ctx) const override { ++run_cnt; } @@ -63,5 +63,5 @@ TEST(OpKernel, all) { ASSERT_EQ(2, infer_shape_cnt); ASSERT_EQ(2, run_cnt); - ASSERT_THROW(net->AddOp(op2), paddle::framework::EnforceNotMet); + ASSERT_THROW(net->AddOp(op2), std::runtime_error); } diff --git a/paddle/framework/op_registry.h b/paddle/framework/op_registry.h index c41fe10729..165a68c1cf 100644 --- a/paddle/framework/op_registry.h +++ b/paddle/framework/op_registry.h @@ -227,10 +227,10 @@ class OpRegistry { } } - static OperatorPtr CreateOp(const std::string& type, - const VarNameList& inputs, - const VarNameList& outputs, - const AttributeMap& attrs) { + static std::shared_ptr CreateOp(const std::string& type, + const VarNameList& inputs, + const VarNameList& outputs, + const AttributeMap& attrs) { auto op_create_it = creators().find(type); PADDLE_ENFORCE(op_create_it != creators().end(), "Operator %s cannot be found", type); @@ -252,10 +252,10 @@ class OpRegistry { } op->Init(); - return OperatorPtr(op); + return std::shared_ptr(op); } - static OperatorPtr CreateOp(const OpDesc& op_desc) { + static std::shared_ptr CreateOp(const OpDesc& op_desc) { std::vector inputs; inputs.reserve((size_t)op_desc.inputs_size()); std::copy(op_desc.inputs().begin(), op_desc.inputs().end(), diff --git a/paddle/framework/op_registry_test.cc b/paddle/framework/op_registry_test.cc index d3a51a361a..05095372d8 100644 --- a/paddle/framework/op_registry_test.cc +++ b/paddle/framework/op_registry_test.cc @@ -7,9 +7,9 @@ namespace paddle { namespace framework { class CosineOp : public OperatorBase { public: - void Run(const ScopePtr& scope, + void Run(const std::shared_ptr& scope, const platform::DeviceContext& dev_ctx) const override {} - void InferShape(const ScopePtr& scope) const override {} + void InferShape(const std::shared_ptr& scope) const override {} }; class CosineOpProtoAndCheckerMaker : public OpProtoAndCheckerMaker { @@ -27,8 +27,8 @@ class CosineOpProtoAndCheckerMaker : public OpProtoAndCheckerMaker { class MyTestOp : public OperatorBase { public: - void InferShape(const ScopePtr& scope) const override {} - void Run(const ScopePtr& scope, + void InferShape(const std::shared_ptr& scope) const override {} + void Run(const std::shared_ptr& scope, const platform::DeviceContext& dev_ctx) const override {} }; @@ -67,7 +67,7 @@ TEST(OpRegistry, CreateOp) { attr->set_type(paddle::framework::AttrType::FLOAT); attr->set_f(scale); - paddle::framework::OperatorPtr op = + std::shared_ptr op = paddle::framework::OpRegistry::CreateOp(op_desc); auto scope = std::make_shared(); paddle::platform::CPUDeviceContext dev_ctx; @@ -89,9 +89,8 @@ TEST(OpRegistry, IllegalAttr) { bool caught = false; try { - paddle::framework::OperatorPtr op __attribute__((unused)) = - paddle::framework::OpRegistry::CreateOp(op_desc); - } catch (paddle::framework::EnforceNotMet err) { + paddle::framework::OpRegistry::CreateOp(op_desc); + } catch (std::runtime_error& err) { caught = true; std::string msg = "larger_than check fail"; const char* err_msg = err.what(); @@ -110,7 +109,7 @@ TEST(OpRegistry, DefaultValue) { ASSERT_TRUE(op_desc.IsInitialized()); - paddle::framework::OperatorPtr op = + std::shared_ptr op = paddle::framework::OpRegistry::CreateOp(op_desc); auto scope = std::make_shared(); paddle::platform::CPUDeviceContext dev_ctx; @@ -136,9 +135,8 @@ TEST(OpRegistry, CustomChecker) { // attr 'test_attr' is not set bool caught = false; try { - paddle::framework::OperatorPtr op __attribute__((unused)) = - paddle::framework::OpRegistry::CreateOp(op_desc); - } catch (paddle::framework::EnforceNotMet err) { + paddle::framework::OpRegistry::CreateOp(op_desc); + } catch (std::runtime_error& err) { caught = true; std::string msg = "Attribute 'test_attr' is required!"; const char* err_msg = err.what(); @@ -155,9 +153,8 @@ TEST(OpRegistry, CustomChecker) { attr->set_i(3); caught = false; try { - paddle::framework::OperatorPtr op __attribute__((unused)) = - paddle::framework::OpRegistry::CreateOp(op_desc); - } catch (paddle::framework::EnforceNotMet err) { + paddle::framework::OpRegistry::CreateOp(op_desc); + } catch (std::runtime_error& err) { caught = true; std::string msg = "'test_attr' must be even!"; const char* err_msg = err.what(); @@ -174,8 +171,7 @@ TEST(OpRegistry, CustomChecker) { attr->set_type(paddle::framework::AttrType::INT); attr->set_i(4); SetInputFormat(&op_desc); - paddle::framework::OperatorPtr op = - paddle::framework::OpRegistry::CreateOp(op_desc); + auto op = paddle::framework::OpRegistry::CreateOp(op_desc); paddle::platform::CPUDeviceContext dev_ctx; auto scope = std::make_shared(); op->Run(scope, dev_ctx); @@ -196,7 +192,7 @@ TEST(ProtoMaker, DuplicatedAttr) { pd::OpProto op_proto; pd::OpAttrChecker op_checker; auto proto_maker = TestAttrProtoMaker(&op_proto, &op_checker); - ASSERT_THROW(proto_maker.Validate(), paddle::framework::EnforceNotMet); + ASSERT_THROW(proto_maker.Validate(), std::runtime_error); } class TestInOutProtoMaker : public pd::OpProtoAndCheckerMaker { @@ -212,5 +208,5 @@ TEST(ProtoMaker, DuplicatedInOut) { pd::OpProto op_proto; pd::OpAttrChecker op_checker; auto proto_maker = TestInOutProtoMaker(&op_proto, &op_checker); - ASSERT_THROW(proto_maker.Validate(), paddle::framework::EnforceNotMet); + ASSERT_THROW(proto_maker.Validate(), std::runtime_error); } diff --git a/paddle/framework/operator.h b/paddle/framework/operator.h index 5f046d6293..6b8dbb39ac 100644 --- a/paddle/framework/operator.h +++ b/paddle/framework/operator.h @@ -47,7 +47,6 @@ struct EigenDeviceConverter { #endif class OperatorBase; -using OperatorPtr = std::shared_ptr; /** * OperatorBase has the basic element that Net will call to do computation. * Only CreateOperator from OpRegistry will new Operator directly. User @@ -80,10 +79,10 @@ class OperatorBase { /// InferShape infer the size of Variables used by this Operator with /// information inside scope - virtual void InferShape(const ScopePtr& scope) const = 0; + virtual void InferShape(const std::shared_ptr& scope) const = 0; /// Net will call this function to Run an op. - virtual void Run(const ScopePtr& scope, + virtual void Run(const std::shared_ptr& scope, const platform::DeviceContext& dev_ctx) const = 0; // Get a input with argument's name described in `op_proto` @@ -208,7 +207,7 @@ class OperatorWithKernel : public OperatorBase { using OpKernelMap = std::unordered_map, OpKernelHash>; - void Run(const ScopePtr& scope, + void Run(const std::shared_ptr& scope, const platform::DeviceContext& dev_ctx) const final { auto& opKernel = AllOpKernels().at(type_).at(OpKernelKey(dev_ctx)); opKernel->Compute(KernelContext(this, scope, dev_ctx)); diff --git a/paddle/framework/operator_test.cc b/paddle/framework/operator_test.cc index 8e55d0111f..3fae356c3e 100644 --- a/paddle/framework/operator_test.cc +++ b/paddle/framework/operator_test.cc @@ -24,8 +24,8 @@ static int op_run_num = 0; class OpWithoutKernelTest : public OperatorBase { public: void Init() override { x = 1; } - void InferShape(const ScopePtr& scope) const override {} - void Run(const ScopePtr& scope, + void InferShape(const std::shared_ptr& scope) const override {} + void Run(const std::shared_ptr& scope, const platform::DeviceContext& dev_ctx) const override { op_run_num++; ASSERT_EQ((int)inputs_.size(), 1); @@ -70,8 +70,7 @@ TEST(OperatorBase, all) { paddle::platform::CPUDeviceContext device_context; auto scope = std::make_shared(); - paddle::framework::OperatorPtr op = - paddle::framework::OpRegistry::CreateOp(op_desc); + auto op = paddle::framework::OpRegistry::CreateOp(op_desc); scope->CreateVariable("OUT1"); ASSERT_EQ(paddle::framework::op_run_num, 0); op->Run(scope, device_context); @@ -189,8 +188,7 @@ TEST(OpKernel, all) { paddle::platform::CPUDeviceContext cpu_device_context; auto scope = std::make_shared(); - paddle::framework::OperatorPtr op = - paddle::framework::OpRegistry::CreateOp(op_desc); + auto op = paddle::framework::OpRegistry::CreateOp(op_desc); ASSERT_EQ(paddle::framework::cpu_kernel_run_num, 0); op->Run(scope, cpu_device_context); ASSERT_EQ(paddle::framework::cpu_kernel_run_num, 1); @@ -236,6 +234,6 @@ TEST(OpKernel, multi_inputs) { paddle::platform::CPUDeviceContext cpu_device_context; auto scope = std::make_shared(); - OperatorPtr op(paddle::framework::OpRegistry::CreateOp(op_desc)); + auto op = paddle::framework::OpRegistry::CreateOp(op_desc); op->Run(scope, cpu_device_context); } diff --git a/paddle/framework/scope.h b/paddle/framework/scope.h index ec62c9189f..79c9ffd1a6 100644 --- a/paddle/framework/scope.h +++ b/paddle/framework/scope.h @@ -24,7 +24,6 @@ namespace paddle { namespace framework { class Scope; -using ScopePtr = std::shared_ptr; /** * @brief Scope that manage all variables. @@ -44,7 +43,7 @@ class Scope { /** * @brief Initialize a Scope with parent. */ - explicit Scope(const ScopePtr& parent) : parent_(parent) {} + explicit Scope(const std::shared_ptr& parent) : parent_(parent) {} /** * @brief Create Variable @@ -91,7 +90,7 @@ class Scope { private: std::unordered_map> vars_; - ScopePtr parent_{nullptr}; + std::shared_ptr parent_{nullptr}; }; } // namespace framework diff --git a/paddle/framework/tensor.h b/paddle/framework/tensor.h index 1dd421cdb6..a36f375d2e 100644 --- a/paddle/framework/tensor.h +++ b/paddle/framework/tensor.h @@ -19,9 +19,8 @@ limitations under the License. */ #include #include #include "paddle/framework/ddim.h" -#include "paddle/framework/enforce.h" -#include "paddle/framework/tensor_types.h" #include "paddle/memory/memory.h" +#include "paddle/platform/enforce.h" #include "paddle/platform/place.h" #include "unsupported/Eigen/CXX11/Tensor" @@ -35,30 +34,41 @@ struct CastToPyBufferImpl; namespace framework { class Tensor { + template + friend struct paddle::pybind::details::CastToPyBufferImpl; + + template + friend struct EigenTensor; + + template + friend struct EigenVector; + public: Tensor() : offset_(0) {} template const T* data() const { - CheckDims(); + EnforceSufficientMemory(); return reinterpret_cast( reinterpret_cast(holder_->ptr()) + offset_); } template - T* raw_data() const { - CheckDims(); + T* data() { + EnforceSufficientMemory(); return reinterpret_cast(reinterpret_cast(holder_->ptr()) + offset_); } - template + template ::value>::type* = nullptr> T* mutable_data(DDim dims, platform::Place place) { - set_dims(dims); + Resize(dims); return mutable_data(place); } - template + template ::value>::type* = nullptr> T* mutable_data(platform::Place place) { PADDLE_ENFORCE(product(dims_) > 0, "Tensor's numel must be larger than zero to call " @@ -86,72 +96,10 @@ class Tensor { offset_); } - template - typename TTypes::Tensor shaped(DDim new_dims) { - Eigen::array dims = - paddle::framework::ToEigenDSizes(new_dims); - return typename TTypes::Tensor(raw_data(), dims); - } - - template - typename TTypes::Tensor tensor() { - return typename TTypes::Tensor( - raw_data(), paddle::framework::ToEigenDSizes(dims_)); - } - - // flat to rank = 1 - template - typename TTypes::Flat flat() { - return shaped(make_ddim({static_cast(product(dims_))})); - } - - // to TensorType Vec - template - typename TTypes::Vec vec() { - return tensor(); - } - - // to TensorType Matrix - template - typename TTypes::Matrix matrix() { - return tensor(); - } - - // const versions of all the methods above. - template - typename TTypes::Tensor shaped(DDim new_dims) const { - Eigen::array dims = - paddle::framework::ToEigenDSizes(new_dims); - return typename TTypes::Tensor(data(), dims); - } - - template - typename TTypes::ConstantTensor tensor() const { - return typename TTypes::Tensor( - data(), paddle::framework::ToEigenDSizes(dims_)); - } - - template - typename TTypes::ConstFlat flat() const { - return shaped(make_ddim({static_cast(product(dims_))})); - } - - template - typename TTypes::ConstVec vec() const { - return tensor(); - } - template - typename TTypes::ConstMatrix matrix() const { - return tensor(); - } - - template - void ShareDataFrom(const Tensor& src) { - src.CheckDims(); - holder_ = src.holder_; - set_dims(src.dims()); - offset_ = src.offset_; + void ShareDataWith(const Tensor& src) { + src.EnforceSufficientMemory(); + *this = src; } template @@ -159,9 +107,9 @@ class Tensor { PADDLE_ENFORCE(platform::is_cpu_place(src.holder_->place()) && platform::is_cpu_place(dst_place), "Tensor::CopyFrom only support CPU now."); - src.CheckDims(); + src.EnforceSufficientMemory(); size_t size = product(src.dims_) * sizeof(T); - set_dims(src.dims()); + Resize(src.dims()); const void* src_ptr = static_cast(src.data()); void* dst_ptr = static_cast(mutable_data(dst_place)); memcpy(dst_ptr, src_ptr, size); @@ -169,34 +117,25 @@ class Tensor { template Tensor Slice(const int& begin_idx, const int& end_idx) const { - CheckDims(); - PADDLE_ENFORCE(begin_idx >= 0 && end_idx <= dims_[0], - "Slice index is less than zero or out of bound."); + EnforceSufficientMemory(); + PADDLE_ENFORCE(begin_idx >= 0, "Slice begin index is less than zero."); + PADDLE_ENFORCE(end_idx <= dims_[0], "Slice end index is out of bound."); PADDLE_ENFORCE(begin_idx < end_idx, "Begin index must be less than end index."); PADDLE_ENFORCE(dims_[0] != 1, "Can not slice a tensor with dims_[0] = 1."); - std::vector d = vectorize(dims_); - int base = 1; - for (size_t i = 1; i < d.size(); ++i) { - base *= d[i]; - } + int base = product(dims_) / dims_[0]; Tensor dst; dst.holder_ = holder_; DDim dst_dims = dims_; dst_dims[0] = end_idx - begin_idx; - dst.set_dims(dst_dims); + dst.Resize(dst_dims); dst.offset_ = offset_ + begin_idx * base * sizeof(T); return dst; } - void set_dims(const DDim& dims) { - if (dims == dims_) { - return; - } - dims_ = dims; - } + void Resize(const DDim& dims) { dims_ = dims; } - DDim dims() const { return dims_; } + const DDim& dims() const { return dims_; } private: // Placeholder hides type T, so it doesn't appear as a template @@ -211,21 +150,9 @@ class Tensor { template struct PlaceholderImpl : public Placeholder { - private: - template - class Deleter { - public: - Deleter(PType place) : place_(place) {} - void operator()(T* ptr) { memory::Free(place_, static_cast(ptr)); } - - private: - PType place_; - }; - - public: PlaceholderImpl(PlaceType place, size_t size) : ptr_(static_cast(memory::Alloc(place, size)), - Deleter(place)), + memory::PODDeleter(place)), place_(place), size_(size) {} @@ -234,13 +161,13 @@ class Tensor { virtual paddle::platform::Place place() const { return place_; } virtual std::type_index type() const { return std::type_index(typeid(T)); } - std::unique_ptr> ptr_; + std::unique_ptr> ptr_; platform::Place place_; // record the place of ptr_. size_t size_; // size of the memory block. }; template - inline void CheckDims() const { + inline void EnforceSufficientMemory() const { PADDLE_ENFORCE(holder_ != nullptr, "Tenosr holds no memory. Call Tensor::mutable_data first."); PADDLE_ENFORCE(holder_->size() >= product(dims_) * sizeof(T) + offset_, @@ -250,9 +177,11 @@ class Tensor { std::shared_ptr holder_; // holds the memory block if allocated. DDim dims_; - size_t offset_; // marks the begin of tensor data area. - template - friend struct paddle::pybind::details::CastToPyBufferImpl; + // A PlaceHolder may be shared by more than one tensor. Some of them may be + // slices of the others. So the offset_ is introduced here to indicate the + // byte offset between PlaceHolder::ptr_ and where tensor's data really + // begins. + size_t offset_; }; } // namespace framework diff --git a/paddle/framework/tensor_test.cc b/paddle/framework/tensor_test.cc index 84c6f0cf65..089844dc01 100644 --- a/paddle/framework/tensor_test.cc +++ b/paddle/framework/tensor_test.cc @@ -19,7 +19,7 @@ TEST(Tensor, Dims) { using namespace paddle::framework; using namespace paddle::platform; Tensor tt; - tt.set_dims(make_ddim({2, 3, 4})); + tt.Resize(make_ddim({2, 3, 4})); DDim dims = tt.dims(); ASSERT_EQ(arity(dims), 3); for (int i = 0; i < 3; ++i) { @@ -33,7 +33,7 @@ TEST(Tensor, DataAssert) { bool caught = false; try { src_tensor.data(); - } catch (paddle::framework::EnforceNotMet err) { + } catch (std::runtime_error& err) { caught = true; std::string msg = "Tenosr holds no memory. Call Tensor::mutable_data first."; @@ -97,7 +97,7 @@ TEST(Tensor, MutableData) { #endif } -TEST(Tensor, ShareDataFrom) { +TEST(Tensor, ShareDataWith) { using namespace paddle::framework; using namespace paddle::platform; { @@ -106,8 +106,8 @@ TEST(Tensor, ShareDataFrom) { // Try to share data form uninitialized tensor bool caught = false; try { - dst_tensor.ShareDataFrom(src_tensor); - } catch (EnforceNotMet err) { + dst_tensor.ShareDataWith(src_tensor); + } catch (std::runtime_error& err) { caught = true; std::string msg = "Tenosr holds no memory. Call Tensor::mutable_data first."; @@ -119,7 +119,7 @@ TEST(Tensor, ShareDataFrom) { ASSERT_TRUE(caught); src_tensor.mutable_data(make_ddim({2, 3, 4}), CPUPlace()); - dst_tensor.ShareDataFrom(src_tensor); + dst_tensor.ShareDataWith(src_tensor); ASSERT_EQ(src_tensor.data(), dst_tensor.data()); } @@ -128,7 +128,7 @@ TEST(Tensor, ShareDataFrom) { Tensor src_tensor; Tensor dst_tensor; src_tensor.mutable_data(make_ddim({2, 3, 4}), GPUPlace()); - dst_tensor.ShareDataFrom(src_tensor); + dst_tensor.ShareDataWith(src_tensor); ASSERT_EQ(src_tensor.data(), dst_tensor.data()); } #endif diff --git a/paddle/framework/tensor_types.h b/paddle/framework/tensor_types.h deleted file mode 100644 index 4bf27a377e..0000000000 --- a/paddle/framework/tensor_types.h +++ /dev/null @@ -1,67 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - -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. */ - -#pragma once - -#include "unsupported/Eigen/CXX11/Tensor" - -namespace paddle { -namespace framework { - -// Helper to define Tensor types given that the scalar is of type T. -template -struct TTypes { - // Rank- tensor of scalar type T. - typedef Eigen::TensorMap, - Eigen::Aligned> - Tensor; - typedef Eigen::TensorMap< - Eigen::Tensor, Eigen::Aligned> - ConstTensor; - - // Scalar tensor (implemented as a rank-0 tensor) of scalar type T. - typedef Eigen::TensorMap< - Eigen::TensorFixedSize, Eigen::RowMajor, IndexType>, - Eigen::Aligned> - Scalar; - typedef Eigen::TensorMap, - Eigen::RowMajor, IndexType>, - Eigen::Aligned> - ConstScalar; - - // Rank-1 tensor (vector) of scalar type T. - typedef Eigen::TensorMap, - Eigen::Aligned> - Flat; - typedef Eigen::TensorMap< - Eigen::Tensor, Eigen::Aligned> - ConstFlat; - typedef Eigen::TensorMap, - Eigen::Aligned> - Vec; - typedef Eigen::TensorMap< - Eigen::Tensor, Eigen::Aligned> - ConstVec; - - // Rank-2 tensor (matrix) of scalar type T. - typedef Eigen::TensorMap, - Eigen::Aligned> - Matrix; - typedef Eigen::TensorMap< - Eigen::Tensor, Eigen::Aligned> - ConstMatrix; -}; - -} // namespace framework -} // namespace paddle diff --git a/paddle/function/CMakeLists.txt b/paddle/function/CMakeLists.txt index a5b14c0c71..2bec00cdb2 100644 --- a/paddle/function/CMakeLists.txt +++ b/paddle/function/CMakeLists.txt @@ -36,6 +36,7 @@ if(WITH_GPU) add_simple_unittest(MulOpTest) add_simple_unittest(CosSimOpTest) add_simple_unittest(RowConvOpTest) + add_simple_unittest(CropOpTest) endif() add_simple_unittest(ConvOpTest) diff --git a/paddle/function/ConvOpTest.cpp b/paddle/function/ConvOpTest.cpp index dfa2f78461..7f32c73479 100644 --- a/paddle/function/ConvOpTest.cpp +++ b/paddle/function/ConvOpTest.cpp @@ -31,13 +31,22 @@ public: ConvolutionTest(const std::string& conv1, const std::string& conv2, TestType type, + bool useGroups = true, std::string algo = "auto") { for (size_t batchSize : {1, 32}) { for (size_t inputSize : {7, 14, 54}) { for (size_t filterSize : {1, 3, 5}) { for (size_t inputChannels : {3, 64}) { - for (size_t outputChannels : {3, 64, 128}) { - if (inputChannels < outputChannels) break; + for (size_t outputChannels : {3, 64}) { + if (inputChannels > outputChannels) break; + size_t groups; + if (!useGroups) { + groups = 1; + } else { + if (outputChannels % inputChannels != 0) continue; + groups = inputChannels; + } + for (size_t stride : {1, 2}) { for (size_t padding : {0, 1}) { if (padding >= filterSize) break; @@ -62,13 +71,24 @@ public: FuncConfig() .set("paddings", paddings) .set("strides", strides) - .set("groups", (size_t)1) + .set("groups", groups) .set("algo", algo)); TensorShape input{ batchSize, inputChannels, inputSize, inputSize}; - TensorShape filter{ - outputChannels, inputChannels, filterSize, filterSize}; + + TensorShape filter; + if (groups > 1) + filter = TensorShape({groups, + outputChannels / groups, + inputChannels / groups, + filterSize, + filterSize}); + else + filter = TensorShape({outputChannels, + inputChannels, + filterSize, + filterSize}); TensorShape output{ batchSize, outputChannels, outputSize, outputSize}; @@ -85,7 +105,8 @@ public: } else if (type == kBackwardFilterTest) { test.addInputs(BufferArg(VALUE_TYPE_FLOAT, output)); test.addInputs(BufferArg(VALUE_TYPE_FLOAT, input)); - test.addOutputs(BufferArg(VALUE_TYPE_FLOAT, filter)); + test.addOutputs(BufferArg(VALUE_TYPE_FLOAT, filter), + ADD_TO); test.run(); } } @@ -106,6 +127,7 @@ public: ConvolutionTest2(const std::string& conv1, const std::string& conv2, TestType type, + bool useGroups = true, std::string algo = "auto") { for (size_t batchSize : {16}) { for (size_t inputHeight : {7, 31}) { @@ -113,7 +135,15 @@ public: for (size_t filterHeight : {1, 5}) { for (size_t filterWidth : {3, 7}) { for (size_t inputChannels : {7}) { - for (size_t outputChannels : {32}) { + for (size_t outputChannels : {7}) { + size_t groups; + if (!useGroups) { + groups = 1; + } else { + if (outputChannels % inputChannels != 0) continue; + groups = inputChannels; + } + size_t stride = 1; size_t padding = 0; size_t outputHeight = @@ -141,13 +171,24 @@ public: FuncConfig() .set("paddings", paddings) .set("strides", strides) - .set("groups", (size_t)1) + .set("groups", groups) .set("algo", algo)); TensorShape input{ batchSize, inputChannels, inputHeight, inputWidth}; - TensorShape filter{ - outputChannels, inputChannels, filterHeight, filterWidth}; + + TensorShape filter; + if (groups > 1) + filter = TensorShape({groups, + outputChannels / groups, + inputChannels / groups, + filterHeight, + filterWidth}); + else + filter = TensorShape({outputChannels, + inputChannels, + filterHeight, + filterWidth}); TensorShape output{ batchSize, outputChannels, outputHeight, outputWidth}; @@ -164,7 +205,8 @@ public: } else if (type == kBackwardFilterTest) { test.addInputs(BufferArg(VALUE_TYPE_FLOAT, output)); test.addInputs(BufferArg(VALUE_TYPE_FLOAT, input)); - test.addOutputs(BufferArg(VALUE_TYPE_FLOAT, filter)); + test.addOutputs(BufferArg(VALUE_TYPE_FLOAT, filter), + ADD_TO); test.run(); } } @@ -177,34 +219,88 @@ public: } }; +// ======Start Convolution TEST====== + TEST(Forward, GEMM) { ConvolutionTest test( - "NaiveConv-CPU", "GemmConv-CPU", kForwardTest); + "NaiveConv-CPU", "GemmConv-CPU", kForwardTest, false); ConvolutionTest2 test2( - "NaiveConv-CPU", "GemmConv-CPU", kForwardTest); + "NaiveConv-CPU", "GemmConv-CPU", kForwardTest, false); } #ifndef PADDLE_ONLY_CPU TEST(Forward, GEMM2) { ConvolutionTest test( - "GemmConv-CPU", "GemmConv-GPU", kForwardTest); + "GemmConv-CPU", "GemmConv-GPU", kForwardTest, false); ConvolutionTest2 test2( - "GemmConv-CPU", "GemmConv-GPU", kForwardTest); + "GemmConv-CPU", "GemmConv-GPU", kForwardTest, false); } TEST(BackwardInput, GEMM) { ConvolutionTest test( - "GemmConvGradInput-CPU", "GemmConvGradInput-GPU", kBackwardInputTest); + "GemmConvGradInput-CPU", + "GemmConvGradInput-GPU", + kBackwardInputTest, + false); ConvolutionTest2 test2( - "GemmConvGradInput-CPU", "GemmConvGradInput-GPU", kBackwardInputTest); + "GemmConvGradInput-CPU", + "GemmConvGradInput-GPU", + kBackwardInputTest, + false); } TEST(BackwardFilter, GEMM) { ConvolutionTest test( - "GemmConvGradFilter-CPU", "GemmConvGradFilter-GPU", kBackwardFilterTest); + "GemmConvGradFilter-CPU", + "GemmConvGradFilter-GPU", + kBackwardFilterTest, + false); ConvolutionTest2 test2( - "GemmConvGradFilter-CPU", "GemmConvGradFilter-GPU", kBackwardFilterTest); + "GemmConvGradFilter-CPU", + "GemmConvGradFilter-GPU", + kBackwardFilterTest, + false); } #endif +// ======End Convolution TEST====== + +// ======Start DepthwiseConvolution TEST====== + +// TODO(zhaolong) The depthwise convolution cpu test will be added when the cpu +// version of depthwiseConv is implemented. + +#ifndef PADDLE_ONLY_CPU + +TEST(DepthwiseConvForward, GEMM2) { + ConvolutionTest test( + "GemmConv-CPU", "DepthwiseConv-GPU", kForwardTest); + ConvolutionTest2 test2( + "GemmConv-CPU", "DepthwiseConv-GPU", kForwardTest); +} + +TEST(DepthwiseConvBackwardInput, GEMM) { + ConvolutionTest test( + "GemmConvGradInput-CPU", + "DepthwiseConvGradInput-GPU", + kBackwardInputTest); + ConvolutionTest2 test2( + "GemmConvGradInput-CPU", + "DepthwiseConvGradInput-GPU", + kBackwardInputTest); +} + +TEST(DepthwiseConvBackwardFilter, GEMM) { + ConvolutionTest test( + "GemmConvGradFilter-CPU", + "DepthwiseConvGradFilter-GPU", + kBackwardFilterTest); + ConvolutionTest2 test2( + "GemmConvGradFilter-CPU", + "DepthwiseConvGradFilter-GPU", + kBackwardFilterTest); +} + +#endif +// ======End DepthwiseConvolution TEST====== } // namespace paddle diff --git a/paddle/function/CropOp.cpp b/paddle/function/CropOp.cpp new file mode 100644 index 0000000000..f12ee43e3d --- /dev/null +++ b/paddle/function/CropOp.cpp @@ -0,0 +1,177 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "CropOp.h" +#include "paddle/function/TensorShape.h" +#include "paddle/math/Vector.h" + +namespace paddle { + +template <> +void Crop(real* outputs, + const real* inputs, + const TensorShape inShape, + const TensorShape outShape, + const FuncConfig& conf) { + std::vector crop_corner = + conf.get>("crop_corner"); + int cCrop = crop_corner[1]; + int hCrop = crop_corner[2]; + int wCrop = crop_corner[3]; + + int num = inShape[0]; + int inC = inShape[1]; + int inH = inShape[2]; + int inW = inShape[3]; + + int outC = outShape[1]; + int outH = outShape[2]; + int outW = outShape[3]; + + for (int n = 0; n < num; n++) { + for (int c = 0; c < outC; c++) { + for (int h = 0; h < outH; h++) { + int outoff = ((n * outC + c) * outH + h) * outW; + int inoff = ((n * inC + c + cCrop) * inH + h + hCrop) * inW + wCrop; + memcpy(outputs + outoff, inputs + inoff, outW * sizeof(real)); + } + } + } +} + +template <> +void CropGrad(const real* inGrad, + real* outGrad, + const TensorShape inShape, + const TensorShape outShape, + const FuncConfig& conf) { + std::vector crop_corner = + conf.get>("crop_corner"); + int cCrop = crop_corner[1]; + int hCrop = crop_corner[2]; + int wCrop = crop_corner[3]; + + int num = outShape[0]; + int outC = outShape[1]; + int outH = outShape[2]; + int outW = outShape[3]; + + int inC = inShape[1]; + int inH = inShape[2]; + int inW = inShape[3]; + + for (int n = 0; n < num; n++) { + for (int c = 0; c < inC; c++) { + for (int h = 0; h < inH; h++) { + int outoff = ((n * outC + c + cCrop) * outH + h + hCrop) * outW + wCrop; + int inoff = ((n * inC + c) * inH + h) * inW; + CpuVector inG = CpuVector(inW, const_cast(inGrad + inoff)); + CpuVector outG = CpuVector(inW, outGrad + outoff); + outG += inG; + } + } + } +} + +/** + * \brief Crop input according to the specify corner and shape. + * The input and output is a 4D tensor. In CropFunc, we only + * crop the 2nd to 4th dimension. + * + * Argument in this Function: + * \param pad_ A struct object contains the cropping corner and shape. + * \param inputs A 4D tensor, only one input. + * \param outputs A 4D tensor, the output value after cropping. + * + * For example, + * Input(2,2,2,3) = [ + * [ [[1,2,3], [3,4,5]], + * [[2,3,5], [1,6,7]] ], + * [ [[4,3,1], [1,8,7]], + * [[3,8,9], [2,3,5]] ] + * ] # the input shape is (2,2,2,3) + * + * pad_: if corner = (0,1,1) and crop_shape = (2,1,2) + * Output(2,2,1,2) = [ + * [ [[4,5]], + * [[6,7]] ], + * [ [[8,7]], + * [[3,5]] ] + * ] # the input shape is (2,2,2,3) + */ +template +class CropFunc : public FunctionBase { +public: + void init(const FuncConfig& config) override { conf_ = config; } + + void calc(const BufferArgs& inputs, const BufferArgs& outputs) override { + CHECK_EQ(1UL, inputs.size()); + CHECK_EQ(1UL, outputs.size()); + CHECK_EQ(outputs[0].getArgType(), ASSIGN_TO); + + TensorShape inShape = inputs[0].shape(); + TensorShape outShape = outputs[0].shape(); + + Crop(outputs[0].data(), + inputs[0].data(), + inShape, + outShape, + conf_); + } + +private: + FuncConfig conf_; +}; + +/** + * \brief The backward propagation of cropping Function. + * + * Argument in this Function: + * \param crop_ The same meaning as it in CropFunc. + * \param inputs The gradient with respect to the output value of CropFunc. + * \param outputs The gradient with respect to the input value of CropFunc. + */ + +template +class CropGradFunc : public FunctionBase { +public: + void init(const FuncConfig& config) override { conf_ = config; } + + void calc(const BufferArgs& inputs, const BufferArgs& outputs) override { + CHECK_EQ(1UL, inputs.size()); + CHECK_EQ(1UL, outputs.size()); + CHECK_EQ(outputs[0].getArgType(), ADD_TO); + + TensorShape outShape = outputs[0].shape(); + TensorShape inShape = inputs[0].shape(); + + CropGrad(inputs[0].data(), + outputs[0].data(), + inShape, + outShape, + conf_); + } + +private: + FuncConfig conf_; +}; + +REGISTER_TYPED_FUNC(Crop, CPU, CropFunc); +REGISTER_TYPED_FUNC(CropGrad, CPU, CropGradFunc); +#ifndef PADDLE_ONLY_CPU +REGISTER_TYPED_FUNC(Crop, GPU, CropFunc); +REGISTER_TYPED_FUNC(CropGrad, GPU, CropGradFunc); +#endif + +} // namespace paddle diff --git a/paddle/function/CropOp.h b/paddle/function/CropOp.h new file mode 100644 index 0000000000..87986fbdc7 --- /dev/null +++ b/paddle/function/CropOp.h @@ -0,0 +1,51 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once + +#include "Function.h" + +namespace paddle { + +/** + * \brief This funtion crops inputs according to the specify start point and + *shape. + * + * \param[out] outputs save results. + * \param[in] inputs input data. + * \param[in] inShape the shape of input tensor. + * \param[in] conf the cropping config + */ +template +void Crop(real* outputs, + const real* inputs, + const TensorShape inShape, + const TensorShape outShape, + const FuncConfig& conf); + +/** + * \brief Cropping operation backward. + * + * \param[out] inGrad gradients of previous layer + * \param[in] outGrad output gradient + * \param[in] inShape the shape of input tensor. + * \param[in] conf the cropping config + */ +template +void CropGrad(const real* inGrad, + real* outGrad, + const TensorShape inShape, + const TensorShape outShape, + const FuncConfig& conf); +} // namespace paddle diff --git a/paddle/function/CropOpGpu.cu b/paddle/function/CropOpGpu.cu new file mode 100644 index 0000000000..786eb268d4 --- /dev/null +++ b/paddle/function/CropOpGpu.cu @@ -0,0 +1,116 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "hl_base.h" +#include "CropOp.h" + +namespace paddle { + +__global__ void KeCrop(real* outputs, const real* inputs, + int inC, int inH, int inW, + int cropC, int cropH, int cropW, + int outC, int outH, int outW, int nthreads) { + const int idx = threadIdx.x + blockIdx.x * blockDim.x; + if (idx < nthreads) { + const int w = idx % outW; + const int h = (idx / outW) % outH; + const int c = (idx / outW / outH) % outC; + const int n = idx / outW / outH / outC; + + const int off = ((n * inC + c + cropC) * inH + h + cropH) * inW + cropW + w; + outputs[idx] = inputs[off]; + } +} + +template <> +void Crop(real* outputs, + const real* inputs, + const TensorShape inShape, + const TensorShape outShape, + const FuncConfig& conf) { + std::vector crop_corner = + conf.get>("crop_corner"); + int cropC = crop_corner[1]; + int cropH = crop_corner[2]; + int cropW = crop_corner[3]; + + int num = inShape[0]; + int inC = inShape[1]; + int inH = inShape[2]; + int inW = inShape[3]; + + int outC = outShape[1]; + int outH = outShape[2]; + int outW = outShape[3]; + + size_t nth = num * outC * outH * outW; + int blockSize = 1024; + int gridSize = (nth + blockSize - 1) / blockSize; + + KeCrop<<>> + (outputs, inputs, inC, inH, inW, cropC, cropH, cropW, + outC, outH, outW, nth); + CHECK_SYNC("Crop"); +} + +__global__ void KeCropDiff(const real* inGrad, real* outGrad, + int inC, int inH, int inW, + int cropC, int cropH, int cropW, + int outC, int outH, int outW, int nthreads) { + const int idx = threadIdx.x + blockIdx.x * blockDim.x; + if (idx < nthreads) { + const int w = idx % inW; + const int h = (idx / inW) % inH; + const int c = (idx / inW / inH) % inC; + const int n = idx / inW / inH / inC; + + const int off = + ((n * outC + c + cropC) * outH + h + cropH) * outW + cropW + w; + + outGrad[off] += inGrad[idx]; + } +} + +template <> +void CropGrad(const real* inGrad, + real* outGrad, + const TensorShape inShape, + const TensorShape outShape, + const FuncConfig& conf) { + std::vector crop_corner = + conf.get>("crop_corner"); + int cropC = crop_corner[1]; + int cropH = crop_corner[2]; + int cropW = crop_corner[3]; + + int num = outShape[0]; + int outC = outShape[1]; + int outH = outShape[2]; + int outW = outShape[3]; + + int inC = inShape[1]; + int inH = inShape[2]; + int inW = inShape[3]; + + size_t nth = num * inC * inH * inW; + int blockSize = 1024; + int gridSize = (nth + blockSize - 1) / blockSize; + + KeCropDiff <<>> + (inGrad, outGrad, inC, inH, inW, cropC, cropH, cropW, + outC, outH, outW, nth); + CHECK_SYNC("CropGrad"); +} + +} // namespace paddle diff --git a/paddle/function/CropOpTest.cpp b/paddle/function/CropOpTest.cpp new file mode 100644 index 0000000000..6f11abfdf6 --- /dev/null +++ b/paddle/function/CropOpTest.cpp @@ -0,0 +1,49 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#include +#include "FunctionTest.h" + +namespace paddle { + +TEST(Crop, real) { + for (size_t numSamples : {5, 32}) { + for (size_t channels : {5, 5, 32}) { + for (size_t imgSizeH : {5, 33, 100}) { + for (size_t imgSizeW : {5, 32, 96}) { + VLOG(3) << " numSamples=" << numSamples << " channels=" << channels + << " imgSizeH=" << imgSizeH << " imgSizeW=" << imgSizeW; + for (bool test_grad : {false, true}) { + CpuGpuFuncCompare compare( + test_grad ? "CropGrad" : "Crop", + FuncConfig() + .set>("crop_corner", {0, 1, 1, 1}) + .set>("crop_shape", {0, 2, 3, 3})); + TensorShape inDims{numSamples, channels, imgSizeH, imgSizeW}; + TensorShape outDims{numSamples, 2, 3, 3}; + compare.addInputs( + BufferArg(VALUE_TYPE_FLOAT, test_grad ? outDims : inDims)); + compare.addOutputs(BufferArg(VALUE_TYPE_FLOAT, + test_grad ? inDims : outDims, + test_grad ? ADD_TO : ASSIGN_TO), + test_grad ? ADD_TO : ASSIGN_TO); + compare.run(); + } + } + } + } + } +} + +} // namespace paddle diff --git a/paddle/function/DepthwiseConvOp.cpp b/paddle/function/DepthwiseConvOp.cpp new file mode 100644 index 0000000000..490e8d546c --- /dev/null +++ b/paddle/function/DepthwiseConvOp.cpp @@ -0,0 +1,306 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "DepthwiseConvOp.h" +#include "ConvOp.h" +#include "GemmFunctor.h" + +namespace paddle { + +template +class DepthwiseConvFunctor { +public: + void operator()(const T* inputData, + const T* filterData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* outputData) { + // TODO(zhaolong) : cpu implementation of depthwise convolution + } +}; + +template +class DepthwiseConvGradInputFunctor { +public: + void operator()(const T* outputGrad, + const T* filterData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* inputGrad) {} + // TODO(zhaolong) : cpu implementation of depthwise convolution +}; + +template +class DepthwiseConvGradFilterFunctor { +public: + void operator()(const T* outputGrad, + const T* inputData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* colData, + T* filterGrad) {} + // TODO(zhaolong) : cpu implementation of depthwise convolution +}; + +/* + * \brief Forward calculation of depthwise convolution. + */ +template +class DepthwiseConvFunction : public ConvFunctionBase { +public: + void init(const FuncConfig& config) override { + ConvFunctionBase::init(config); + } + + void check(const BufferArgs& inputs, const BufferArgs& outputs) override { + const TensorShape& input = inputs[0].shape(); + const TensorShape& filter = inputs[1].shape(); + const TensorShape& output = outputs[0].shape(); + checkShape(input, filter, output); + } + + void calc(const BufferArgs& inputs, const BufferArgs& outputs) override { + CHECK_EQ(numInputs_, inputs.size()); + CHECK_EQ(numOutputs_, outputs.size()); + check(inputs, outputs); + + const TensorShape& input = inputs[0].shape(); + const TensorShape& filter = inputs[1].shape(); + const TensorShape& output = outputs[0].shape(); + + size_t batchSize = input[0]; + size_t inputChannels = input[1]; + size_t inputHeight = input[2]; + size_t inputWidth = input[3]; + size_t filterHeight = getFilterHeight(filter); + size_t filterWidth = getFilterWidth(filter); + size_t outputChannels = output[1]; + size_t outputHeight = output[2]; + size_t outputWidth = output[3]; + size_t filterMultiplier = outputChannels / groups_; + CHECK_EQ(inputChannels, groups_); + + real* inputData = inputs[0].data(); + real* filterData = inputs[1].data(); + real* outputData = outputs[0].data(); + + DepthwiseConvFunctor depthwiseConv; + depthwiseConv(inputData, + filterData, + batchSize, + outputChannels, + outputHeight, + outputWidth, + inputChannels, + inputHeight, + inputWidth, + filterMultiplier, + filterHeight, + filterWidth, + strideH(), + strideW(), + paddingH(), + paddingW(), + outputData); + } +}; + +/* + * \brief Backward input calculation of depthwise convolution. + */ +template +class DepthwiseConvGradInputFunction : public ConvFunctionBase { +public: + void init(const FuncConfig& config) override { + ConvFunctionBase::init(config); + } + + void check(const BufferArgs& inputs, const BufferArgs& outputs) override { + const TensorShape& output = inputs[0].shape(); + const TensorShape& filter = inputs[1].shape(); + const TensorShape& input = outputs[0].shape(); + checkShape(input, filter, output); + } + + void calc(const BufferArgs& inputs, const BufferArgs& outputs) override { + CHECK_EQ(numInputs_, inputs.size()); + CHECK_EQ(numOutputs_, outputs.size()); + CHECK_EQ(outputs[0].getArgType(), ADD_TO); + check(inputs, outputs); + CHECK_EQ(outputs[0].getArgType(), ADD_TO); + const TensorShape& output = inputs[0].shape(); + const TensorShape& filter = inputs[1].shape(); + const TensorShape& input = outputs[0].shape(); + + size_t batchSize = input[0]; + size_t inputChannels = input[1]; + size_t inputHeight = input[2]; + size_t inputWidth = input[3]; + size_t filterHeight = getFilterHeight(filter); + size_t filterWidth = getFilterWidth(filter); + size_t outputChannels = output[1]; + size_t outputHeight = output[2]; + size_t outputWidth = output[3]; + size_t filterMultiplier = outputChannels / groups_; + CHECK_EQ(inputChannels, groups_); + + real* outputGrad = inputs[0].data(); + real* filterData = inputs[1].data(); + real* inputGrad = outputs[0].data(); + + DepthwiseConvGradInputFunctor depthwiseConvGradInput; + depthwiseConvGradInput(outputGrad, + filterData, + batchSize, + outputChannels, + outputHeight, + outputWidth, + inputChannels, + inputHeight, + inputWidth, + filterMultiplier, + filterHeight, + filterWidth, + strideH(), + strideW(), + paddingH(), + paddingW(), + inputGrad); + } +}; + +/* + * \brief Backward filter calculation of depthwise convolution. + */ +template +class DepthwiseConvGradFilterFunction : public ConvFunctionBase { +public: + void init(const FuncConfig& config) override { + ConvFunctionBase::init(config); + } + + void check(const BufferArgs& inputs, const BufferArgs& outputs) override { + const TensorShape& output = inputs[0].shape(); + const TensorShape& input = inputs[1].shape(); + const TensorShape& filter = outputs[0].shape(); + checkShape(input, filter, output); + } + + void calc(const BufferArgs& inputs, const BufferArgs& outputs) override { + CHECK_EQ(numInputs_, inputs.size()); + CHECK_EQ(numOutputs_, outputs.size()); + CHECK_EQ(outputs[0].getArgType(), ADD_TO); + check(inputs, outputs); + const TensorShape& output = inputs[0].shape(); + const TensorShape& input = inputs[1].shape(); + const TensorShape& filter = outputs[0].shape(); + + size_t batchSize = input[0]; + size_t inputChannels = input[1]; + size_t inputHeight = input[2]; + size_t inputWidth = input[3]; + size_t filterHeight = getFilterHeight(filter); + size_t filterWidth = getFilterWidth(filter); + size_t outputChannels = output[1]; + size_t outputHeight = output[2]; + size_t outputWidth = output[3]; + size_t filterMultiplier = outputChannels / groups_; + CHECK_EQ(inputChannels, groups_); + + real* outputGrad = inputs[0].data(); + real* inputData = inputs[1].data(); + real* filterGrad = outputs[0].data(); + + int size = outputChannels * filterHeight * filterWidth * outputHeight * + outputWidth; + resizeBuffer(size); + real* colData = reinterpret_cast(memory_->getBuf()); + + DepthwiseConvGradFilterFunctor depthwiseConvGradFilter; + + depthwiseConvGradFilter(outputGrad, + inputData, + batchSize, + outputChannels, + outputHeight, + outputWidth, + inputChannels, + inputHeight, + inputWidth, + filterMultiplier, + filterHeight, + filterWidth, + strideH(), + strideW(), + paddingH(), + paddingW(), + colData, + filterGrad); + } +}; + +REGISTER_TYPED_FUNC(DepthwiseConv, CPU, DepthwiseConvFunction); +REGISTER_TYPED_FUNC(DepthwiseConvGradInput, + CPU, + DepthwiseConvGradInputFunction); +REGISTER_TYPED_FUNC(DepthwiseConvGradFilter, + CPU, + DepthwiseConvGradFilterFunction); +#ifndef PADDLE_ONLY_CPU +REGISTER_TYPED_FUNC(DepthwiseConv, GPU, DepthwiseConvFunction); +REGISTER_TYPED_FUNC(DepthwiseConvGradInput, + GPU, + DepthwiseConvGradInputFunction); +REGISTER_TYPED_FUNC(DepthwiseConvGradFilter, + GPU, + DepthwiseConvGradFilterFunction); +#endif + +} // namespace paddle diff --git a/paddle/function/DepthwiseConvOp.h b/paddle/function/DepthwiseConvOp.h new file mode 100644 index 0000000000..1bf70e52f3 --- /dev/null +++ b/paddle/function/DepthwiseConvOp.h @@ -0,0 +1,159 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once + +#include "TensorType.h" + +namespace paddle { + +/** + *\brief Depthwise convolution forward. The outputData + * of depthwise convolution is same with ExpandConvLayer + * when groups equals inputChannels in ExpandConvLayer. + * + * \param[in] inputData input data. + * \param[in] filterData the Paramters of the depthwise conv layer.. + * \param[in] batchSize batch size of input data. + * \param[in] outputChannels channels of outputData. + * \param[in] outputHeight height of outputData. + * \param[in] outputWidth width of outputData. + * \param[in] inputChannels channels of inputData. + * \param[in] inputHeight height of inputData. + * \param[in] inputWidth width of inputData.. + * \param[in] filterMultiplier equals to outputChannels/groups_. + * \param[in] filterHeight height of filter. + * \param[in] filterWidth widht of filter. + * \param[in] strideH stride size in height direction. + * \param[in] strideW stride size in width direction. + * \param[in] paddingH padding size in height direction. + * \param[in] paddingW padding size in width direction. + * \param[out] outputData outputData. + * + */ +template +class DepthwiseConvFunctor { +public: + void operator()(const T* inputData, + const T* filterData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* outputData); +}; + +/** + *\brief Functor tot compute the depthwise convolution backprop w.r.t input. + * + * + * \param[in] outputGradData the grad data of output. + * \param[in] filterData the Paramters of the depthwise conv layer.. + * \param[in] batchSize batch size of input data. + * \param[in] outputChannels channels of outputData. + * \param[in] outputHeight height of outputData. + * \param[in] outputWidth width of outputData. + * \param[in] inputChannels channels of input data. + * \param[in] inputHeight height of inputData. + * \param[in] inputWidth width of inputData. + * \param[in] filterMultiplier equals to outputChannels/groups_. + * \param[in] filterHeight height of filter. + * \param[in] filterWidth widht of filter. + * \param[in] strideH stride size in height direction. + * \param[in] strideW stride size in width direction. + * \param[in] paddingH padding size in height direction. + * \param[in] paddingW padding size in width direction. + * \param[out] inputGrad the grad data of input. + * + */ +template +class DepthwiseConvGradInputFunctor { +public: + void operator()(const T* outputGrad, + const T* filterData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* inputGrad); +}; + +/** + *\brief Functor tot compute the depthwise convolution backprop w.r.t filter. + * + * \param[in] outputGradData the grad data of output. + * \param[in] inputData inputData. + * \param[in] batchSize batch size of input data. + * \param[in] outputChannels channels of outputData. + * \param[in] outputHeight height of outputData. + * \param[in] outputWidth width of outputData. + * \param[in] inputChannels channels of input data. + * \param[in] inputHeight height of inputData. + * \param[in] inputWidth width of inputData. + * \param[in] filterMultiplier equals to outputChannels/groups_. + * \param[in] filterHeight height of filter. + * \param[in] filterWidth widht of filter. + * \param[in] strideH stride size in height direction. + * \param[in] strideW stride size in width direction. + * \param[in] paddingH padding size in height direction. + * \param[in] paddingW padding size in width direction. + * \param[in] colData Auxiliary data when calculating filterGrad. + * \param[in] multiplierData Auxiliary data when calculating filterGrad. + * \param[out] filterGrad the grad data of filter. + * + */ +template +class DepthwiseConvGradFilterFunctor { +public: + void operator()(const T* outputGrad, + const T* inputData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* colData, + T* filterGrad); +}; + +} // namespace paddle diff --git a/paddle/function/DepthwiseConvOpGpu.cu b/paddle/function/DepthwiseConvOpGpu.cu new file mode 100644 index 0000000000..ede0d27aa8 --- /dev/null +++ b/paddle/function/DepthwiseConvOpGpu.cu @@ -0,0 +1,342 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "DepthwiseConvOp.h" +#include "GemmFunctor.h" +#include "paddle/math/BaseMatrix.h" + +namespace paddle { + +// CUDA kernel to compute the depthwise convolution forward pass +template +__global__ +void ConvolutionDepthwiseForward(const int nthreads, + const T* const inputData, const T* const filterData, + const int batchSize, const int outputChannels, const int outputHeight, + const int outputWidth, const int inputChannels, const int inputHeight, + const int inputWidth, const int filterMultiplier, const int filterHeight, + const int filterWidth, const int strideH, const int strideW, + const int paddingH, const int paddingW, T* const outputData) { + + int index = + (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x; + + if (index < nthreads) { + const int batch = index / outputChannels / outputHeight / outputWidth; + const int c_out = (index / outputHeight / outputWidth) % outputChannels; + const int h_out = (index / outputWidth) % outputHeight; + const int w_out = index % outputWidth; + + const int c_in = c_out / filterMultiplier; + const T* weight = filterData + c_out * filterHeight * filterWidth; + T value = 0; + const int h_in_start = -paddingH + h_out * strideH; + const int w_in_start = -paddingW + w_out * strideW; + const int h_in_end = -paddingH + h_out * strideH + filterHeight - 1; + const int w_in_end = -paddingW + w_out * strideW + filterWidth - 1; + if ((h_in_start >= 0) && (h_in_end < inputHeight) + && (w_in_start >= 0) && (w_in_end < inputWidth)) { + for (int kh = 0; kh < filterHeight; ++kh) { + for (int kw = 0; kw < filterWidth; ++kw) { + const int h_in = -paddingH + h_out * strideH + kh; + const int w_in = -paddingW + w_out * strideW + kw; + const int offset = ((batch * inputChannels + c_in) + * inputHeight + h_in) * inputWidth + w_in; + value += (*weight) * inputData[offset]; + ++weight; + } + } + } else { + for (int kh = 0; kh < filterHeight; ++kh) { + for (int kw = 0; kw < filterWidth; ++kw) { + const int h_in = -paddingH + h_out * strideH + kh; + const int w_in = -paddingW + w_out * strideW + kw; + if ((h_in >= 0) && (h_in < inputHeight) + && (w_in >= 0) && (w_in < inputWidth)) { + const int offset = ((batch * inputChannels + c_in) + * inputHeight + h_in) * inputWidth + w_in; + value += (*weight) * inputData[offset]; + } + ++weight; + } + } + } + outputData[index] = value; + } +} + +// CUDA kernel to compute the depthwise convolution backprop w.r.t input. +template +__global__ +void ConvolutionDepthwiseInputBackward(const int nthreads, + const T* const top_diff, const T* const weight_data, + const int num, const int outputChannels, const int outputHeight, + const int outputWidth, const int inputChannels, const int inputHeight, + const int inputWidth, const int filterMultiplier, const int filterHeight, + const int filterWidth, const int strideH, const int strideW, + const int paddingH, const int paddingW, T* const bottom_diff) { + int index = + (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x; + if (index < nthreads) { + const int batch = index / inputChannels / inputHeight / inputWidth; + const int c_in = (index / inputHeight / inputWidth) % inputChannels; + const int h_in = (index / inputWidth) % inputHeight; + const int w_in = index % inputWidth; + + const int c_out_start = c_in * filterMultiplier; + + int h_out_start = (h_in - filterHeight + paddingH + strideH)/strideH; + h_out_start = 0 > h_out_start ? 0 : h_out_start; + int h_out_end = (h_in + paddingH)/strideH; + h_out_end = outputHeight - 1 < h_out_end? outputHeight - 1 : h_out_end; + int w_out_start = (w_in - filterWidth + paddingW + strideW)/strideW; + w_out_start = 0 > w_out_start ? 0 : w_out_start; + int w_out_end = (w_in + paddingW)/strideW; + w_out_end = outputWidth - 1 < w_out_end? outputWidth - 1 : w_out_end; + + T value = 0; + + for (int c_out = c_out_start; + c_out < c_out_start + filterMultiplier; c_out ++) { + for (int h_out = h_out_start; h_out <= h_out_end; ++h_out) { + const int filter_h = h_in + paddingH - h_out * strideH; + for (int w_out = w_out_start; w_out <= w_out_end; ++w_out) { + const int filter_w = w_in + paddingW - w_out * strideW; + const int filter_offset = c_out * filterHeight * filterWidth + + filter_h * filterWidth + filter_w; + const int top_diff_offset = ((batch * outputChannels + c_out) * + outputHeight + h_out)* outputWidth + w_out; + value += top_diff[top_diff_offset] * weight_data[filter_offset]; + } + } + } + bottom_diff[index] += value; + } +} + +// CUDA kernel to compute the depthwise convolution backprop w.r.t filter. +template +__global__ +void ConvolutionDepthwiseFilterBackward(const int num_i, const int nthreads, + const T* const top_diff, const T* const inputData, + const int num, const int outputChannels, const int outputHeight, + const int outputWidth, const int inputChannels, const int inputHeight, + const int inputWidth, const int filterMultiplier, const int filterHeight, + const int filterWidth, const int strideH, const int strideW, + const int paddingH, const int paddingW, T* const buffer_data) { + int index = + (blockIdx.x * gridDim.y + blockIdx.y) * blockDim.x + threadIdx.x; + if (index < nthreads) { + const int h_out = (index / outputWidth) % outputHeight; + const int w_out = index % outputWidth; + const int kh = (index / filterWidth / outputHeight / outputWidth) + % filterHeight; + const int kw = (index / outputHeight / outputWidth) % filterWidth; + const int h_in = -paddingH + h_out * strideH + kh; + const int w_in = -paddingW + w_out * strideW + kw; + if ((h_in >= 0) && (h_in < inputHeight) + && (w_in >= 0) && (w_in < inputWidth)) { + const int c_out = index / + (filterHeight * filterWidth * outputHeight * outputWidth); + const int c_in = c_out / filterMultiplier; + const int batch = num_i; + const int top_offset = ((batch * outputChannels + c_out) * + outputHeight + h_out) * outputWidth + w_out; + const int bottom_offset = ((batch * inputChannels + c_in) + * inputHeight + h_in) * inputWidth + w_in; + buffer_data[index] = top_diff[top_offset] * inputData[bottom_offset]; + } else { + buffer_data[index] = 0; + } + } +} + +template +class DepthwiseConvFunctor{ +public: + void operator()(const T* inputData, + const T* filterData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* outputData){ + int outputSize = batchSize * outputChannels * outputHeight * outputWidth; + + size_t blocks = (outputSize + 1024 -1) / 1024; + size_t blockX = 512; + size_t blockY = (blocks+512-1)/512; + dim3 threads(1024, 1); + dim3 grid(blockX, blockY); + + ConvolutionDepthwiseForward + <<< grid, threads, 0, STREAM_DEFAULT >>>( + outputSize, + inputData, + filterData, + batchSize, + outputChannels, + outputHeight, + outputWidth, + inputChannels, + inputHeight, + inputWidth, + filterMultiplier, + filterHeight, + filterWidth, + strideH, + strideW, + paddingH, + paddingW, + outputData); + } +}; + +template +class DepthwiseConvGradInputFunctor{ +public: + void operator()(const T* outputGrad, + const T* filterData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* inputGrad){ + int inputSize = batchSize * inputChannels * inputHeight * inputWidth; + + size_t blocks = (inputSize + 1024 -1) / 1024; + size_t blockX = 512; + size_t blockY = (blocks+512-1)/512; + dim3 threads(1024, 1); + dim3 grid(blockX, blockY); + + + ConvolutionDepthwiseInputBackward + // NOLINT_NEXT_LINE(whitespace/operators) + <<< grid, threads, 0, STREAM_DEFAULT >>>( + inputSize, + outputGrad, + filterData, + batchSize, + outputChannels, + outputHeight, + outputWidth, + inputChannels, + inputHeight, + inputWidth, + filterMultiplier, + filterHeight, + filterWidth, + strideH, + strideW, + paddingH, + paddingW, + inputGrad); + } +}; + +template +class DepthwiseConvGradFilterFunctor { +public: + void operator()(const T* outputGrad, + const T* inputData, + int batchSize, + int outputChannels, + int outputHeight, + int outputWidth, + int inputChannels, + int inputHeight, + int inputWidth, + int filterMultiplier, + int filterHeight, + int filterWidth, + int strideH, + int strideW, + int paddingH, + int paddingW, + T* colData, + T* filterGrad){ + int colDataSize = outputChannels * filterHeight * filterWidth + * outputHeight * outputWidth; + + size_t blocks = (colDataSize + 1024 -1) / 1024; + size_t blockX = 512; + size_t blockY = (blocks+512-1)/512; + dim3 threads(1024, 1); + dim3 grid(blockX, blockY); + BaseMatrix filterGradMatrix(outputChannels * filterHeight * filterWidth, + 1, filterGrad, false, true); + + for (int i = 0; i < batchSize; i++) { + ConvolutionDepthwiseFilterBackward + <<< grid, threads, 0, STREAM_DEFAULT >>>( + i, + colDataSize, + outputGrad, + inputData, + batchSize, + outputChannels, + outputHeight, + outputWidth, + inputChannels, + inputHeight, + inputWidth, + filterMultiplier, + filterHeight, + filterWidth, + strideH, + strideW, + paddingH, + paddingW, + colData); + int K = outputHeight * outputWidth; + int M = colDataSize / K; + + BaseMatrix colMatrix(M, K, colData, false, true); + filterGradMatrix.sumRows(colMatrix, (T)1.0, (T)1.0); + } + } +}; + +#ifdef PADDLE_TYPE_DOUBLE +template class DepthwiseConvGradInputFunctor; +template class DepthwiseConvFunctor; +template class DepthwiseConvGradFilterFunctor; +#else +template class DepthwiseConvGradInputFunctor; +template class DepthwiseConvFunctor; +template class DepthwiseConvGradFilterFunctor; +#endif + +} // namespace paddle diff --git a/paddle/gserver/layers/CropLayer.cpp b/paddle/gserver/layers/CropLayer.cpp new file mode 100644 index 0000000000..69ad913420 --- /dev/null +++ b/paddle/gserver/layers/CropLayer.cpp @@ -0,0 +1,146 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "CropLayer.h" +#include "paddle/utils/Stat.h" +namespace paddle { + +REGISTER_LAYER(crop, CropLayer); + +bool CropLayer::init(const LayerMap& layerMap, + const ParameterMap& parameterMap) { + /* Initialize the basic parent class */ + Layer::init(layerMap, parameterMap); + CHECK_LE(static_cast(inputLayers_.size()), 2); + CHECK_GE(static_cast(inputLayers_.size()), 1); + crop_axis_ = config_.axis(); + for (int i = 0; i < config_.offset_size(); i++) { + crop_offsets_.push_back(config_.offset(i)); + } + + // 1. get input_0 shape + auto& input0_img_conf = config_.inputs(0).image_conf(); + inDims_ = TensorShape({0, + input0_img_conf.channels(), + input0_img_conf.has_img_size_y() + ? input0_img_conf.img_size_y() + : input0_img_conf.img_size(), + input0_img_conf.img_size()}); + // 2. get target dims from config + if (config_.inputs_size() == 1) { + targetDims_ = TensorShape({config_.shape(0), + config_.shape(1), + config_.shape(2), + config_.shape(3)}); + } else { + // 2. get input_1 shape + auto& input1_img_conf = config_.inputs(1).image_conf(); + targetDims_ = TensorShape({0, + input1_img_conf.channels(), + input1_img_conf.has_img_size_y() + ? input1_img_conf.img_size_y() + : input1_img_conf.img_size(), + input1_img_conf.img_size()}); + } + + // 3. get final crop corner + int dimSize = 4; + crop_corner_ = {0, 0, 0, 0}; + for (int i = 0; i < dimSize; i++) { + if (i >= crop_axis_) { + if (crop_offsets_.size() > 1) { + crop_corner_[i] = crop_offsets_[i - crop_axis_]; + } else { + crop_corner_[i] = crop_offsets_[0]; + } + } + } + + outDims_ = TensorShape(4); + + createFunction( + forward_, "Crop", FuncConfig().set("crop_corner", crop_corner_)); + createFunction( + backward_, "CropGrad", FuncConfig().set("crop_corner", crop_corner_)); + + return true; +} + +void CropLayer::setOutDims() { + MatrixPtr input = inputLayers_[1]->getOutputValue(); + size_t batchSize = input->getHeight(); + // get target dims from input_1 + if (config_.inputs_size() == 2) { + targetDims_.setDim(0, batchSize); + int ch = config_.inputs(0).image_conf().channels(); + if (ch != 0) targetDims_.setDim(1, ch); + int h = inputLayers_[1]->getOutput().getFrameHeight(); + if (h != 0) targetDims_.setDim(2, h); + int w = inputLayers_[1]->getOutput().getFrameWidth(); + if (w != 0) targetDims_.setDim(3, w); + } + // get final crop shape from target dims and crop axis + std::vector crop_shape; + int dimSize = 4; + for (int i = 0; i < dimSize; i++) { + if (i >= crop_axis_) { + crop_shape.push_back(targetDims_[i]); + } else { + crop_shape.push_back(inDims_[i]); + } + } + + outDims_.reshape( + {crop_shape[0], crop_shape[1], crop_shape[2], crop_shape[3]}); + output_.setFrameHeight(crop_shape[2]); + output_.setFrameWidth(crop_shape[3]); +} + +void CropLayer::setInDims() { + MatrixPtr input = inputLayers_[0]->getOutputValue(); + size_t batchSize = input->getHeight(); + inDims_.setDim(0, batchSize); + int h = inputLayers_[0]->getOutput().getFrameHeight(); + if (h != 0) inDims_.setDim(2, h); + int w = inputLayers_[0]->getOutput().getFrameWidth(); + if (w != 0) inDims_.setDim(3, w); +} + +void CropLayer::forward(PassType passType) { + Layer::forward(passType); + setInDims(); + setOutDims(); + int size = outDims_[1] * outDims_[2] * outDims_[3]; + resetOutput(outDims_[0], size); + MatrixPtr outV = getOutputValue(); + REGISTER_TIMER_INFO("CropForward", getName().c_str()); + + BufferArgs inputs; + BufferArgs outputs; + inputs.addArg(*getInputValue(0), inDims_); + outputs.addArg(*getOutputValue(), outDims_, ASSIGN_TO); + forward_[0]->calc(inputs, outputs); +} + +void CropLayer::backward(const UpdateCallback& callback) { + (void)callback; + REGISTER_TIMER_INFO("CropBackward", getName().c_str()); + + BufferArgs inputs; + BufferArgs outputs; + inputs.addArg(*getOutputGrad(), outDims_); + outputs.addArg(*getInputGrad(0), inDims_, ADD_TO); + backward_[0]->calc(inputs, outputs); +} +} // namespace paddle diff --git a/paddle/gserver/layers/CropLayer.h b/paddle/gserver/layers/CropLayer.h new file mode 100644 index 0000000000..6b62026210 --- /dev/null +++ b/paddle/gserver/layers/CropLayer.h @@ -0,0 +1,52 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once + +#include "Layer.h" + +namespace paddle { + +/** + * \brief This layer crop input according to the specify conf. + * input_0: input to be cropped + * input_1: optional reference input + * axis: start dimension to be croped + * offset: offset of cropping in each dimension + * shape: if reference input layer was not setted, + * crop input as this shape conf + */ +class CropLayer : public Layer { +public: + explicit CropLayer(const LayerConfig& config) : Layer(config) {} + + ~CropLayer() {} + + bool init(const LayerMap& layerMap, + const ParameterMap& parameterMap) override; + void forward(PassType passType) override; + void backward(const UpdateCallback& callback = nullptr) override; + +protected: + void setOutDims(); + void setInDims(); + + int32_t crop_axis_; + std::vector crop_offsets_; + std::vector crop_corner_; + TensorShape inDims_; + TensorShape targetDims_; + TensorShape outDims_; +}; +} // namespace paddle diff --git a/paddle/gserver/layers/ExpandConvLayer.cpp b/paddle/gserver/layers/ExpandConvLayer.cpp index af79e65a7c..783e02e47c 100644 --- a/paddle/gserver/layers/ExpandConvLayer.cpp +++ b/paddle/gserver/layers/ExpandConvLayer.cpp @@ -38,10 +38,25 @@ bool ExpandConvLayer::init(const LayerMap &layerMap, inputShape_.resize(numInputs); filterShape_.resize(numInputs); outputShape_.resize(numInputs); + + std::string convType; + std::string convGradInputType; + std::string convGradFilterType; + for (int i = 0; i < config_.inputs_size(); i++) { std::vector paddings = {(size_t)paddingY_[i], (size_t)padding_[i]}; std::vector strides = {(size_t)strideY_[i], (size_t)stride_[i]}; + if (useGpu_ && (size_t)groups_[i] == (size_t)channels_[i] && !isDeconv_) { + convType = "DepthwiseConv"; + convGradInputType = "DepthwiseConvGradInput"; + convGradFilterType = "DepthwiseConvGradFilter"; + } else { + convType = "GemmConv"; + convGradInputType = "GemmConvGradInput"; + convGradFilterType = "GemmConvGradFilter"; + } + if (FLAGS_use_nnpack) { CHECK_EQ(isDeconv_, false); createFunction(forward_, @@ -53,21 +68,21 @@ bool ExpandConvLayer::init(const LayerMap &layerMap, .set("algo", std::string("auto"))); } else { createFunction(forward_, - !isDeconv_ ? "GemmConv" : "GemmConvGradInput", + !isDeconv_ ? convType : convGradInputType, FuncConfig() .set("paddings", paddings) .set("strides", strides) .set("groups", (size_t)groups_[i])); createFunction(backward_, - !isDeconv_ ? "GemmConvGradInput" : "GemmConv", + !isDeconv_ ? convGradInputType : convType, FuncConfig() .set("paddings", paddings) .set("strides", strides) .set("groups", (size_t)groups_[i])); createFunction(backward_, - "GemmConvGradFilter", + convGradFilterType, FuncConfig() .set("paddings", paddings) .set("strides", strides) diff --git a/paddle/gserver/tests/CMakeLists.txt b/paddle/gserver/tests/CMakeLists.txt index 92f6cbcfe5..a43adc7ce7 100644 --- a/paddle/gserver/tests/CMakeLists.txt +++ b/paddle/gserver/tests/CMakeLists.txt @@ -56,7 +56,7 @@ add_test(NAME test_DetectionOutput add_unittest_without_exec(test_ConvUnify test_ConvUnify.cpp LayerGradUtil.cpp) - + add_test(NAME test_ConvUnify COMMAND test_ConvUnify) ################# test_BatchNorm ####################### diff --git a/paddle/gserver/tests/test_LayerGrad.cpp b/paddle/gserver/tests/test_LayerGrad.cpp index 67251f08e3..0975c3bc95 100644 --- a/paddle/gserver/tests/test_LayerGrad.cpp +++ b/paddle/gserver/tests/test_LayerGrad.cpp @@ -347,6 +347,55 @@ TEST(Layer, CosSimVecMatLayer) { } } +void testDepthwiseConvLayer(const string& type, bool useGpu) { + TestConfig config; + config.biasSize = 32; + config.layerConfig.set_type(type); + config.layerConfig.set_num_filters(32); + config.layerConfig.set_partial_sum(1); + config.layerConfig.set_shared_biases(true); + + config.inputDefs.push_back({INPUT_DATA, "layer_0", 2048, 192}); + LayerInputConfig* input = config.layerConfig.add_inputs(); + ConvConfig* conv = input->mutable_conv_conf(); + conv->set_filter_size(2); + conv->set_filter_size_y(3); + conv->set_channels(16); + conv->set_padding(0); + conv->set_padding_y(1); + conv->set_stride(2); + conv->set_stride_y(2); + conv->set_groups(16); + conv->set_filter_channels(conv->channels() / conv->groups()); + conv->set_img_size(16); + conv->set_img_size_y(8); + conv->set_output_x(outputSize(conv->img_size(), + conv->filter_size(), + conv->padding(), + conv->stride(), + /* caffeMode */ true)); + conv->set_output_y(outputSize(conv->img_size_y(), + conv->filter_size_y(), + conv->padding_y(), + conv->stride_y(), + /* caffeMode */ true)); + config.layerConfig.set_size(conv->output_x() * conv->output_y() * + config.layerConfig.num_filters()); + + testLayerGrad(config, "depthwise_conv", 100, false, useGpu); + // Use small batch_size and useWeight=true to test biasGrad + testLayerGrad(config, "depthwise_conv", 2, false, useGpu, true, 0.02); +} + +TEST(Layer, depthwiseConvLayer) { + // 'depthwise_conv' is a sepecial case of 'exconv' whose + // groups size equals to the input channels size. + testDepthwiseConvLayer("exconv", /* useGpu= */ false); +#ifndef PADDLE_ONLY_CPU + testDepthwiseConvLayer("exconv", /* useGpu= */ true); +#endif +} + void testConvLayer(const string& type, bool trans, bool useGpu) { TestConfig config; config.biasSize = 16; @@ -1802,6 +1851,34 @@ TEST(Layer, RowConvLayer) { } } +TEST(Layer, CropLayer) { + TestConfig config; + // config input_0 + config.inputDefs.push_back({INPUT_DATA, "layer_0", 1024, 0}); + LayerInputConfig* input = config.layerConfig.add_inputs(); + ImageConfig* img = input->mutable_image_conf(); + img->set_channels(4); + img->set_img_size(16); + config.layerConfig.set_axis(2); + config.layerConfig.add_offset(0); + config.layerConfig.add_offset(0); + + // config input_1 + config.inputDefs.push_back({INPUT_DATA, "layer_1", 128, 0}); + input = config.layerConfig.add_inputs(); + img = input->mutable_image_conf(); + img->set_channels(2); + img->set_img_size(8); + + // config crop layer + config.layerConfig.set_type("crop"); + config.layerConfig.set_name("cropLayer"); + + for (auto useGpu : {false, true}) { + testLayerGrad(config, "crop", 100, false, useGpu, false); + } +} + int main(int argc, char** argv) { testing::InitGoogleTest(&argc, argv); initMain(argc, argv); diff --git a/paddle/math/MathFunctions.cpp b/paddle/math/MathFunctions.cpp index 7045562dd4..c8ba1074a1 100644 --- a/paddle/math/MathFunctions.cpp +++ b/paddle/math/MathFunctions.cpp @@ -202,7 +202,7 @@ double dotProduct(const int n, const double* x, const double* y) { return cblas_ddot(n, x, 1, y, 1); } -#ifdef PADDLE_USE_MKL +#if defined(PADDLE_USE_MKL) || defined(PADDLE_USE_MKLML) template <> void vExp(const int n, const float* a, float* r) { @@ -243,7 +243,55 @@ template <> void vAdd(const int n, const double* a, const double* b, double* r) { vdAdd(n, a, b, r); } +#else + +DEFINE_MATRIX_BINARY_OP(vExp, b = std::exp(a)); +template +void vExp(const int n, const T* a, T* r) { + hl_cpu_apply_binary_op, 0, 0>( + binary::vExp(), const_cast(a), r, 1, n, n, n); +} + +DEFINE_MATRIX_BINARY_OP(vLog, b = std::log(a)); +template +void vLog(const int n, const T* a, T* r) { + hl_cpu_apply_binary_op, 0, 0>( + binary::vLog(), const_cast(a), r, 1, n, n, n); +} + +DEFINE_MATRIX_BINARY_PARAMETER_OP(vPow, ONE_PARAMETER, b = std::pow(a, p)); +template +void vPow(const int n, const T* a, const T b, T* r) { + hl_cpu_apply_binary_op, 0, 0>( + binary::vPow(b), const_cast(a), r, 1, n, n, n); +} + +DEFINE_MATRIX_TERNARY_OP(vAdd, c = a + b); +template +void vAdd(const int n, const T* a, const T* b, T* r) { + hl_cpu_apply_ternary_op, 0, 0>(ternary::vAdd(), + const_cast(a), + const_cast(b), + r, + 1, + n, + n, + n, + n); +} + +template void vExp(const int n, const float* a, float* r); +template void vExp(const int n, const double* a, double* r); +template void vLog(const int n, const float* a, float* r); +template void vLog(const int n, const double* a, double* r); +template void vPow(const int n, const float* a, const float b, float* r); +template void vPow(const int n, const double* a, const double b, double* r); +template void vAdd(const int n, const float* a, const float* b, float* r); +template void vAdd(const int n, const double* a, const double* b, double* r); +#endif + +#ifdef PADDLE_USE_MKL template <> void vInvSqrt(const int n, const float* a, float* r) { vsInvSqrt(n, a, r); @@ -275,20 +323,6 @@ void vTanh(const int n, const double* a, double* r) { } #else -DEFINE_MATRIX_BINARY_OP(vExp, b = std::exp(a)); -template -void vExp(const int n, const T* a, T* r) { - hl_cpu_apply_binary_op, 0, 0>( - binary::vExp(), const_cast(a), r, 1, n, n, n); -} - -DEFINE_MATRIX_BINARY_OP(vLog, b = std::log(a)); -template -void vLog(const int n, const T* a, T* r) { - hl_cpu_apply_binary_op, 0, 0>( - binary::vLog(), const_cast(a), r, 1, n, n, n); -} - DEFINE_MATRIX_BINARY_OP(vInvSqrt, b = 1.0f / std::sqrt(a)); template void vInvSqrt(const int n, const T* a, T* r) { @@ -312,41 +346,12 @@ void vTanh(const int n, const T* a, T* r) { binary::vTanh(), const_cast(a), r, 1, n, n, n); } -DEFINE_MATRIX_BINARY_PARAMETER_OP(vPow, ONE_PARAMETER, b = std::pow(a, p)); -template -void vPow(const int n, const T* a, const T b, T* r) { - hl_cpu_apply_binary_op, 0, 0>( - binary::vPow(b), const_cast(a), r, 1, n, n, n); -} - -DEFINE_MATRIX_TERNARY_OP(vAdd, c = a + b); -template -void vAdd(const int n, const T* a, const T* b, T* r) { - hl_cpu_apply_ternary_op, 0, 0>(ternary::vAdd(), - const_cast(a), - const_cast(b), - r, - 1, - n, - n, - n, - n); -} - -template void vExp(const int n, const float* a, float* r); -template void vExp(const int n, const double* a, double* r); -template void vLog(const int n, const float* a, float* r); -template void vLog(const int n, const double* a, double* r); template void vInvSqrt(const int n, const double* a, double* r); template void vInvSqrt(const int n, const float* a, float* r); template void vLog1p(const int n, const float* a, float* r); template void vLog1p(const int n, const double* a, double* r); template void vTanh(const int n, const float* a, float* r); template void vTanh(const int n, const double* a, double* r); -template void vPow(const int n, const float* a, const float b, float* r); -template void vPow(const int n, const double* a, const double b, double* r); -template void vAdd(const int n, const float* a, const float* b, float* r); -template void vAdd(const int n, const double* a, const double* b, double* r); #endif diff --git a/paddle/math/MathFunctions.h b/paddle/math/MathFunctions.h index 8ada0d34c6..637643838f 100644 --- a/paddle/math/MathFunctions.h +++ b/paddle/math/MathFunctions.h @@ -15,6 +15,12 @@ limitations under the License. */ #ifndef MATHFUNCTIONS_H_ #define MATHFUNCTIONS_H_ +#ifdef PADDLE_USE_MKLML +#include +#include +#include +#endif + #ifdef PADDLE_USE_MKL #include #include diff --git a/paddle/memory/CMakeLists.txt b/paddle/memory/CMakeLists.txt index fac442cca5..8035d93bfe 100644 --- a/paddle/memory/CMakeLists.txt +++ b/paddle/memory/CMakeLists.txt @@ -1,11 +1,16 @@ add_subdirectory(detail) cc_library(memory SRCS memory.cc) +cc_library(memcpy SRCS memcpy.cc DEPS device_context) cc_library(paddle_memory DEPS - memory meta_data - meta_cache memory_block - buddy_allocator system_allocator) + memory + memcpy + meta_data + meta_cache + memory_block + buddy_allocator + system_allocator) cc_test(memory_test SRCS memory_test.cc DEPS place paddle_memory) diff --git a/paddle/memory/README.md b/paddle/memory/README.md index 96a331a486..7f95e80f98 100644 --- a/paddle/memory/README.md +++ b/paddle/memory/README.md @@ -1,140 +1,4 @@ -## Design +# Region-based Heterogeneous Memory Management -### Usage - -To allocate 4KB CPU memory: - -```cpp -p = memory::Alloc(platform::CPUPlace(), 4*1024); -``` - -To allocate 4KB memory on the 3rd GPU: - -```cpp -p = memory::Alloc(platform::GPUPlace(2), 4*1024); -``` - -To free memory and check the so-far used amount of memory on a place: - -```cpp -auto pl = platform::GPUPlace(0); -p = memory::Alloc(pl, 4*1024); -cout << memory::Used(pl); -memory::Free(pl, p); -``` - -### API - -In `paddle/memory/memory.h` we have: - -```cpp -namespace memory { -template void* Alloc(Place, size_t); -template void Free(Place, void*); -template size_t Used(Place); -} // namespace memory -``` - -These function templates have specializations on either `platform::CPUPlace` or `platform::GPUPlace`: - -```cpp -template<> -void* Alloc(CPUPlace p, size_t size) { - return GetCPUBuddyAllocator()->Alloc(size); -} -``` - -and - -```cpp -template<> -void Alloc(GPUPlace p, size_t size) { - return GetGPUBuddyAllocator(p.id)->Alloc(size); -} -``` - -Similar specializations exist for `Free` and `Used`. - -### Implementation - -`GetCPUBuddyAllocator` and `GetGPUBuddyAllocator` are singletions. - -```cpp -BuddyAllocator* GetCPUBuddyAllocator() { - static BuddyAllocator* a = NULL; - if (a == NULL) { - a = new BuddyAllocator(new CPUAllocator /*backup allocator*/, ...); - } - return a; -} - -BuddyAllocator* GetGPUBuddyAllocator(int gpu_id) { - static BuddyAllocator* as = NULL; - if (as == NULL) { - as = new BuddyAllocator*[platform::NumGPUs()]; - for (int gpu = 0; gpu < platform::NumGPUs(); gpu++) { - as[gpu] = new BuddyAllocator(new GPUAllocator(gpu) /* backup allocator */, ...); - } - } - return as[gpu_id); -``` - -#### `BuddyAllocator` - -`BuddyAllocator` implements the buddy allocation algorithm. Its constructor takes parameters only related with the algorithm: - -```cpp -BuddyAllocator::BuddyAllocator(initial_pool_size, max_pool_size) { - ... -} -``` - -Please be aware that **`BuddyAllocator` always allocate aligned memory**, aligned on 32-bytes, which can hold a `BuddyAllocator::Block` object: - -```cpp -class BuddyAllocator { - private: - struct Block { - size_t size; - Block* left, right; - size_t index; // allocator id - }; - ... -}; -``` - -Because BuddyAllocator has the meta-data of each block, it can trace the used memory -- record the amount returned by `Alloc` freed in `Free`. Instead, `CPUAllocator` and `GPUAllocator` doesn't know the size of freed memory block and cannot do the trace. - -#### System Allocators - -The `GPUAllocator` and `CPUAllocator` are calls *system allocators*. They work as the fallback allocators of `BuddyAllocator`. - -## Justification - -I got inspiration from Majel and Caffe2, though above design look different from both. - -### Caffe2 - -In Caffe2, `Tensor::mutable_data()` allocates the memroy. In particular, [`Tensor::mutable_data`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/tensor.h#L523) calls [`Tensor::raw_mutable_data`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/tensor.h#L459), which in turn calls [`Context::New`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/tensor.h#L479). - -There are two implementations of `Context`: - -1. [`CPUContext`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/context.h#L105), whose [`New` method](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/context.h#L131) calls [`g_cpu_allocator.get()->New(size_t)`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/context.cc#L15) to allocate the memory. - -1. [`CUDAContext`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/context_gpu.h#L99), which has a data member [`int gpu_id_`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/context_gpu.h#L202). This looks very similar to class `majel::GPUPlace`, who also has an `int id_` data member. `CUDAContext::New(size_t)` calls [`g_cub_allocator->DeviceAllocate(&ptr, nbytes)`](https://github.com/caffe2/caffe2/blob/v0.7.0/caffe2/core/context_gpu.cu#L355) to allocate the memory. - -### Majel - -In Majel, there are basically two allocator types: - -1. `cpu::SystemAllocator`, which has similar functionality to `caffe2::CPUContext::New/Delete`. -1. `gpu::SystemAllocator`, which has similar functionality to `caffe2::CUDAContext::New/Delete`. - -However, memory allocation is not via these two allocators. Instead, these two allocators are defined in hidden namespaces. - -In Majel there are hidden global variables like: - -1. `cpu::SystemAllocator g_cpu_allocator`, and -1. `vector g_gpu_allocators(NUM_GPUS)`. - -Programs allocate memory via a BuddyAllocator, which can take the `g_cpu_allocator` or a `g_gpu_allocators[gpu_id]` as its *fallback allocator*, so that if BuddyAllocator cannot find a block in its memory pool, it extends its memory pool by calling the fallback allocator's `New(size_t)`. +Please check out the [design documentation](http://gangliao.me) to find out more details about +buddy memory allocator for both CPU and GPU. diff --git a/paddle/memory/detail/system_allocator.cc b/paddle/memory/detail/system_allocator.cc index 1579174b1a..f61e67a329 100644 --- a/paddle/memory/detail/system_allocator.cc +++ b/paddle/memory/detail/system_allocator.cc @@ -14,7 +14,7 @@ limitations under the License. */ #include "paddle/memory/detail/system_allocator.h" #include "paddle/platform/assert.h" -#include "paddle/platform/error.h" +#include "paddle/platform/enforce.h" #include "paddle/platform/gpu_info.h" #include // for malloc and free @@ -128,8 +128,7 @@ void GPUAllocator::Free(void* p, size_t size, size_t index) { // process is terminating, in which case we don't care if // cudaFree succeeds. if (err != cudaErrorCudartUnloading) { - platform::throw_on_error(err, - "cudaFree{Host} failed in GPUAllocator::Free."); + PADDLE_ENFORCE(err, "cudaFree{Host} failed in GPUAllocator::Free."); } } diff --git a/paddle/memory/memcpy.cc b/paddle/memory/memcpy.cc new file mode 100644 index 0000000000..098931c887 --- /dev/null +++ b/paddle/memory/memcpy.cc @@ -0,0 +1,70 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "paddle/memory/memcpy.h" + +#include // for memcpy + +#include "paddle/platform/device_context.h" + +namespace paddle { +namespace memory { + +template <> +void Copy(platform::CPUPlace, void* dst, + platform::CPUPlace, + const void* src, size_t num) { + std::memcpy(dst, src, num); +} + +#ifndef PADDLE_ONLY_CPU +template <> +void Copy(platform::CPUPlace dst_place, + void* dst, + platform::GPUPlace src_place, + const void* src, size_t num, + cudaStream_t stream) { + platform::GPUPlaceGuard g(src_place.device); + platform::GpuMemcpyAsync(dst, src, num, cudaMemcpyDeviceToHost, stream); +} + +template <> +void Copy(platform::GPUPlace dst_place, + void* dst, + platform::CPUPlace src_place, + const void* src, size_t num, + cudaStream_t stream) { + platform::GPUPlaceGuard g(dst_place.device); + platform::GpuMemcpyAsync(dst, src, num, cudaMemcpyHostToDevice, stream); +} + +template <> +void Copy(platform::GPUPlace dst_place, + void* dst, + platform::GPUPlace src_place, + const void* src, size_t num, + cudaStream_t stream) { + if (dst_place == src_place) { + platform::GPUPlaceGuard g(src_place.device); + platform::GpuMemcpyAsync(dst, src, num, cudaMemcpyDeviceToDevice, stream); + } else { + platform::GpuMemcpyPeer(dst, dst_place.device, src, src_place.device, num, + stream); + } +} + +#endif // PADDLE_ONLY_CPU + +} // namespace memory +} // namespace paddle diff --git a/paddle/memory/memcpy.h b/paddle/memory/memcpy.h new file mode 100644 index 0000000000..99b1c2e1c3 --- /dev/null +++ b/paddle/memory/memcpy.h @@ -0,0 +1,33 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once + +#include "paddle/platform/gpu_info.h" +#include "paddle/platform/place.h" + +namespace paddle { +namespace memory { + +template +void Copy(DstPlace, void* dst, SrcPlace, const void* src, size_t num); + +#ifndef PADDLE_ONLY_CPU +template +void Copy(DstPlace, void* dst, SrcPlace, const void* src, size_t num, + cudaStream_t stream); +#endif // PADDLE_ONLY_CPU + +} // namespace memory +} // namespace paddle diff --git a/paddle/memory/memory.cc b/paddle/memory/memory.cc index df3d57d629..c2e046926f 100644 --- a/paddle/memory/memory.cc +++ b/paddle/memory/memory.cc @@ -15,7 +15,8 @@ limitations under the License. */ #include "paddle/memory/memory.h" #include "paddle/memory/detail/buddy_allocator.h" #include "paddle/memory/detail/system_allocator.h" -#include "paddle/platform/assert.h" + +#include // for memcpy namespace paddle { namespace memory { diff --git a/paddle/memory/memory.h b/paddle/memory/memory.h index 2d6f4fd2a0..5e0d647072 100644 --- a/paddle/memory/memory.h +++ b/paddle/memory/memory.h @@ -14,19 +14,32 @@ limitations under the License. */ #pragma once +#include "paddle/platform/gpu_info.h" #include "paddle/platform/place.h" namespace paddle { namespace memory { -template +template void* Alloc(Place, size_t); -template +template void Free(Place, void*); -template +template size_t Used(Place); +template ::value>::type* = nullptr> +class PODDeleter { + public: + PODDeleter(Place place) : place_(place) {} + void operator()(T* ptr) { Free(place_, static_cast(ptr)); } + + private: + Place place_; +}; + } // namespace memory } // namespace paddle diff --git a/paddle/operators/CMakeLists.txt b/paddle/operators/CMakeLists.txt index bc64bfd7ec..0a14dc2114 100644 --- a/paddle/operators/CMakeLists.txt +++ b/paddle/operators/CMakeLists.txt @@ -48,6 +48,9 @@ op_library(mul_op SRCS mul_op.cc mul_op.cu) op_library(rowwise_add_op SRCS rowwise_add_op.cu rowwise_add_op.cc) op_library(sigmoid_op SRCS sigmoid_op.cu sigmoid_op.cc) op_library(softmax_op SRCS softmax_op.cc softmax_op.cu) +op_library(cross_entropy_op SRCS cross_entropy_op.cc cross_entropy_op.cu) op_library(fc_op SRCS fc_op.cc DEPS mul_op rowwise_add_op sigmoid_op softmax_op net) + +op_library(sgd_op SRCS sgd_op.cc sgd_op.cu) diff --git a/paddle/operators/add_op.cc b/paddle/operators/add_op.cc index 41d044cdb7..ebe9ceebe4 100644 --- a/paddle/operators/add_op.cc +++ b/paddle/operators/add_op.cc @@ -31,7 +31,7 @@ protected: "Inputs/Outputs of AddOp must all be set"); PADDLE_ENFORCE(inputs[0]->dims() == inputs[1]->dims(), "Two input of Add Op's dimension must be same."); - outputs[0]->set_dims(inputs[0]->dims()); + outputs[0]->Resize(inputs[0]->dims()); } }; @@ -53,6 +53,5 @@ The equation is: Out = X + Y } // namespace paddle REGISTER_OP(add_two, paddle::operators::AddOp, paddle::operators::AddOpMaker); -typedef paddle::operators::AddKernel<::paddle::platform::CPUPlace, float> - AddKernel_CPU_float; -REGISTER_OP_CPU_KERNEL(add_two, AddKernel_CPU_float); +REGISTER_OP_CPU_KERNEL( + add_two, paddle::operators::AddKernel); diff --git a/paddle/operators/add_op.cu b/paddle/operators/add_op.cu index 0edf142ee4..2e5a755f92 100644 --- a/paddle/operators/add_op.cu +++ b/paddle/operators/add_op.cu @@ -1,6 +1,5 @@ #include "paddle/operators/add_op.h" #include "paddle/framework/op_registry.h" -typedef paddle::operators::AddKernel<::paddle::platform::GPUPlace, float> AddKernel_GPU_float; REGISTER_OP_GPU_KERNEL(add_two, - AddKernel_GPU_float); \ No newline at end of file + paddle::operators::AddKernel); \ No newline at end of file diff --git a/paddle/operators/add_op.h b/paddle/operators/add_op.h index e08b3fb187..39d54a63bd 100644 --- a/paddle/operators/add_op.h +++ b/paddle/operators/add_op.h @@ -14,6 +14,7 @@ limitations under the License. */ #pragma once #include "glog/logging.h" +#include "paddle/framework/eigen.h" #include "paddle/framework/operator.h" namespace paddle { @@ -29,8 +30,10 @@ public: output->mutable_data(context.GetPlace()); - output->flat().device(*(context.GetEigenDevice())) = - input0.flat() + input1.flat(); + framework::EigenVector::Flatten(*output).device( + *(context.GetEigenDevice())) = + framework::EigenVector::Flatten(input0) + + framework::EigenVector::Flatten(input1); } }; diff --git a/paddle/operators/cross_entropy_op.cc b/paddle/operators/cross_entropy_op.cc new file mode 100644 index 0000000000..7d7bb09f3d --- /dev/null +++ b/paddle/operators/cross_entropy_op.cc @@ -0,0 +1,67 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "paddle/operators/cross_entropy_op.h" +#include "paddle/framework/op_registry.h" +#include "paddle/framework/tensor.h" + +namespace paddle { +namespace operators { + +class OnehotCrossEntropyOp : public framework::OperatorWithKernel { +protected: + void InferShape( + const std::vector &inputs, + const std::vector &outputs) const override { + PADDLE_ENFORCE(inputs.size() == 2, + "Input size of OnehotCrossEntropyOp must be two"); + PADDLE_ENFORCE(outputs.size() == 1, + "Output size of OnehotCrossEntropyOp must be one"); + PADDLE_ENFORCE(inputs[0] != nullptr && inputs[1] != nullptr, + "Inputs of OnehotCrossEntropyOp must all be set"); + PADDLE_ENFORCE(outputs[0] != nullptr, + "Outputs of OnehotCrossEntropyOp must all be set"); + PADDLE_ENFORCE(inputs[0]->dims().size() == 2, "X's dimension must be 2."); + PADDLE_ENFORCE(outputs[0]->dims().size() == 1, + "label's dimension must be 1."); + outputs[0]->Resize(framework::make_ddim({inputs[0]->dims()[0]})); + } +}; + +class OnehotCrossEntropyOpMaker : public framework::OpProtoAndCheckerMaker { +public: + OnehotCrossEntropyOpMaker(framework::OpProto *proto, + framework::OpAttrChecker *op_checker) + : framework::OpProtoAndCheckerMaker(proto, op_checker) { + AddInput("X", "The first input of OnehotCrossEntropyOp"); + AddInput("label", "The second input of OnehotCrossEntropyOp"); + AddOutput("Y", "The output of OnehotCrossEntropyOp"); + AddComment(R"DOC( +OnehotCrossEntropy Operator. + + Y[i] = -log(X[i][j]) + +)DOC"); + } +}; +} // namespace operators +} // namespace paddle + +REGISTER_OP(onehot_cross_entropy, + paddle::operators::OnehotCrossEntropyOp, + paddle::operators::OnehotCrossEntropyOpMaker); +REGISTER_OP_CPU_KERNEL( + onehot_cross_entropy, + paddle::operators::OnehotCrossEntropyOpKernel<::paddle::platform::CPUPlace, + float>); diff --git a/paddle/operators/cross_entropy_op.cu b/paddle/operators/cross_entropy_op.cu new file mode 100644 index 0000000000..1bcdcb7ea6 --- /dev/null +++ b/paddle/operators/cross_entropy_op.cu @@ -0,0 +1,6 @@ +#include "paddle/operators/cross_entropy_op.h" +#include "paddle/framework/op_registry.h" + +REGISTER_OP_GPU_KERNEL(onehot_cross_entropy, + paddle::operators::OnehotCrossEntropyOpKernel< + ::paddle::platform::GPUPlace, float>); \ No newline at end of file diff --git a/paddle/operators/cross_entropy_op.h b/paddle/operators/cross_entropy_op.h new file mode 100644 index 0000000000..ad2c7f34e1 --- /dev/null +++ b/paddle/operators/cross_entropy_op.h @@ -0,0 +1,50 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once +#include "glog/logging.h" +#include "paddle/framework/operator.h" + +namespace paddle { +namespace operators { + +template +class OnehotCrossEntropyOpKernel : public framework::OpKernel { +public: + constexpr T LOG_THRESHOLD() const { return static_cast(1e-20); } + + void Compute(const framework::KernelContext& context) const override { + auto X = context.Input(0)->Get(); + const T* X_data = X.data(); + const int* label_data = + context.Input(1)->Get().data(); + auto* Y = context.Output(0)->GetMutable(); + + Y->mutable_data(context.GetPlace()); + + T* Y_data = Y->data(); + + int batch_size = X.dims()[0]; + int class_num = X.dims()[1]; + + // Y[i] = -log(X[i][j]) + for (int i = 0; i < batch_size; ++i) { + Y_data[i] = -std::log( + std::max(X_data[i * class_num + label_data[i]], LOG_THRESHOLD())); + } + } +}; + +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/mul_op.cc b/paddle/operators/mul_op.cc index 713b2a5dc8..079a580080 100644 --- a/paddle/operators/mul_op.cc +++ b/paddle/operators/mul_op.cc @@ -12,9 +12,9 @@ See the License for the specific language governing permissions and limitations under the License. */ -#include -#include -#include +#include "paddle/operators/mul_op.h" +#include "paddle/framework/op_registry.h" +#include "paddle/framework/tensor.h" namespace paddle { namespace operators { @@ -33,7 +33,7 @@ protected: dim0[1] == dim1[0], "First matrix's width must be equal with second matrix's height."); PADDLE_ENFORCE(outputs.size() == 1, "The mul op must take one output"); - outputs[0]->set_dims({dim0[0], dim1[1]}); + outputs[0]->Resize({dim0[0], dim1[1]}); } }; @@ -57,4 +57,4 @@ The equation is: Out = X * Y REGISTER_OP(mul, paddle::operators::MulOp, paddle::operators::MulOpMaker); REGISTER_OP_CPU_KERNEL( - mul, paddle::operators::MulKernel); + mul, paddle::operators::MulKernel); diff --git a/paddle/operators/mul_op.cu b/paddle/operators/mul_op.cu index 201723df24..3ee581dc77 100644 --- a/paddle/operators/mul_op.cu +++ b/paddle/operators/mul_op.cu @@ -12,9 +12,9 @@ See the License for the specific language governing permissions and limitations under the License. */ -#include -#include +#include "paddle/operators/mul_op.h" +#include "paddle/framework/op_registry.h" REGISTER_OP_GPU_KERNEL(mul, paddle::operators::MulKernel); \ No newline at end of file + ::GPUPlace, float>); \ No newline at end of file diff --git a/paddle/operators/mul_op.h b/paddle/operators/mul_op.h index ce8a0169e0..e6bad7fb9d 100644 --- a/paddle/operators/mul_op.h +++ b/paddle/operators/mul_op.h @@ -14,17 +14,30 @@ #pragma once -#include -#include +#include "glog/logging.h" +#include "paddle/framework/eigen.h" +#include "paddle/framework/operator.h" namespace paddle { namespace operators { -template +template class MulKernel : public framework::OpKernel { public: - void Compute(const framework::KernelContext &context) const override { - LOG(INFO) << "Mul kernel in " << typeid(Place).name(); + void Compute(const framework::KernelContext& context) const override { + Eigen::array, 1> dim_pair = { + {Eigen::IndexPair(1, 0)}}; + + auto input0 = context.Input(0)->Get(); + auto input1 = context.Input(1)->Get(); + auto* output = context.Output(0)->GetMutable(); + + output->mutable_data(context.GetPlace()); + + framework::EigenMatrix::From(*output).device( + *(context.GetEigenDevice())) = + framework::EigenMatrix::From(input0).contract( + framework::EigenMatrix::From(input1), dim_pair); } }; } // namespace operators diff --git a/paddle/operators/rowwise_add_op.cc b/paddle/operators/rowwise_add_op.cc index 414bafd046..e04d69fa72 100644 --- a/paddle/operators/rowwise_add_op.cc +++ b/paddle/operators/rowwise_add_op.cc @@ -12,8 +12,8 @@ See the License for the specific language governing permissions and limitations under the License. */ -#include -#include +#include "paddle/operators/rowwise_add_op.h" +#include "paddle/framework/op_registry.h" namespace paddle { namespace operators { @@ -30,7 +30,7 @@ protected: PADDLE_ENFORCE(dim1.size() == 1, "The second input must be vector"); PADDLE_ENFORCE(dim0[1] == dim1[0], "The width of two input must be same"); PADDLE_ENFORCE(outputs.size() == 1, "The output size must be 1"); - outputs[0]->set_dims(inputs[0]->dims()); + outputs[0]->Resize(inputs[0]->dims()); } }; @@ -58,4 +58,4 @@ REGISTER_OP(rowwise_add, paddle::operators::RowWiseAddOpMaker); REGISTER_OP_CPU_KERNEL( rowwise_add, - paddle::operators::RowWiseAddKernel); + paddle::operators::RowWiseAddKernel); diff --git a/paddle/operators/rowwise_add_op.cu b/paddle/operators/rowwise_add_op.cu index 2c4bfbf93a..5dfac4fd2c 100644 --- a/paddle/operators/rowwise_add_op.cu +++ b/paddle/operators/rowwise_add_op.cu @@ -1,6 +1,6 @@ -#include -#include +#include "paddle/framework/op_registry.h" +#include "paddle/operators/rowwise_add_op.h" REGISTER_OP_GPU_KERNEL( rowwise_add, - paddle::operators::RowWiseAddKernel); + paddle::operators::RowWiseAddKernel); diff --git a/paddle/operators/rowwise_add_op.h b/paddle/operators/rowwise_add_op.h index 35f43e6376..dc47fe7c84 100644 --- a/paddle/operators/rowwise_add_op.h +++ b/paddle/operators/rowwise_add_op.h @@ -13,17 +13,32 @@ limitations under the License. */ #pragma once -#include -#include +#include "glog/logging.h" +#include "paddle/framework/eigen.h" +#include "paddle/framework/operator.h" namespace paddle { namespace operators { -template +template class RowWiseAddKernel : public framework::OpKernel { public: - void Compute(const framework::KernelContext &context) const override { - LOG(INFO) << "RowWiseAdd kernel in " << typeid(Place).name(); + void Compute(const framework::KernelContext& context) const override { + auto in0 = context.Input(0)->Get(); + auto in1 = context.Input(1)->Get(); + auto* out = context.Output(0)->GetMutable(); + out->mutable_data(context.GetPlace()); + + auto input = framework::EigenMatrix::From(in0); + auto bias = framework::EigenVector::From(in1); + auto output = framework::EigenMatrix::From(*out); + + const int bias_size = bias.dimension(0); + const int rest_size = input.size() / bias_size; + Eigen::DSizes one_d(input.size()); + Eigen::DSizes bcast(rest_size); + output.reshape(one_d).device(*(context.GetEigenDevice())) = + input.reshape(one_d) + bias.broadcast(bcast).reshape(one_d); } }; diff --git a/paddle/operators/sgd_op.cc b/paddle/operators/sgd_op.cc new file mode 100644 index 0000000000..66ab1e0011 --- /dev/null +++ b/paddle/operators/sgd_op.cc @@ -0,0 +1,61 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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 "paddle/operators/sgd_op.h" +#include "paddle/framework/op_registry.h" +#include "paddle/framework/tensor.h" + +namespace paddle { +namespace operators { + +class SGDOp : public framework::OperatorWithKernel { +protected: + void InferShape( + const std::vector &inputs, + const std::vector &outputs) const override { + PADDLE_ENFORCE(inputs.size() == 2, "Input size of SGDOp must be two"); + PADDLE_ENFORCE(outputs.size() == 1, "Output size of SGDOp must be one"); + PADDLE_ENFORCE(inputs[0] != nullptr, "inputs[0] mast be set"); + PADDLE_ENFORCE(inputs[1] != nullptr, "inputs[1] mast be set"); + PADDLE_ENFORCE(outputs[0] != nullptr, "outputs[0] mast be set"); + PADDLE_ENFORCE(inputs[0]->dims() == inputs[1]->dims(), + "Two input of SGD Op's dimension must be same."); + outputs[0]->Resize(inputs[0]->dims()); + } +}; + +class SGDOpMaker : public framework::OpProtoAndCheckerMaker { +public: + SGDOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker) + : framework::OpProtoAndCheckerMaker(proto, op_checker) { + AddInput("param", "input parameter"); + AddInput("grad", "input gradient"); + AddOutput("param_out", "output parameter"); + AddAttr("learning_rate", "learning rate of sgd"); + AddComment(R"DOC( + +Simplest sgd algorithm. + +param_out = param - learning_rate * grad; + +)DOC"); + } +}; +} // namespace operators +} // namespace paddle + +REGISTER_OP(sgd, paddle::operators::SGDOp, paddle::operators::SGDOpMaker); +typedef paddle::operators::SGDOpKernel<::paddle::platform::CPUPlace, float> + SGDOpKernel_CPU_float; +REGISTER_OP_CPU_KERNEL(sgd, SGDOpKernel_CPU_float); diff --git a/paddle/operators/sgd_op.cu b/paddle/operators/sgd_op.cu new file mode 100644 index 0000000000..400425db10 --- /dev/null +++ b/paddle/operators/sgd_op.cu @@ -0,0 +1,5 @@ +#include "paddle/operators/sgd_op.h" +#include "paddle/framework/op_registry.h" + +typedef paddle::operators::SGDOpKernel<::paddle::platform::GPUPlace, float> SGDOpKernel_GPU_float; +REGISTER_OP_GPU_KERNEL(sgd, SGDOpKernel_GPU_float); \ No newline at end of file diff --git a/paddle/operators/sgd_op.h b/paddle/operators/sgd_op.h new file mode 100644 index 0000000000..4b2d214618 --- /dev/null +++ b/paddle/operators/sgd_op.h @@ -0,0 +1,42 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once +#include "glog/logging.h" +#include "paddle/framework/eigen.h" +#include "paddle/framework/operator.h" + +namespace paddle { +namespace operators { + +template +class SGDOpKernel : public framework::OpKernel { +public: + void Compute(const framework::KernelContext& ctx) const override { + auto param = ctx.Input("param")->Get(); + auto grad = ctx.Input("grad")->Get(); + auto* param_out = ctx.Output(0)->GetMutable(); + float lr = ctx.op_.GetAttr("learning_rate"); + + param_out->mutable_data(ctx.GetPlace()); + + framework::EigenVector::Flatten(*param_out) + .device(*(ctx.GetEigenDevice())) = + framework::EigenVector::Flatten(param) - + lr * framework::EigenVector::Flatten(grad); + } +}; + +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/sgd_op_test.cc b/paddle/operators/sgd_op_test.cc new file mode 100644 index 0000000000..75137259f5 --- /dev/null +++ b/paddle/operators/sgd_op_test.cc @@ -0,0 +1,22 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#include +#include +USE_OP(sgd); +TEST(SGDOp, GetOpProto) { + auto& protos = paddle::framework::OpRegistry::protos(); + auto it = protos.find("sgd"); + ASSERT_NE(it, protos.end()); +} diff --git a/paddle/operators/sigmoid_op.cc b/paddle/operators/sigmoid_op.cc index 45ae277c53..91f7d86aeb 100644 --- a/paddle/operators/sigmoid_op.cc +++ b/paddle/operators/sigmoid_op.cc @@ -12,8 +12,8 @@ See the License for the specific language governing permissions and limitations under the License. */ -#include -#include +#include "paddle/operators/sigmoid_op.h" +#include "paddle/framework/op_registry.h" namespace paddle { namespace operators { @@ -24,7 +24,7 @@ protected: const std::vector &outputs) const override { PADDLE_ENFORCE(inputs.size() == 1, "Sigmoid Op only have one input"); PADDLE_ENFORCE(outputs.size() == 1, "Sigmoid Op only have one output"); - outputs[0]->set_dims(inputs[0]->dims()); + outputs[0]->Resize(inputs[0]->dims()); } }; @@ -34,7 +34,7 @@ public: framework::OpAttrChecker *op_checker) : framework::OpProtoAndCheckerMaker(proto, op_checker) { AddInput("X", "sigmoid input"); - AddInput("Y", "sigmoid output"); + AddOutput("Y", "sigmoid output"); AddComment("Sigmoid function"); } }; @@ -46,4 +46,5 @@ REGISTER_OP(sigmoid, paddle::operators::SigmoidOp, paddle::operators::SigmoidOpMaker); REGISTER_OP_CPU_KERNEL( - sigmoid, paddle::operators::SigmoidKernel); + sigmoid, + paddle::operators::SigmoidKernel); diff --git a/paddle/operators/sigmoid_op.cu b/paddle/operators/sigmoid_op.cu index 79d5222348..ed344b2bfd 100644 --- a/paddle/operators/sigmoid_op.cu +++ b/paddle/operators/sigmoid_op.cu @@ -1,5 +1,5 @@ -#include -#include +#include "paddle/operators/sigmoid_op.h" +#include "paddle/framework/op_registry.h" REGISTER_OP_GPU_KERNEL( - sigmoid, paddle::operators::SigmoidKernel); + sigmoid, paddle::operators::SigmoidKernel); diff --git a/paddle/operators/sigmoid_op.h b/paddle/operators/sigmoid_op.h index 42173343f3..2b9356246c 100644 --- a/paddle/operators/sigmoid_op.h +++ b/paddle/operators/sigmoid_op.h @@ -14,17 +14,25 @@ #pragma once -#include -#include +#include "glog/logging.h" +#include "paddle/framework/eigen.h" +#include "paddle/framework/operator.h" namespace paddle { namespace operators { -template +template class SigmoidKernel : public framework::OpKernel { public: - void Compute(const framework::KernelContext &context) const override { - LOG(INFO) << "Sigmoid kernel in " << typeid(Place).name(); + void Compute(const framework::KernelContext& context) const override { + auto input = context.Input(0)->Get(); + auto* output = context.Output(0)->GetMutable(); + + output->mutable_data(context.GetPlace()); + + framework::EigenVector::Flatten(*output).device( + *(context.GetEigenDevice())) = + 1.0 / (1.0 + (-1.0 * framework::EigenVector::Flatten(input)).exp()); } }; } // namespace operators diff --git a/paddle/operators/softmax_op.cc b/paddle/operators/softmax_op.cc index 4ca7be359e..cf5e273de6 100644 --- a/paddle/operators/softmax_op.cc +++ b/paddle/operators/softmax_op.cc @@ -11,8 +11,8 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. */ -#include -#include +#include "paddle/operators/softmax_op.h" +#include "paddle/framework/op_registry.h" namespace paddle { namespace operators { @@ -23,9 +23,11 @@ protected: const std::vector &inputs, const std::vector &outputs) const override { PADDLE_ENFORCE(inputs.size() == 1, "Only one input is need for softmax"); + PADDLE_ENFORCE(inputs[0]->dims().size() == 2, + "The input of softmax op must be matrix"); PADDLE_ENFORCE(outputs.size() == 1, "Only one output is need for softmax"); - outputs[0]->set_dims(inputs[0]->dims()); + outputs[0]->Resize(inputs[0]->dims()); } }; @@ -46,4 +48,5 @@ public: namespace ops = paddle::operators; REGISTER_OP(softmax, ops::SoftmaxOp, ops::SoftmaxOpMaker); -REGISTER_OP_CPU_KERNEL(softmax, ops::SoftmaxKernel); +REGISTER_OP_CPU_KERNEL(softmax, + ops::SoftmaxKernel); diff --git a/paddle/operators/softmax_op.cu b/paddle/operators/softmax_op.cu index 903eef1b62..60676191eb 100644 --- a/paddle/operators/softmax_op.cu +++ b/paddle/operators/softmax_op.cu @@ -1,5 +1,5 @@ -#include -#include +#include "paddle/framework/op_registry.h" +#include "paddle/operators/softmax_op.h" REGISTER_OP_GPU_KERNEL( - softmax, paddle::operators::SoftmaxKernel); + softmax, paddle::operators::SoftmaxKernel); diff --git a/paddle/operators/softmax_op.h b/paddle/operators/softmax_op.h index 74e9e2786b..500c188dbf 100644 --- a/paddle/operators/softmax_op.h +++ b/paddle/operators/softmax_op.h @@ -14,17 +14,49 @@ #pragma once -#include -#include +#include "glog/logging.h" +#include "paddle/framework/eigen.h" +#include "paddle/framework/operator.h" namespace paddle { namespace operators { -template +template class SoftmaxKernel : public framework::OpKernel { public: - void Compute(const framework::KernelContext &context) const override { - LOG(INFO) << "Softmax kernel in " << typeid(Place).name(); + void Compute(const framework::KernelContext& context) const override { + auto input = context.Input(0)->Get(); + auto* output = context.Output(0)->GetMutable(); + output->mutable_data(context.GetPlace()); + + auto logits = framework::EigenMatrix::From(input); + auto softmax = framework::EigenMatrix::From(*output); + + const int kBatchDim = 0; + const int kClassDim = 1; + + const int batch_size = logits.dimension(kBatchDim); + const int num_classes = logits.dimension(kClassDim); + + Eigen::DSizes along_class(kClassDim); + Eigen::DSizes batch_by_one(batch_size, 1); + Eigen::DSizes one_by_class(1, num_classes); + + auto shifted_logits = (logits - + logits.maximum(along_class) + .eval() + .reshape(batch_by_one) + .broadcast(one_by_class)); + + softmax.device(*(context.GetEigenDevice())) = shifted_logits.exp(); + + softmax.device(*(context.GetEigenDevice())) = + (softmax * + softmax.sum(along_class) + .inverse() + .eval() + .reshape(batch_by_one) + .broadcast(one_by_class)); } }; } // namespace operators diff --git a/paddle/platform/CMakeLists.txt b/paddle/platform/CMakeLists.txt index 6ac4035c0f..bd77bb7daa 100644 --- a/paddle/platform/CMakeLists.txt +++ b/paddle/platform/CMakeLists.txt @@ -8,6 +8,8 @@ cc_test(place_test SRCS place_test.cc DEPS place glog gflags) add_subdirectory(dynload) +cc_test(enforce_test SRCS enforce_test.cc) + IF(WITH_GPU) set(GPU_CTX_DEPS dynload_cuda dynamic_loader) ELSE() diff --git a/paddle/platform/cpu_info.cc b/paddle/platform/cpu_info.cc index dfab391cfb..78e1fa9df5 100644 --- a/paddle/platform/cpu_info.cc +++ b/paddle/platform/cpu_info.cc @@ -22,7 +22,6 @@ limitations under the License. */ #endif #include "gflags/gflags.h" -#include "paddle/platform/error.h" DEFINE_double(fraction_of_cpu_memory_to_use, 1, "Default use 100% of CPU memory for PaddlePaddle," diff --git a/paddle/platform/device_context.h b/paddle/platform/device_context.h index f226a75c20..fe6f13e399 100644 --- a/paddle/platform/device_context.h +++ b/paddle/platform/device_context.h @@ -11,12 +11,13 @@ limitations under the License. */ #pragma once -#include "paddle/framework/enforce.h" +#include "paddle/platform/enforce.h" +#include "paddle/platform/place.h" + #ifndef PADDLE_ONLY_CPU #include "paddle/platform/dynload/cublas.h" #include "paddle/platform/dynload/cudnn.h" #include "paddle/platform/dynload/curand.h" -#include "paddle/platform/error.h" #include "paddle/platform/gpu_info.h" #define EIGEN_USE_GPU #endif @@ -71,8 +72,7 @@ class CUDADeviceContext : public DeviceContext { public: explicit CUDADeviceContext(const GPUPlace gpu_place) : gpu_place_(gpu_place) { GPUPlaceGuard guard(gpu_place_); - paddle::platform::throw_on_error(cudaStreamCreate(&stream_), - "cudaStreamCreate failed"); + PADDLE_ENFORCE(cudaStreamCreate(&stream_), "cudaStreamCreate failed"); eigen_stream_.reset(new Eigen::CudaStreamDevice(&stream_)); eigen_device_.reset(new Eigen::GpuDevice(eigen_stream_.get())); } @@ -83,8 +83,8 @@ class CUDADeviceContext : public DeviceContext { } void Wait() { - paddle::platform::throw_on_error(cudaStreamSynchronize(stream_), - "cudaStreamSynchronize failed"); + PADDLE_ENFORCE(cudaStreamSynchronize(stream_), + "cudaStreamSynchronize failed"); } cudaStream_t stream() { return stream_; } @@ -94,12 +94,11 @@ class CUDADeviceContext : public DeviceContext { cublasHandle_t cublas_handle() { if (!blas_handle_) { GPUPlaceGuard guard(gpu_place_); - PADDLE_ENFORCE(paddle::platform::dynload::cublasCreate(&blas_handle_) == - CUBLAS_STATUS_SUCCESS, + PADDLE_ENFORCE(paddle::platform::dynload::cublasCreate(&blas_handle_), "cublasCreate failed"); - PADDLE_ENFORCE(paddle::platform::dynload::cublasSetStream( - blas_handle_, stream_) == CUBLAS_STATUS_SUCCESS, - "cublasSetStream failed"); + PADDLE_ENFORCE( + paddle::platform::dynload::cublasSetStream(blas_handle_, stream_), + "cublasSetStream failed"); } return blas_handle_; } @@ -107,12 +106,11 @@ class CUDADeviceContext : public DeviceContext { cudnnHandle_t cudnn_handle() { if (!dnn_handle_) { GPUPlaceGuard guard(gpu_place_); - PADDLE_ENFORCE(paddle::platform::dynload::cudnnCreate(&dnn_handle_) == - CUDNN_STATUS_SUCCESS, + PADDLE_ENFORCE(paddle::platform::dynload::cudnnCreate(&dnn_handle_), "cudnnCreate failed"); - PADDLE_ENFORCE(paddle::platform::dynload::cudnnSetStream( - dnn_handle_, stream_) == CUDNN_STATUS_SUCCESS, - "cudnnSetStream failed"); + PADDLE_ENFORCE( + paddle::platform::dynload::cudnnSetStream(dnn_handle_, stream_), + "cudnnSetStream failed"); } return dnn_handle_; } @@ -121,16 +119,15 @@ class CUDADeviceContext : public DeviceContext { if (!rand_generator_) { GPUPlaceGuard guard(gpu_place_); PADDLE_ENFORCE(paddle::platform::dynload::curandCreateGenerator( - &rand_generator_, CURAND_RNG_PSEUDO_DEFAULT) == - CURAND_STATUS_SUCCESS, + &rand_generator_, CURAND_RNG_PSEUDO_DEFAULT), "curandCreateGenerator failed"); PADDLE_ENFORCE( paddle::platform::dynload::curandSetPseudoRandomGeneratorSeed( - rand_generator_, random_seed_) == CURAND_STATUS_SUCCESS, + rand_generator_, random_seed_), "curandSetPseudoRandomGeneratorSeed failed"); - PADDLE_ENFORCE(paddle::platform::dynload::curandSetStream( - rand_generator_, stream_) == CURAND_STATUS_SUCCESS, - "curandSetStream failed"); + PADDLE_ENFORCE( + paddle::platform::dynload::curandSetStream(rand_generator_, stream_), + "curandSetStream failed"); } return rand_generator_; } @@ -138,26 +135,23 @@ class CUDADeviceContext : public DeviceContext { ~CUDADeviceContext() { Wait(); if (blas_handle_) { - PADDLE_ENFORCE(paddle::platform::dynload::cublasDestroy(blas_handle_) == - CUBLAS_STATUS_SUCCESS, + PADDLE_ENFORCE(paddle::platform::dynload::cublasDestroy(blas_handle_), "cublasDestroy failed"); } if (dnn_handle_) { - PADDLE_ENFORCE(paddle::platform::dynload::cudnnDestroy(dnn_handle_) == - CUDNN_STATUS_SUCCESS, + PADDLE_ENFORCE(paddle::platform::dynload::cudnnDestroy(dnn_handle_), "cudnnDestroy failed"); } if (rand_generator_) { - PADDLE_ENFORCE(paddle::platform::dynload::curandDestroyGenerator( - rand_generator_) == CURAND_STATUS_SUCCESS, - "curandDestroyGenerator failed"); + PADDLE_ENFORCE( + paddle::platform::dynload::curandDestroyGenerator(rand_generator_), + "curandDestroyGenerator failed"); } eigen_stream_.reset(); eigen_device_.reset(); - paddle::platform::throw_on_error(cudaStreamDestroy(stream_), - "cudaStreamDestroy failed"); + PADDLE_ENFORCE(cudaStreamDestroy(stream_), "cudaStreamDestroy failed"); } private: diff --git a/paddle/platform/dynload/dynamic_loader.cc b/paddle/platform/dynload/dynamic_loader.cc index dd914e006d..ae9a0a982c 100644 --- a/paddle/platform/dynload/dynamic_loader.cc +++ b/paddle/platform/dynload/dynamic_loader.cc @@ -19,7 +19,7 @@ limitations under the License. */ #include #include "gflags/gflags.h" #include "glog/logging.h" -#include "paddle/framework/enforce.h" +#include "paddle/platform/enforce.h" DEFINE_string(cudnn_dir, "", "Specify path for loading libcudnn.so. For instance, " diff --git a/paddle/platform/enforce.h b/paddle/platform/enforce.h new file mode 100644 index 0000000000..b06ab8a2f1 --- /dev/null +++ b/paddle/platform/enforce.h @@ -0,0 +1,148 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +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. */ + +#pragma once + +#include +#include +#include +#include + +#ifndef PADDLE_ONLY_CPU + +#include "paddle/platform/dynload/cublas.h" +#include "paddle/platform/dynload/cudnn.h" +#include "paddle/platform/dynload/curand.h" + +#include +#include +#include +#include +#include + +#endif // PADDLE_ONLY_CPU + +namespace paddle { +namespace platform { + +// Because most enforce conditions would evaluate to true, we can use +// __builtin_expect to instruct the C++ compiler to generate code that +// always forces branch prediction of true. +// This generates faster binary code. __builtin_expect is since C++11. +// For more details, please check https://stackoverflow.com/a/43870188/724872. +#define UNLIKELY(condition) __builtin_expect(static_cast(condition), 0) + +template +inline void throw_on_error(T e) { + throw_on_error(e, ""); +} + +template +inline typename std::enable_if::type throw_on_error( + int stat, const Args&... args) { + if (UNLIKELY(!(stat))) { + throw std::runtime_error( + string::Sprintf(args...) + + string::Sprintf(" at [%s:%s];", __FILE__, __LINE__)); + } +} + +#ifndef PADDLE_ONLY_CPU + +template +inline typename std::enable_if::type throw_on_error( + cudaError_t e, const Args&... args) { + if (UNLIKELY(e)) { + // clang-format off + throw thrust::system_error( + e, thrust::cuda_category(), + string::Sprintf(args...) + + string::Sprintf(" at [%s:%s];", __FILE__, __LINE__)); + // clang-format on + } +} + +template +inline typename std::enable_if::type throw_on_error( + curandStatus_t stat, const Args&... args) { + if (stat != CURAND_STATUS_SUCCESS) { + // clang-format off + throw thrust::system_error( + cudaErrorLaunchFailure, thrust::cuda_category(), + string::Sprintf(args...) + + string::Sprintf(" at [%s:%s];", __FILE__, __LINE__)); + // clang-format on + } +} + +template +inline typename std::enable_if::type throw_on_error( + cudnnStatus_t stat, const Args&... args) { + if (stat == CUDNN_STATUS_SUCCESS) { + return; + } else { + // clang-format off + throw std::runtime_error( + platform::dynload::cudnnGetErrorString(stat) + + string::Sprintf(args...) + + string::Sprintf(" at [%s:%s];", __FILE__, __LINE__)); + // clang-format on + } +} + +template +inline typename std::enable_if::type throw_on_error( + cublasStatus_t stat, const Args&... args) { + std::string err; + if (stat == CUBLAS_STATUS_SUCCESS) { + return; + } else if (stat == CUBLAS_STATUS_NOT_INITIALIZED) { + err = "CUBLAS: not initialized, "; + } else if (stat == CUBLAS_STATUS_ALLOC_FAILED) { + err = "CUBLAS: alloc failed, "; + } else if (stat == CUBLAS_STATUS_INVALID_VALUE) { + err = "CUBLAS: invalid value, "; + } else if (stat == CUBLAS_STATUS_ARCH_MISMATCH) { + err = "CUBLAS: arch mismatch, "; + } else if (stat == CUBLAS_STATUS_MAPPING_ERROR) { + err = "CUBLAS: mapping error, "; + } else if (stat == CUBLAS_STATUS_EXECUTION_FAILED) { + err = "CUBLAS: execution failed, "; + } else if (stat == CUBLAS_STATUS_INTERNAL_ERROR) { + err = "CUBLAS: internal error, "; + } else if (stat == CUBLAS_STATUS_NOT_SUPPORTED) { + err = "CUBLAS: not supported, "; + } else if (stat == CUBLAS_STATUS_LICENSE_ERROR) { + err = "CUBLAS: license error, "; + } + throw std::runtime_error(err + string::Sprintf(args...) + + string::Sprintf(" at [%s:%s];", __FILE__, __LINE__)); +} + +#endif // PADDLE_ONLY_CPU + +#define PADDLE_THROW(...) \ + do { \ + throw std::runtime_error( \ + string::Sprintf(__VA_ARGS__) + \ + string::Sprintf(" at [%s:%s];", __FILE__, __LINE__)); \ + } while (0) + +#define PADDLE_ENFORCE(...) \ + do { \ + ::paddle::platform::throw_on_error(__VA_ARGS__); \ + } while (0) + +} // namespace platform +} // namespace paddle diff --git a/paddle/framework/enforce_test.cc b/paddle/platform/enforce_test.cc similarity index 85% rename from paddle/framework/enforce_test.cc rename to paddle/platform/enforce_test.cc index f8da1a192f..d7152f8150 100644 --- a/paddle/framework/enforce_test.cc +++ b/paddle/platform/enforce_test.cc @@ -9,8 +9,8 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. */ -#include -#include +#include "paddle/platform/enforce.h" +#include "gtest/gtest.h" TEST(ENFORCE, OK) { PADDLE_ENFORCE(true, "Enforce is ok %d now %f", 123, 0.345); @@ -23,13 +23,14 @@ TEST(ENFORCE, FAILED) { bool in_catch = false; try { PADDLE_ENFORCE(false, "Enforce is not ok %d at all", 123); - } catch (paddle::framework::EnforceNotMet err) { + } catch (const std::runtime_error& error) { + // your error handling code here in_catch = true; std::string msg = "Enforce is not ok 123 at all"; - const char* what = err.what(); + const char* what = error.what(); for (size_t i = 0; i < msg.length(); ++i) { ASSERT_EQ(what[i], msg[i]); } } ASSERT_TRUE(in_catch); -} \ No newline at end of file +} diff --git a/paddle/platform/error.h b/paddle/platform/error.h deleted file mode 100644 index 93424bb610..0000000000 --- a/paddle/platform/error.h +++ /dev/null @@ -1,87 +0,0 @@ -#pragma once - -#include -#include -#include - -#ifndef PADDLE_ONLY_CPU - -#include -#include -#include -#include -#include - -#endif // PADDLE_ONLY_CPU - -namespace paddle { -namespace platform { - -#ifndef PADDLE_ONLY_CPU - -inline void throw_on_error(cudaError_t e, const char* message) { - if (e) { - throw thrust::system_error(e, thrust::cuda_category(), message); - } -} - -inline void throw_on_error(curandStatus_t stat, const char* message) { - if (stat != CURAND_STATUS_SUCCESS) { - throw thrust::system_error(cudaErrorLaunchFailure, thrust::cuda_category(), - message); - } -} - -inline void throw_on_error(cudnnStatus_t stat, const char* message) { - std::stringstream ss; - if (stat == CUDNN_STATUS_SUCCESS) { - return; - } else { - ss << cudnnGetErrorString(stat); - ss << ", " << message; - throw std::runtime_error(ss.str()); - } -} - -inline void throw_on_error(cublasStatus_t stat, const char* message) { - std::stringstream ss; - if (stat == CUBLAS_STATUS_SUCCESS) { - return; - } else if (stat == CUBLAS_STATUS_NOT_INITIALIZED) { - ss << "CUBLAS: not initialized"; - } else if (stat == CUBLAS_STATUS_ALLOC_FAILED) { - ss << "CUBLAS: alloc failed"; - } else if (stat == CUBLAS_STATUS_INVALID_VALUE) { - ss << "CUBLAS: invalid value"; - } else if (stat == CUBLAS_STATUS_ARCH_MISMATCH) { - ss << "CUBLAS: arch mismatch"; - } else if (stat == CUBLAS_STATUS_MAPPING_ERROR) { - ss << "CUBLAS: mapping error"; - } else if (stat == CUBLAS_STATUS_EXECUTION_FAILED) { - ss << "CUBLAS: execution failed"; - } else if (stat == CUBLAS_STATUS_INTERNAL_ERROR) { - ss << "CUBLAS: internal error"; - } else if (stat == CUBLAS_STATUS_NOT_SUPPORTED) { - ss << "CUBLAS: not supported"; - } else if (stat == CUBLAS_STATUS_LICENSE_ERROR) { - ss << "CUBLAS: license error"; - } - ss << ", " << message; - throw std::runtime_error(ss.str()); -} - -inline void throw_on_error(cublasStatus_t stat) { - const char* message = ""; - throw_on_error(stat, message); -} - -#endif // PADDLE_ONLY_CPU - -inline void throw_on_error(int stat, const char* message) { - if (stat) { - throw std::runtime_error(message + (", stat = " + std::to_string(stat))); - } -} - -} // namespace platform -} // namespace paddle diff --git a/paddle/platform/gpu_info.cc b/paddle/platform/gpu_info.cc index a1383d3524..edeb3ecd7b 100644 --- a/paddle/platform/gpu_info.cc +++ b/paddle/platform/gpu_info.cc @@ -14,7 +14,7 @@ limitations under the License. */ #include "paddle/platform/gpu_info.h" #include "gflags/gflags.h" -#include "paddle/platform/error.h" +#include "paddle/platform/enforce.h" DEFINE_double(fraction_of_gpu_memory_to_use, 0.95, "Default use 95% of GPU memory for PaddlePaddle," @@ -25,7 +25,7 @@ namespace platform { int GetDeviceCount() { int count; - throw_on_error( + PADDLE_ENFORCE( cudaGetDeviceCount(&count), "cudaGetDeviceCount failed in paddle::platform::GetDeviceCount"); return count; @@ -33,19 +33,19 @@ int GetDeviceCount() { int GetCurrentDeviceId() { int device_id; - throw_on_error( + PADDLE_ENFORCE( cudaGetDevice(&device_id), "cudaGetDevice failed in paddle::platform::GetCurrentDeviceId"); return device_id; } void SetDeviceId(int id) { - throw_on_error(cudaSetDevice(id), + PADDLE_ENFORCE(cudaSetDevice(id), "cudaSetDevice failed in paddle::platform::SetDeviceId"); } -void GpuMemoryUsage(size_t& available, size_t& total) { - throw_on_error(cudaMemGetInfo(&available, &total), +void GpuMemoryUsage(size_t &available, size_t &total) { + PADDLE_ENFORCE(cudaMemGetInfo(&available, &total), "cudaMemGetInfo failed in paddle::platform::GetMemoryUsage"); } @@ -82,5 +82,28 @@ size_t GpuMaxChunkSize() { return usable; } +void GpuMemcpyAsync(void *dst, const void *src, size_t count, + enum cudaMemcpyKind kind, cudaStream_t stream) { + PADDLE_ENFORCE(cudaMemcpyAsync(dst, src, count, kind, stream), + "cudaMemcpyAsync failed in paddle::platform::GpuMemcpyAsync"); +} + +void GpuMemcpySync(void *dst, const void *src, size_t count, + enum cudaMemcpyKind kind) { + PADDLE_ENFORCE(cudaMemcpy(dst, src, count, kind), + "cudaMemcpy failed in paddle::platform::GpuMemcpySync"); + // note: cudaMemcpy may actually be asynchronous with respect to the caller, + // block on stream 0 to make sure the copy has completed + PADDLE_ENFORCE( + cudaStreamSynchronize(0), + "cudaStreamSynchronize failed in paddle::platform::GpuMemcpySync"); +} + +void GpuMemcpyPeer(void *dst, int dst_device, const void *src, int src_device, + size_t count, cudaStream_t stream) { + PADDLE_ENFORCE( + cudaMemcpyPeerAsync(dst, dst_device, src, src_device, count, stream), + "cudaMemcpyPeerAsync failed in paddle::platform::GpuMemcpyPeer"); +} } // namespace platform } // namespace paddle diff --git a/paddle/platform/gpu_info.h b/paddle/platform/gpu_info.h index 79e71956bd..d3a5f5f13f 100644 --- a/paddle/platform/gpu_info.h +++ b/paddle/platform/gpu_info.h @@ -16,6 +16,7 @@ limitations under the License. */ #ifndef PADDLE_ONLY_CPU +#include #include namespace paddle { @@ -31,7 +32,7 @@ int GetCurrentDeviceId(); void SetDeviceId(int device_id); //！Get the memory usage of current GPU device. -void GpuMemoryUsage(size_t& available, size_t& total); +void GpuMemoryUsage(size_t &available, size_t &total); //! Get the maximum allocation size of current GPU device. size_t GpuMaxAllocSize(); @@ -42,6 +43,18 @@ size_t GpuMinChunkSize(); //! Get the maximum chunk size for GPU buddy allocator. size_t GpuMaxChunkSize(); +//! Copy memory from address src to dst asynchronously. +void GpuMemcpyAsync(void *dst, const void *src, size_t count, + enum cudaMemcpyKind kind, cudaStream_t stream); + +//! Copy memory from address src to dst synchronously. +void GpuMemcpySync(void *dst, const void *src, size_t count, + enum cudaMemcpyKind kind); + +//! Copy memory from one device to another device. +void GpuMemcpyPeer(void *dst, int dst_device, const void *src, int src_device, + size_t count, cudaStream_t stream); + } // namespace platform } // namespace paddle diff --git a/paddle/pybind/CMakeLists.txt b/paddle/pybind/CMakeLists.txt index 29fb29c7c1..fd1a142b40 100644 --- a/paddle/pybind/CMakeLists.txt +++ b/paddle/pybind/CMakeLists.txt @@ -1,2 +1,2 @@ cc_library(paddle_pybind SHARED SRCS pybind.cc DEPS pybind python - add_op fc_op) + add_op fc_op sgd_op cross_entropy_op) diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index bd126f0e97..62539c1076 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -24,11 +24,37 @@ limitations under the License. */ #include "pybind11/pybind11.h" #include "pybind11/stl.h" +#include "paddle/framework/net.h" +#include "paddle/framework/op_registry.h" +#include "paddle/framework/operator.h" +#include "paddle/framework/scope.h" +#include "paddle/pybind/tensor_bind.h" +#include "pybind11/numpy.h" +#include "pybind11/pybind11.h" +#include "pybind11/stl.h" + namespace py = pybind11; namespace pd = paddle::framework; USE_OP(add_two); +USE_OP(onehot_cross_entropy); USE_OP_WITHOUT_KERNEL(fc); +USE_OP(sgd); +USE_OP(mul); +USE_OP(sigmoid); +USE_OP(softmax); +USE_OP(rowwise_add); + +template +void ExposeOperator(ClassType& m) { + m.def("infer_shape", &ClassType::type::InferShape) + .def("run", &ClassType::type::Run) + .def("outputs", + [](const typename ClassType::type& op) -> std::vector { + return op.outputs_; + }) + .def("__str__", &ClassType::type::DebugString); +} template void ExposeOperator(ClassType& m) { @@ -42,7 +68,7 @@ void ExposeOperator(ClassType& m) { } PYBIND11_PLUGIN(core) { - py::module m("core", "C++ core of Paddle Paddle"); + py::module m("core", "C++ core of PaddlePaddle"); py::class_(m, "Tensor", py::buffer_protocol()) .def_buffer([](pd::Tensor& self) -> py::buffer_info { @@ -52,7 +78,7 @@ PYBIND11_PLUGIN(core) { [](const pd::Tensor& self) { return pd::vectorize(self.dims()); }) .def("set_dims", [](pd::Tensor& self, const std::vector& dim) { - self.set_dims(pd::make_ddim(dim)); + self.Resize(pd::make_ddim(dim)); }) .def("alloc_float", [](pd::Tensor& self) { @@ -119,9 +145,10 @@ All parameter, weight, gradient are variables in Paddle. return new paddle::platform::CPUDeviceContext(); }); - py::class_ operator_base(m, "Operator"); + py::class_> operator_base( + m, "Operator"); - operator_base.def_static("create", [](py::bytes protobin) -> pd::OperatorPtr { + operator_base.def_static("create", [](py::bytes protobin) { pd::OpDesc desc; PADDLE_ENFORCE(desc.ParsePartialFromString(protobin), "Cannot parse user input to OpDesc"); diff --git a/paddle/pybind/tensor_bind.h b/paddle/pybind/tensor_bind.h index b96516643a..995e102bf9 100644 --- a/paddle/pybind/tensor_bind.h +++ b/paddle/pybind/tensor_bind.h @@ -86,7 +86,7 @@ void PyTensorSetFromArray( dims.push_back((int)array.shape()[i]); } - self.set_dims(framework::make_ddim(dims)); + self.Resize(framework::make_ddim(dims)); auto *dst = self.mutable_data(paddle::platform::CPUPlace()); std::memcpy(dst, array.data(), sizeof(T) * array.size()); } diff --git a/proto/ModelConfig.proto b/proto/ModelConfig.proto index 37cd16c798..83f72c137b 100644 --- a/proto/ModelConfig.proto +++ b/proto/ModelConfig.proto @@ -472,10 +472,16 @@ message LayerConfig { // blank label used in ctc loss optional uint32 blank = 52 [default = 0]; - // stride parameter for seqlastins layer, AverageLayer, MaxLayer, which + // stride parameter for seqlastins layer, AverageLayer, MaxLayer, which // controls the scope of pooling operation. can be set > 0. // leave empty or set to -1 to disable this stride pooling. optional int32 seq_pool_stride = 53 [default = -1]; + + // for crop layer + optional int32 axis = 54 [default = 2]; + repeated uint32 offset = 55; + repeated uint32 shape = 56; + } message EvaluatorConfig { diff --git a/python/paddle/trainer/config_parser.py b/python/paddle/trainer/config_parser.py index ef3d81e4c0..fc112f1327 100644 --- a/python/paddle/trainer/config_parser.py +++ b/python/paddle/trainer/config_parser.py @@ -1998,6 +1998,23 @@ class PadLayer(LayerBase): self.config.size = out_ch * out_h * out_w +@config_layer('crop') +class CropLayer(LayerBase): + def __init__(self, name, inputs, axis, offset, shape, **xargs): + super(CropLayer, self).__init__(name, 'crop', 0, inputs=inputs, **xargs) + self.config.axis = axis + self.config.offset.extend(offset) + self.config.shape.extend(shape) + + # get channel, width and height from input_0 layer + input_layer = self.get_input_layer(0) + image_conf = self.config.inputs[0].image_conf + image_conf.img_size = input_layer.width + image_conf.img_size_y = input_layer.height + image_conf.channels = input_layer.size / (input_layer.width * + input_layer.height) + + @config_layer('batch_norm') class BatchNormLayer(LayerBase): layer_type = 'batch_norm' @@ -3202,6 +3219,10 @@ def ParameterHook(type, **kwargs): if sparsity_ratio is not None: hook.sparsity_ratio = sparsity_ratio return hook + elif type == 'dpruning': + hook = ParameterUpdaterHookConfig() + hook.type = type + return hook else: return None diff --git a/python/paddle/trainer_config_helpers/layers.py b/python/paddle/trainer_config_helpers/layers.py index 78aa0778f8..21eba71527 100755 --- a/python/paddle/trainer_config_helpers/layers.py +++ b/python/paddle/trainer_config_helpers/layers.py @@ -127,6 +127,7 @@ __all__ = [ 'dropout_layer', 'prelu_layer', 'gated_unit_layer', + 'crop_layer', ] @@ -218,6 +219,7 @@ class LayerType(object): SMOOTH_L1 = 'smooth_l1' PRELU = 'prelu' + CROP_LAYER = 'crop' @staticmethod def is_layer_type(type_name): @@ -3171,11 +3173,11 @@ def memory(name, @wrap_bias_attr_default() -@wrap_act_default( - param_names=['gate_act', 'state_act'], act=SigmoidActivation()) +@wrap_act_default(param_names=['gate_act'], act=SigmoidActivation()) +@wrap_act_default(param_names=['state_act'], act=TanhActivation()) @wrap_act_default(act=TanhActivation()) @wrap_name_default('lstm_step') -@layer_support() +@layer_support(ERROR_CLIPPING, DROPOUT) def lstm_step_layer(input, state, size=None, @@ -3529,12 +3531,7 @@ def SubsequenceInput(input): @wrap_name_default("recurrent_group") -def recurrent_group(step, - input, - reverse=False, - name=None, - targetInlink=None, - is_generating=False): +def recurrent_group(step, input, reverse=False, name=None, targetInlink=None): """ Recurrent layer group is an extremely flexible recurrent unit in PaddlePaddle. As long as the user defines the calculation done within a @@ -3600,21 +3597,12 @@ def recurrent_group(step, :type targetInlink: LayerOutput|SubsequenceInput - :param is_generating: If is generating, none of input type should be LayerOutput; - else, for training or testing, one of the input type must - be LayerOutput. - - :type is_generating: bool - :return: LayerOutput object. :rtype: LayerOutput """ model_type('recurrent_nn') - def is_single_input(x): - return isinstance(x, LayerOutput) or isinstance(x, StaticInput) - - if is_single_input(input): + if isinstance(input, LayerOutput) or isinstance(input, StaticInput): input = [input] assert isinstance(input, collections.Sequence) @@ -3628,13 +3616,8 @@ def recurrent_group(step, in_links=map(lambda x: x.name, in_links), seq_reversed=reverse) in_args = [] - has_LayerOutput = False for each_input in input: - assert is_single_input(each_input) - if isinstance(each_input, LayerOutput): - in_args.append(each_input) - has_LayerOutput = True - else: # StaticInput + if isinstance(each_input, StaticInput): # StaticInput mem_name = "__%s_memory__" % each_input.input.name mem = memory( name=None, @@ -3642,24 +3625,26 @@ def recurrent_group(step, boot_layer=each_input.input) mem.set_input(mem) in_args.append(mem) - - assert (is_generating != has_LayerOutput) + else: + in_args.append(each_input) layer_outs = step(*in_args) if isinstance(layer_outs, LayerOutput): layer_outs = [layer_outs] - for ot in layer_outs: - assert isinstance(ot, LayerOutput) - ot.reverse = reverse - RecurrentLayerGroupSetOutLink(ot.name) + for layer_out in layer_outs: + assert isinstance( + layer_out, LayerOutput + ), "Type of step function's return value must be LayerOutput." + layer_out.reverse = reverse + RecurrentLayerGroupSetOutLink(layer_out.name) RecurrentLayerGroupEnd(name=name) for layer_out in layer_outs: - # Thee previous full_name is the name is the rnn group - # We need a full_name outside the rnn group + # The previous full_name is the name inside the recurrent group. + # We need a full_name outside the recurrent group. layer_out.full_name = MakeLayerNameInSubmodel(layer_out.name) if len(layer_outs) == 1: @@ -3682,7 +3667,20 @@ class BaseGeneratedInput(object): class GeneratedInput(BaseGeneratedInput): def after_real_step(self, input): - return maxid_layer(input=input, name='__beam_search_predict__') + if isinstance(input, LayerOutput): + input = [input] + elif isinstance(input, collections.Sequence): + input = list(input) + if len(input) > 1: + logger.info( + ("More than one layers inside the recurrent_group " + "are returned as outputs of the entire recurrent_group " + "PLEASE garantee the first output is probability of " + "the predicted next word.")) + + return [maxid_layer( + input=input[0], name='__beam_search_predict__')] + ( + input[1:] if len(input) > 1 else []) def before_real_step(self): predict_id = memory( @@ -3869,6 +3867,7 @@ def beam_search(step, :type step: callable :param input: Input data for the recurrent unit, which should include the previously generated words as a GeneratedInput object. + In beam_search, none of the input's type should be LayerOutput. :type input: list :param bos_id: Index of the start symbol in the dictionary. The start symbol is a special token for NLP task, which indicates the @@ -3910,15 +3909,18 @@ def beam_search(step, real_input = [] for i, each_input in enumerate(input): - assert isinstance(each_input, StaticInput) or isinstance( - each_input, BaseGeneratedInput) + assert not isinstance(each_input, LayerOutput), ( + "in beam_search, " + "none of the input should has a type of LayerOutput.") if isinstance(each_input, BaseGeneratedInput): - assert generated_input_index == -1 + assert generated_input_index == -1, ("recurrent_group accepts " + "only one GeneratedInput.") generated_input_index = i + else: real_input.append(each_input) - assert generated_input_index != -1 + assert generated_input_index != -1, "No GeneratedInput is given." gipt = input[generated_input_index] @@ -3939,17 +3941,11 @@ def beam_search(step, predict = gipt.after_real_step(step(*args)) - eos_layer(input=predict, eos_id=eos_id, name=eos_name) + eos_layer(input=predict[0], eos_id=eos_id, name=eos_name) return predict - tmp = recurrent_group( - step=__real_step__, - input=real_input, - reverse=False, - name=name, - is_generating=True) - - return tmp + return recurrent_group( + step=__real_step__, input=real_input, reverse=False, name=name) def __cost_input__(input, label, weight=None): @@ -5970,3 +5966,52 @@ def gated_unit_layer(input, name="%s_gated_act" % name, input=dotmul_operator(input_proj, gate), layer_attr=layer_attr) + + +@wrap_name_default() +@layer_support() +def crop_layer(input, offset, axis=2, shape=None, name=None, layer_attr=None): + """ + The crop layer crops images by offset and shape. User can set crop shape by + args 'shape' explicitly or by reference input layer. + + The example usage is: + + .. code-block:: python + crop = crop_layer(input=[image_input, reference_input], axis=2, offset=[2, 3]) + + :param input: The input layer.If two inputs were setted, + the second input will be regarded as reference input + :type input: LayerOutput or Sequence + :param offset: The crop offset + :type offset: Sequence + :param axis: start axis to be cropped. To image input layer: + - 0: batch size + - 1: channels + - 2: height + - 3: width + :type partial_sum: int + :param shape: The shape to be cropped. Default is None. + :type shape: Sequence | None + :param name: Name of this layer. + :type name: basestring + :return: LayerOutput object. + :rtype: LayerOutput + """ + if isinstance(input, LayerOutput): + input = [input] + else: + assert isinstance(input, collections.Sequence) + l = Layer( + inputs=[x.name for x in input], + axis=axis, + offset=offset, + shape=shape, + name=name, + type=LayerType.CROP_LAYER, + **ExtraLayerAttribute.to_kwargs(layer_attr)) + return LayerOutput( + name=name, + layer_type=LayerType.CROP_LAYER, + parents=input, + size=l.config.size) diff --git a/python/paddle/trainer_config_helpers/networks.py b/python/paddle/trainer_config_helpers/networks.py index 810bea913e..dcc4fec4f3 100755 --- a/python/paddle/trainer_config_helpers/networks.py +++ b/python/paddle/trainer_config_helpers/networks.py @@ -614,18 +614,17 @@ def simple_lstm(input, @wrap_name_default('lstm_unit') def lstmemory_unit(input, - memory_boot=None, + out_memory=None, name=None, size=None, param_attr=None, act=None, gate_act=None, state_act=None, - mixed_bias_attr=None, + input_proj_bias_attr=None, + input_proj_layer_attr=None, lstm_bias_attr=None, - mixed_layer_attr=None, - lstm_layer_attr=None, - get_output_layer_attr=None): + lstm_layer_attr=None): """ Define calculations that a LSTM unit performs during a single time step. This function itself is not a recurrent layer, so it can not be @@ -662,8 +661,8 @@ def lstmemory_unit(input, :param input: input layer name. :type input: LayerOutput - :param memory_boot: the initialization state of the LSTM cell. - :type memory_boot: LayerOutput | None + :param out_memory: output of previous time step + :type out_memory: LayerOutput | None :param name: lstmemory unit name. :type name: basestring :param size: lstmemory unit size. @@ -676,33 +675,35 @@ def lstmemory_unit(input, :type gate_act: BaseActivation :param state_act: lstm state activiation type. :type state_act: BaseActivation - :param mixed_bias_attr: bias parameter attribute of mixed layer. - False means no bias, None means default bias. - :type mixed_bias_attr: ParameterAttribute|False + :param input_proj_bias_attr: bias attribute for input-to-hidden projection. + False means no bias, None means default bias. + :type input_proj_bias_attr: ParameterAttribute|False|None + :param input_proj_layer_attr: extra layer attribute for input to hidden + projection of the LSTM unit, such as dropout, error clipping. + :type input_proj_layer_attr: ExtraLayerAttribute :param lstm_bias_attr: bias parameter attribute of lstm layer. - False means no bias, None means default bias. + False means no bias, None means default bias. :type lstm_bias_attr: ParameterAttribute|False - :param mixed_layer_attr: mixed layer's extra attribute. - :type mixed_layer_attr: ExtraLayerAttribute :param lstm_layer_attr: lstm layer's extra attribute. :type lstm_layer_attr: ExtraLayerAttribute - :param get_output_layer_attr: get output layer's extra attribute. - :type get_output_layer_attr: ExtraLayerAttribute :return: lstmemory unit name. :rtype: LayerOutput """ if size is None: assert input.size % 4 == 0 size = input.size / 4 - out_mem = memory(name=name, size=size) - state_mem = memory( - name="%s_state" % name, size=size, boot_layer=memory_boot) + if out_memory is None: + out_mem = memory(name=name, size=size) + else: + out_mem = out_memory + + state_mem = memory(name="%s_state" % name, size=size) with mixed_layer( name="%s_input_recurrent" % name, size=size * 4, - bias_attr=mixed_bias_attr, - layer_attr=mixed_layer_attr, + bias_attr=input_proj_bias_attr, + layer_attr=input_proj_layer_attr, act=IdentityActivation()) as m: m += identity_projection(input=input) m += full_matrix_projection(input=out_mem, param_attr=param_attr) @@ -717,11 +718,7 @@ def lstmemory_unit(input, gate_act=gate_act, state_act=state_act, layer_attr=lstm_layer_attr) - get_output_layer( - name='%s_state' % name, - input=lstm_out, - arg_name='state', - layer_attr=get_output_layer_attr) + get_output_layer(name='%s_state' % name, input=lstm_out, arg_name='state') return lstm_out @@ -730,17 +727,16 @@ def lstmemory_unit(input, def lstmemory_group(input, size=None, name=None, - memory_boot=None, + out_memory=None, reverse=False, param_attr=None, act=None, gate_act=None, state_act=None, - mixed_bias_attr=None, + input_proj_bias_attr=None, + input_proj_layer_attr=None, lstm_bias_attr=None, - mixed_layer_attr=None, - lstm_layer_attr=None, - get_output_layer_attr=None): + lstm_layer_attr=None): """ lstm_group is a recurrent_group version of Long Short Term Memory. It does exactly the same calculation as the lstmemory layer (see lstmemory in @@ -774,8 +770,8 @@ def lstmemory_group(input, :type size: int :param name: name of the lstmemory group. :type name: basestring - :param memory_boot: the initialization state of LSTM cell. - :type memory_boot: LayerOutput | None + :param out_memory: output of previous time step + :type out_memory: LayerOutput | None :param reverse: is lstm reversed :type reverse: bool :param param_attr: Parameter config, None if use default. @@ -786,18 +782,17 @@ def lstmemory_group(input, :type gate_act: BaseActivation :param state_act: lstm state activiation type. :type state_act: BaseActivation - :param mixed_bias_attr: bias parameter attribute of mixed layer. - False means no bias, None means default bias. - :type mixed_bias_attr: ParameterAttribute|False :param lstm_bias_attr: bias parameter attribute of lstm layer. False means no bias, None means default bias. :type lstm_bias_attr: ParameterAttribute|False - :param mixed_layer_attr: mixed layer's extra attribute. - :type mixed_layer_attr: ExtraLayerAttribute + :param input_proj_bias_attr: bias attribute for input-to-hidden projection. + False means no bias, None means default bias. + :type input_proj_bias_attr: ParameterAttribute|False|None + :param input_proj_layer_attr: extra layer attribute for input to hidden + projection of the LSTM unit, such as dropout, error clipping. + :type input_proj_layer_attr: ExtraLayerAttribute :param lstm_layer_attr: lstm layer's extra attribute. :type lstm_layer_attr: ExtraLayerAttribute - :param get_output_layer_attr: get output layer's extra attribute. - :type get_output_layer_attr: ExtraLayerAttribute :return: the lstmemory group. :rtype: LayerOutput """ @@ -805,18 +800,17 @@ def lstmemory_group(input, def __lstm_step__(ipt): return lstmemory_unit( input=ipt, - memory_boot=memory_boot, name=name, size=size, - mixed_bias_attr=mixed_bias_attr, - mixed_layer_attr=mixed_layer_attr, - param_attr=param_attr, - lstm_bias_attr=lstm_bias_attr, act=act, gate_act=gate_act, state_act=state_act, + out_memory=out_memory, + input_proj_bias_attr=input_proj_bias_attr, + input_proj_layer_attr=input_proj_layer_attr, + param_attr=param_attr, lstm_layer_attr=lstm_layer_attr, - get_output_layer_attr=get_output_layer_attr) + lstm_bias_attr=lstm_bias_attr) return recurrent_group( name='%s_recurrent_group' % name, diff --git a/python/paddle/trainer_config_helpers/tests/configs/protostr/shared_lstm.protostr b/python/paddle/trainer_config_helpers/tests/configs/protostr/shared_lstm.protostr index 7f2aa5a0fe..75cf231203 100644 --- a/python/paddle/trainer_config_helpers/tests/configs/protostr/shared_lstm.protostr +++ b/python/paddle/trainer_config_helpers/tests/configs/protostr/shared_lstm.protostr @@ -104,7 +104,7 @@ layers { } bias_parameter_name: "lstm_bias" active_gate_type: "sigmoid" - active_state_type: "sigmoid" + active_state_type: "tanh" } layers { name: "__lstm_group_0___state@__lstm_group_0___recurrent_group" @@ -183,7 +183,7 @@ layers { } bias_parameter_name: "lstm_bias" active_gate_type: "sigmoid" - active_state_type: "sigmoid" + active_state_type: "tanh" } layers { name: "__lstm_group_1___state@__lstm_group_1___recurrent_group" diff --git a/python/paddle/trainer_config_helpers/tests/configs/protostr/test_rnn_group.protostr b/python/paddle/trainer_config_helpers/tests/configs/protostr/test_rnn_group.protostr index af1b63c5df..711785be37 100644 --- a/python/paddle/trainer_config_helpers/tests/configs/protostr/test_rnn_group.protostr +++ b/python/paddle/trainer_config_helpers/tests/configs/protostr/test_rnn_group.protostr @@ -258,7 +258,7 @@ layers { } bias_parameter_name: "___lstm_group_0__@__lstm_group_0___recurrent_group.wbias" active_gate_type: "sigmoid" - active_state_type: "sigmoid" + active_state_type: "tanh" } layers { name: "__lstm_group_0___state@__lstm_group_0___recurrent_group" diff --git a/python/paddle/trainer_config_helpers/tests/configs/shared_lstm.py b/python/paddle/trainer_config_helpers/tests/configs/shared_lstm.py index 05810597b3..565e281a6e 100644 --- a/python/paddle/trainer_config_helpers/tests/configs/shared_lstm.py +++ b/python/paddle/trainer_config_helpers/tests/configs/shared_lstm.py @@ -20,12 +20,13 @@ lstm1 = lstmemory_group( input=m1, param_attr=lstm_param, lstm_bias_attr=lstm_bias, - mixed_bias_attr=False) + input_proj_bias_attr=False) + lstm2 = lstmemory_group( input=m2, param_attr=lstm_param, lstm_bias_attr=lstm_bias, - mixed_bias_attr=False) + input_proj_bias_attr=False) softmax_param = ParamAttr(name='softmax_param') diff --git a/python/paddle/trainer_config_helpers/tests/configs/test_crop.py b/python/paddle/trainer_config_helpers/tests/configs/test_crop.py new file mode 100644 index 0000000000..8314a7e9a5 --- /dev/null +++ b/python/paddle/trainer_config_helpers/tests/configs/test_crop.py @@ -0,0 +1,21 @@ +from paddle.trainer_config_helpers import * + +settings(batch_size=1000, learning_rate=1e-5) + +data = data_layer(name='data', size=2016, height=48, width=42) +refernce_data = data_layer(name='data', size=768, height=16, width=16) + +conv = img_conv_layer( + input=data, + filter_size=3, + num_channels=1, + num_filters=16, + padding=1, + act=LinearActivation(), + bias_attr=True) + +pool = img_pool_layer(input=conv, pool_size=2, stride=2, pool_type=MaxPooling()) + +crop = crop_layer(input=[pool, refernce_data], axis=2) + +outputs(pad) diff --git a/python/paddle/v2/dataset/flowers.py b/python/paddle/v2/dataset/flowers.py index e2a21e6e3e..634388094c 100644 --- a/python/paddle/v2/dataset/flowers.py +++ b/python/paddle/v2/dataset/flowers.py @@ -116,7 +116,7 @@ def reader_creator(data_file, data = batch['data'] labels = batch['label'] for sample, label in itertools.izip(data, batch['label']): - yield sample, int(label) + yield sample, int(label) - 1 if use_xmap: return xmap_readers(mapper, reader, cpu_count(), buffered_size) diff --git a/python/paddle/v2/framework/tests/CMakeLists.txt b/python/paddle/v2/framework/tests/CMakeLists.txt index f71009aa85..b3eb2ef8a8 100644 --- a/python/paddle/v2/framework/tests/CMakeLists.txt +++ b/python/paddle/v2/framework/tests/CMakeLists.txt @@ -1,3 +1,15 @@ -add_python_test(test_framework test_protobuf.py test_scope.py - test_default_scope_funcs.py test_op_creation_methods.py - test_tensor.py test_fc_op.py test_add_two_op.py) +add_python_test(test_framework + test_protobuf.py + test_scope.py + test_default_scope_funcs.py + test_op_creation_methods.py + test_plain_net.py + test_tensor.py + test_fc_op.py + test_add_two_op.py + test_sgd_op.py + test_cross_entropy_op.py + test_mul_op.py + test_sigmoid_op.py + test_softmax_op.py + test_rowwise_add_op.py) diff --git a/python/paddle/v2/framework/tests/op_test_util.py b/python/paddle/v2/framework/tests/op_test_util.py index 237f9b7eb0..7b62313f8a 100644 --- a/python/paddle/v2/framework/tests/op_test_util.py +++ b/python/paddle/v2/framework/tests/op_test_util.py @@ -5,6 +5,18 @@ import paddle.v2.framework.create_op_creation_methods as creation class OpTestMeta(type): + """ + Operator Test ClassMeta. + + It injects `test_all` method into user's OperatorTest class, to make Python + unittest module run that method. + + The `test_all` read what value is stored in `self`. It use self's values to + create and run a operator, and check whether that op is OK or not. + + See `test_add_two_op` for example usage. + """ + def __new__(cls, name, bases, attrs): obj = super(OpTestMeta, cls).__new__(cls, name, bases, attrs) @@ -44,7 +56,10 @@ class OpTestMeta(type): for out_name in func.all_output_args: actual = numpy.array(scope.get_var(out_name).get_tensor()) expect = getattr(self, out_name) - numpy.testing.assert_almost_equal(actual, expect) + # TODO(qijun) The default decimal is 7, but numpy.dot and eigen.mul + # has some diff, and could not pass unittest. So I set decimal 3 here. + # And I will check this in future. + numpy.testing.assert_almost_equal(actual, expect, decimal=3) obj.test_all = test_all return obj diff --git a/python/paddle/v2/framework/tests/test_cross_entropy_op.py b/python/paddle/v2/framework/tests/test_cross_entropy_op.py new file mode 100644 index 0000000000..609c56535e --- /dev/null +++ b/python/paddle/v2/framework/tests/test_cross_entropy_op.py @@ -0,0 +1,22 @@ +import unittest +import numpy +from op_test_util import OpTestMeta + + +class TestSGD(unittest.TestCase): + __metaclass__ = OpTestMeta + + def setUp(self): + self.type = "onehot_cross_entropy" + batch_size = 100 + class_num = 10 + self.X = numpy.random.random((batch_size, class_num)).astype("float32") + self.label = 5 * numpy.ones(batch_size).astype("int32") + Y = [] + for i in range(0, batch_size): + Y.append(-numpy.log(self.X[i][self.label[i]])) + self.Y = numpy.array(Y).astype("float32") + + +if __name__ == "__main__": + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_mul_op.py b/python/paddle/v2/framework/tests/test_mul_op.py new file mode 100644 index 0000000000..0a87e66cd0 --- /dev/null +++ b/python/paddle/v2/framework/tests/test_mul_op.py @@ -0,0 +1,17 @@ +import unittest +from op_test_util import OpTestMeta +import numpy as np + + +class TestMulOp(unittest.TestCase): + __metaclass__ = OpTestMeta + + def setUp(self): + self.type = "mul" + self.X = np.random.random((32, 784)).astype("float32") + self.Y = np.random.random((784, 100)).astype("float32") + self.Out = np.dot(self.X, self.Y) + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_plain_net.py b/python/paddle/v2/framework/tests/test_plain_net.py new file mode 100644 index 0000000000..2b919aca28 --- /dev/null +++ b/python/paddle/v2/framework/tests/test_plain_net.py @@ -0,0 +1,30 @@ +import paddle.v2.framework.core as core +from paddle.v2.framework.create_op_creation_methods import op_creations +import unittest + + +class TestNet(unittest.TestCase): + def test_net_all(self): + net = core.PlainNet.create() + op1 = op_creations.add_two(X="X", Y="Y", Out="Out") + net.add_op(op1) + + net2 = core.PlainNet.create() + net2.add_op(op_creations.fc(X="X", W="w", Y="fc.out")) + net2.complete_add_op(True) + net.add_op(net2) + net.complete_add_op(True) + + expected = ''' +Op(plain_net), inputs:(@EMPTY@, X, Y, w), outputs:(@TEMP@fc@0, Out, fc.out). + Op(add_two), inputs:(X, Y), outputs:(Out). + Op(plain_net), inputs:(@EMPTY@, X, w), outputs:(@TEMP@fc@0, fc.out). + Op(fc), inputs:(X, w, @EMPTY@), outputs:(fc.out, @TEMP@fc@0). + Op(mul), inputs:(X, w), outputs:(@TEMP@fc@0). + Op(sigmoid), inputs:(@TEMP@fc@0), outputs:(fc.out). +''' + self.assertEqual(expected, "\n" + str(net)) + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_rowwise_add_op.py b/python/paddle/v2/framework/tests/test_rowwise_add_op.py new file mode 100644 index 0000000000..ef1514983c --- /dev/null +++ b/python/paddle/v2/framework/tests/test_rowwise_add_op.py @@ -0,0 +1,17 @@ +import unittest +from op_test_util import OpTestMeta +import numpy as np + + +class TestRowwiseAddOp(unittest.TestCase): + __metaclass__ = OpTestMeta + + def setUp(self): + self.type = "rowwise_add" + self.X = np.random.random((32, 784)).astype("float32") + self.b = np.random.random(784).astype("float32") + self.Out = np.add(self.X, self.b) + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_sgd_op.py b/python/paddle/v2/framework/tests/test_sgd_op.py new file mode 100644 index 0000000000..405d73b224 --- /dev/null +++ b/python/paddle/v2/framework/tests/test_sgd_op.py @@ -0,0 +1,18 @@ +import unittest +import numpy +from op_test_util import OpTestMeta + + +class TestSGD(unittest.TestCase): + __metaclass__ = OpTestMeta + + def setUp(self): + self.type = "sgd" + self.param = numpy.random.random((342, 345)).astype("float32") + self.grad = numpy.random.random((342, 345)).astype("float32") + self.learning_rate = 0.1 + self.param_out = self.param - self.learning_rate * self.grad + + +if __name__ == "__main__": + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_sigmoid_op.py b/python/paddle/v2/framework/tests/test_sigmoid_op.py new file mode 100644 index 0000000000..50044a122f --- /dev/null +++ b/python/paddle/v2/framework/tests/test_sigmoid_op.py @@ -0,0 +1,16 @@ +import unittest +from op_test_util import OpTestMeta +import numpy as np + + +class TestSigmoidOp(unittest.TestCase): + __metaclass__ = OpTestMeta + + def setUp(self): + self.type = "sigmoid" + self.X = np.random.random((32, 100)).astype("float32") + self.Y = 1 / (1 + np.exp(-self.X)) + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_softmax_op.py b/python/paddle/v2/framework/tests/test_softmax_op.py new file mode 100644 index 0000000000..191b698c1c --- /dev/null +++ b/python/paddle/v2/framework/tests/test_softmax_op.py @@ -0,0 +1,23 @@ +import unittest +from op_test_util import OpTestMeta +import numpy as np + + +def stable_softmax(x): + """Compute the softmax of vector x in a numerically stable way.""" + shiftx = x - np.max(x) + exps = np.exp(shiftx) + return exps / np.sum(exps) + + +class TestSoftmaxOp(unittest.TestCase): + __metaclass__ = OpTestMeta + + def setUp(self): + self.type = "softmax" + self.X = np.random.random((32, 100)).astype("float32") + self.Y = np.apply_along_axis(stable_softmax, 1, self.X) + + +if __name__ == '__main__': + unittest.main()