diff --git a/.github/ISSUE_TEMPLATE.md b/.github/ISSUE_TEMPLATE.md index d78a3dc3..c981f62f 100644 --- a/.github/ISSUE_TEMPLATE.md +++ b/.github/ISSUE_TEMPLATE.md @@ -1,19 +1,34 @@ -Please use the [caffe-users list](https://groups.google.com/forum/#!forum/caffe-users) for usage, installation, or modeling questions, or other requests for help. -_Do not post such requests to Issues._ Doing so interferes with the development of Caffe. +## Important - read before submitting -Please read the [guidelines for contributing](https://github.com/BVLC/caffe/blob/master/CONTRIBUTING.md) before submitting this issue. +*Please read the [guidelines for contributing](https://github.com/BVLC/caffe/blob/master/CONTRIBUTING.md) before submitting this issue!* + +*Please do not post installation, build, usage, or modeling questions, or other requests for help to Issues.* +Use the [caffe-users list](https://groups.google.com/forum/#!forum/caffe-users) instead. +This helps developers maintain a clear, uncluttered, and efficient view of the state of Caffe. ### Issue summary ### Steps to reproduce -If you are having difficulty building Caffe or training a model, please ask the caffe-users mailing list. If you are reporting a build error that seems to be due to a bug in Caffe, please attach your build configuration (either Makefile.config or CMakeCache.txt) and the output of the make (or cmake) command. -### Your system configuration -Operating system: -Compiler: -CUDA version (if applicable): -CUDNN version (if applicable): -BLAS: -Python or MATLAB version (for pycaffe and matcaffe respectively): +### Tried solutions + + +### System configuration + +* Operating system: +* Compiler: +* CUDA version (if applicable): +* CUDNN version (if applicable): +* BLAS: +* Python version (if using pycaffe): +* MATLAB version (if using matcaffe): + +### Issue checklist + +- [ ] read the guidelines and removed the first paragraph +- [ ] written a short summary and detailed steps to reproduce +- [ ] explained how solutions to related problems failed (tick if found none) +- [ ] filled system configuration +- [ ] attached relevant logs/config files (tick if not applicable) diff --git a/CMakeLists.txt b/CMakeLists.txt index 08f56a33..27d172f9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -42,6 +42,9 @@ caffe_option(USE_LMDB "Build with lmdb" ON) caffe_option(ALLOW_LMDB_NOLOCK "Allow MDB_NOLOCK when reading LMDB files (only if necessary)" OFF) caffe_option(USE_OPENMP "Link with OpenMP (when your BLAS wants OpenMP and you get linker errors)" OFF) +# This code is taken from https://github.com/sh1r0/caffe-android-lib +caffe_option(USE_HDF5 "Build with hdf5" ON) + # ---[ Dependencies include(cmake/Dependencies.cmake) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 8cd5e56c..45f7e186 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -1,21 +1,63 @@ # Contributing +Below you will find a collection of guidelines for submitting issues as well as contributing code to the Caffe repository. +Please read those before starting an issue or a pull request. + ## Issues Specific Caffe design and development issues, bugs, and feature requests are maintained by GitHub Issues. -_Please do not post usage, installation, or modeling questions, or other requests for help to Issues._ -Use the [caffe-users list](https://groups.google.com/forum/#!forum/caffe-users) instead. This helps developers maintain a clear, uncluttered, and efficient view of the state of Caffe. - -When reporting a bug, it's most helpful to provide the following information, where applicable: +*Please do not post installation, build, usage, or modeling questions, or other requests for help to Issues.* +Use the [caffe-users list](https://groups.google.com/forum/#!forum/caffe-users) instead. +This helps developers maintain a clear, uncluttered, and efficient view of the state of Caffe. +See the chapter [caffe-users](#caffe-users) below for guidance on posting to the users list. -* What steps reproduce the bug? -* Can you reproduce the bug using the latest [master](https://github.com/BVLC/caffe/tree/master), compiled with the `DEBUG` make option? -* What hardware and operating system/distribution are you running? +When reporting an issue, it's most helpful to provide the following information, where applicable: +* How does the problem look like and what steps reproduce it? +* Can you reproduce it using the latest [master](https://github.com/BVLC/caffe/tree/master), compiled with the `DEBUG` make option? +* What hardware and software are you running? In particular: + * GPU make and model, if relevant, + * operating system/distribution, + * compiler; please also post which version (for example, with GCC run `gcc --version` to check), + * CUDA version, if applicable (run `nvcc --version` to check), + * cuDNN version, if applicable (version number is stored in `cudnn.h`, look for lines containing `CUDNN_MAJOR`, `CUDNN_MINOR` and `CUDNN_PATCHLEVEL`), + * BLAS library, + * Python version, if relevant, + * MATLAB version, if relevant. +* **What have you already tried** to solve the problem? How did it fail? Are there any other issues related to yours? +* If this is not a build-related issue, does your installation pass `make runtest`? * If the bug is a crash, provide the backtrace (usually printed by Caffe; always obtainable with `gdb`). +* If you are reporting a build error that seems to be due to a bug in Caffe, please attach your build configuration (either Makefile.config or CMakeCache.txt) and the output of the make (or cmake) command. + +If only a small portion of the code/log is relevant to your issue, you may paste it directly into the post, preferably using Markdown syntax for code block: triple backtick ( \`\`\` ) to open/close a block. +In other cases (multiple files, or long files), please **attach** them to the post - this greatly improves readability. + +If the problem arises during a complex operation (e.g. large script using pycaffe, long network prototxt), please reduce the example to the minimal size that still causes the error. +Also, minimize influence of external modules, data etc. - this way it will be easier for others to understand and reproduce your issue, and eventually help you. +Sometimes you will find the root cause yourself in the process. Try to give your issue a title that is succinct and specific. The devs will rename issues as needed to keep track of them. +## Caffe-users + +Before you post to the [caffe-users list](https://groups.google.com/forum/#!forum/caffe-users), make sure you look for existing solutions. +The Caffe community has encountered and found solutions to countless problems - benefit from the collective experience. +Recommended places to look: +* the [users list](https://groups.google.com/forum/#!forum/caffe-users) itself, +* [`caffe`](https://stackoverflow.com/questions/tagged/caffe) tag on StackOverflow, +* [GitHub issues](https://github.com/BVLC/caffe/issues) tracker (some problems have been answered there), +* the public [wiki](https://github.com/BVLC/caffe/wiki), +* the official [documentation](http://caffe.berkeleyvision.org/). + +Found a post/issue with your exact problem, but with no answer? +Don't just leave a "me too" message - provide the details of your case. +Problems with more available information are easier to solve and attract good attention. + +When posting to the list, make sure you provide as much relevant information as possible - recommendations for an issue report (see above) are a good starting point. +*Please make it very clear which version of Caffe you are using, especially if it is a fork not maintained by BVLC.* + +Formatting recommendations hold: paste short logs/code fragments into the post (use fixed-width text for them), **attach** long logs or multiple files. + ## Pull Requests Caffe welcomes all contributions. diff --git a/Makefile b/Makefile index c3bc54c3..609bf885 100644 --- a/Makefile +++ b/Makefile @@ -178,11 +178,13 @@ ifneq ($(CPU_ONLY), 1) LIBRARIES := cudart cublas curand endif -LIBRARIES += glog gflags protobuf boost_system boost_filesystem m hdf5_hl hdf5 +LIBRARIES += glog gflags protobuf boost_system boost_filesystem m # handle IO dependencies USE_LEVELDB ?= 1 USE_LMDB ?= 1 +# This code is taken from https://github.com/sh1r0/caffe-android-lib +USE_HDF5 ?= 1 USE_OPENCV ?= 1 ifeq ($(USE_LEVELDB), 1) @@ -191,6 +193,10 @@ endif ifeq ($(USE_LMDB), 1) LIBRARIES += lmdb endif +# This code is taken from https://github.com/sh1r0/caffe-android-lib +ifeq ($(USE_HDF5), 1) + LIBRARIES += hdf5_hl hdf5 +endif ifeq ($(USE_OPENCV), 1) LIBRARIES += opencv_core opencv_highgui opencv_imgproc @@ -347,6 +353,10 @@ ifeq ($(ALLOW_LMDB_NOLOCK), 1) COMMON_FLAGS += -DALLOW_LMDB_NOLOCK endif endif +# This code is taken from https://github.com/sh1r0/caffe-android-lib +ifeq ($(USE_HDF5), 1) + COMMON_FLAGS += -DUSE_HDF5 +endif # CPU-only configuration ifeq ($(CPU_ONLY), 1) @@ -577,7 +587,7 @@ $(STATIC_NAME): $(OBJS) | $(LIB_BUILD_DIR) @ echo AR -o $@ $(Q)ar rcs $@ $(OBJS) -$(BUILD_DIR)/%.o: %.cpp | $(ALL_BUILD_DIRS) +$(BUILD_DIR)/%.o: %.cpp $(PROTO_GEN_HEADER) | $(ALL_BUILD_DIRS) @ echo CXX $< $(Q)$(CXX) $< $(CXXFLAGS) -c -o $@ 2> $@.$(WARNS_EXT) \ || (cat $@.$(WARNS_EXT); exit 1) @@ -641,7 +651,7 @@ $(PROTO_BUILD_DIR)/%.pb.cc $(PROTO_BUILD_DIR)/%.pb.h : \ $(PY_PROTO_BUILD_DIR)/%_pb2.py : $(PROTO_SRC_DIR)/%.proto \ $(PY_PROTO_INIT) | $(PY_PROTO_BUILD_DIR) @ echo PROTOC \(python\) $< - $(Q)protoc --proto_path=$(PROTO_SRC_DIR) --python_out=$(PY_PROTO_BUILD_DIR) $< + $(Q)protoc --proto_path=src --python_out=python $< $(PY_PROTO_INIT): | $(PY_PROTO_BUILD_DIR) touch $(PY_PROTO_INIT) @@ -694,6 +704,6 @@ $(DISTRIBUTE_DIR): all py | $(DISTRIBUTE_SUBDIRS) install -m 644 $(DYNAMIC_NAME) $(DISTRIBUTE_DIR)/lib cd $(DISTRIBUTE_DIR)/lib; rm -f $(DYNAMIC_NAME_SHORT); ln -s $(DYNAMIC_VERSIONED_NAME_SHORT) $(DYNAMIC_NAME_SHORT) # add python - it's not the standard way, indeed... - cp -r python $(DISTRIBUTE_DIR)/python + cp -r python $(DISTRIBUTE_DIR)/ -include $(DEPS) diff --git a/Makefile.config.example b/Makefile.config.example index d552b38a..24ca6327 100644 --- a/Makefile.config.example +++ b/Makefile.config.example @@ -11,6 +11,8 @@ # USE_OPENCV := 0 # USE_LEVELDB := 0 # USE_LMDB := 0 +# This code is taken from https://github.com/sh1r0/caffe-android-lib +# USE_HDF5 := 0 # uncomment to allow MDB_NOLOCK when reading LMDB files (only if necessary) # You should not set this flag if you will be reading LMDBs with any @@ -33,6 +35,7 @@ CUDA_DIR := /usr/local/cuda # CUDA architecture setting: going with all of them. # For CUDA < 6.0, comment the *_50 through *_61 lines for compatibility. # For CUDA < 8.0, comment the *_60 and *_61 lines for compatibility. +# For CUDA >= 9.0, comment the *_20 and *_21 lines for compatibility. CUDA_ARCH := -gencode arch=compute_20,code=sm_20 \ -gencode arch=compute_20,code=sm_21 \ -gencode arch=compute_30,code=sm_30 \ diff --git a/cmake/ConfigGen.cmake b/cmake/ConfigGen.cmake index 09bb09b4..69889c24 100644 --- a/cmake/ConfigGen.cmake +++ b/cmake/ConfigGen.cmake @@ -24,6 +24,18 @@ function(caffe_generate_export_configs) set(HAVE_CUDA FALSE) endif() + set(HDF5_IMPORTED OFF) + foreach(_lib ${HDF5_LIBRARIES} ${HDF5_HL_LIBRARIES}) + if(TARGET ${_lib}) + set(HDF5_IMPORTED ON) + endif() + endforeach() + + # This code is taken from https://github.com/sh1r0/caffe-android-lib + if(USE_HDF5) + list(APPEND Caffe_DEFINITIONS -DUSE_HDF5) + endif() + if(NOT HAVE_CUDNN) set(HAVE_CUDNN FALSE) endif() diff --git a/cmake/Cuda.cmake b/cmake/Cuda.cmake index b2b19e8b..e03feabf 100644 --- a/cmake/Cuda.cmake +++ b/cmake/Cuda.cmake @@ -109,6 +109,12 @@ function(caffe_select_nvcc_arch_flags out_variable) set(__nvcc_flags "") set(__nvcc_archs_readable "") + string(COMPARE LESS "${CUDA_VERSION}" "9.0" iscudaolderthan90) + if(NOT iscudaolderthan90) + string(REPLACE "21(20)" "" __cuda_arch_bin "${__cuda_arch_bin}") + string(REPLACE "20" "" __cuda_arch_bin "${__cuda_arch_bin}") + endif() + # Tell NVCC to add binaries for the specified GPUs foreach(__arch ${__cuda_arch_bin}) if(__arch MATCHES "([0-9]+)\\(([0-9]+)\\)") @@ -232,7 +238,7 @@ endfunction() ################################################################################################ find_package(CUDA 5.5 QUIET) -find_cuda_helper_libs(curand) # cmake 2.8.7 compartibility which doesn't search for curand +find_cuda_helper_libs(curand) # cmake 2.8.7 compatibility which doesn't search for curand if(NOT CUDA_FOUND) return() diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index f48a0ebe..da1b984f 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -47,6 +47,14 @@ find_package(HDF5 COMPONENTS HL REQUIRED) list(APPEND Caffe_INCLUDE_DIRS PUBLIC ${HDF5_INCLUDE_DIRS}) list(APPEND Caffe_LINKER_LIBS PUBLIC ${HDF5_LIBRARIES} ${HDF5_HL_LIBRARIES}) +# This code is taken from https://github.com/sh1r0/caffe-android-lib +if(USE_HDF5) + find_package(HDF5 COMPONENTS HL REQUIRED) + include_directories(SYSTEM ${HDF5_INCLUDE_DIRS} ${HDF5_HL_INCLUDE_DIR}) + list(APPEND Caffe_LINKER_LIBS ${HDF5_LIBRARIES} ${HDF5_HL_LIBRARIES}) + add_definitions(-DUSE_HDF5) +endif() + # ---[ LMDB if(USE_LMDB) find_package(LMDB REQUIRED) diff --git a/cmake/Misc.cmake b/cmake/Misc.cmake index 9dd2609b..fcb24647 100644 --- a/cmake/Misc.cmake +++ b/cmake/Misc.cmake @@ -32,9 +32,10 @@ endif() set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE CACHE BOOLEAN "Use link paths for shared library rpath") set(CMAKE_MACOSX_RPATH TRUE) -list(FIND CMAKE_PLATFORM_IMPLICIT_LINK_DIRECTORIES ${CMAKE_INSTALL_PREFIX}/lib __is_systtem_dir) +list(FIND CMAKE_PLATFORM_IMPLICIT_LINK_DIRECTORIES + ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_LIBDIR} __is_systtem_dir) if(${__is_systtem_dir} STREQUAL -1) - set(CMAKE_INSTALL_RPATH ${CMAKE_INSTALL_PREFIX}/lib) + set(CMAKE_INSTALL_RPATH ${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_LIBDIR}) endif() # ---[ Funny target diff --git a/cmake/Modules/FindMKL.cmake b/cmake/Modules/FindMKL.cmake index 8ac6fc0c..150b3cb6 100644 --- a/cmake/Modules/FindMKL.cmake +++ b/cmake/Modules/FindMKL.cmake @@ -9,7 +9,7 @@ # This module defines the following variables: # # MKL_FOUND : True mkl is found -# MKL_INCLUDE_DIR : unclude directory +# MKL_INCLUDE_DIR : include directory # MKL_LIBRARIES : the libraries to link against. diff --git a/cmake/ProtoBuf.cmake b/cmake/ProtoBuf.cmake index 8005b448..72ea3230 100644 --- a/cmake/ProtoBuf.cmake +++ b/cmake/ProtoBuf.cmake @@ -78,7 +78,7 @@ function(caffe_protobuf_generate_cpp_py output_dir srcs_var hdrs_var python_var) "${output_dir}/${fil_we}_pb2.py" COMMAND ${CMAKE_COMMAND} -E make_directory "${output_dir}" COMMAND ${PROTOBUF_PROTOC_EXECUTABLE} --cpp_out ${output_dir} ${_protoc_include} ${abs_fil} - COMMAND ${PROTOBUF_PROTOC_EXECUTABLE} --python_out ${output_dir} ${_protoc_include} ${abs_fil} + COMMAND ${PROTOBUF_PROTOC_EXECUTABLE} --python_out ${PROJECT_BINARY_DIR}/include --proto_path ${PROJECT_SOURCE_DIR}/src ${_protoc_include} ${abs_fil} DEPENDS ${abs_fil} COMMENT "Running C++/Python protocol buffer compiler on ${fil}" VERBATIM ) endforeach() diff --git a/cmake/Summary.cmake b/cmake/Summary.cmake index ed8c2526..40b8c2f2 100644 --- a/cmake/Summary.cmake +++ b/cmake/Summary.cmake @@ -119,6 +119,8 @@ function(caffe_print_configuration_summary) caffe_status(" USE_LMDB : ${USE_LMDB}") caffe_status(" USE_NCCL : ${USE_NCCL}") caffe_status(" ALLOW_LMDB_NOLOCK : ${ALLOW_LMDB_NOLOCK}") + # This code is taken from https://github.com/sh1r0/caffe-android-lib + caffe_status(" USE_HDF5 : ${USE_HDF5}") caffe_status("") caffe_status("Dependencies:") caffe_status(" BLAS : " APPLE THEN "Yes (vecLib)" ELSE "Yes (${BLAS})") diff --git a/docs/development.md b/docs/development.md index ec05bbee..36cd3995 100644 --- a/docs/development.md +++ b/docs/development.md @@ -116,5 +116,5 @@ To get a list of all options `googletest` provides, simply pass the `--help` fla - **Run `make lint` to check C++ code.** - Wrap lines at 80 chars. -- Follow [Google C++ style](http://google-styleguide.googlecode.com/svn/trunk/cppguide.xml) and [Google python style](http://google-styleguide.googlecode.com/svn/trunk/pyguide.html) + [PEP 8](http://legacy.python.org/dev/peps/pep-0008/). +- Follow [Google C++ style](https://google.github.io/styleguide/cppguide.html) and [Google python style](https://google.github.io/styleguide/pyguide.html) + [PEP 8](http://legacy.python.org/dev/peps/pep-0008/). - Remember that “a foolish consistency is the hobgoblin of little minds,” so use your best judgement to write the clearest code for your particular case. diff --git a/docs/install_apt.md b/docs/install_apt.md index b6cb1c2d..e361a92d 100644 --- a/docs/install_apt.md +++ b/docs/install_apt.md @@ -40,6 +40,7 @@ Continue with [compilation](installation.html#compilation). sudo apt-get install libprotobuf-dev libleveldb-dev libsnappy-dev libopencv-dev libhdf5-serial-dev protobuf-compiler sudo apt-get install --no-install-recommends libboost-all-dev + sudo apt-get install libgflags-dev libgoogle-glog-dev liblmdb-dev **CUDA**: Install by `apt-get` or the NVIDIA `.run` package. The NVIDIA package tends to follow more recent library and driver versions, but the installation is more manual. @@ -54,12 +55,6 @@ This can be skipped for CPU-only installation. CUDA 8 is required on Ubuntu 16.04. -**Remaining dependencies, 14.04** - -Everything is packaged in 14.04. - - sudo apt-get install libgflags-dev libgoogle-glog-dev liblmdb-dev - **Remaining dependencies, 12.04** These dependencies need manual installation in 12.04. diff --git a/docs/installation.md b/docs/installation.md index 42f1d0ce..c4822853 100644 --- a/docs/installation.md +++ b/docs/installation.md @@ -17,7 +17,7 @@ The official Makefile and `Makefile.config` build are complemented by a [communi - [RHEL / CentOS / Fedora installation](install_yum.html) - [Windows](https://github.com/BVLC/caffe/tree/windows) *see the Windows branch led by Guillaume Dumont* - [OpenCL](https://github.com/BVLC/caffe/tree/opencl) *see the OpenCL branch led by Fabian Tschopp* -- [AWS AMI](https://github.com/bitfusionio/amis/tree/master/awsmrkt-bfboost-ubuntu14-cuda75-caffe) *pre-configured for AWS* +- [AWS AMI](https://aws.amazon.com/marketplace/pp/B01M0AXXQB) *official deep learning amazon machine image from AWS* **Overview**: @@ -42,14 +42,14 @@ Optional dependencies: * [OpenCV](http://opencv.org/) >= 2.4 including 3.0 * IO libraries: `lmdb`, `leveldb` (note: leveldb requires `snappy`) -* cuDNN for GPU acceleration (v6) +* cuDNN for GPU acceleration (v7) Pycaffe and Matcaffe interfaces have their own natural needs. * For Python Caffe: `Python 2.7` or `Python 3.3+`, `numpy (>= 1.7)`, boost-provided `boost.python` * For MATLAB Caffe: MATLAB with the `mex` compiler. -**cuDNN Caffe**: for fastest operation Caffe is accelerated by drop-in integration of [NVIDIA cuDNN](https://developer.nvidia.com/cudnn). To speed up your Caffe models, install cuDNN then uncomment the `USE_CUDNN := 1` flag in `Makefile.config` when installing Caffe. Acceleration is automatic. The current version is cuDNN v6; older versions are supported in older Caffe. +**cuDNN Caffe**: for fastest operation Caffe is accelerated by drop-in integration of [NVIDIA cuDNN](https://developer.nvidia.com/cudnn). To speed up your Caffe models, install cuDNN then uncomment the `USE_CUDNN := 1` flag in `Makefile.config` when installing Caffe. Acceleration is automatic. The current version is cuDNN v7; older versions are supported in older Caffe. **CPU-only Caffe**: for cold-brewed CPU-only Caffe uncomment the `CPU_ONLY := 1` flag in `Makefile.config` to configure and build Caffe without CUDA. This is helpful for cloud or cluster deployment. @@ -80,7 +80,7 @@ The main requirements are `numpy` and `boost.python` (provided by boost). `panda You can install the dependencies with - for req in $(cat requirements.txt); do pip install $req; done + pip install -r requirements.txt but we suggest first installing the [Anaconda](https://store.continuum.io/cshop/anaconda/) Python distribution, which provides most of the necessary packages, as well as the `hdf5` library dependency. diff --git a/docs/tutorial/interfaces.md b/docs/tutorial/interfaces.md index b5a4f1ad..2578af5d 100644 --- a/docs/tutorial/interfaces.md +++ b/docs/tutorial/interfaces.md @@ -129,8 +129,8 @@ Use CPU: Use GPU and specify its gpu_id: - caffe.set_mode_gpu(); caffe.set_device(gpu_id); + caffe.set_mode_gpu(); #### Create a network and access its layers and blobs diff --git a/docs/tutorial/layers.md b/docs/tutorial/layers.md index 2faacc58..5036d4fd 100644 --- a/docs/tutorial/layers.md +++ b/docs/tutorial/layers.md @@ -87,12 +87,13 @@ Layers: * [ELU](layers/elu.html) - exponential linear rectification. * [Sigmoid](layers/sigmoid.html) * [TanH](layers/tanh.html) -* [Absolute Value](layers/abs.html) +* [Absolute Value](layers/absval.html) * [Power](layers/power.html) - f(x) = (shift + scale * x) ^ power. * [Exp](layers/exp.html) - f(x) = base ^ (shift + scale * x). * [Log](layers/log.html) - f(x) = log(x). * [BNLL](layers/bnll.html) - f(x) = log(1 + exp(x)). * [Threshold](layers/threshold.html) - performs step function at user defined threshold. +* [Clip](layers/clip.html) - clips a blob between a fixed minimum and maximum value. * [Bias](layers/bias.html) - adds a bias to a blob that can either be learned or fixed. * [Scale](layers/scale.html) - scales a blob by an amount that can either be learned or fixed. diff --git a/docs/tutorial/layers/clip.md b/docs/tutorial/layers/clip.md new file mode 100644 index 00000000..d6a20f5f --- /dev/null +++ b/docs/tutorial/layers/clip.md @@ -0,0 +1,20 @@ +--- +title: Clip Layer +--- + +# Clip Layer + +* Layer type: `Clip` +* [Doxygen Documentation](http://caffe.berkeleyvision.org/doxygen/classcaffe_1_1ClipLayer.html) +* Header: [`./include/caffe/layers/clip_layer.hpp`](https://github.com/BVLC/caffe/blob/master/include/caffe/layers/clip_layer.hpp) +* CPU implementation: [`./src/caffe/layers/clip_layer.cpp`](https://github.com/BVLC/caffe/blob/master/src/caffe/layers/clip_layer.cpp) +* CUDA GPU implementation: [`./src/caffe/layers/clip_layer.cu`](https://github.com/BVLC/caffe/blob/master/src/caffe/layers/clip_layer.cu) + +## Parameters + +* Parameters (`ClipParameter clip_param`) +* From [`./src/caffe/proto/caffe.proto`](https://github.com/BVLC/caffe/blob/master/src/caffe/proto/caffe.proto): + +{% highlight Protobuf %} +{% include proto/ClipParameter.txt %} +{% endhighlight %} diff --git a/examples/web_demo/readme.md b/examples/web_demo/readme.md index fe74b9ef..e50c4f10 100644 --- a/examples/web_demo/readme.md +++ b/examples/web_demo/readme.md @@ -11,7 +11,7 @@ priority: 10 ## Requirements The demo server requires Python with some dependencies. -To make sure you have the dependencies, please run `pip install -r examples/web_demo/requirements.txt`, and also make sure that you've compiled the Python Caffe interface and that it is on your `PYTHONPATH` (see [installation instructions](/installation.html)). +To make sure you have the dependencies, please run `pip install -r examples/web_demo/requirements.txt`, and also make sure that you've compiled the Python Caffe interface and that it is on your `PYTHONPATH` (see [installation instructions](http://caffe.berkeleyvision.org/installation.html)). Make sure that you have obtained the Reference CaffeNet Model and the ImageNet Auxiliary Data: diff --git a/include/caffe/filler.hpp b/include/caffe/filler.hpp index dad9ad46..a4477361 100644 --- a/include/caffe/filler.hpp +++ b/include/caffe/filler.hpp @@ -108,9 +108,9 @@ class PositiveUnitballFiller : public Filler { caffe_rng_uniform(blob->count(), 0, 1, blob->mutable_cpu_data()); // We expect the filler to not be called very frequently, so we will // just use a simple implementation - int dim = blob->count() / blob->num(); + int dim = blob->count() / blob->shape(0); CHECK(dim); - for (int i = 0; i < blob->num(); ++i) { + for (int i = 0; i < blob->shape(0); ++i) { Dtype sum = 0; for (int j = 0; j < dim; ++j) { sum += data[i * dim + j]; @@ -147,8 +147,11 @@ class XavierFiller : public Filler { : Filler(param) {} virtual void Fill(Blob* blob) { CHECK(blob->count()); - int fan_in = blob->count() / blob->num(); - int fan_out = blob->count() / blob->channels(); + int fan_in = blob->count() / blob->shape(0); + // Compatibility with ND blobs + int fan_out = blob->num_axes() > 1 ? + blob->count() / blob->shape(1) : + blob->count(); Dtype n = fan_in; // default to fan_in if (this->filler_param_.variance_norm() == FillerParameter_VarianceNorm_AVERAGE) { @@ -189,8 +192,11 @@ class MSRAFiller : public Filler { : Filler(param) {} virtual void Fill(Blob* blob) { CHECK(blob->count()); - int fan_in = blob->count() / blob->num(); - int fan_out = blob->count() / blob->channels(); + int fan_in = blob->count() / blob->shape(0); + // Compatibility with ND blobs + int fan_out = blob->num_axes() > 1 ? + blob->count() / blob->shape(1) : + blob->count(); Dtype n = fan_in; // default to fan_in if (this->filler_param_.variance_norm() == FillerParameter_VarianceNorm_AVERAGE) { @@ -250,10 +256,10 @@ class BilinearFiller : public Filler { CHECK_EQ(blob->width(), blob->height()) << "Filter must be square"; Dtype* data = blob->mutable_cpu_data(); int f = ceil(blob->width() / 2.); - float c = (2 * f - 1 - f % 2) / (2. * f); + Dtype c = (blob->width() - 1) / (2. * f); for (int i = 0; i < blob->count(); ++i) { - float x = i % blob->width(); - float y = (i / blob->width()) % blob->height(); + Dtype x = i % blob->width(); + Dtype y = (i / blob->width()) % blob->height(); data[i] = (1 - fabs(x / f - c)) * (1 - fabs(y / f - c)); } CHECK_EQ(this->filler_param_.sparse(), -1) diff --git a/include/caffe/layers/accuracy_layer.hpp b/include/caffe/layers/accuracy_layer.hpp index a9ad3225..dd2247b9 100644 --- a/include/caffe/layers/accuracy_layer.hpp +++ b/include/caffe/layers/accuracy_layer.hpp @@ -68,6 +68,8 @@ class AccuracyLayer : public Layer { */ virtual void Forward_cpu(const vector*>& bottom, const vector*>& top); + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); /// @brief Not implemented -- AccuracyLayer cannot be used as a loss. @@ -77,6 +79,8 @@ class AccuracyLayer : public Layer { if (propagate_down[i]) { NOT_IMPLEMENTED; } } } + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); int label_axis_, outer_num_, inner_num_; diff --git a/include/caffe/layers/clip_layer.hpp b/include/caffe/layers/clip_layer.hpp new file mode 100644 index 00000000..2788193e --- /dev/null +++ b/include/caffe/layers/clip_layer.hpp @@ -0,0 +1,75 @@ +#ifndef CAFFE_CLIP_LAYER_HPP_ +#define CAFFE_CLIP_LAYER_HPP_ + +#include + +#include "caffe/blob.hpp" +#include "caffe/layer.hpp" +#include "caffe/proto/caffe.pb.h" + +#include "caffe/layers/neuron_layer.hpp" + +namespace caffe { + +/** + * @brief Clip: @f$ y = \max(min, \min(max, x)) @f$. + */ +template +class ClipLayer : public NeuronLayer { + public: + /** + * @param param provides ClipParameter clip_param, + * with ClipLayer options: + * - min + * - max + */ + explicit ClipLayer(const LayerParameter& param) + : NeuronLayer(param) {} + + virtual inline const char* type() const { return "Clip"; } + + protected: + /** + * @param bottom input Blob vector (length 1) + * -# @f$ (N \times C \times H \times W) @f$ + * the inputs @f$ x @f$ + * @param top output Blob vector (length 1) + * -# @f$ (N \times C \times H \times W) @f$ + * the computed outputs @f$ + * y = \max(min, \min(max, x)) + * @f$ + */ + virtual void Forward_cpu(const vector*>& bottom, + const vector*>& top); + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); + + /** + * @brief Computes the error gradient w.r.t. the clipped inputs. + * + * @param top output Blob vector (length 1), providing the error gradient with + * respect to the outputs + * -# @f$ (N \times C \times H \times W) @f$ + * containing error gradients @f$ \frac{\partial E}{\partial y} @f$ + * with respect to computed outputs @f$ y @f$ + * @param propagate_down see Layer::Backward. + * @param bottom input Blob vector (length 1) + * -# @f$ (N \times C \times H \times W) @f$ + * the inputs @f$ x @f$; Backward fills their diff with + * gradients @f$ + * \frac{\partial E}{\partial x} = \left\{ + * \begin{array}{lr} + * 0 & \mathrm{if} \; x < min \vee x > max \\ + * \frac{\partial E}{\partial y} & \mathrm{if} \; x \ge min \wedge x \le max + * \end{array} \right. + * @f$ + */ + virtual void Backward_cpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); +}; + +} // namespace caffe + +#endif // CAFFE_CLIP_LAYER_HPP_ diff --git a/include/caffe/layers/cudnn_deconv_layer.hpp b/include/caffe/layers/cudnn_deconv_layer.hpp new file mode 100644 index 00000000..12799e5b --- /dev/null +++ b/include/caffe/layers/cudnn_deconv_layer.hpp @@ -0,0 +1,68 @@ +#ifndef CAFFE_CUDNN_DECONV_LAYER_HPP_ +#define CAFFE_CUDNN_DECONV_LAYER_HPP_ + +#include + +#include "caffe/blob.hpp" +#include "caffe/layer.hpp" +#include "caffe/proto/caffe.pb.h" + +#include "caffe/layers/deconv_layer.hpp" + +namespace caffe { + +#ifdef USE_CUDNN +/* + * @brief cuDNN implementation of DeConvolutionLayer. + * Fallback to DeConvolutionLayer for CPU mode. + * + * cuDNN accelerates deconvolution through forward kernels for filtering and + * bias plus backward kernels for the gradient w.r.t. the filters, biases, and + * inputs. Caffe + cuDNN further speeds up the computation through forward + * parallelism across groups and backward parallelism across gradients. +*/ +template +class CuDNNDeconvolutionLayer : public DeconvolutionLayer { + public: + explicit CuDNNDeconvolutionLayer(const LayerParameter& param) + : DeconvolutionLayer(param), handles_setup_(false) {} + virtual void LayerSetUp(const vector*>& bottom, + const vector*>& top); + virtual void Reshape(const vector*>& bottom, + const vector*>& top); + virtual ~CuDNNDeconvolutionLayer(); + + protected: + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom); + + bool handles_setup_; + cudnnHandle_t* handle_; + cudaStream_t* stream_; + + // algorithms for forward and backwards convolutions + cudnnConvolutionFwdAlgo_t *fwd_algo_; + cudnnConvolutionBwdFilterAlgo_t *bwd_filter_algo_; + cudnnConvolutionBwdDataAlgo_t *bwd_data_algo_; + + vector bottom_descs_, top_descs_; + cudnnTensorDescriptor_t bias_desc_; + cudnnFilterDescriptor_t filter_desc_; + vector conv_descs_; + int bottom_offset_, top_offset_, bias_offset_; + + size_t *workspace_fwd_sizes_; + size_t *workspace_bwd_data_sizes_; + size_t *workspace_bwd_filter_sizes_; + size_t workspaceSizeInBytes; // size of underlying storage + void *workspaceData; // underlying storage + void **workspace; // aliases into workspaceData +}; +#endif + +} // namespace caffe + +#endif // CAFFE_CUDNN_DECONV_LAYER_HPP_ diff --git a/include/caffe/layers/euclidean_loss_layer.hpp b/include/caffe/layers/euclidean_loss_layer.hpp index f564569e..24568c54 100644 --- a/include/caffe/layers/euclidean_loss_layer.hpp +++ b/include/caffe/layers/euclidean_loss_layer.hpp @@ -30,7 +30,7 @@ namespace caffe { * This can be used for least-squares regression tasks. An InnerProductLayer * input to a EuclideanLossLayer exactly formulates a linear least squares * regression problem. With non-zero weight decay the problem becomes one of - * ridge regression -- see src/caffe/test/test_sgd_solver.cpp for a concrete + * ridge regression -- see src/caffe/test/test_gradient_based_solver.cpp for a concrete * example wherein we check that the gradients computed for a Net with exactly * this structure match hand-computed gradient formulas for ridge regression. * diff --git a/include/caffe/layers/infogain_loss_layer.hpp b/include/caffe/layers/infogain_loss_layer.hpp index edecde82..3b3caa27 100644 --- a/include/caffe/layers/infogain_loss_layer.hpp +++ b/include/caffe/layers/infogain_loss_layer.hpp @@ -13,20 +13,21 @@ namespace caffe { /** - * @brief A generalization of MultinomialLogisticLossLayer that takes an + * @brief A generalization of SoftmaxWithLossLayer that takes an * "information gain" (infogain) matrix specifying the "value" of all label * pairs. * - * Equivalent to the MultinomialLogisticLossLayer if the infogain matrix is the + * Equivalent to the SoftmaxWithLossLayer if the infogain matrix is the * identity. * * @param bottom input Blob vector (length 2-3) * -# @f$ (N \times C \times H \times W) @f$ - * the predictions @f$ \hat{p} @f$, a Blob with values in - * @f$ [0, 1] @f$ indicating the predicted probability of each of the - * @f$ K = CHW @f$ classes. Each prediction vector @f$ \hat{p}_n @f$ - * should sum to 1 as in a probability distribution: @f$ - * \forall n \sum\limits_{k=1}^K \hat{p}_{nk} = 1 @f$. + * the predictions @f$ x @f$, a Blob with values in + * @f$ [-\infty, +\infty] @f$ indicating the predicted score for each of + * the @f$ K = CHW @f$ classes. This layer maps these scores to a + * probability distribution over classes using the softmax function + * @f$ \hat{p}_{nk} = \exp(x_{nk}) / + * \left[\sum_{k'} \exp(x_{nk'})\right] @f$ (see SoftmaxLayer). * -# @f$ (N \times 1 \times 1 \times 1) @f$ * the labels @f$ l @f$, an integer-valued Blob with values * @f$ l_n \in [0, 1, 2, ..., K - 1] @f$ @@ -35,7 +36,7 @@ namespace caffe { * (\b optional) the infogain matrix @f$ H @f$. This must be provided as * the third bottom blob input if not provided as the infogain_mat in the * InfogainLossParameter. If @f$ H = I @f$, this layer is equivalent to the - * MultinomialLogisticLossLayer. + * SoftmaxWithLossLayer. * @param top output Blob vector (length 1) * -# @f$ (1 \times 1 \times 1 \times 1) @f$ * the computed infogain multinomial logistic loss: @f$ E = @@ -98,8 +99,8 @@ class InfogainLossLayer : public LossLayer { * infogain matrix, if provided as bottom[2]) * @param bottom input Blob vector (length 2-3) * -# @f$ (N \times C \times H \times W) @f$ - * the predictions @f$ \hat{p} @f$; Backward computes diff - * @f$ \frac{\partial E}{\partial \hat{p}} @f$ + * the predictions @f$ x @f$; Backward computes diff + * @f$ \frac{\partial E}{\partial x} @f$ * -# @f$ (N \times 1 \times 1 \times 1) @f$ * the labels -- ignored as we can't compute their error gradients * -# @f$ (1 \times 1 \times K \times K) @f$ diff --git a/include/caffe/layers/pooling_layer.hpp b/include/caffe/layers/pooling_layer.hpp index f4d6803b..38a43283 100644 --- a/include/caffe/layers/pooling_layer.hpp +++ b/include/caffe/layers/pooling_layer.hpp @@ -51,6 +51,7 @@ class PoolingLayer : public Layer { int height_, width_; int pooled_height_, pooled_width_; bool global_pooling_; + PoolingParameter_RoundMode round_mode_; Blob rand_idx_; Blob max_idx_; }; diff --git a/include/caffe/layers/swish_layer.hpp b/include/caffe/layers/swish_layer.hpp new file mode 100644 index 00000000..d538ff6d --- /dev/null +++ b/include/caffe/layers/swish_layer.hpp @@ -0,0 +1,96 @@ +#ifndef CAFFE_SWISH_LAYER_HPP_ +#define CAFFE_SWISH_LAYER_HPP_ + +#include + +#include "caffe/blob.hpp" +#include "caffe/layer.hpp" +#include "caffe/proto/caffe.pb.h" + +#include "caffe/layers/neuron_layer.hpp" +#include "caffe/layers/sigmoid_layer.hpp" + +namespace caffe { + +/** + * @brief Swish non-linearity @f$ y = x \sigma (\beta x) @f$. + * A novel activation function that tends to work better than ReLU [1]. + * + * [1] Prajit Ramachandran, Barret Zoph, Quoc V. Le. "Searching for + * Activation Functions". arXiv preprint arXiv:1710.05941v2 (2017). + */ +template +class SwishLayer : public NeuronLayer { + public: + /** + * @param param provides SwishParameter swish_param, + * with SwishLayer options: + * - beta (\b optional, default 1). + * the value @f$ \beta @f$ in the @f$ y = x \sigma (\beta x) @f$. + */ + explicit SwishLayer(const LayerParameter& param) + : NeuronLayer(param), + sigmoid_layer_(new SigmoidLayer(param)), + sigmoid_input_(new Blob()), + sigmoid_output_(new Blob()) {} + virtual void LayerSetUp(const vector*>& bottom, + const vector*>& top); + virtual void Reshape(const vector*>& bottom, + const vector*>& top); + + virtual inline const char* type() const { return "Swish"; } + + protected: + /** + * @param bottom input Blob vector (length 1) + * -# @f$ (N \times C \times H \times W) @f$ + * the inputs @f$ x @f$ + * @param top output Blob vector (length 1) + * -# @f$ (N \times C \times H \times W) @f$ + * the computed outputs @f$ + * y = x \sigma (\beta x) + * @f$. + */ + virtual void Forward_cpu(const vector*>& bottom, + const vector*>& top); + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); + + /** + * @brief Computes the error gradient w.r.t. the sigmoid inputs. + * + * @param top output Blob vector (length 1), providing the error gradient with + * respect to the outputs + * -# @f$ (N \times C \times H \times W) @f$ + * containing error gradients @f$ \frac{\partial E}{\partial y} @f$ + * with respect to computed outputs @f$ y @f$ + * @param propagate_down see Layer::Backward. + * @param bottom input Blob vector (length 1) + * -# @f$ (N \times C \times H \times W) @f$ + * the inputs @f$ x @f$; Backward fills their diff with + * gradients @f$ + * \frac{\partial E}{\partial x} + * = \frac{\partial E}{\partial y}(\beta y + + * \sigma (\beta x)(1 - \beta y)) + * @f$ if propagate_down[0] + */ + virtual void Backward_cpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + + /// The internal SigmoidLayer + shared_ptr > sigmoid_layer_; + /// sigmoid_input_ stores the input of the SigmoidLayer. + shared_ptr > sigmoid_input_; + /// sigmoid_output_ stores the output of the SigmoidLayer. + shared_ptr > sigmoid_output_; + /// bottom vector holder to call the underlying SigmoidLayer::Forward + vector*> sigmoid_bottom_vec_; + /// top vector holder to call the underlying SigmoidLayer::Forward + vector*> sigmoid_top_vec_; +}; + +} // namespace caffe + +#endif // CAFFE_SWISH_LAYER_HPP_ diff --git a/include/caffe/net.hpp b/include/caffe/net.hpp index d3c9306e..143d5d28 100644 --- a/include/caffe/net.hpp +++ b/include/caffe/net.hpp @@ -111,9 +111,9 @@ class Net { * another Net. */ void CopyTrainedLayersFrom(const NetParameter& param); - void CopyTrainedLayersFrom(const string trained_filename); - void CopyTrainedLayersFromBinaryProto(const string trained_filename); - void CopyTrainedLayersFromHDF5(const string trained_filename); + void CopyTrainedLayersFrom(const string& trained_filename); + void CopyTrainedLayersFromBinaryProto(const string& trained_filename); + void CopyTrainedLayersFromHDF5(const string& trained_filename); /// @brief Writes the net to a proto. void ToProto(NetParameter* param, bool write_diff = false) const; /// @brief Writes the net to an HDF5 file. diff --git a/include/caffe/sgd_solvers.hpp b/include/caffe/sgd_solvers.hpp index 1fc52d87..925ff783 100644 --- a/include/caffe/sgd_solvers.hpp +++ b/include/caffe/sgd_solvers.hpp @@ -23,10 +23,11 @@ class SGDSolver : public Solver { const vector > >& history() { return history_; } + virtual void ApplyUpdate(); + Dtype GetLearningRate(); + protected: void PreSolve(); - Dtype GetLearningRate(); - virtual void ApplyUpdate(); virtual void Normalize(int param_id); virtual void Regularize(int param_id); virtual void ComputeUpdateValue(int param_id, Dtype rate); diff --git a/include/caffe/solver.hpp b/include/caffe/solver.hpp index a28d8cb8..7a0d7777 100644 --- a/include/caffe/solver.hpp +++ b/include/caffe/solver.hpp @@ -55,7 +55,7 @@ class Solver { // The main entry of the solver function. In default, iter will be zero. Pass // in a non-zero iter number to resume training for a pre-trained net. virtual void Solve(const char* resume_file = NULL); - inline void Solve(const string resume_file) { Solve(resume_file.c_str()); } + inline void Solve(const string& resume_file) { Solve(resume_file.c_str()); } void Step(int iters); // The Restore method simply dispatches to one of the // RestoreSolverStateFrom___ protected methods. You should implement these @@ -94,10 +94,11 @@ class Solver { */ virtual inline const char* type() const { return ""; } - protected: // Make and apply the update value for the current iteration. virtual void ApplyUpdate() = 0; - string SnapshotFilename(const string extension); + + protected: + string SnapshotFilename(const string& extension); string SnapshotToBinaryProto(); string SnapshotToHDF5(); // The test routine diff --git a/include/caffe/syncedmem.hpp b/include/caffe/syncedmem.hpp index 317ce29a..8d650a34 100644 --- a/include/caffe/syncedmem.hpp +++ b/include/caffe/syncedmem.hpp @@ -66,8 +66,8 @@ class SyncedMemory { void* mutable_cpu_data(); void* mutable_gpu_data(); enum SyncedHead { UNINITIALIZED, HEAD_AT_CPU, HEAD_AT_GPU, SYNCED }; - SyncedHead head() { return head_; } - size_t size() { return size_; } + SyncedHead head() const { return head_; } + size_t size() const { return size_; } #ifndef CPU_ONLY void async_gpu_push(const cudaStream_t& stream); diff --git a/include/caffe/util/cudnn.hpp b/include/caffe/util/cudnn.hpp index 7b798fe0..68edd406 100644 --- a/include/caffe/util/cudnn.hpp +++ b/include/caffe/util/cudnn.hpp @@ -45,6 +45,12 @@ inline const char* cudnnGetErrorString(cudnnStatus_t status) { #if CUDNN_VERSION_MIN(6, 0, 0) case CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING: return "CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING"; +#endif +#if CUDNN_VERSION_MIN(7, 0, 0) + case CUDNN_STATUS_RUNTIME_IN_PROGRESS: + return "CUDNN_STATUS_RUNTIME_IN_PROGRESS"; + case CUDNN_STATUS_RUNTIME_FP_OVERFLOW: + return "CUDNN_STATUS_RUNTIME_FP_OVERFLOW"; #endif } return "Unknown cudnn status"; diff --git a/include/caffe/util/hdf5.hpp b/include/caffe/util/hdf5.hpp index 71549c1c..dbd8bb6c 100644 --- a/include/caffe/util/hdf5.hpp +++ b/include/caffe/util/hdf5.hpp @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 #ifndef CAFFE_UTIL_HDF5_H_ #define CAFFE_UTIL_HDF5_H_ @@ -37,3 +38,4 @@ string hdf5_get_name_by_idx(hid_t loc_id, int idx); } // namespace caffe #endif // CAFFE_UTIL_HDF5_H_ +#endif // USE_HDF5 diff --git a/include/caffe/util/signal_handler.h b/include/caffe/util/signal_handler.h index fb84c65b..52463325 100644 --- a/include/caffe/util/signal_handler.h +++ b/include/caffe/util/signal_handler.h @@ -8,7 +8,7 @@ namespace caffe { class SignalHandler { public: - // Contructor. Specify what action to take when a signal is received. + // Constructor. Specify what action to take when a signal is received. SignalHandler(SolverAction::Enum SIGINT_action, SolverAction::Enum SIGHUP_action); ~SignalHandler(); diff --git a/python/caffe/_caffe.cpp b/python/caffe/_caffe.cpp index 6ddf4621..e5d32a31 100644 --- a/python/caffe/_caffe.cpp +++ b/python/caffe/_caffe.cpp @@ -419,7 +419,7 @@ BOOST_PYTHON_MODULE(_caffe) { .def("reshape", &Net::Reshape) .def("clear_param_diffs", &Net::ClearParamDiffs) // The cast is to select a particular overload. - .def("copy_from", static_cast::*)(const string)>( + .def("copy_from", static_cast::*)(const string&)>( &Net::CopyTrainedLayersFrom)) .def("share_with", &Net::ShareTrainedLayersWith) .add_property("_blob_loss_weights", bp::make_function( @@ -494,7 +494,9 @@ BOOST_PYTHON_MODULE(_caffe) { bp::class_("SolverParameter", bp::no_init) .add_property("max_iter", &SolverParameter::max_iter) .add_property("display", &SolverParameter::display) - .add_property("layer_wise_reduce", &SolverParameter::layer_wise_reduce); + .add_property("layer_wise_reduce", &SolverParameter::layer_wise_reduce) + .add_property("base_lr", &SolverParameter::base_lr, + &SolverParameter::set_base_lr); bp::class_("LayerParameter", bp::no_init); bp::class_, shared_ptr >, boost::noncopyable>( @@ -511,26 +513,28 @@ BOOST_PYTHON_MODULE(_caffe) { .def("restore", &Solver::Restore) .def("snapshot", &Solver::Snapshot) .def("share_weights", &share_weights) + .def("apply_update", &Solver::ApplyUpdate) .add_property("param", bp::make_function(&Solver::param, - bp::return_value_policy())); + bp::return_internal_reference<>())); BP_REGISTER_SHARED_PTR_TO_PYTHON(Solver); bp::class_, bp::bases >, shared_ptr >, boost::noncopyable>( - "SGDSolver", bp::init()); - bp::class_, bp::bases >, + "SGDSolver", bp::init()) + .add_property("lr", &SGDSolver::GetLearningRate); + bp::class_, bp::bases >, shared_ptr >, boost::noncopyable>( "NesterovSolver", bp::init()); - bp::class_, bp::bases >, + bp::class_, bp::bases >, shared_ptr >, boost::noncopyable>( "AdaGradSolver", bp::init()); - bp::class_, bp::bases >, + bp::class_, bp::bases >, shared_ptr >, boost::noncopyable>( "RMSPropSolver", bp::init()); - bp::class_, bp::bases >, + bp::class_, bp::bases >, shared_ptr >, boost::noncopyable>( "AdaDeltaSolver", bp::init()); - bp::class_, bp::bases >, + bp::class_, bp::bases >, shared_ptr >, boost::noncopyable>( "AdamSolver", bp::init()); diff --git a/python/caffe/classifier.py b/python/caffe/classifier.py index 983760a7..64d804be 100644 --- a/python/caffe/classifier.py +++ b/python/caffe/classifier.py @@ -23,7 +23,7 @@ class Classifier(caffe.Net): def __init__(self, model_file, pretrained_file, image_dims=None, mean=None, input_scale=None, raw_scale=None, channel_swap=None): - caffe.Net.__init__(self, model_file, pretrained_file, caffe.TEST) + caffe.Net.__init__(self, model_file, caffe.TEST, weights=pretrained_file) # configure pre-processing in_ = self.inputs[0] diff --git a/python/caffe/detector.py b/python/caffe/detector.py index ef1f9173..ceee5d36 100644 --- a/python/caffe/detector.py +++ b/python/caffe/detector.py @@ -35,7 +35,7 @@ class Detector(caffe.Net): def __init__(self, model_file, pretrained_file, mean=None, input_scale=None, raw_scale=None, channel_swap=None, context_pad=None): - caffe.Net.__init__(self, model_file, pretrained_file, caffe.TEST) + caffe.Net.__init__(self, model_file, caffe.TEST, weights=pretrained_file) # configure pre-processing in_ = self.inputs[0] diff --git a/python/caffe/draw.py b/python/caffe/draw.py index 8411a41d..0061f490 100644 --- a/python/caffe/draw.py +++ b/python/caffe/draw.py @@ -59,18 +59,60 @@ def get_edge_label(layer): return edge_label -def get_layer_label(layer, rankdir): +def get_layer_lr_mult(layer): + """Get the learning rate multipliers. + + Get the learning rate multipliers for the given layer. Assumes a + Convolution/Deconvolution/InnerProduct layer. + + Parameters + ---------- + layer : caffe_pb2.LayerParameter + A Convolution, Deconvolution, or InnerProduct layer. + + Returns + ------- + learning_rates : tuple of floats + the learning rate multipliers for the weights and biases. + """ + if layer.type not in ['Convolution', 'Deconvolution', 'InnerProduct']: + raise ValueError("%s layers do not have a " + "learning rate multiplier" % layer.type) + + if not hasattr(layer, 'param'): + return (1.0, 1.0) + + params = getattr(layer, 'param') + + if len(params) == 0: + return (1.0, 1.0) + + if len(params) == 1: + lrm0 = getattr(params[0],'lr_mult', 1.0) + return (lrm0, 1.0) + + if len(params) == 2: + lrm0, lrm1 = [getattr(p,'lr_mult', 1.0) for p in params] + return (lrm0, lrm1) + + raise ValueError("Could not parse the learning rate multiplier") + + +def get_layer_label(layer, rankdir, display_lrm=False): """Define node label based on layer type. Parameters ---------- - layer : ? + layer : caffe_pb2.LayerParameter rankdir : {'LR', 'TB', 'BT'} Direction of graph layout. + display_lrm : boolean, optional + If True include the learning rate multipliers in the label (default is + False). Returns ------- - string : + node_label : string A label for the current layer """ @@ -81,36 +123,54 @@ def get_layer_label(layer, rankdir): else: # If graph orientation is horizontal, vertical space is free and # horizontal space is not; separate words with newlines - separator = '\\n' - - if layer.type == 'Convolution' or layer.type == 'Deconvolution': - # Outer double quotes needed or else colon characters don't parse - # properly - node_label = '"%s%s(%s)%skernel size: %d%sstride: %d%spad: %d"' %\ - (layer.name, - separator, - layer.type, - separator, - layer.convolution_param.kernel_size[0] if len(layer.convolution_param.kernel_size) else 1, - separator, - layer.convolution_param.stride[0] if len(layer.convolution_param.stride) else 1, - separator, - layer.convolution_param.pad[0] if len(layer.convolution_param.pad) else 0) - elif layer.type == 'Pooling': + separator = r'\n' + + # Initializes a list of descriptors that will be concatenated into the + # `node_label` + descriptors_list = [] + # Add the layer's name + descriptors_list.append(layer.name) + # Add layer's type + if layer.type == 'Pooling': pooling_types_dict = get_pooling_types_dict() - node_label = '"%s%s(%s %s)%skernel size: %d%sstride: %d%spad: %d"' %\ - (layer.name, - separator, - pooling_types_dict[layer.pooling_param.pool], - layer.type, - separator, - layer.pooling_param.kernel_size, - separator, - layer.pooling_param.stride, - separator, - layer.pooling_param.pad) + layer_type = '(%s %s)' % (layer.type, + pooling_types_dict[layer.pooling_param.pool]) else: - node_label = '"%s%s(%s)"' % (layer.name, separator, layer.type) + layer_type = '(%s)' % layer.type + descriptors_list.append(layer_type) + + # Describe parameters for spatial operation layers + if layer.type in ['Convolution', 'Deconvolution', 'Pooling']: + if layer.type == 'Pooling': + kernel_size = layer.pooling_param.kernel_size + stride = layer.pooling_param.stride + padding = layer.pooling_param.pad + else: + kernel_size = layer.convolution_param.kernel_size[0] if \ + len(layer.convolution_param.kernel_size) else 1 + stride = layer.convolution_param.stride[0] if \ + len(layer.convolution_param.stride) else 1 + padding = layer.convolution_param.pad[0] if \ + len(layer.convolution_param.pad) else 0 + spatial_descriptor = separator.join([ + "kernel size: %d" % kernel_size, + "stride: %d" % stride, + "pad: %d" % padding, + ]) + descriptors_list.append(spatial_descriptor) + + # Add LR multiplier for learning layers + if display_lrm and layer.type in ['Convolution', 'Deconvolution', 'InnerProduct']: + lrm0, lrm1 = get_layer_lr_mult(layer) + if any([lrm0, lrm1]): + lr_mult = "lr mult: %.1f, %.1f" % (lrm0, lrm1) + descriptors_list.append(lr_mult) + + # Concatenate the descriptors into one label + node_label = separator.join(descriptors_list) + # Outer double quotes needed or else colon characters don't parse + # properly + node_label = '"%s"' % node_label return node_label @@ -127,7 +187,7 @@ def choose_color_by_layertype(layertype): return color -def get_pydot_graph(caffe_net, rankdir, label_edges=True, phase=None): +def get_pydot_graph(caffe_net, rankdir, label_edges=True, phase=None, display_lrm=False): """Create a data structure which represents the `caffe_net`. Parameters @@ -140,6 +200,9 @@ def get_pydot_graph(caffe_net, rankdir, label_edges=True, phase=None): phase : {caffe_pb2.Phase.TRAIN, caffe_pb2.Phase.TEST, None} optional Include layers from this network phase. If None, include all layers. (the default is None) + display_lrm : boolean, optional + If True display the learning rate multipliers when relevant (default is + False). Returns ------- @@ -164,7 +227,7 @@ def get_pydot_graph(caffe_net, rankdir, label_edges=True, phase=None): included = included and not layer_phase.phase == phase if not included: continue - node_label = get_layer_label(layer, rankdir) + node_label = get_layer_label(layer, rankdir, display_lrm=display_lrm) node_name = "%s_%s" % (layer.name, layer.type) if (len(layer.bottom) == 1 and len(layer.top) == 1 and layer.bottom[0] == layer.top[0]): @@ -202,7 +265,7 @@ def get_pydot_graph(caffe_net, rankdir, label_edges=True, phase=None): return pydot_graph -def draw_net(caffe_net, rankdir, ext='png', phase=None): +def draw_net(caffe_net, rankdir, ext='png', phase=None, display_lrm=False): """Draws a caffe net and returns the image string encoded using the given extension. @@ -214,16 +277,20 @@ def draw_net(caffe_net, rankdir, ext='png', phase=None): phase : {caffe_pb2.Phase.TRAIN, caffe_pb2.Phase.TEST, None} optional Include layers from this network phase. If None, include all layers. (the default is None) + display_lrm : boolean, optional + If True display the learning rate multipliers for the learning layers + (default is False). Returns ------- string : Postscript representation of the graph. """ - return get_pydot_graph(caffe_net, rankdir, phase=phase).create(format=ext) + return get_pydot_graph(caffe_net, rankdir, phase=phase, + display_lrm=display_lrm).create(format=ext) -def draw_net_to_file(caffe_net, filename, rankdir='LR', phase=None): +def draw_net_to_file(caffe_net, filename, rankdir='LR', phase=None, display_lrm=False): """Draws a caffe net, and saves it to file using the format given as the file extension. Use '.raw' to output raw text that you can manually feed to graphviz to draw graphs. @@ -238,7 +305,10 @@ def draw_net_to_file(caffe_net, filename, rankdir='LR', phase=None): phase : {caffe_pb2.Phase.TRAIN, caffe_pb2.Phase.TEST, None} optional Include layers from this network phase. If None, include all layers. (the default is None) + display_lrm : boolean, optional + If True display the learning rate multipliers for the learning layers + (default is False). """ ext = filename[filename.rfind('.')+1:] with open(filename, 'wb') as fid: - fid.write(draw_net(caffe_net, rankdir, ext, phase)) + fid.write(draw_net(caffe_net, rankdir, ext, phase, display_lrm)) diff --git a/python/caffe/io.py b/python/caffe/io.py index eb7aa34c..a2ac61fb 100644 --- a/python/caffe/io.py +++ b/python/caffe/io.py @@ -188,13 +188,14 @@ def deprocess(self, in_, data): def set_transpose(self, in_, order): """ - Set the input channel order for e.g. RGB to BGR conversion - as needed for the reference ImageNet model. + Set the order of dimensions, e.g. to convert OpenCV's HxWxC images + into CxHxW. Parameters ---------- - in_ : which input to assign this channel order + in_ : which input to assign this dimension order order : the order to transpose the dimensions + for example (2,0,1) changes HxWxC into CxHxW and (1,2,0) reverts """ self.__check_input(in_) if len(order) != len(self.inputs[in_]) - 1: @@ -258,7 +259,12 @@ def set_mean(self, in_, mean): if len(ms) != 3: raise ValueError('Mean shape invalid') if ms != self.inputs[in_][1:]: - raise ValueError('Mean shape incompatible with input shape.') + in_shape = self.inputs[in_][1:] + m_min, m_max = mean.min(), mean.max() + normal_mean = (mean - m_min) / (m_max - m_min) + mean = resize_image(normal_mean.transpose((1,2,0)), + in_shape[1:]).transpose((2,0,1)) * \ + (m_max - m_min) + m_min # repeat mean over all length (C3D) if len(self.inputs[in_]) == 5: diff --git a/python/caffe/test/test_net.py b/python/caffe/test/test_net.py index afd27690..ee1d38c3 100644 --- a/python/caffe/test/test_net.py +++ b/python/caffe/test/test_net.py @@ -72,41 +72,41 @@ def test_forward_backward(self): self.net.backward() def test_forward_start_end(self): - conv_blob=self.net.blobs['conv']; - ip_blob=self.net.blobs['ip_blob']; - sample_data=np.random.uniform(size=conv_blob.data.shape); - sample_data=sample_data.astype(np.float32); - conv_blob.data[:]=sample_data; - forward_blob=self.net.forward(start='ip',end='ip'); - self.assertIn('ip_blob',forward_blob); - - manual_forward=[]; + conv_blob=self.net.blobs['conv'] + ip_blob=self.net.blobs['ip_blob'] + sample_data=np.random.uniform(size=conv_blob.data.shape) + sample_data=sample_data.astype(np.float32) + conv_blob.data[:]=sample_data + forward_blob=self.net.forward(start='ip',end='ip') + self.assertIn('ip_blob',forward_blob) + + manual_forward=[] for i in range(0,conv_blob.data.shape[0]): dot=np.dot(self.net.params['ip'][0].data, - conv_blob.data[i].reshape(-1)); - manual_forward.append(dot+self.net.params['ip'][1].data); - manual_forward=np.array(manual_forward); + conv_blob.data[i].reshape(-1)) + manual_forward.append(dot+self.net.params['ip'][1].data) + manual_forward=np.array(manual_forward) - np.testing.assert_allclose(ip_blob.data,manual_forward,rtol=1e-3); + np.testing.assert_allclose(ip_blob.data,manual_forward,rtol=1e-3,atol=1e-5) def test_backward_start_end(self): - conv_blob=self.net.blobs['conv']; - ip_blob=self.net.blobs['ip_blob']; + conv_blob=self.net.blobs['conv'] + ip_blob=self.net.blobs['ip_blob'] sample_data=np.random.uniform(size=ip_blob.data.shape) - sample_data=sample_data.astype(np.float32); - ip_blob.diff[:]=sample_data; - backward_blob=self.net.backward(start='ip',end='ip'); - self.assertIn('conv',backward_blob); + sample_data=sample_data.astype(np.float32) + ip_blob.diff[:]=sample_data + backward_blob=self.net.backward(start='ip',end='ip') + self.assertIn('conv',backward_blob) - manual_backward=[]; + manual_backward=[] for i in range(0,conv_blob.data.shape[0]): dot=np.dot(self.net.params['ip'][0].data.transpose(), - sample_data[i].reshape(-1)); - manual_backward.append(dot); - manual_backward=np.array(manual_backward); - manual_backward=manual_backward.reshape(conv_blob.data.shape); + sample_data[i].reshape(-1)) + manual_backward.append(dot) + manual_backward=np.array(manual_backward) + manual_backward=manual_backward.reshape(conv_blob.data.shape) - np.testing.assert_allclose(conv_blob.diff,manual_backward,rtol=1e-3); + np.testing.assert_allclose(conv_blob.diff,manual_backward,rtol=1e-3,atol=1e-5) def test_clear_param_diffs(self): # Run a forward/backward step to have non-zero diffs diff --git a/python/caffe/test/test_solver.py b/python/caffe/test/test_solver.py index f618fded..50c9d541 100644 --- a/python/caffe/test/test_solver.py +++ b/python/caffe/test/test_solver.py @@ -38,6 +38,17 @@ def test_solve(self): self.solver.solve() self.assertEqual(self.solver.iter, 100) + def test_apply_update(self): + net = self.solver.net + data = net.layers[1].blobs[0].data[...] + # Reset the weights of that layer to 0 + data[...] = 0 + net.layers[1].blobs[0].diff[...] = 1 + # Apply the update, the initial learning rate should be 0.01 + self.solver.apply_update() + # Check that the new weights are -0.01, with a precision of 1e-7 + self.assertTrue((data - -0.01 * np.ones(data.shape)).max() < 1e-7) + def test_net_memory(self): """Check that nets survive after the solver is destroyed.""" diff --git a/python/draw_net.py b/python/draw_net.py index dfe70d26..23cae30a 100755 --- a/python/draw_net.py +++ b/python/draw_net.py @@ -33,6 +33,10 @@ def parse_args(): 'TEST, or ALL. If ALL, then all layers are drawn ' 'regardless of phase.'), default="ALL") + parser.add_argument('--display_lrm', action='store_true', + help=('Use this flag to visualize the learning rate ' + 'multiplier, when non-zero, for the learning ' + 'layers (Convolution, Deconvolution, InnerProduct).')) args = parser.parse_args() return args @@ -51,7 +55,7 @@ def main(): elif args.phase != "ALL": raise ValueError("Unknown phase: " + args.phase) caffe.draw.draw_net_to_file(net, args.output_image_file, args.rankdir, - phase) + phase, args.display_lrm) if __name__ == '__main__': diff --git a/python/train.py b/python/train.py index 5897f5dc..14a38b8c 100644 --- a/python/train.py +++ b/python/train.py @@ -63,8 +63,8 @@ def show_time(): def solve(proto, snapshot, gpus, timing, uid, rank): - caffe.set_mode_gpu() caffe.set_device(gpus[rank]) + caffe.set_mode_gpu() caffe.set_solver_count(len(gpus)) caffe.set_solver_rank(rank) caffe.set_multiprocess(True) diff --git a/scripts/travis/install-deps.sh b/scripts/travis/install-deps.sh index 2fa2a74a..abf9cf1c 100755 --- a/scripts/travis/install-deps.sh +++ b/scripts/travis/install-deps.sh @@ -106,7 +106,7 @@ if $WITH_CUDA ; then ln -s /usr/local/cuda-$CUDA_VERSION /usr/local/cuda if $WITH_CUDNN ; then - apt-get install -y --no-install-recommends libcudnn6-dev + apt-get install -y --no-install-recommends libcudnn7-dev fi fi diff --git a/src/caffe/layer_factory.cpp b/src/caffe/layer_factory.cpp index 3e1198c9..a8d75040 100644 --- a/src/caffe/layer_factory.cpp +++ b/src/caffe/layer_factory.cpp @@ -7,7 +7,9 @@ #include "caffe/layer.hpp" #include "caffe/layer_factory.hpp" +#include "caffe/layers/clip_layer.hpp" #include "caffe/layers/conv_layer.hpp" +#include "caffe/layers/deconv_layer.hpp" #include "caffe/layers/lrn_layer.hpp" #include "caffe/layers/pooling_layer.hpp" #include "caffe/layers/relu_layer.hpp" @@ -18,6 +20,7 @@ #ifdef USE_CUDNN #include "caffe/layers/cudnn_conv_layer.hpp" +#include "caffe/layers/cudnn_deconv_layer.hpp" #include "caffe/layers/cudnn_lcn_layer.hpp" #include "caffe/layers/cudnn_lrn_layer.hpp" #include "caffe/layers/cudnn_ndconv_layer.hpp" @@ -37,7 +40,7 @@ namespace caffe { // Get convolution layer according to engine. template shared_ptr > GetConvolutionLayer( - const LayerParameter& param) { + const LayerParameter& param) { ConvolutionParameter conv_param = param.convolution_param(); ConvolutionParameter_Engine engine = conv_param.engine(); #ifdef USE_CUDNN @@ -96,6 +99,45 @@ shared_ptr > GetNdConvolutionLayer( REGISTER_LAYER_CREATOR(NdConvolution, GetNdConvolutionLayer); #endif +// Get deconvolution layer according to engine. +template +shared_ptr > GetDeconvolutionLayer(const LayerParameter& param) { + ConvolutionParameter conv_param = param.convolution_param(); + ConvolutionParameter_Engine engine = conv_param.engine(); +#ifdef USE_CUDNN + bool use_dilation = false; + for (int i = 0; i < conv_param.dilation_size(); ++i) { + if (conv_param.dilation(i) > 1) { + use_dilation = true; + } + } +#endif + if (engine == ConvolutionParameter_Engine_DEFAULT) { + engine = ConvolutionParameter_Engine_CAFFE; +#ifdef USE_CUDNN + if (!use_dilation) { + engine = ConvolutionParameter_Engine_CUDNN; + } +#endif + } + if (engine == ConvolutionParameter_Engine_CAFFE) { + return shared_ptr >(new DeconvolutionLayer(param)); +#ifdef USE_CUDNN + } else if (engine == ConvolutionParameter_Engine_CUDNN) { + if (use_dilation) { + LOG(FATAL) << "CuDNN doesn't support the dilated deconvolution at Layer " + << param.name(); + } + return shared_ptr >(new CuDNNDeconvolutionLayer(param)); +#endif + } else { + LOG(FATAL) << "Layer " << param.name() << " has unknown engine."; + throw; // Avoids missing return warning + } +} + +REGISTER_LAYER_CREATOR(Deconvolution, GetDeconvolutionLayer); + // Get pooling layer according to engine. template shared_ptr > GetPoolingLayer(const LayerParameter& param) { diff --git a/src/caffe/layers/accuracy_layer.cpp b/src/caffe/layers/accuracy_layer.cpp index 4eddbb5c..b6d95b54 100644 --- a/src/caffe/layers/accuracy_layer.cpp +++ b/src/caffe/layers/accuracy_layer.cpp @@ -52,8 +52,6 @@ void AccuracyLayer::Forward_cpu(const vector*>& bottom, const Dtype* bottom_label = bottom[1]->cpu_data(); const int dim = bottom[0]->count() / outer_num_; const int num_labels = bottom[0]->shape(label_axis_); - vector maxval(top_k_+1); - vector max_id(top_k_+1); if (top.size() > 1) { caffe_set(nums_buffer_.count(), Dtype(0), nums_buffer_.mutable_cpu_data()); caffe_set(top[1]->count(), Dtype(0), top[1]->mutable_cpu_data()); @@ -66,32 +64,29 @@ void AccuracyLayer::Forward_cpu(const vector*>& bottom, if (has_ignore_label_ && label_value == ignore_label_) { continue; } - if (top.size() > 1) ++nums_buffer_.mutable_cpu_data()[label_value]; DCHECK_GE(label_value, 0); DCHECK_LT(label_value, num_labels); + if (top.size() > 1) ++nums_buffer_.mutable_cpu_data()[label_value]; + const Dtype prob_of_true_class = bottom_data[i * dim + + label_value * inner_num_ + + j]; + int num_better_predictions = -1; // true_class also counts as "better" // Top-k accuracy - std::vector > bottom_data_vector; - for (int k = 0; k < num_labels; ++k) { - bottom_data_vector.push_back(std::make_pair( - bottom_data[i * dim + k * inner_num_ + j], k)); + for (int k = 0; k < num_labels && num_better_predictions < top_k_; ++k) { + num_better_predictions += + (bottom_data[i * dim + k * inner_num_ + j] >= prob_of_true_class); } - std::partial_sort( - bottom_data_vector.begin(), bottom_data_vector.begin() + top_k_, - bottom_data_vector.end(), std::greater >()); - // check if true label is in top k predictions - for (int k = 0; k < top_k_; k++) { - if (bottom_data_vector[k].second == label_value) { - ++accuracy; - if (top.size() > 1) ++top[1]->mutable_cpu_data()[label_value]; - break; - } + // check if there are less than top_k_ predictions + if (num_better_predictions < top_k_) { + ++accuracy; + if (top.size() > 1) ++top[1]->mutable_cpu_data()[label_value]; } ++count; } } // LOG(INFO) << "Accuracy: " << accuracy; - top[0]->mutable_cpu_data()[0] = accuracy / count; + top[0]->mutable_cpu_data()[0] = (count == 0) ? 0 : (accuracy / count); if (top.size() > 1) { for (int i = 0; i < top[1]->count(); ++i) { top[1]->mutable_cpu_data()[i] = @@ -102,6 +97,10 @@ void AccuracyLayer::Forward_cpu(const vector*>& bottom, // Accuracy layer should not be used as a loss function. } +#ifdef CPU_ONLY +STUB_GPU(AccuracyLayer); +#endif + INSTANTIATE_CLASS(AccuracyLayer); REGISTER_LAYER_CLASS(Accuracy); diff --git a/src/caffe/layers/accuracy_layer.cu b/src/caffe/layers/accuracy_layer.cu new file mode 100644 index 00000000..904aab42 --- /dev/null +++ b/src/caffe/layers/accuracy_layer.cu @@ -0,0 +1,148 @@ +#include + +#include "caffe/layers/accuracy_layer.hpp" +#include "caffe/util/math_functions.hpp" + + +namespace caffe { + +template +__global__ void AccuracyForwardGPU(const int nthreads, + const Dtype* bottom_data, const Dtype* label, Dtype* acc, + const int num, const int dim, const int spatial_dim, + const int num_labels, const int top_k, + const bool has_ignore_label_, const int ignore_label_, + Dtype* counts) { + CUDA_KERNEL_LOOP(index, nthreads) { + const int n = index / spatial_dim; + const int s = index % spatial_dim; + const int label_value = static_cast(label[n * spatial_dim + s]); + const Dtype prob_of_true_class = bottom_data[n * dim + + label_value * spatial_dim + + s]; + int num_better_predictions = -1; // true_class also counts as "better" + if (has_ignore_label_ && label_value == ignore_label_) { + acc[index] = 0; + counts[index] = 0; + } else { + for (int k = 0; k < num_labels & num_better_predictions < top_k; k++) { + num_better_predictions += + (bottom_data[n * dim + k * spatial_dim + s] >= prob_of_true_class); + } + acc[index] = (num_better_predictions < top_k); + counts[index] = 1; + } + } +} + +template +__global__ void AccuracyForwardWithPerClassGPU(const int nthreads, + const Dtype* bottom_data, const Dtype* label, + Dtype* acc, Dtype* counts, + const int num, const int dim, const int spatial_dim, + const int num_labels, const int top_k, + const bool has_ignore_label_, const int ignore_label_) { + CUDA_KERNEL_LOOP(index, nthreads) { + const int n = index / spatial_dim; + const int s = index % spatial_dim; + const int label_value = static_cast(label[n * spatial_dim + s]); + const Dtype prob_of_true_class = bottom_data[n * dim + + label_value * spatial_dim + + s]; + if (has_ignore_label_ && label_value == ignore_label_) { + // nothing to be done. + } else { + int num_better_predictions = -1; // true_class also counts as "better" + for (int k = 0; k < num_labels & num_better_predictions < top_k; k++) { + num_better_predictions += + (bottom_data[n * dim + k * spatial_dim + s] >= prob_of_true_class); + } + acc[label_value*nthreads + index] += (num_better_predictions < top_k); + counts[label_value*nthreads + index] = 1; + } + } +} + +template +void AccuracyLayer::Forward_gpu( + const vector*>& bottom, const vector*>& top) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + const Dtype* bottom_label = bottom[1]->gpu_data(); + const int dim = bottom[0]->count() / outer_num_; + const int num_labels = bottom[0]->shape(label_axis_); + const int nthreads = outer_num_ * inner_num_; + // Since this memory is not used for anything, we use it here to avoid having + // to allocate new GPU memory to accumulate intermediate results. + Dtype* acc_data = bottom[0]->mutable_gpu_diff(); + if (top.size() == 1) { + // simple case - report only global accuracy. + + // Similarly, this memory is never used elsewhere, and thus we can use it + // to avoid having to allocate additional GPU memory. + Dtype* counts = bottom[1]->mutable_gpu_diff(); + // NOLINT_NEXT_LINE(whitespace/operators) + AccuracyForwardGPU<<>>(nthreads, bottom_data, bottom_label, + acc_data, outer_num_, dim, inner_num_, num_labels, top_k_, + has_ignore_label_, ignore_label_, counts); + Dtype acc; + caffe_gpu_asum(nthreads, acc_data, &acc); + Dtype valid_count; + caffe_gpu_asum(nthreads, counts, &valid_count); + if (valid_count > 0) { + top[0]->mutable_cpu_data()[0] = acc / valid_count; + } else { + top[0]->mutable_cpu_data()[0] = 0; + } + } else { + // need to report per-class accuracy as well + + // allocate space for more detailed "counts" + nums_buffer_.ReshapeLike(*bottom[0]); + Dtype* counts = nums_buffer_.mutable_gpu_data(); + + caffe_gpu_set(bottom[0]->count(), Dtype(0), acc_data); + caffe_gpu_set(nums_buffer_.count(), Dtype(0), counts); + + // NOLINT_NEXT_LINE(whitespace/operators) + AccuracyForwardWithPerClassGPU<<>>(nthreads, bottom_data, bottom_label, + acc_data, counts, outer_num_, dim, inner_num_, num_labels, top_k_, + has_ignore_label_, ignore_label_); + + // get the overall accuracy + Dtype acc; + caffe_gpu_asum(bottom[0]->count(), acc_data, &acc); + Dtype valid_count; + caffe_gpu_asum(nums_buffer_.count(), counts, &valid_count); + if (valid_count > 0) { + top[0]->mutable_cpu_data()[0] = acc / valid_count; + } else { + top[0]->mutable_cpu_data()[0] = 0; + } + + // get per-class accuracy + Dtype* per_class_acc = top[1]->mutable_cpu_data(); + for (int l = 0; l < num_labels; l++) { + caffe_gpu_asum(nthreads, acc_data + l*nthreads, per_class_acc+l); + caffe_gpu_asum(nthreads, counts + l*nthreads, &valid_count); + if (valid_count > 0) { + per_class_acc[l] /= valid_count; + } else { + per_class_acc[l] = 0; + } + } + } + // Clear scratch memory to prevent interfering with backward (see #6202). + caffe_gpu_set(bottom[0]->count(), Dtype(0), bottom[0]->mutable_gpu_diff()); +} + + +template +void AccuracyLayer::Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom) { + if (propagate_down[1]) { NOT_IMPLEMENTED; } +} + +INSTANTIATE_LAYER_GPU_FUNCS(AccuracyLayer); +} // namespace caffe diff --git a/src/caffe/layers/base_conv_layer.cpp b/src/caffe/layers/base_conv_layer.cpp index a204b70d..6ae49f3e 100644 --- a/src/caffe/layers/base_conv_layer.cpp +++ b/src/caffe/layers/base_conv_layer.cpp @@ -210,7 +210,9 @@ void BaseConvolutionLayer::Reshape(const vector*>& bottom, // TODO: generalize to handle inputs of different shapes. for (int bottom_id = 1; bottom_id < bottom.size(); ++bottom_id) { CHECK(bottom[0]->shape() == bottom[bottom_id]->shape()) - << "All inputs must have the same shape."; + << "shape mismatch - bottom[0]: " << bottom[0]->shape_string() + << " vs. bottom[" << bottom_id << "]: " + << bottom[bottom_id]->shape_string(); } // Shape the tops. bottom_shape_ = &bottom[0]->shape(); diff --git a/src/caffe/layers/clip_layer.cpp b/src/caffe/layers/clip_layer.cpp new file mode 100644 index 00000000..9d9a5967 --- /dev/null +++ b/src/caffe/layers/clip_layer.cpp @@ -0,0 +1,51 @@ +#include +#include + +#include "caffe/layers/clip_layer.hpp" + +namespace caffe { + +template +void ClipLayer::Forward_cpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + Dtype* top_data = top[0]->mutable_cpu_data(); + const int count = bottom[0]->count(); + + Dtype min = this->layer_param_.clip_param().min(); + Dtype max = this->layer_param_.clip_param().max(); + + for (int i = 0; i < count; ++i) { + top_data[i] = std::max(min, std::min(bottom_data[i], max)); + } +} + +template +void ClipLayer::Backward_cpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + if (propagate_down[0]) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + const Dtype* top_diff = top[0]->cpu_diff(); + Dtype* bottom_diff = bottom[0]->mutable_cpu_diff(); + const int count = bottom[0]->count(); + + Dtype min = this->layer_param_.clip_param().min(); + Dtype max = this->layer_param_.clip_param().max(); + + for (int i = 0; i < count; ++i) { + bottom_diff[i] = top_diff[i] * ( + bottom_data[i] >= min && bottom_data[i] <= max); + } + } +} + + +#ifdef CPU_ONLY +STUB_GPU(ClipLayer); +#endif + +INSTANTIATE_CLASS(ClipLayer); +REGISTER_LAYER_CLASS(Clip); + +} // namespace caffe diff --git a/src/caffe/layers/clip_layer.cu b/src/caffe/layers/clip_layer.cu new file mode 100644 index 00000000..56f3be32 --- /dev/null +++ b/src/caffe/layers/clip_layer.cu @@ -0,0 +1,67 @@ +#include + +#include "caffe/layers/clip_layer.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +__global__ void ClipForward(const int n, const float* in, float* out, + float p_min, float p_max) { + CUDA_KERNEL_LOOP(index, n) { + out[index] = fmaxf(p_min, fminf(in[index], p_max)); + } +} + +__global__ void ClipForward(const int n, const double* in, double* out, + double p_min, double p_max) { + CUDA_KERNEL_LOOP(index, n) { + out[index] = fmax(p_min, fmin(in[index], p_max)); + } +} + +template +void ClipLayer::Forward_gpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + Dtype* top_data = top[0]->mutable_gpu_data(); + const int count = bottom[0]->count(); + Dtype p_min = this->layer_param_.clip_param().min(); + Dtype p_max = this->layer_param_.clip_param().max(); + // NOLINT_NEXT_LINE(whitespace/operators) + ClipForward<<>>( + count, bottom_data, top_data, p_min, p_max); + CUDA_POST_KERNEL_CHECK; +} + +template +__global__ void ClipBackward(const int n, const Dtype* in_diff, + const Dtype* in_data, Dtype* out_diff, Dtype p_min, Dtype p_max) { + CUDA_KERNEL_LOOP(index, n) { + out_diff[index] = in_diff[index] * ( + in_data[index] >= p_min && in_data[index] <= p_max); + } +} + +template +void ClipLayer::Backward_gpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + if (propagate_down[0]) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + const Dtype* top_diff = top[0]->gpu_diff(); + Dtype* bottom_diff = bottom[0]->mutable_gpu_diff(); + const int count = bottom[0]->count(); + Dtype p_min = this->layer_param_.clip_param().min(); + Dtype p_max = this->layer_param_.clip_param().max(); + // NOLINT_NEXT_LINE(whitespace/operators) + ClipBackward<<>>( + count, top_diff, bottom_data, bottom_diff, p_min, p_max); + CUDA_POST_KERNEL_CHECK; + } +} + + +INSTANTIATE_LAYER_GPU_FUNCS(ClipLayer); + + +} // namespace caffe diff --git a/src/caffe/layers/cudnn_deconv_layer.cpp b/src/caffe/layers/cudnn_deconv_layer.cpp new file mode 100644 index 00000000..260da5c1 --- /dev/null +++ b/src/caffe/layers/cudnn_deconv_layer.cpp @@ -0,0 +1,327 @@ +#ifdef USE_CUDNN +#include +#include + +#include "caffe/layers/cudnn_deconv_layer.hpp" + +namespace caffe { + +// Set to three for the benefit of the backward pass, which +// can use separate streams for calculating the gradient w.r.t. +// bias, filter weights, and bottom data for each group independently +#define CUDNN_STREAMS_PER_GROUP 3 + +/** + * TODO(dox) explain cuDNN interface + */ +template +void CuDNNDeconvolutionLayer::LayerSetUp( + const vector*>& bottom, const vector*>& top) { + DeconvolutionLayer::LayerSetUp(bottom, top); + // Initialize CUDA streams and cuDNN. + stream_ = new cudaStream_t[this->group_ * CUDNN_STREAMS_PER_GROUP]; + handle_ = new cudnnHandle_t[this->group_ * CUDNN_STREAMS_PER_GROUP]; + + // Initialize algorithm arrays + fwd_algo_ = new cudnnConvolutionFwdAlgo_t[bottom.size()]; + bwd_filter_algo_= new cudnnConvolutionBwdFilterAlgo_t[bottom.size()]; + bwd_data_algo_ = new cudnnConvolutionBwdDataAlgo_t[bottom.size()]; + + // initialize size arrays + workspace_fwd_sizes_ = new size_t[bottom.size()]; + workspace_bwd_filter_sizes_ = new size_t[bottom.size()]; + workspace_bwd_data_sizes_ = new size_t[bottom.size()]; + + // workspace data + workspaceSizeInBytes = 0; + workspaceData = NULL; + workspace = new void*[this->group_ * CUDNN_STREAMS_PER_GROUP]; + + for (size_t i = 0; i < bottom.size(); ++i) { + // initialize all to default algorithms + fwd_algo_[i] = (cudnnConvolutionFwdAlgo_t)0; + bwd_filter_algo_[i] = (cudnnConvolutionBwdFilterAlgo_t)0; + bwd_data_algo_[i] = (cudnnConvolutionBwdDataAlgo_t)0; + // default algorithms don't require workspace + workspace_fwd_sizes_[i] = 0; + workspace_bwd_data_sizes_[i] = 0; + workspace_bwd_filter_sizes_[i] = 0; + } + + for (int g = 0; g < this->group_ * CUDNN_STREAMS_PER_GROUP; g++) { + CUDA_CHECK(cudaStreamCreate(&stream_[g])); + CUDNN_CHECK(cudnnCreate(&handle_[g])); + CUDNN_CHECK(cudnnSetStream(handle_[g], stream_[g])); + workspace[g] = NULL; + } + + // Set the indexing parameters. + bias_offset_ = (this->num_output_ / this->group_); + + // Create filter descriptor. + const int* kernel_shape_data = this->kernel_shape_.cpu_data(); + const int kernel_h = kernel_shape_data[0]; + const int kernel_w = kernel_shape_data[1]; + cudnn::createFilterDesc(&filter_desc_, + this->channels_ / this->group_, + this->num_output_ / this->group_, + kernel_h, + kernel_w); + + // Create tensor descriptor(s) for data and corresponding convolution(s). + for (int i = 0; i < bottom.size(); i++) { + cudnnTensorDescriptor_t bottom_desc; + cudnn::createTensor4dDesc(&bottom_desc); + bottom_descs_.push_back(bottom_desc); + cudnnTensorDescriptor_t top_desc; + cudnn::createTensor4dDesc(&top_desc); + top_descs_.push_back(top_desc); + cudnnConvolutionDescriptor_t conv_desc; + cudnn::createConvolutionDesc(&conv_desc); + conv_descs_.push_back(conv_desc); + } + + // Tensor descriptor for bias. + if (this->bias_term_) { + cudnn::createTensor4dDesc(&bias_desc_); + } + + handles_setup_ = true; +} + +template +void CuDNNDeconvolutionLayer::Reshape( + const vector*>& bottom, const vector*>& top) { + DeconvolutionLayer::Reshape(bottom, top); + CHECK_EQ(2, this->num_spatial_axes_) + << "CuDNNDeconvolutionLayer input must have 2 spatial axes " + << "(e.g., height and width). " + << "Use 'engine: CAFFE' for general ND convolution."; + bottom_offset_ = this->bottom_dim_ / this->group_; + top_offset_ = this->top_dim_ / this->group_; + const int height = bottom[0]->shape(this->channel_axis_ + 1); + const int width = bottom[0]->shape(this->channel_axis_ + 2); + const int height_out = top[0]->shape(this->channel_axis_ + 1); + const int width_out = top[0]->shape(this->channel_axis_ + 2); + const int* pad_data = this->pad_.cpu_data(); + const int pad_h = pad_data[0]; + const int pad_w = pad_data[1]; + const int* stride_data = this->stride_.cpu_data(); + const int stride_h = stride_data[0]; + const int stride_w = stride_data[1]; + + // Specify workspace limit for kernels directly until we have a + // planning strategy and a rewrite of Caffe's GPU memory mangagement + size_t workspace_limit_bytes = 8*1024*1024; + + for (int i = 0; i < bottom.size(); i++) { + cudnn::setTensor4dDesc(&bottom_descs_[i], + this->num_, + this->channels_ / this->group_, + height, + width, + this->channels_ * height * width, + height * width, + width, + 1); + cudnn::setTensor4dDesc(&top_descs_[i], + this->num_, + this->num_output_ / this->group_, + height_out, + width_out, + this->num_output_ * height_out * width_out, + height_out * width_out, + width_out, + 1); + cudnn::setConvolutionDesc(&conv_descs_[i], + top_descs_[i], + filter_desc_, + pad_h, + pad_w, + stride_h, + stride_w); + + // choose forward and backward algorithms + workspace(s) + CUDNN_CHECK(cudnnGetConvolutionForwardAlgorithm( + handle_[0], + top_descs_[i], + filter_desc_, + conv_descs_[i], + bottom_descs_[i], + CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT, + workspace_limit_bytes, + &fwd_algo_[i])); + + // We have found that CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM is + // buggy. Thus, if this algo was chosen, choose winograd instead. If + // winograd is not supported or workspace is larger than threshold, choose + // implicit_gemm instead. + if (fwd_algo_[i] == CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM) { + size_t winograd_workspace_size; + cudnnStatus_t status = cudnnGetConvolutionForwardWorkspaceSize( + handle_[0], + top_descs_[i], + filter_desc_, + conv_descs_[i], + bottom_descs_[i], + CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD, + &winograd_workspace_size); + if (status != CUDNN_STATUS_SUCCESS || + winograd_workspace_size >= workspace_limit_bytes) { + fwd_algo_[i] = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM; + } else { + fwd_algo_[i] = CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD; + } + } + + CUDNN_CHECK(cudnnGetConvolutionForwardWorkspaceSize( + handle_[0], + top_descs_[i], + filter_desc_, + conv_descs_[i], + bottom_descs_[i], + fwd_algo_[i], + &(workspace_fwd_sizes_[i]))); + + // choose backward algorithm for filter + CUDNN_CHECK(cudnnGetConvolutionBackwardFilterAlgorithm( + handle_[0], + top_descs_[i], + bottom_descs_[i], + conv_descs_[i], + filter_desc_, + CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT, + workspace_limit_bytes, + &bwd_filter_algo_[i])); + + // get workspace for backwards filter algorithm + CUDNN_CHECK(cudnnGetConvolutionBackwardFilterWorkspaceSize( + handle_[0], + top_descs_[i], + bottom_descs_[i], + conv_descs_[i], + filter_desc_, + bwd_filter_algo_[i], + &workspace_bwd_filter_sizes_[i])); + + // choose backward algo for data + CUDNN_CHECK(cudnnGetConvolutionBackwardDataAlgorithm( + handle_[0], + filter_desc_, + bottom_descs_[i], + conv_descs_[i], + top_descs_[i], + CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT, + workspace_limit_bytes, + &bwd_data_algo_[i])); + + // get workspace size + CUDNN_CHECK(cudnnGetConvolutionBackwardDataWorkspaceSize( + handle_[0], + filter_desc_, + bottom_descs_[i], + conv_descs_[i], + top_descs_[i], + bwd_data_algo_[i], + &workspace_bwd_data_sizes_[i])); + } + + // reduce over all workspace sizes to get a maximum to allocate / reallocate + size_t total_workspace_fwd = 0; + size_t total_workspace_bwd_data = 0; + size_t total_workspace_bwd_filter = 0; + + for (size_t i = 0; i < bottom.size(); i++) { + total_workspace_fwd = std::max(total_workspace_fwd, + workspace_fwd_sizes_[i]); + total_workspace_bwd_data = std::max(total_workspace_bwd_data, + workspace_bwd_data_sizes_[i]); + total_workspace_bwd_filter = std::max(total_workspace_bwd_filter, + workspace_bwd_filter_sizes_[i]); + } + // get max over all operations + size_t max_workspace = std::max(total_workspace_fwd, + total_workspace_bwd_data); + max_workspace = std::max(max_workspace, total_workspace_bwd_filter); + // ensure all groups have enough workspace + size_t total_max_workspace = max_workspace * + (this->group_ * CUDNN_STREAMS_PER_GROUP); + + // this is the total amount of storage needed over all groups + streams + if (total_max_workspace > workspaceSizeInBytes) { + DLOG(INFO) << "Reallocating workspace storage: " << total_max_workspace; + workspaceSizeInBytes = total_max_workspace; + + // free the existing workspace and allocate a new (larger) one + cudaFree(this->workspaceData); + + cudaError_t err = cudaMalloc(&(this->workspaceData), workspaceSizeInBytes); + if (err != cudaSuccess) { + // force zero memory path + for (int i = 0; i < bottom.size(); i++) { + workspace_fwd_sizes_[i] = 0; + workspace_bwd_filter_sizes_[i] = 0; + workspace_bwd_data_sizes_[i] = 0; + fwd_algo_[i] = CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING; + bwd_filter_algo_[i] = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0; + bwd_data_algo_[i] = CUDNN_CONVOLUTION_BWD_DATA_ALGO_0; + } + + // NULL out all workspace pointers + for (int g = 0; g < (this->group_ * CUDNN_STREAMS_PER_GROUP); g++) { + workspace[g] = NULL; + } + // NULL out underlying data + workspaceData = NULL; + workspaceSizeInBytes = 0; + } + + // if we succeed in the allocation, set pointer aliases for workspaces + for (int g = 0; g < (this->group_ * CUDNN_STREAMS_PER_GROUP); g++) { + workspace[g] = reinterpret_cast(workspaceData) + g*max_workspace; + } + } + + // Tensor descriptor for bias. + if (this->bias_term_) { + cudnn::setTensor4dDesc( + &bias_desc_, 1, this->num_output_ / this->group_, 1, 1); + } +} + +template +CuDNNDeconvolutionLayer::~CuDNNDeconvolutionLayer() { + // Check that handles have been setup before destroying. + if (!handles_setup_) { return; } + + for (int i = 0; i < bottom_descs_.size(); i++) { + cudnnDestroyTensorDescriptor(bottom_descs_[i]); + cudnnDestroyTensorDescriptor(top_descs_[i]); + cudnnDestroyConvolutionDescriptor(conv_descs_[i]); + } + if (this->bias_term_) { + cudnnDestroyTensorDescriptor(bias_desc_); + } + cudnnDestroyFilterDescriptor(filter_desc_); + + for (int g = 0; g < this->group_ * CUDNN_STREAMS_PER_GROUP; g++) { + cudaStreamDestroy(stream_[g]); + cudnnDestroy(handle_[g]); + } + + cudaFree(workspaceData); + delete [] workspace; + delete [] stream_; + delete [] handle_; + delete [] fwd_algo_; + delete [] bwd_filter_algo_; + delete [] bwd_data_algo_; + delete [] workspace_fwd_sizes_; + delete [] workspace_bwd_data_sizes_; + delete [] workspace_bwd_filter_sizes_; +} + +INSTANTIATE_CLASS(CuDNNDeconvolutionLayer); + +} // namespace caffe +#endif diff --git a/src/caffe/layers/cudnn_deconv_layer.cu b/src/caffe/layers/cudnn_deconv_layer.cu new file mode 100644 index 00000000..eb1df329 --- /dev/null +++ b/src/caffe/layers/cudnn_deconv_layer.cu @@ -0,0 +1,138 @@ +#ifdef USE_CUDNN +#include + +#include "caffe/layers/cudnn_deconv_layer.hpp" + +namespace caffe { + +__global__ void sync_deconv_groups() {} + +template +void CuDNNDeconvolutionLayer::Forward_gpu( + const vector*>& bottom, const vector*>& top) { + const Dtype* weight = this->blobs_[0]->gpu_data(); + for (int i = 0; i < bottom.size(); ++i) { + const Dtype* bottom_data = bottom[i]->gpu_data(); + Dtype* top_data = top[i]->mutable_gpu_data(); + + // Forward through cuDNN in parallel over groups. + for (int g = 0; g < this->group_; g++) { + // Filters. + CUDNN_CHECK(cudnnConvolutionBackwardData( + handle_[g], + cudnn::dataType::one, + filter_desc_, + weight + this->weight_offset_ * g, + bottom_descs_[i], + bottom_data + bottom_offset_ * g, + conv_descs_[i], + bwd_data_algo_[i], + workspace[g], + workspace_bwd_data_sizes_[i], + cudnn::dataType::zero, + top_descs_[i], + top_data + top_offset_ * g)); + + // Bias. + if (this->bias_term_) { + const Dtype* bias_data = this->blobs_[1]->gpu_data(); + CUDNN_CHECK(cudnnAddTensor(handle_[g], + cudnn::dataType::one, + bias_desc_, + bias_data + bias_offset_ * g, + cudnn::dataType::one, + top_descs_[i], + top_data + top_offset_ * g)); + } + } + + // Synchronize the work across groups, each of which went into its own + // stream, by launching an empty kernel into the default (null) stream. + // NOLINT_NEXT_LINE(whitespace/operators) + sync_deconv_groups<<<1, 1>>>(); + } +} + +template +void CuDNNDeconvolutionLayer::Backward_gpu( + const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + const Dtype* weight = NULL; + Dtype* weight_diff = NULL; + if (this->param_propagate_down_[0]) { + weight = this->blobs_[0]->gpu_data(); + weight_diff = this->blobs_[0]->mutable_gpu_diff(); + } + Dtype* bias_diff = NULL; + if (this->bias_term_ && this->param_propagate_down_[1]) { + bias_diff = this->blobs_[1]->mutable_gpu_diff(); + } + for (int i = 0; i < top.size(); ++i) { + const Dtype* top_diff = top[i]->gpu_diff(); + // Backward through cuDNN in parallel over groups and gradients. + for (int g = 0; g < this->group_; g++) { + // Gradient w.r.t. bias. + if (this->bias_term_ && this->param_propagate_down_[1]) { + CUDNN_CHECK(cudnnConvolutionBackwardBias(handle_[0 * this->group_ + g], + cudnn::dataType::one, + top_descs_[i], + top_diff + top_offset_ * g, + cudnn::dataType::one, + bias_desc_, + bias_diff + bias_offset_ * g)); + } + + // Gradient w.r.t. weights. + if (this->param_propagate_down_[0]) { + const Dtype* bottom_data = bottom[i]->gpu_data(); + CUDNN_CHECK(cudnnConvolutionBackwardFilter( + handle_[1 * this->group_ + g], + cudnn::dataType::one, + top_descs_[i], + top_diff + top_offset_ * g, + bottom_descs_[i], + bottom_data + bottom_offset_ * g, + conv_descs_[i], + bwd_filter_algo_[i], + workspace[1 * this->group_ + g], + workspace_bwd_filter_sizes_[i], + cudnn::dataType::one, + filter_desc_, + weight_diff + this->weight_offset_ * g)); + } + + // Gradient w.r.t. bottom data. + if (propagate_down[i]) { + if (weight == NULL) { + weight = this->blobs_[0]->gpu_data(); + } + Dtype* bottom_diff = bottom[i]->mutable_gpu_diff(); + CUDNN_CHECK( + cudnnConvolutionForward(handle_[2 * this->group_ + g], + cudnn::dataType::one, + top_descs_[i], + top_diff + top_offset_ * g, + filter_desc_, + weight + this->weight_offset_ * g, + conv_descs_[i], + fwd_algo_[i], + workspace[2 * this->group_ + g], + workspace_fwd_sizes_[i], + cudnn::dataType::zero, + bottom_descs_[i], + bottom_diff + bottom_offset_ * g)); + } + } + + // Synchronize the work across groups, each of which went into its own + // stream, by launching an empty kernel into the default (null) stream. + // NOLINT_NEXT_LINE(whitespace/operators) + sync_deconv_groups<<<1, 1>>>(); + } +} + +INSTANTIATE_LAYER_GPU_FUNCS(CuDNNDeconvolutionLayer); + +} // namespace caffe +#endif diff --git a/src/caffe/layers/deconv_layer.cpp b/src/caffe/layers/deconv_layer.cpp index 20a460fb..b86472b3 100644 --- a/src/caffe/layers/deconv_layer.cpp +++ b/src/caffe/layers/deconv_layer.cpp @@ -79,6 +79,5 @@ STUB_GPU(DeconvolutionLayer); #endif INSTANTIATE_CLASS(DeconvolutionLayer); -REGISTER_LAYER_CLASS(Deconvolution); } // namespace caffe diff --git a/src/caffe/layers/embed_layer.cu b/src/caffe/layers/embed_layer.cu index 6324a3a8..3cf39fd9 100644 --- a/src/caffe/layers/embed_layer.cu +++ b/src/caffe/layers/embed_layer.cu @@ -15,6 +15,11 @@ __global__ void EmbedForward(const int nthreads, const Dtype* bottom_data, const int n = top_index / N; const int d = top_index % N; const int index = static_cast(bottom_data[n]); + #ifdef DEBUG + assert(index >= 0); + assert(index < K); + assert(static_cast(index) == bottom_data[n]); + #endif const int weight_index = index * N + d; top_data[top_index] = weight[weight_index]; } diff --git a/src/caffe/layers/hdf5_data_layer.cpp b/src/caffe/layers/hdf5_data_layer.cpp index 00716a92..7668854c 100644 --- a/src/caffe/layers/hdf5_data_layer.cpp +++ b/src/caffe/layers/hdf5_data_layer.cpp @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 /* TODO: - load file in a separate thread ("prefetch") @@ -184,3 +185,4 @@ INSTANTIATE_CLASS(HDF5DataLayer); REGISTER_LAYER_CLASS(HDF5Data); } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/layers/hdf5_data_layer.cu b/src/caffe/layers/hdf5_data_layer.cu index 33eebd41..70cd9f32 100644 --- a/src/caffe/layers/hdf5_data_layer.cu +++ b/src/caffe/layers/hdf5_data_layer.cu @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 /* TODO: - only load parts of the file, in accordance with a prototxt param "max_mem" @@ -34,3 +35,4 @@ void HDF5DataLayer::Forward_gpu(const vector*>& bottom, INSTANTIATE_LAYER_GPU_FUNCS(HDF5DataLayer); } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/layers/hdf5_output_layer.cpp b/src/caffe/layers/hdf5_output_layer.cpp index f8f1edcd..28c453a2 100644 --- a/src/caffe/layers/hdf5_output_layer.cpp +++ b/src/caffe/layers/hdf5_output_layer.cpp @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 #include #include "hdf5.h" @@ -72,3 +73,4 @@ INSTANTIATE_CLASS(HDF5OutputLayer); REGISTER_LAYER_CLASS(HDF5Output); } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/layers/hdf5_output_layer.cu b/src/caffe/layers/hdf5_output_layer.cu index c1685cd3..891aea03 100644 --- a/src/caffe/layers/hdf5_output_layer.cu +++ b/src/caffe/layers/hdf5_output_layer.cu @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 #include #include "hdf5.h" @@ -37,3 +38,4 @@ void HDF5OutputLayer::Backward_gpu(const vector*>& top, INSTANTIATE_LAYER_GPU_FUNCS(HDF5OutputLayer); } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/layers/inner_product_layer.cpp b/src/caffe/layers/inner_product_layer.cpp index e65349f0..57fdbe1f 100644 --- a/src/caffe/layers/inner_product_layer.cpp +++ b/src/caffe/layers/inner_product_layer.cpp @@ -42,7 +42,7 @@ void InnerProductLayer::LayerSetUp(const vector*>& bottom, shared_ptr > weight_filler(GetFiller( this->layer_param_.inner_product_param().weight_filler())); weight_filler->Fill(this->blobs_[0].get()); - // If necessary, intiialize and fill the bias term + // If necessary, initialize and fill the bias term if (bias_term_) { vector bias_shape(1, N_); this->blobs_[1].reset(new Blob(bias_shape)); diff --git a/src/caffe/layers/pooling_layer.cpp b/src/caffe/layers/pooling_layer.cpp index 4740262b..a4306afa 100644 --- a/src/caffe/layers/pooling_layer.cpp +++ b/src/caffe/layers/pooling_layer.cpp @@ -35,6 +35,7 @@ void PoolingLayer::LayerSetUp(const vector*>& bottom, || (!pool_param.has_stride_h() && !pool_param.has_stride_w())) << "Stride is stride OR stride_h and stride_w are required."; global_pooling_ = pool_param.global_pooling(); + round_mode_ = pool_param.round_mode(); if (global_pooling_) { kernel_h_ = bottom[0]->height(); kernel_w_ = bottom[0]->width(); @@ -89,10 +90,22 @@ void PoolingLayer::Reshape(const vector*>& bottom, kernel_h_ = bottom[0]->height(); kernel_w_ = bottom[0]->width(); } - pooled_height_ = static_cast(ceil(static_cast( - height_ + 2 * pad_h_ - kernel_h_) / stride_h_)) + 1; - pooled_width_ = static_cast(ceil(static_cast( - width_ + 2 * pad_w_ - kernel_w_) / stride_w_)) + 1; + switch (round_mode_) { + case PoolingParameter_RoundMode_CEIL: + pooled_height_ = static_cast(ceil(static_cast( + height_ + 2 * pad_h_ - kernel_h_) / stride_h_)) + 1; + pooled_width_ = static_cast(ceil(static_cast( + width_ + 2 * pad_w_ - kernel_w_) / stride_w_)) + 1; + break; + case PoolingParameter_RoundMode_FLOOR: + pooled_height_ = static_cast(floor(static_cast( + height_ + 2 * pad_h_ - kernel_h_) / stride_h_)) + 1; + pooled_width_ = static_cast(floor(static_cast( + width_ + 2 * pad_w_ - kernel_w_) / stride_w_)) + 1; + break; + default: + LOG(FATAL) << "Unknown rounding mode."; + } if (pad_h_ || pad_w_) { // If we have padding, ensure that the last pooling starts strictly // inside the image (instead of at the padding); otherwise clip the last. @@ -134,7 +147,7 @@ void PoolingLayer::Forward_cpu(const vector*>& bottom, const int top_count = top[0]->count(); // We'll output the mask to top[1] if it's of size >1. const bool use_top_mask = top.size() > 1; - int* mask = NULL; // suppress warnings about uninitalized variables + int* mask = NULL; // suppress warnings about uninitialized variables Dtype* top_mask = NULL; // Different pooling methods. We explicitly do the switch outside the for // loop to save time, although this results in more code. diff --git a/src/caffe/layers/recurrent_layer.cpp b/src/caffe/layers/recurrent_layer.cpp index e0c82773..9cd3206f 100644 --- a/src/caffe/layers/recurrent_layer.cpp +++ b/src/caffe/layers/recurrent_layer.cpp @@ -214,8 +214,9 @@ void RecurrentLayer::Reshape(const vector*>& bottom, const int bottom_offset = 2 + static_input_; for (int i = bottom_offset, j = 0; i < bottom.size(); ++i, ++j) { CHECK(recur_input_blobs_[j]->shape() == bottom[i]->shape()) - << "bottom[" << i << "] shape must match hidden state input shape: " - << recur_input_blobs_[j]->shape_string(); + << "shape mismatch - recur_input_blobs_[" << j << "]: " + << recur_input_blobs_[j]->shape_string() + << " vs. bottom[" << i << "]: " << bottom[i]->shape_string(); recur_input_blobs_[j]->ShareData(*bottom[i]); } } diff --git a/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu b/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu index b9877e6a..7497e4aa 100644 --- a/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu +++ b/src/caffe/layers/sigmoid_cross_entropy_loss_layer.cu @@ -48,9 +48,8 @@ void SigmoidCrossEntropyLossLayer::Forward_gpu( // Stable version of loss computation from input data const Dtype* input_data = bottom[0]->gpu_data(); const Dtype* target = bottom[1]->gpu_data(); - // Since this memory is not used for anything until it is overwritten - // on the backward pass, we use it here to avoid having to allocate new GPU - // memory to accumulate intermediate results in the kernel. + // Since this memory is not used for anything, we use it here to avoid having + // to allocate new GPU memory to accumulate intermediate results. Dtype* loss_data = bottom[0]->mutable_gpu_diff(); Dtype* count_data = bottom[1]->mutable_gpu_diff(); Dtype valid_count; @@ -69,6 +68,10 @@ void SigmoidCrossEntropyLossLayer::Forward_gpu( caffe_gpu_asum(count, loss_data, &loss); normalizer_ = get_normalizer(normalization_, valid_count); top[0]->mutable_cpu_data()[0] = loss / normalizer_; + + // Clear scratch memory to prevent interfering with backward (see #6202). + caffe_gpu_set(bottom[0]->count(), Dtype(0), bottom[0]->mutable_gpu_diff()); + caffe_gpu_set(bottom[1]->count(), Dtype(0), bottom[1]->mutable_gpu_diff()); } template diff --git a/src/caffe/layers/slice_layer.cpp b/src/caffe/layers/slice_layer.cpp index 759beafe..64de0964 100644 --- a/src/caffe/layers/slice_layer.cpp +++ b/src/caffe/layers/slice_layer.cpp @@ -41,7 +41,9 @@ void SliceLayer::Reshape(const vector*>& bottom, int count = 0; if (slice_point_.size() != 0) { CHECK_EQ(slice_point_.size(), top.size() - 1); - CHECK_LE(top.size(), bottom_slice_axis); + CHECK_LE(top.size(), bottom_slice_axis) + << "slice axis: " << slice_axis_ + << ", bottom[0] shape: " << bottom[0]->shape_string(); int prev = 0; vector slices; for (int i = 0; i < slice_point_.size(); ++i) { diff --git a/src/caffe/layers/softmax_loss_layer.cu b/src/caffe/layers/softmax_loss_layer.cu index 660e1b39..b3c8ffa6 100644 --- a/src/caffe/layers/softmax_loss_layer.cu +++ b/src/caffe/layers/softmax_loss_layer.cu @@ -36,9 +36,8 @@ void SoftmaxWithLossLayer::Forward_gpu( const Dtype* label = bottom[1]->gpu_data(); const int dim = prob_.count() / outer_num_; const int nthreads = outer_num_ * inner_num_; - // Since this memory is not used for anything until it is overwritten - // on the backward pass, we use it here to avoid having to allocate new GPU - // memory to accumulate intermediate results in the kernel. + // Since this memory is not used for anything, we use it here to avoid having + // to allocate new GPU memory to accumulate intermediate results. Dtype* loss_data = bottom[0]->mutable_gpu_diff(); // Similarly, this memory is never used elsewhere, and thus we can use it // to avoid having to allocate additional GPU memory. @@ -61,6 +60,9 @@ void SoftmaxWithLossLayer::Forward_gpu( if (top.size() == 2) { top[1]->ShareData(prob_); } + + // Clear scratch memory to prevent interfering with backward (see #6202). + caffe_gpu_set(bottom[0]->count(), Dtype(0), bottom[0]->mutable_gpu_diff()); } template diff --git a/src/caffe/layers/swish_layer.cpp b/src/caffe/layers/swish_layer.cpp new file mode 100644 index 00000000..28935679 --- /dev/null +++ b/src/caffe/layers/swish_layer.cpp @@ -0,0 +1,68 @@ +#include +#include + +#include "caffe/layers/swish_layer.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +template +void SwishLayer::LayerSetUp(const vector*>& bottom, + const vector*>& top) { + NeuronLayer::LayerSetUp(bottom, top); + sigmoid_bottom_vec_.clear(); + sigmoid_bottom_vec_.push_back(sigmoid_input_.get()); + sigmoid_top_vec_.clear(); + sigmoid_top_vec_.push_back(sigmoid_output_.get()); + sigmoid_layer_->SetUp(sigmoid_bottom_vec_, sigmoid_top_vec_); +} + +template +void SwishLayer::Reshape(const vector*>& bottom, + const vector*>& top) { + NeuronLayer::Reshape(bottom, top); + sigmoid_input_->ReshapeLike(*bottom[0]); + sigmoid_layer_->Reshape(sigmoid_bottom_vec_, sigmoid_top_vec_); +} + +template +void SwishLayer::Forward_cpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + Dtype* sigmoid_input_data = sigmoid_input_->mutable_cpu_data(); + Dtype* top_data = top[0]->mutable_cpu_data(); + const int count = bottom[0]->count(); + Dtype beta = this->layer_param_.swish_param().beta(); + caffe_copy(count, bottom_data, sigmoid_input_data); + caffe_scal(count, beta, sigmoid_input_data); + sigmoid_layer_->Forward(sigmoid_bottom_vec_, sigmoid_top_vec_); + caffe_mul(count, bottom_data, sigmoid_output_->cpu_data(), top_data); +} + +template +void SwishLayer::Backward_cpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + if (propagate_down[0]) { + const Dtype* top_data = top[0]->cpu_data(); + const Dtype* top_diff = top[0]->cpu_diff(); + const Dtype* sigmoid_output_data = sigmoid_output_->cpu_data(); + Dtype* bottom_diff = bottom[0]->mutable_cpu_diff(); + const int count = bottom[0]->count(); + Dtype beta = this->layer_param_.swish_param().beta(); + for (int i = 0; i < count; ++i) { + const Dtype swish_x = top_data[i]; + bottom_diff[i] = top_diff[i] * (beta * swish_x + sigmoid_output_data[i] + * (1. - beta * swish_x)); + } + } +} + +#ifdef CPU_ONLY +STUB_GPU(SwishLayer); +#endif + +INSTANTIATE_CLASS(SwishLayer); +REGISTER_LAYER_CLASS(Swish); + +} // namespace caffe diff --git a/src/caffe/layers/swish_layer.cu b/src/caffe/layers/swish_layer.cu new file mode 100644 index 00000000..c4fef53b --- /dev/null +++ b/src/caffe/layers/swish_layer.cu @@ -0,0 +1,54 @@ +#include +#include + +#include "caffe/layers/swish_layer.hpp" +#include "caffe/util/math_functions.hpp" + +namespace caffe { + +template +void SwishLayer::Forward_gpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->gpu_data(); + Dtype* sigmoid_input_data = sigmoid_input_->mutable_gpu_data(); + Dtype* top_data = top[0]->mutable_gpu_data(); + const int count = bottom[0]->count(); + Dtype beta = this->layer_param_.swish_param().beta(); + caffe_copy(count, bottom_data, sigmoid_input_data); + caffe_gpu_scal(count, beta, sigmoid_input_data); + sigmoid_layer_->Forward(sigmoid_bottom_vec_, sigmoid_top_vec_); + caffe_gpu_mul(count, bottom_data, sigmoid_output_->gpu_data(), top_data); +} + +template +__global__ void SwishBackward(const int n, const Dtype* in_diff, + const Dtype* out_data, const Dtype* sigmoid_output_data, Dtype* out_diff, + const Dtype beta) { + CUDA_KERNEL_LOOP(index, n) { + const Dtype swish_x = out_data[index]; + out_diff[index] = in_diff[index] * (beta * swish_x + + sigmoid_output_data[index] * (1 - beta * swish_x)); + } +} + +template +void SwishLayer::Backward_gpu(const vector*>& top, + const vector& propagate_down, + const vector*>& bottom) { + if (propagate_down[0]) { + const Dtype* top_data = top[0]->gpu_data(); + const Dtype* top_diff = top[0]->gpu_diff(); + const Dtype* sigmoid_output_data = sigmoid_output_->gpu_data(); + Dtype* bottom_diff = bottom[0]->mutable_gpu_diff(); + const int count = bottom[0]->count(); + Dtype beta = this->layer_param_.swish_param().beta(); + // NOLINT_NEXT_LINE(whitespace/operators) + SwishBackward<<>>( + count, top_diff, top_data, sigmoid_output_data, bottom_diff, beta); + CUDA_POST_KERNEL_CHECK; + } +} + +INSTANTIATE_LAYER_GPU_FUNCS(SwishLayer); + +} // namespace caffe diff --git a/src/caffe/net.cpp b/src/caffe/net.cpp index 353c2f95..5e844b03 100644 --- a/src/caffe/net.cpp +++ b/src/caffe/net.cpp @@ -5,7 +5,9 @@ #include #include +#ifdef USE_HDF5 #include "hdf5.h" +#endif // USE_HDF5 #include "caffe/common.hpp" #include "caffe/layer.hpp" @@ -164,7 +166,7 @@ void Net::Init(const NetParameter& in_param) { // loss. We can skip backward computation for blobs that don't contribute // to the loss. // Also checks if all bottom blobs don't need backward computation (possible - // because the skip_propagate_down param) and so we can skip bacward + // because the skip_propagate_down param) and so we can skip backward // computation for the entire layer set blobs_under_loss; set blobs_skip_backp; @@ -768,7 +770,7 @@ void Net::CopyTrainedLayersFrom(const NetParameter& param) { } template -void Net::CopyTrainedLayersFrom(const string trained_filename) { +void Net::CopyTrainedLayersFrom(const string& trained_filename) { if (H5Fis_hdf5(trained_filename.c_str())) { CopyTrainedLayersFromHDF5(trained_filename); } else { @@ -778,14 +780,15 @@ void Net::CopyTrainedLayersFrom(const string trained_filename) { template void Net::CopyTrainedLayersFromBinaryProto( - const string trained_filename) { + const string& trained_filename) { NetParameter param; ReadNetParamsFromBinaryFileOrDie(trained_filename, ¶m); CopyTrainedLayersFrom(param); } template -void Net::CopyTrainedLayersFromHDF5(const string trained_filename) { +void Net::CopyTrainedLayersFromHDF5(const string& trained_filename) { +#ifdef USE_HDF5 hid_t file_hid = H5Fopen(trained_filename.c_str(), H5F_ACC_RDONLY, H5P_DEFAULT); CHECK_GE(file_hid, 0) << "Couldn't open " << trained_filename; @@ -832,6 +835,10 @@ void Net::CopyTrainedLayersFromHDF5(const string trained_filename) { } H5Gclose(data_hid); H5Fclose(file_hid); +#else + LOG(FATAL) << "CopyTrainedLayersFromHDF5 requires hdf5;" + << " compile with USE_HDF5."; +#endif // USE_HDF5 } template @@ -848,6 +855,8 @@ void Net::ToProto(NetParameter* param, bool write_diff) const { template void Net::ToHDF5(const string& filename, bool write_diff) const { +// This code is taken from https://github.com/sh1r0/caffe-android-lib +#ifdef USE_HDF5 hid_t file_hid = H5Fcreate(filename.c_str(), H5F_ACC_TRUNC, H5P_DEFAULT, H5P_DEFAULT); CHECK_GE(file_hid, 0) @@ -901,6 +910,10 @@ void Net::ToHDF5(const string& filename, bool write_diff) const { H5Gclose(diff_hid); } H5Fclose(file_hid); +// This code is taken from https://github.com/sh1r0/caffe-android-lib +#else + LOG(FATAL) << "ToHDF5 requires hdf5; compile with USE_HDF5."; +#endif // USE_HDF5 } template diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 9db6c964..8ad70e6a 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -101,7 +101,7 @@ message NetParameter { // NOTE // Update the next available ID when you add a new SolverParameter field. // -// SolverParameter next available ID: 42 (last added: layer_wise_reduce) +// SolverParameter next available ID: 43 (last added: weights) message SolverParameter { ////////////////////////////////////////////////////////////////////////////// // Specifying the train and test networks @@ -189,7 +189,11 @@ message SolverParameter { optional float clip_gradients = 35 [default = -1]; optional int32 snapshot = 14 [default = 0]; // The snapshot interval - optional string snapshot_prefix = 15; // The prefix for the snapshot. + // The prefix for the snapshot. + // If not set then is replaced by prototxt file path without extension. + // If is set to directory then is augmented by prototxt file name + // without extention. + optional string snapshot_prefix = 15; // whether to snapshot diff in the results or not. Snapshotting diff will help // debugging but the final protocol buffer size will be much larger. optional bool snapshot_diff = 16 [default = false]; @@ -244,6 +248,16 @@ message SolverParameter { // Overlap compute and communication for data parallel training optional bool layer_wise_reduce = 41 [default = true]; + + // Path to caffemodel file(s) with pretrained weights to initialize finetuning. + // Tha same as command line --weights parameter for caffe train command. + // If command line --weights parameter is specified, it has higher priority + // and overwrites this one(s). + // If --snapshot command line parameter is specified, this one(s) are ignored. + // If several model files are expected, they can be listed in a one + // weights parameter separated by ',' (like in a command string) or + // in repeated weights parameters separately. + repeated string weights = 42; } // A message that stores the solver snapshots @@ -311,7 +325,7 @@ message ParamSpec { // NOTE // Update the next available ID when you add a new LayerParameter field. // -// LayerParameter next available layer-specific ID: 147 (last added: recurrent_param) +// LayerParameter next available layer-specific ID: 149 (last added: clip_param) // video-caffe custom layers start with 7777 // Next available video-caffe layer ID: 7778 (last added: video_data_layer) message LayerParameter { @@ -369,6 +383,7 @@ message LayerParameter { optional ArgMaxParameter argmax_param = 103; optional BatchNormParameter batch_norm_param = 139; optional BiasParameter bias_param = 141; + optional ClipParameter clip_param = 148; optional ConcatParameter concat_param = 104; optional ContrastiveLossParameter contrastive_loss_param = 105; optional ConvolutionParameter convolution_param = 106; @@ -406,6 +421,7 @@ message LayerParameter { optional SoftmaxParameter softmax_param = 125; optional SPPParameter spp_param = 132; optional SliceParameter slice_param = 126; + optional SwishParameter swish_param = 147; optional TanHParameter tanh_param = 127; optional ThresholdParameter threshold_param = 128; optional TileParameter tile_param = 138; @@ -496,6 +512,12 @@ message ArgMaxParameter { optional int32 axis = 3; } +// Message that stores parameters used by ClipLayer +message ClipParameter { + required float min = 1; + required float max = 2; +} + message ConcatParameter { // The axis along which to concatenate -- may be negative to index from the // end (e.g., -1 for the last axis). Other axes must have the @@ -941,6 +963,12 @@ message PoolingParameter { // If global_pooling then it will pool over the size of the bottom by doing // kernel_h = bottom->height and kernel_w = bottom->width optional bool global_pooling = 12 [default = false]; + // How to calculate the output size - using ceil (default) or floor rounding. + enum RoundMode { + CEIL = 0; + FLOOR = 1; + } + optional RoundMode round_mode = 13 [default = CEIL]; } message PowerParameter { @@ -1162,6 +1190,15 @@ message SoftmaxParameter { optional int32 axis = 2 [default = 1]; } +// Message that stores parameters used by SwishLayer +message SwishParameter { + // Beta parameter for the Swish activation function + // Described in: + // Prajit Ramachandran, Barret Zoph, Quoc V. Le. (2017). Searching for + // Activation Functions. https://arxiv.org/abs/1710.05941v2 + optional float beta = 1 [default = 1]; +} + message TanHParameter { enum Engine { DEFAULT = 0; diff --git a/src/caffe/solver.cpp b/src/caffe/solver.cpp index 04426937..842312e0 100644 --- a/src/caffe/solver.cpp +++ b/src/caffe/solver.cpp @@ -3,6 +3,7 @@ #include #include +#include "boost/algorithm/string.hpp" #include "caffe/solver.hpp" #include "caffe/util/format.hpp" #include "caffe/util/hdf5.hpp" @@ -59,11 +60,25 @@ void Solver::Init(const SolverParameter& param) { current_step_ = 0; } +// Load weights from the caffemodel(s) specified in "weights" solver parameter +// into the train and test nets. +template +void LoadNetWeights(shared_ptr > net, + const std::string& model_list) { + std::vector model_names; + boost::split(model_names, model_list, boost::is_any_of(",")); + for (int i = 0; i < model_names.size(); ++i) { + boost::trim(model_names[i]); + LOG(INFO) << "Finetuning from " << model_names[i]; + net->CopyTrainedLayersFrom(model_names[i]); + } +} + template void Solver::InitTrainNet() { const int num_train_nets = param_.has_net() + param_.has_net_param() + param_.has_train_net() + param_.has_train_net_param(); - const string& field_names = "net, net_param, train_net, train_net_param"; + const string field_names = "net, net_param, train_net, train_net_param"; CHECK_GE(num_train_nets, 1) << "SolverParameter must specify a train net " << "using one of these fields: " << field_names; CHECK_LE(num_train_nets, 1) << "SolverParameter must not contain more than " @@ -98,6 +113,9 @@ void Solver::InitTrainNet() { net_state.MergeFrom(param_.train_state()); net_param.mutable_state()->CopyFrom(net_state); net_.reset(new Net(net_param)); + for (int w_idx = 0; w_idx < param_.weights_size(); ++w_idx) { + LoadNetWeights(net_, param_.weights(w_idx)); + } } template @@ -173,6 +191,9 @@ void Solver::InitTestNets() { << "Creating test net (#" << i << ") specified by " << sources[i]; test_nets_[i].reset(new Net(net_params[i])); test_nets_[i]->set_debug_info(param_.debug_info()); + for (int w_idx = 0; w_idx < param_.weights_size(); ++w_idx) { + LoadNetWeights(test_nets_[i], param_.weights(w_idx)); + } } } @@ -245,10 +266,6 @@ void Solver::Step(int iters) { } ApplyUpdate(); - // Increment the internal iter_ counter -- its value should always indicate - // the number of times the weights have been updated. - ++iter_; - SolverAction::Enum request = GetRequestedAction(); // Save a snapshot if needed. @@ -430,13 +447,13 @@ void Solver::CheckSnapshotWritePermissions() { } else { LOG(FATAL) << "Cannot write to snapshot prefix '" << param_.snapshot_prefix() << "'. Make sure " - << "that the directory exists and is writeable."; + << "that the directory exists and is writable."; } } } template -string Solver::SnapshotFilename(const string extension) { +string Solver::SnapshotFilename(const string& extension) { return param_.snapshot_prefix() + "_iter_" + caffe::format_int(iter_) + extension; } diff --git a/src/caffe/solvers/sgd_solver.cpp b/src/caffe/solvers/sgd_solver.cpp index ad6abe54..081c47eb 100644 --- a/src/caffe/solvers/sgd_solver.cpp +++ b/src/caffe/solvers/sgd_solver.cpp @@ -30,12 +30,16 @@ Dtype SGDSolver::GetLearningRate() { if (lr_policy == "fixed") { rate = this->param_.base_lr(); } else if (lr_policy == "step") { + CHECK_GT(this->param_.stepsize(), 0); this->current_step_ = this->iter_ / this->param_.stepsize(); + CHECK_GE(this->param_.gamma(), 0); rate = this->param_.base_lr() * pow(this->param_.gamma(), this->current_step_); } else if (lr_policy == "exp") { + CHECK_GE(this->param_.gamma(), 0); rate = this->param_.base_lr() * pow(this->param_.gamma(), this->iter_); } else if (lr_policy == "inv") { + CHECK_GE(this->param_.gamma(), 0); rate = this->param_.base_lr() * pow(Dtype(1) + this->param_.gamma() * this->iter_, - this->param_.power()); @@ -46,6 +50,7 @@ Dtype SGDSolver::GetLearningRate() { LOG(INFO) << "MultiStep Status: Iteration " << this->iter_ << ", step = " << this->current_step_; } + CHECK_GE(this->param_.gamma(), 0); rate = this->param_.base_lr() * pow(this->param_.gamma(), this->current_step_); } else if (lr_policy == "poly") { @@ -53,6 +58,8 @@ Dtype SGDSolver::GetLearningRate() { (Dtype(this->iter_) / Dtype(this->param_.max_iter())), this->param_.power()); } else if (lr_policy == "sigmoid") { + CHECK_GE(this->param_.gamma(), 0); + CHECK_GT(this->param_.stepsize(), 0); rate = this->param_.base_lr() * (Dtype(1.) / (Dtype(1.) + exp(-this->param_.gamma() * (Dtype(this->iter_) - Dtype(this->param_.stepsize()))))); @@ -113,6 +120,10 @@ void SGDSolver::ApplyUpdate() { ComputeUpdateValue(param_id, rate); } this->net_->Update(); + + // Increment the internal iter_ counter -- its value should always indicate + // the number of times the weights have been updated. + ++this->iter_; } template @@ -278,6 +289,8 @@ void SGDSolver::SnapshotSolverStateToBinaryProto( template void SGDSolver::SnapshotSolverStateToHDF5( const string& model_filename) { +// This code is taken from https://github.com/sh1r0/caffe-android-lib +#ifdef USE_HDF5 string snapshot_filename = Solver::SnapshotFilename(".solverstate.h5"); LOG(INFO) << "Snapshotting solver state to HDF5 file " << snapshot_filename; @@ -299,6 +312,11 @@ void SGDSolver::SnapshotSolverStateToHDF5( } H5Gclose(history_hid); H5Fclose(file_hid); +// This code is taken from https://github.com/sh1r0/caffe-android-lib +#else + LOG(FATAL) << "SnapshotSolverStateToHDF5 requires hdf5;" + << " compile with USE_HDF5."; +#endif // USE_HDF5 } template @@ -323,6 +341,7 @@ void SGDSolver::RestoreSolverStateFromBinaryProto( template void SGDSolver::RestoreSolverStateFromHDF5(const string& state_file) { +#ifdef USE_HDF5 hid_t file_hid = H5Fopen(state_file.c_str(), H5F_ACC_RDONLY, H5P_DEFAULT); CHECK_GE(file_hid, 0) << "Couldn't open solver state file " << state_file; this->iter_ = hdf5_load_int(file_hid, "iter"); @@ -344,6 +363,10 @@ void SGDSolver::RestoreSolverStateFromHDF5(const string& state_file) { } H5Gclose(history_hid); H5Fclose(file_hid); +#else + LOG(FATAL) << "RestoreSolverStateFromHDF5 requires hdf5;" + << " compile with USE_HDF5."; +#endif // USE_HDF5 } INSTANTIATE_CLASS(SGDSolver); diff --git a/src/caffe/test/test_accuracy_layer.cpp b/src/caffe/test/test_accuracy_layer.cpp index 6fe808bd..e5cc9d5e 100644 --- a/src/caffe/test/test_accuracy_layer.cpp +++ b/src/caffe/test/test_accuracy_layer.cpp @@ -13,8 +13,10 @@ namespace caffe { -template -class AccuracyLayerTest : public CPUDeviceTest { +template +class AccuracyLayerTest : public MultiDeviceTest { + typedef typename TypeParam::Dtype Dtype; + protected: AccuracyLayerTest() : blob_bottom_data_(new Blob()), @@ -69,11 +71,12 @@ class AccuracyLayerTest : public CPUDeviceTest { int top_k_; }; -TYPED_TEST_CASE(AccuracyLayerTest, TestDtypes); +TYPED_TEST_CASE(AccuracyLayerTest, TestDtypesAndDevices); TYPED_TEST(AccuracyLayerTest, TestSetup) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); EXPECT_EQ(this->blob_top_->num(), 1); EXPECT_EQ(this->blob_top_->channels(), 1); @@ -82,11 +85,12 @@ TYPED_TEST(AccuracyLayerTest, TestSetup) { } TYPED_TEST(AccuracyLayerTest, TestSetupTopK) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; AccuracyParameter* accuracy_param = layer_param.mutable_accuracy_param(); accuracy_param->set_top_k(5); - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); EXPECT_EQ(this->blob_top_->num(), 1); EXPECT_EQ(this->blob_top_->channels(), 1); @@ -95,8 +99,9 @@ TYPED_TEST(AccuracyLayerTest, TestSetupTopK) { } TYPED_TEST(AccuracyLayerTest, TestSetupOutputPerClass) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_per_class_vec_); EXPECT_EQ(this->blob_top_->num(), 1); EXPECT_EQ(this->blob_top_->channels(), 1); @@ -108,33 +113,39 @@ TYPED_TEST(AccuracyLayerTest, TestSetupOutputPerClass) { EXPECT_EQ(this->blob_top_per_class_->width(), 1); } -TYPED_TEST(AccuracyLayerTest, TestForwardCPU) { +TYPED_TEST(AccuracyLayerTest, TestForward) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); - - TypeParam max_value; - int max_id; - int num_correct_labels = 0; - for (int i = 0; i < 100; ++i) { - max_value = -FLT_MAX; - max_id = 0; - for (int j = 0; j < 10; ++j) { - if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { - max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); - max_id = j; + + // repeat the forward + for (int iter = 0; iter < 3; iter++) { + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + + Dtype max_value; + int max_id; + int num_correct_labels = 0; + for (int i = 0; i < 100; ++i) { + max_value = -FLT_MAX; + max_id = 0; + for (int j = 0; j < 10; ++j) { + if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { + max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); + max_id = j; + } + } + if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + ++num_correct_labels; } } - if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - ++num_correct_labels; - } + EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), + num_correct_labels / Dtype(100.0), 1e-4); } - EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), - num_correct_labels / 100.0, 1e-4); } TYPED_TEST(AccuracyLayerTest, TestForwardWithSpatialAxes) { + typedef typename TypeParam::Dtype Dtype; this->blob_bottom_data_->Reshape(2, 10, 4, 5); vector label_shape(3); label_shape[0] = 2; label_shape[1] = 4; label_shape[2] = 5; @@ -142,195 +153,218 @@ TYPED_TEST(AccuracyLayerTest, TestForwardWithSpatialAxes) { this->FillBottoms(); LayerParameter layer_param; layer_param.mutable_accuracy_param()->set_axis(1); - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); - - TypeParam max_value; - const int num_labels = this->blob_bottom_label_->count(); - int max_id; - int num_correct_labels = 0; - vector label_offset(3); - for (int n = 0; n < this->blob_bottom_data_->num(); ++n) { - for (int h = 0; h < this->blob_bottom_data_->height(); ++h) { - for (int w = 0; w < this->blob_bottom_data_->width(); ++w) { - max_value = -FLT_MAX; - max_id = 0; - for (int c = 0; c < this->blob_bottom_data_->channels(); ++c) { - const TypeParam pred_value = - this->blob_bottom_data_->data_at(n, c, h, w); - if (pred_value > max_value) { - max_value = pred_value; - max_id = c; + + // repeat the forward + for (int iter = 0; iter < 3; iter++) { + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + + Dtype max_value; + const int num_labels = this->blob_bottom_label_->count(); + int max_id; + int num_correct_labels = 0; + vector label_offset(3); + for (int n = 0; n < this->blob_bottom_data_->num(); ++n) { + for (int h = 0; h < this->blob_bottom_data_->height(); ++h) { + for (int w = 0; w < this->blob_bottom_data_->width(); ++w) { + max_value = -FLT_MAX; + max_id = 0; + for (int c = 0; c < this->blob_bottom_data_->channels(); ++c) { + const Dtype pred_value = + this->blob_bottom_data_->data_at(n, c, h, w); + if (pred_value > max_value) { + max_value = pred_value; + max_id = c; + } + } + label_offset[0] = n; label_offset[1] = h; label_offset[2] = w; + const int correct_label = + static_cast(this->blob_bottom_label_->data_at(label_offset)); + if (max_id == correct_label) { + ++num_correct_labels; } - } - label_offset[0] = n; label_offset[1] = h; label_offset[2] = w; - const int correct_label = - static_cast(this->blob_bottom_label_->data_at(label_offset)); - if (max_id == correct_label) { - ++num_correct_labels; } } } + EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), + num_correct_labels / Dtype(num_labels), 1e-4); } - EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), - num_correct_labels / TypeParam(num_labels), 1e-4); } TYPED_TEST(AccuracyLayerTest, TestForwardIgnoreLabel) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - const TypeParam kIgnoreLabelValue = -1; + const Dtype kIgnoreLabelValue = -1; layer_param.mutable_accuracy_param()->set_ignore_label(kIgnoreLabelValue); - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); // Manually set some labels to the ignore label value (-1). this->blob_bottom_label_->mutable_cpu_data()[2] = kIgnoreLabelValue; this->blob_bottom_label_->mutable_cpu_data()[5] = kIgnoreLabelValue; this->blob_bottom_label_->mutable_cpu_data()[32] = kIgnoreLabelValue; layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); - - TypeParam max_value; - int max_id; - int num_correct_labels = 0; - int count = 0; - for (int i = 0; i < 100; ++i) { - if (kIgnoreLabelValue == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - continue; - } - ++count; - max_value = -FLT_MAX; - max_id = 0; - for (int j = 0; j < 10; ++j) { - if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { - max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); - max_id = j; + + // repeat the forward + for (int iter = 0; iter < 3; iter++) { + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + + Dtype max_value; + int max_id; + int num_correct_labels = 0; + int count = 0; + for (int i = 0; i < 100; ++i) { + if (kIgnoreLabelValue == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + continue; + } + ++count; + max_value = -FLT_MAX; + max_id = 0; + for (int j = 0; j < 10; ++j) { + if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { + max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); + max_id = j; + } + } + if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + ++num_correct_labels; } } - if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - ++num_correct_labels; - } + EXPECT_EQ(count, 97); // We set 3 out of 100 labels to kIgnoreLabelValue. + EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), + num_correct_labels / Dtype(count), 1e-4); } - EXPECT_EQ(count, 97); // We set 3 out of 100 labels to kIgnoreLabelValue. - EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), - num_correct_labels / TypeParam(count), 1e-4); } -TYPED_TEST(AccuracyLayerTest, TestForwardCPUTopK) { +TYPED_TEST(AccuracyLayerTest, TestForwardTopK) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; AccuracyParameter* accuracy_param = layer_param.mutable_accuracy_param(); accuracy_param->set_top_k(this->top_k_); - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); - - TypeParam current_value; - int current_rank; - int num_correct_labels = 0; - for (int i = 0; i < 100; ++i) { - for (int j = 0; j < 10; ++j) { - current_value = this->blob_bottom_data_->data_at(i, j, 0, 0); - current_rank = 0; - for (int k = 0; k < 10; ++k) { - if (this->blob_bottom_data_->data_at(i, k, 0, 0) > current_value) { - ++current_rank; + + // repeat the forward + for (int iter = 0; iter < 3; iter++) { + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + + Dtype current_value; + int current_rank; + int num_correct_labels = 0; + for (int i = 0; i < 100; ++i) { + for (int j = 0; j < 10; ++j) { + current_value = this->blob_bottom_data_->data_at(i, j, 0, 0); + current_rank = 0; + for (int k = 0; k < 10; ++k) { + if (this->blob_bottom_data_->data_at(i, k, 0, 0) > current_value) { + ++current_rank; + } + } + if (current_rank < this->top_k_ && + j == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + ++num_correct_labels; } - } - if (current_rank < this->top_k_ && - j == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - ++num_correct_labels; } } - } - EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), - num_correct_labels / 100.0, 1e-4); + EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), + num_correct_labels / Dtype(100.0), 1e-4); + } } -TYPED_TEST(AccuracyLayerTest, TestForwardCPUPerClass) { +TYPED_TEST(AccuracyLayerTest, TestForwardPerClass) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_per_class_vec_); - layer.Forward(this->blob_bottom_vec_, this->blob_top_per_class_vec_); - - TypeParam max_value; - int max_id; - int num_correct_labels = 0; - const int num_class = this->blob_top_per_class_->num(); - vector correct_per_class(num_class, 0); - vector num_per_class(num_class, 0); - for (int i = 0; i < 100; ++i) { - max_value = -FLT_MAX; - max_id = 0; - for (int j = 0; j < 10; ++j) { - if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { - max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); - max_id = j; + // repeat the forward + for (int iter = 0; iter < 3; iter++) { + layer.Forward(this->blob_bottom_vec_, this->blob_top_per_class_vec_); + + Dtype max_value; + int max_id; + int num_correct_labels = 0; + const int num_class = this->blob_top_per_class_->num(); + vector correct_per_class(num_class, 0); + vector num_per_class(num_class, 0); + for (int i = 0; i < 100; ++i) { + max_value = -FLT_MAX; + max_id = 0; + for (int j = 0; j < 10; ++j) { + if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { + max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); + max_id = j; + } + } + ++num_per_class[this->blob_bottom_label_->data_at(i, 0, 0, 0)]; + if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + ++num_correct_labels; + ++correct_per_class[max_id]; } } - ++num_per_class[this->blob_bottom_label_->data_at(i, 0, 0, 0)]; - if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - ++num_correct_labels; - ++correct_per_class[max_id]; + EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), + num_correct_labels / 100.0, 1e-4); + for (int i = 0; i < num_class; ++i) { + Dtype accuracy_per_class = (num_per_class[i] > 0 ? + static_cast(correct_per_class[i]) / num_per_class[i] : 0); + EXPECT_NEAR(this->blob_top_per_class_->data_at(i, 0, 0, 0), + accuracy_per_class, 1e-4); } } - EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), - num_correct_labels / 100.0, 1e-4); - for (int i = 0; i < num_class; ++i) { - TypeParam accuracy_per_class = (num_per_class[i] > 0 ? - static_cast(correct_per_class[i]) / num_per_class[i] : 0); - EXPECT_NEAR(this->blob_top_per_class_->data_at(i, 0, 0, 0), - accuracy_per_class, 1e-4); - } } -TYPED_TEST(AccuracyLayerTest, TestForwardCPUPerClassWithIgnoreLabel) { +TYPED_TEST(AccuracyLayerTest, TestForwardPerClassWithIgnoreLabel) { + typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - const TypeParam kIgnoreLabelValue = -1; + const Dtype kIgnoreLabelValue = -1; layer_param.mutable_accuracy_param()->set_ignore_label(kIgnoreLabelValue); - AccuracyLayer layer(layer_param); + AccuracyLayer layer(layer_param); // Manually set some labels to the ignore label value (-1). this->blob_bottom_label_->mutable_cpu_data()[2] = kIgnoreLabelValue; this->blob_bottom_label_->mutable_cpu_data()[5] = kIgnoreLabelValue; this->blob_bottom_label_->mutable_cpu_data()[32] = kIgnoreLabelValue; layer.SetUp(this->blob_bottom_vec_, this->blob_top_per_class_vec_); - layer.Forward(this->blob_bottom_vec_, this->blob_top_per_class_vec_); - - TypeParam max_value; - int max_id; - int num_correct_labels = 0; - const int num_class = this->blob_top_per_class_->num(); - vector correct_per_class(num_class, 0); - vector num_per_class(num_class, 0); - int count = 0; - for (int i = 0; i < 100; ++i) { - if (kIgnoreLabelValue == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - continue; - } - ++count; - max_value = -FLT_MAX; - max_id = 0; - for (int j = 0; j < 10; ++j) { - if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { - max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); - max_id = j; + + // repeat the forward + for (int iter = 0; iter < 3; iter++) { + layer.Forward(this->blob_bottom_vec_, this->blob_top_per_class_vec_); + + Dtype max_value; + int max_id; + int num_correct_labels = 0; + const int num_class = this->blob_top_per_class_->num(); + vector correct_per_class(num_class, 0); + vector num_per_class(num_class, 0); + int count = 0; + for (int i = 0; i < 100; ++i) { + if (kIgnoreLabelValue == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + continue; + } + ++count; + max_value = -FLT_MAX; + max_id = 0; + for (int j = 0; j < 10; ++j) { + if (this->blob_bottom_data_->data_at(i, j, 0, 0) > max_value) { + max_value = this->blob_bottom_data_->data_at(i, j, 0, 0); + max_id = j; + } + } + ++num_per_class[this->blob_bottom_label_->data_at(i, 0, 0, 0)]; + if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { + ++num_correct_labels; + ++correct_per_class[max_id]; } } - ++num_per_class[this->blob_bottom_label_->data_at(i, 0, 0, 0)]; - if (max_id == this->blob_bottom_label_->data_at(i, 0, 0, 0)) { - ++num_correct_labels; - ++correct_per_class[max_id]; + EXPECT_EQ(count, 97); + EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), + num_correct_labels / Dtype(count), 1e-4); + for (int i = 0; i < 10; ++i) { + Dtype accuracy_per_class = (num_per_class[i] > 0 ? + static_cast(correct_per_class[i]) / num_per_class[i] : 0); + EXPECT_NEAR(this->blob_top_per_class_->data_at(i, 0, 0, 0), + accuracy_per_class, 1e-4); } } - EXPECT_EQ(count, 97); - EXPECT_NEAR(this->blob_top_->data_at(0, 0, 0, 0), - num_correct_labels / TypeParam(count), 1e-4); - for (int i = 0; i < 10; ++i) { - TypeParam accuracy_per_class = (num_per_class[i] > 0 ? - static_cast(correct_per_class[i]) / num_per_class[i] : 0); - EXPECT_NEAR(this->blob_top_per_class_->data_at(i, 0, 0, 0), - accuracy_per_class, 1e-4); - } } } // namespace caffe diff --git a/src/caffe/test/test_deconvolution_layer.cpp b/src/caffe/test/test_deconvolution_layer.cpp index 18f17dc1..eb37d08d 100644 --- a/src/caffe/test/test_deconvolution_layer.cpp +++ b/src/caffe/test/test_deconvolution_layer.cpp @@ -5,6 +5,7 @@ #include "caffe/blob.hpp" #include "caffe/common.hpp" #include "caffe/filler.hpp" +#include "caffe/layers/cudnn_deconv_layer.hpp" #include "caffe/layers/deconv_layer.hpp" #include "caffe/test/test_caffe_main.hpp" @@ -302,4 +303,268 @@ TYPED_TEST(DeconvolutionLayerTest, TestGradient3D) { this->blob_top_vec_); } +#ifdef USE_CUDNN + +// Since ConvolutionLayerTest checks the shared conv/deconv code in detail, +// we'll just do a simple forward test and a gradient check. +template +class CuDNNDeconvolutionLayerTest : public MultiDeviceTest { + typedef typename TypeParam::Dtype Dtype; + + protected: + CuDNNDeconvolutionLayerTest() + : blob_bottom_(new Blob(2, 3, 6, 4)), + blob_bottom_2_(new Blob(2, 3, 6, 4)), + blob_top_(new Blob()), + blob_top_2_(new Blob()) {} + virtual void SetUp() { + // fill the values + FillerParameter filler_param; + filler_param.set_value(1.); + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + filler.Fill(this->blob_bottom_2_); + blob_bottom_vec_.push_back(blob_bottom_); + blob_top_vec_.push_back(blob_top_); + } + + virtual ~CuDNNDeconvolutionLayerTest() { + delete blob_bottom_; + delete blob_bottom_2_; + delete blob_top_; + delete blob_top_2_; + } + + Blob* const blob_bottom_; + Blob* const blob_bottom_2_; + Blob* const blob_top_; + Blob* const blob_top_2_; + vector*> blob_bottom_vec_; + vector*> blob_top_vec_; +}; + +TYPED_TEST_CASE(CuDNNDeconvolutionLayerTest, TestDtypesAndDevices); + +TYPED_TEST(CuDNNDeconvolutionLayerTest, TestSetup) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + ConvolutionParameter* convolution_param = + layer_param.mutable_convolution_param(); + convolution_param->add_kernel_size(3); + convolution_param->add_stride(2); + convolution_param->set_num_output(4); + this->blob_bottom_vec_.push_back(this->blob_bottom_2_); + this->blob_top_vec_.push_back(this->blob_top_2_); + shared_ptr > layer( + new CuDNNDeconvolutionLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->num(), 2); + EXPECT_EQ(this->blob_top_->channels(), 4); + EXPECT_EQ(this->blob_top_->height(), 13); + EXPECT_EQ(this->blob_top_->width(), 9); + EXPECT_EQ(this->blob_top_2_->num(), 2); + EXPECT_EQ(this->blob_top_2_->channels(), 4); + EXPECT_EQ(this->blob_top_2_->height(), 13); + EXPECT_EQ(this->blob_top_2_->width(), 9); + // setting group should not change the shape + convolution_param->set_num_output(3); + convolution_param->set_group(3); + layer.reset(new CuDNNDeconvolutionLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->num(), 2); + EXPECT_EQ(this->blob_top_->channels(), 3); + EXPECT_EQ(this->blob_top_->height(), 13); + EXPECT_EQ(this->blob_top_->width(), 9); + EXPECT_EQ(this->blob_top_2_->num(), 2); + EXPECT_EQ(this->blob_top_2_->channels(), 3); + EXPECT_EQ(this->blob_top_2_->height(), 13); + EXPECT_EQ(this->blob_top_2_->width(), 9); +} + +TYPED_TEST(CuDNNDeconvolutionLayerTest, TestSimpleCuDNNDeconvolution) { + typedef typename TypeParam::Dtype Dtype; + this->blob_bottom_vec_.push_back(this->blob_bottom_2_); + this->blob_top_vec_.push_back(this->blob_top_2_); + LayerParameter layer_param; + ConvolutionParameter* convolution_param = + layer_param.mutable_convolution_param(); + convolution_param->add_kernel_size(3); + convolution_param->add_stride(2); + convolution_param->set_num_output(4); + convolution_param->mutable_weight_filler()->set_type("constant"); + convolution_param->mutable_weight_filler()->set_value(1); + convolution_param->mutable_bias_filler()->set_type("constant"); + convolution_param->mutable_bias_filler()->set_value(0.1); + shared_ptr > layer( + new CuDNNDeconvolutionLayer(layer_param)); + layer->SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + // constant-fill the bottom blobs + FillerParameter filler_param; + filler_param.set_value(1.); + ConstantFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + filler.Fill(this->blob_bottom_2_); + layer->Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // simply check that accumulation works with overlapping filters + const Dtype* top_data = this->blob_top_->cpu_data(); + for (int n = 0; n < this->blob_top_->num(); ++n) { + for (int c = 0; c < this->blob_top_->channels(); ++c) { + for (int h = 0; h < this->blob_top_->height(); ++h) { + for (int w = 0; w < this->blob_top_->width(); ++w) { + Dtype expected = 3.1; + bool h_overlap = h % 2 == 0 && h > 0 + && h < this->blob_top_->height() - 1; + bool w_overlap = w % 2 == 0 && w > 0 + && w < this->blob_top_->width() - 1; + if (h_overlap && w_overlap) { + expected += 9; + } else if (h_overlap || w_overlap) { + expected += 3; + } + EXPECT_NEAR(top_data[this->blob_top_->offset(n, c, h, w)], + expected, 1e-4); + } + } + } + } +} + +TYPED_TEST(CuDNNDeconvolutionLayerTest, TestGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + ConvolutionParameter* convolution_param = + layer_param.mutable_convolution_param(); + this->blob_bottom_vec_.push_back(this->blob_bottom_2_); + this->blob_top_vec_.push_back(this->blob_top_2_); + convolution_param->add_kernel_size(2); + convolution_param->add_stride(1); + convolution_param->set_num_output(1); + convolution_param->mutable_weight_filler()->set_type("gaussian"); + convolution_param->mutable_bias_filler()->set_type("gaussian"); + CuDNNDeconvolutionLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(CuDNNDeconvolutionLayerTest, TestNDAgainst2D) { + typedef typename TypeParam::Dtype Dtype; + const int kernel_h = 11; + const int kernel_w = 13; + vector bottom_shape(4); + bottom_shape[0] = 15; + bottom_shape[1] = 12; + bottom_shape[2] = kernel_h * 2; + bottom_shape[3] = kernel_w * 2; + FillerParameter filler_param; + GaussianFiller filler(filler_param); + for (int i = 0; i < this->blob_bottom_vec_.size(); ++i) { + this->blob_bottom_vec_[i]->Reshape(bottom_shape); + filler.Fill(this->blob_bottom_vec_[i]); + } + LayerParameter layer_param; + ConvolutionParameter* convolution_param = + layer_param.mutable_convolution_param(); + convolution_param->set_num_output(18); + convolution_param->set_bias_term(false); + convolution_param->set_group(6); + convolution_param->set_kernel_h(kernel_h); + convolution_param->set_kernel_w(kernel_w); + convolution_param->mutable_weight_filler()->set_type("gaussian"); + Blob weights; + Blob top_diff; + // Shape and fill weights and top_diff. + bool copy_diff; + bool reshape; + { + CuDNNDeconvolutionLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + top_diff.ReshapeLike(*this->blob_top_); + filler.Fill(&top_diff); + ASSERT_EQ(1, layer.blobs().size()); + copy_diff = false; reshape = true; + weights.CopyFrom(*layer.blobs()[0], copy_diff, reshape); + } + vector propagate_down(1, true); + Blob result_2d; + Blob backward_result_2d; + Blob backward_weight_result_2d; + // Test with 2D im2col + { + caffe_set(this->blob_top_->count(), Dtype(0), + this->blob_top_->mutable_cpu_data()); + caffe_set(this->blob_bottom_->count(), Dtype(0), + this->blob_bottom_->mutable_cpu_diff()); + caffe_set(weights.count(), Dtype(0), weights.mutable_cpu_diff()); + // Do SetUp and Forward; save Forward result in result_2d. + convolution_param->set_force_nd_im2col(false); + CuDNNDeconvolutionLayer layer_2d(layer_param); + layer_2d.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + ASSERT_EQ(1, layer_2d.blobs().size()); + copy_diff = false; reshape = false; + layer_2d.blobs()[0]->CopyFrom(weights, copy_diff, reshape); + layer_2d.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + copy_diff = false; reshape = true; + result_2d.CopyFrom(*this->blob_top_, copy_diff, reshape); + // Copy pre-generated top diff into actual top diff; + // do Backward and save result in backward_result_2d. + ASSERT_EQ(this->blob_top_->shape(), top_diff.shape()); + caffe_copy(top_diff.count(), top_diff.cpu_data(), + this->blob_top_->mutable_cpu_diff()); + layer_2d.Backward(this->blob_top_vec_, propagate_down, + this->blob_bottom_vec_); + copy_diff = true; reshape = true; + backward_result_2d.CopyFrom(*this->blob_bottom_, copy_diff, reshape); + backward_weight_result_2d.CopyFrom(weights, copy_diff, reshape); + } + Blob result_nd; + Blob backward_result_nd; + Blob backward_weight_result_nd; + // Test with ND im2col + { + caffe_set(this->blob_top_->count(), Dtype(0), + this->blob_top_->mutable_cpu_data()); + caffe_set(this->blob_bottom_->count(), Dtype(0), + this->blob_bottom_->mutable_cpu_diff()); + caffe_set(weights.count(), Dtype(0), weights.mutable_cpu_diff()); + // Do SetUp and Forward; save Forward result in result_nd. + convolution_param->set_force_nd_im2col(true); + CuDNNDeconvolutionLayer layer_nd(layer_param); + layer_nd.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + ASSERT_EQ(1, layer_nd.blobs().size()); + copy_diff = false; reshape = false; + layer_nd.blobs()[0]->CopyFrom(weights, copy_diff, reshape); + layer_nd.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + copy_diff = false; reshape = true; + result_nd.CopyFrom(*this->blob_top_, copy_diff, reshape); + // Copy pre-generated top diff into actual top diff; + // do Backward and save result in backward_result_nd. + ASSERT_EQ(this->blob_top_->shape(), top_diff.shape()); + caffe_copy(top_diff.count(), top_diff.cpu_data(), + this->blob_top_->mutable_cpu_diff()); + layer_nd.Backward(this->blob_top_vec_, propagate_down, + this->blob_bottom_vec_); + copy_diff = true; reshape = true; + backward_result_nd.CopyFrom(*this->blob_bottom_, copy_diff, reshape); + backward_weight_result_nd.CopyFrom(weights, copy_diff, reshape); + } + ASSERT_EQ(result_nd.count(), result_2d.count()); + for (int i = 0; i < result_2d.count(); ++i) { + EXPECT_NEAR(result_2d.cpu_data()[i], result_nd.cpu_data()[i], 1e-4); + } + ASSERT_EQ(backward_result_nd.count(), backward_result_2d.count()); + for (int i = 0; i < backward_result_2d.count(); ++i) { + EXPECT_EQ(backward_result_2d.cpu_diff()[i], + backward_result_nd.cpu_diff()[i]); + } + ASSERT_EQ(backward_weight_result_nd.count(), + backward_weight_result_2d.count()); + for (int i = 0; i < backward_weight_result_2d.count(); ++i) { + EXPECT_EQ(backward_weight_result_2d.cpu_diff()[i], + backward_weight_result_nd.cpu_diff()[i]); + } +} + +#endif + } // namespace caffe diff --git a/src/caffe/test/test_filler.cpp b/src/caffe/test/test_filler.cpp index 26e9b217..34f7007d 100644 --- a/src/caffe/test/test_filler.cpp +++ b/src/caffe/test/test_filler.cpp @@ -1,3 +1,5 @@ +#include + #include "gtest/gtest.h" #include "caffe/filler.hpp" @@ -10,11 +12,20 @@ template class ConstantFillerTest : public ::testing::Test { protected: ConstantFillerTest() - : blob_(new Blob(2, 3, 4, 5)), + : blob_(new Blob()), filler_param_() { filler_param_.set_value(10.); filler_.reset(new ConstantFiller(filler_param_)); + } + virtual void test_params(const vector& shape) { + EXPECT_TRUE(blob_); + blob_->Reshape(shape); filler_->Fill(blob_); + const int count = blob_->count(); + const Dtype* data = blob_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_EQ(data[i], filler_param_.value()); + } } virtual ~ConstantFillerTest() { delete blob_; } Blob* const blob_; @@ -25,12 +36,34 @@ class ConstantFillerTest : public ::testing::Test { TYPED_TEST_CASE(ConstantFillerTest, TestDtypes); TYPED_TEST(ConstantFillerTest, TestFill) { - EXPECT_TRUE(this->blob_); - const int count = this->blob_->count(); - const TypeParam* data = this->blob_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_GE(data[i], this->filler_param_.value()); - } + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + this->test_params(blob_shape); +} + +TYPED_TEST(ConstantFillerTest, TestFill1D) { + vector blob_shape(1, 15); + this->test_params(blob_shape); +} + +TYPED_TEST(ConstantFillerTest, TestFill2D) { + vector blob_shape; + blob_shape.push_back(8); + blob_shape.push_back(3); + this->test_params(blob_shape); +} + +TYPED_TEST(ConstantFillerTest, TestFill5D) { + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + blob_shape.push_back(2); + this->test_params(blob_shape); } @@ -38,12 +71,22 @@ template class UniformFillerTest : public ::testing::Test { protected: UniformFillerTest() - : blob_(new Blob(2, 3, 4, 5)), + : blob_(new Blob()), filler_param_() { filler_param_.set_min(1.); filler_param_.set_max(2.); filler_.reset(new UniformFiller(filler_param_)); + } + virtual void test_params(const vector& shape) { + EXPECT_TRUE(blob_); + blob_->Reshape(shape); filler_->Fill(blob_); + const int count = blob_->count(); + const Dtype* data = blob_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_GE(data[i], filler_param_.min()); + EXPECT_LE(data[i], filler_param_.max()); + } } virtual ~UniformFillerTest() { delete blob_; } Blob* const blob_; @@ -54,23 +97,64 @@ class UniformFillerTest : public ::testing::Test { TYPED_TEST_CASE(UniformFillerTest, TestDtypes); TYPED_TEST(UniformFillerTest, TestFill) { - EXPECT_TRUE(this->blob_); - const int count = this->blob_->count(); - const TypeParam* data = this->blob_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_GE(data[i], this->filler_param_.min()); - EXPECT_LE(data[i], this->filler_param_.max()); - } + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + this->test_params(blob_shape); +} + +TYPED_TEST(UniformFillerTest, TestFill1D) { + vector blob_shape(1, 15); + this->test_params(blob_shape); +} + +TYPED_TEST(UniformFillerTest, TestFill2D) { + vector blob_shape; + blob_shape.push_back(8); + blob_shape.push_back(3); + this->test_params(blob_shape); +} + +TYPED_TEST(UniformFillerTest, TestFill5D) { + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + blob_shape.push_back(2); + this->test_params(blob_shape); } template class PositiveUnitballFillerTest : public ::testing::Test { protected: PositiveUnitballFillerTest() - : blob_(new Blob(2, 3, 4, 5)), + : blob_(new Blob()), filler_param_() { filler_.reset(new PositiveUnitballFiller(filler_param_)); + } + virtual void test_params(const vector& shape) { + EXPECT_TRUE(blob_); + blob_->Reshape(shape); filler_->Fill(blob_); + const int num = blob_->shape(0); + const int count = blob_->count(); + const int dim = count / num; + const Dtype* data = blob_->cpu_data(); + for (int i = 0; i < count; ++i) { + EXPECT_GE(data[i], 0); + EXPECT_LE(data[i], 1); + } + for (int i = 0; i < num; ++i) { + Dtype sum = Dtype(0); + for (int j = 0; j < dim; ++j) { + sum += data[i * dim + j]; + } + EXPECT_GE(sum, 0.999); + EXPECT_LE(sum, 1.001); + } } virtual ~PositiveUnitballFillerTest() { delete blob_; } Blob* const blob_; @@ -81,35 +165,78 @@ class PositiveUnitballFillerTest : public ::testing::Test { TYPED_TEST_CASE(PositiveUnitballFillerTest, TestDtypes); TYPED_TEST(PositiveUnitballFillerTest, TestFill) { - EXPECT_TRUE(this->blob_); - const int num = this->blob_->num(); - const int count = this->blob_->count(); - const int dim = count / num; - const TypeParam* data = this->blob_->cpu_data(); - for (int i = 0; i < count; ++i) { - EXPECT_GE(data[i], 0); - EXPECT_LE(data[i], 1); - } - for (int i = 0; i < num; ++i) { - TypeParam sum = 0; - for (int j = 0; j < dim; ++j) { - sum += data[i * dim + j]; - } - EXPECT_GE(sum, 0.999); - EXPECT_LE(sum, 1.001); - } + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + this->test_params(blob_shape); +} + +TYPED_TEST(PositiveUnitballFillerTest, TestFill1D) { + vector blob_shape(1, 15); + this->test_params(blob_shape); +} + +TYPED_TEST(PositiveUnitballFillerTest, TestFill2D) { + vector blob_shape; + blob_shape.push_back(8); + blob_shape.push_back(3); + this->test_params(blob_shape); +} + +TYPED_TEST(PositiveUnitballFillerTest, TestFill5D) { + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + blob_shape.push_back(2); + this->test_params(blob_shape); } template class GaussianFillerTest : public ::testing::Test { protected: GaussianFillerTest() - : blob_(new Blob(2, 3, 4, 5)), + : blob_(new Blob()), filler_param_() { filler_param_.set_mean(10.); filler_param_.set_std(0.1); filler_.reset(new GaussianFiller(filler_param_)); + } + virtual void test_params(const vector& shape, + const Dtype tolerance = Dtype(5), const int repetitions = 100) { + // Tests for statistical properties should be ran multiple times. + EXPECT_TRUE(blob_); + blob_->Reshape(shape); + for (int i = 0; i < repetitions; ++i) { + test_params_iter(shape, tolerance); + } + } + virtual void test_params_iter(const vector& shape, + const Dtype tolerance) { + // This test has a configurable tolerance parameter - by default it was + // equal to 5.0 which is very loose - allowing some tuning (e.g. for tests + // on smaller blobs the actual variance will be larger than desired, so the + // tolerance can be increased to account for that). filler_->Fill(blob_); + const int count = blob_->count(); + const Dtype* data = blob_->cpu_data(); + Dtype mean = Dtype(0); + Dtype var = Dtype(0); + for (int i = 0; i < count; ++i) { + mean += data[i]; + var += data[i] * data[i]; + } + mean /= count; + var /= count; + var -= mean*mean; + EXPECT_GE(mean, filler_param_.mean() - filler_param_.std() * tolerance); + EXPECT_LE(mean, filler_param_.mean() + filler_param_.std() * tolerance); + Dtype target_var = filler_param_.std() * filler_param_.std(); + EXPECT_GE(var, target_var / tolerance); + EXPECT_LE(var, target_var * tolerance); } virtual ~GaussianFillerTest() { delete blob_; } Blob* const blob_; @@ -120,41 +247,62 @@ class GaussianFillerTest : public ::testing::Test { TYPED_TEST_CASE(GaussianFillerTest, TestDtypes); TYPED_TEST(GaussianFillerTest, TestFill) { - EXPECT_TRUE(this->blob_); - const int count = this->blob_->count(); - const TypeParam* data = this->blob_->cpu_data(); - TypeParam mean = 0.; - TypeParam var = 0.; - for (int i = 0; i < count; ++i) { - mean += data[i]; - var += (data[i] - this->filler_param_.mean()) * - (data[i] - this->filler_param_.mean()); - } - mean /= count; - var /= count; - // Very loose test. - EXPECT_GE(mean, this->filler_param_.mean() - this->filler_param_.std() * 5); - EXPECT_LE(mean, this->filler_param_.mean() + this->filler_param_.std() * 5); - TypeParam target_var = this->filler_param_.std() * this->filler_param_.std(); - EXPECT_GE(var, target_var / 5.); - EXPECT_LE(var, target_var * 5.); + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + const TypeParam tolerance = TypeParam(3); // enough for a 120-element blob + this->test_params(blob_shape, tolerance); +} + +TYPED_TEST(GaussianFillerTest, TestFill1D) { + vector blob_shape(1, 125); + const TypeParam tolerance = TypeParam(3); + this->test_params(blob_shape, tolerance); +} + +TYPED_TEST(GaussianFillerTest, TestFill2D) { + vector blob_shape; + blob_shape.push_back(8); + blob_shape.push_back(15); + const TypeParam tolerance = TypeParam(3); + this->test_params(blob_shape, tolerance); +} + +TYPED_TEST(GaussianFillerTest, TestFill5D) { + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + blob_shape.push_back(2); + const TypeParam tolerance = TypeParam(2); + this->test_params(blob_shape, tolerance); } template class XavierFillerTest : public ::testing::Test { protected: XavierFillerTest() - : blob_(new Blob(1000, 2, 4, 5)), + : blob_(new Blob()), filler_param_() { } virtual void test_params(FillerParameter_VarianceNorm variance_norm, + Dtype n, const vector& shape, const int repetitions = 100) { + EXPECT_TRUE(blob_); + blob_->Reshape(shape); + for (int i = 0; i < repetitions; ++i) { + test_params_iter(variance_norm, n); + } + } + virtual void test_params_iter(FillerParameter_VarianceNorm variance_norm, Dtype n) { - this->filler_param_.set_variance_norm(variance_norm); - this->filler_.reset(new XavierFiller(this->filler_param_)); - this->filler_->Fill(blob_); - EXPECT_TRUE(this->blob_); - const int count = this->blob_->count(); - const Dtype* data = this->blob_->cpu_data(); + filler_param_.set_variance_norm(variance_norm); + filler_.reset(new XavierFiller(filler_param_)); + filler_->Fill(blob_); + const int count = blob_->count(); + const Dtype* data = blob_->cpu_data(); Dtype mean = 0.; Dtype ex2 = 0.; for (int i = 0; i < count; ++i) { @@ -177,33 +325,92 @@ class XavierFillerTest : public ::testing::Test { TYPED_TEST_CASE(XavierFillerTest, TestDtypes); TYPED_TEST(XavierFillerTest, TestFillFanIn) { + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(4); + blob_shape.push_back(5); TypeParam n = 2*4*5; - this->test_params(FillerParameter_VarianceNorm_FAN_IN, n); + this->test_params(FillerParameter_VarianceNorm_FAN_IN, n, blob_shape); } + TYPED_TEST(XavierFillerTest, TestFillFanOut) { + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(4); + blob_shape.push_back(5); TypeParam n = 1000*4*5; - this->test_params(FillerParameter_VarianceNorm_FAN_OUT, n); + this->test_params(FillerParameter_VarianceNorm_FAN_OUT, n, blob_shape); } + TYPED_TEST(XavierFillerTest, TestFillAverage) { + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(4); + blob_shape.push_back(5); TypeParam n = (2*4*5 + 1000*4*5) / 2.0; - this->test_params(FillerParameter_VarianceNorm_AVERAGE, n); + this->test_params(FillerParameter_VarianceNorm_AVERAGE, n, blob_shape); +} + +TYPED_TEST(XavierFillerTest, TestFill1D) { + // This makes little sense but at least we will know that we can fill it + EXPECT_TRUE(this->blob_); + vector blob_shape(1, 25); + this->blob_->Reshape(blob_shape); + this->filler_param_.set_variance_norm(FillerParameter_VarianceNorm_AVERAGE); + this->filler_.reset(new XavierFiller(this->filler_param_)); + this->filler_->Fill(this->blob_); +} + +TYPED_TEST(XavierFillerTest, TestFill2D) { + EXPECT_TRUE(this->blob_); + vector blob_shape; + blob_shape.push_back(8); + blob_shape.push_back(3); + this->blob_->Reshape(blob_shape); + this->filler_param_.set_variance_norm(FillerParameter_VarianceNorm_AVERAGE); + this->filler_.reset(new XavierFiller(this->filler_param_)); + this->filler_->Fill(this->blob_); +} + +TYPED_TEST(XavierFillerTest, TestFill5D) { + EXPECT_TRUE(this->blob_); + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + blob_shape.push_back(2); + this->blob_->Reshape(blob_shape); + this->filler_param_.set_variance_norm(FillerParameter_VarianceNorm_AVERAGE); + this->filler_.reset(new XavierFiller(this->filler_param_)); + this->filler_->Fill(this->blob_); } template class MSRAFillerTest : public ::testing::Test { protected: MSRAFillerTest() - : blob_(new Blob(1000, 2, 4, 5)), + : blob_(new Blob()), filler_param_() { } virtual void test_params(FillerParameter_VarianceNorm variance_norm, + Dtype n, const vector& shape, const int repetitions = 100) { + EXPECT_TRUE(blob_); + blob_->Reshape(shape); + for (int i = 0; i < repetitions; ++i) { + test_params_iter(variance_norm, n); + } + } + virtual void test_params_iter(FillerParameter_VarianceNorm variance_norm, Dtype n) { - this->filler_param_.set_variance_norm(variance_norm); - this->filler_.reset(new MSRAFiller(this->filler_param_)); - this->filler_->Fill(blob_); - EXPECT_TRUE(this->blob_); - const int count = this->blob_->count(); - const Dtype* data = this->blob_->cpu_data(); + filler_param_.set_variance_norm(variance_norm); + filler_.reset(new MSRAFiller(filler_param_)); + filler_->Fill(blob_); + const int count = blob_->count(); + const Dtype* data = blob_->cpu_data(); Dtype mean = 0.; Dtype ex2 = 0.; for (int i = 0; i < count; ++i) { @@ -226,16 +433,123 @@ class MSRAFillerTest : public ::testing::Test { TYPED_TEST_CASE(MSRAFillerTest, TestDtypes); TYPED_TEST(MSRAFillerTest, TestFillFanIn) { + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(4); + blob_shape.push_back(5); TypeParam n = 2*4*5; - this->test_params(FillerParameter_VarianceNorm_FAN_IN, n); + this->test_params(FillerParameter_VarianceNorm_FAN_IN, n, blob_shape); } + TYPED_TEST(MSRAFillerTest, TestFillFanOut) { + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(4); + blob_shape.push_back(5); TypeParam n = 1000*4*5; - this->test_params(FillerParameter_VarianceNorm_FAN_OUT, n); + this->test_params(FillerParameter_VarianceNorm_FAN_OUT, n, blob_shape); } + TYPED_TEST(MSRAFillerTest, TestFillAverage) { + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(4); + blob_shape.push_back(5); TypeParam n = (2*4*5 + 1000*4*5) / 2.0; - this->test_params(FillerParameter_VarianceNorm_AVERAGE, n); + this->test_params(FillerParameter_VarianceNorm_AVERAGE, n, blob_shape); +} + +TYPED_TEST(MSRAFillerTest, TestFill1D) { + // Like with Xavier - no checking for correctness, just if it can be filled. + EXPECT_TRUE(this->blob_); + vector blob_shape(1, 25); + this->blob_->Reshape(blob_shape); + this->filler_param_.set_variance_norm(FillerParameter_VarianceNorm_AVERAGE); + this->filler_.reset(new MSRAFiller(this->filler_param_)); + this->filler_->Fill(this->blob_); +} + +TYPED_TEST(MSRAFillerTest, TestFill2D) { + EXPECT_TRUE(this->blob_); + vector blob_shape; + blob_shape.push_back(8); + blob_shape.push_back(3); + this->blob_->Reshape(blob_shape); + this->filler_param_.set_variance_norm(FillerParameter_VarianceNorm_AVERAGE); + this->filler_.reset(new MSRAFiller(this->filler_param_)); + this->filler_->Fill(this->blob_); +} + +TYPED_TEST(MSRAFillerTest, TestFill5D) { + EXPECT_TRUE(this->blob_); + vector blob_shape; + blob_shape.push_back(2); + blob_shape.push_back(3); + blob_shape.push_back(4); + blob_shape.push_back(5); + blob_shape.push_back(2); + this->blob_->Reshape(blob_shape); + this->filler_param_.set_variance_norm(FillerParameter_VarianceNorm_AVERAGE); + this->filler_.reset(new MSRAFiller(this->filler_param_)); + this->filler_->Fill(this->blob_); +} + +template +class BilinearFillerTest : public ::testing::Test { + protected: + BilinearFillerTest() + : blob_(new Blob()), + filler_param_() { + } + virtual void test_params(const vector& shape) { + EXPECT_TRUE(blob_); + blob_->Reshape(shape); + filler_.reset(new BilinearFiller(filler_param_)); + filler_->Fill(blob_); + CHECK_EQ(blob_->num_axes(), 4); + const int outer_num = blob_->count(0, 2); + const int inner_num = blob_->count(2, 4); + const Dtype* data = blob_->cpu_data(); + int f = ceil(blob_->shape(3) / 2.); + Dtype c = (blob_->shape(3) - 1) / (2. * f); + for (int i = 0; i < outer_num; ++i) { + for (int j = 0; j < inner_num; ++j) { + Dtype x = j % blob_->shape(3); + Dtype y = (j / blob_->shape(3)) % blob_->shape(2); + Dtype expected_value = (1 - fabs(x / f - c)) * (1 - fabs(y / f - c)); + const Dtype actual_value = data[i * inner_num + j]; + EXPECT_NEAR(expected_value, actual_value, 0.01); + } + } + } + virtual ~BilinearFillerTest() { delete blob_; } + Blob* blob_; + FillerParameter filler_param_; + shared_ptr > filler_; +}; + +TYPED_TEST_CASE(BilinearFillerTest, TestDtypes); + +TYPED_TEST(BilinearFillerTest, TestFillOdd) { + const int n = 7; + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(n); + blob_shape.push_back(n); + this->test_params(blob_shape); +} +TYPED_TEST(BilinearFillerTest, TestFillEven) { + const int n = 6; + vector blob_shape; + blob_shape.push_back(1000); + blob_shape.push_back(2); + blob_shape.push_back(n); + blob_shape.push_back(n); + this->test_params(blob_shape); } } // namespace caffe diff --git a/src/caffe/test/test_hdf5_output_layer.cpp b/src/caffe/test/test_hdf5_output_layer.cpp index 55739153..3b19a8cf 100644 --- a/src/caffe/test/test_hdf5_output_layer.cpp +++ b/src/caffe/test/test_hdf5_output_layer.cpp @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 #include #include @@ -120,3 +121,4 @@ TYPED_TEST(HDF5OutputLayerTest, TestForward) { } } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/test/test_hdf5data_layer.cpp b/src/caffe/test/test_hdf5data_layer.cpp index 3977c486..0e5c398f 100644 --- a/src/caffe/test/test_hdf5data_layer.cpp +++ b/src/caffe/test/test_hdf5data_layer.cpp @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 #include #include @@ -163,3 +164,4 @@ TYPED_TEST(HDF5DataLayerTest, TestSkip) { } } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/test/test_neuron_layer.cpp b/src/caffe/test/test_neuron_layer.cpp index 180871a2..d1ecc37b 100644 --- a/src/caffe/test/test_neuron_layer.cpp +++ b/src/caffe/test/test_neuron_layer.cpp @@ -10,6 +10,7 @@ #include "caffe/layers/absval_layer.hpp" #include "caffe/layers/bnll_layer.hpp" +#include "caffe/layers/clip_layer.hpp" #include "caffe/layers/dropout_layer.hpp" #include "caffe/layers/elu_layer.hpp" #include "caffe/layers/exp_layer.hpp" @@ -19,6 +20,7 @@ #include "caffe/layers/prelu_layer.hpp" #include "caffe/layers/relu_layer.hpp" #include "caffe/layers/sigmoid_layer.hpp" +#include "caffe/layers/swish_layer.hpp" #include "caffe/layers/tanh_layer.hpp" #include "caffe/layers/threshold_layer.hpp" @@ -205,6 +207,66 @@ TYPED_TEST(NeuronLayerTest, TestAbsGradient) { this->blob_top_vec_); } +TYPED_TEST(NeuronLayerTest, TestClip) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CHECK(google::protobuf::TextFormat::ParseFromString( + "clip_param { min: -1, max: 2 }", &layer_param)); + ClipLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // Now, check values + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + const Dtype* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_GE(top_data[i], -1); + EXPECT_LE(top_data[i], 2); + EXPECT_TRUE(bottom_data[i] > -1 || top_data[i] == -1); + EXPECT_TRUE(bottom_data[i] < 2 || top_data[i] == 2); + EXPECT_TRUE(!(bottom_data[i] >= -1 && bottom_data[i] <= 2) + || top_data[i] == bottom_data[i]); + } +} + +TYPED_TEST(NeuronLayerTest, TestClipGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CHECK(google::protobuf::TextFormat::ParseFromString( + "clip_param { min: -1, max: 2 }", &layer_param)); + ClipLayer layer(layer_param); + // Unfortunately, it might happen that an input value lands exactly within + // the discontinuity region of the Clip function. In this case the numeric + // gradient is likely to differ significantly (i.e. by a value larger than + // checker tolerance) from the computed gradient. To handle such cases, we + // eliminate such values from the input blob before the gradient check. + const Dtype epsilon = 1e-2; + const Dtype min_range_start = layer_param.clip_param().min() - epsilon; + const Dtype min_range_end = layer_param.clip_param().min() + epsilon; + const Dtype max_range_start = layer_param.clip_param().max() - epsilon; + const Dtype max_range_end = layer_param.clip_param().max() + epsilon; + // The input blob is owned by the NeuronLayerTest object, so we begin with + // creating a temporary blob and copying the input data there. + Blob temp_bottom; + temp_bottom.ReshapeLike(*this->blob_bottom_); + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + Dtype* temp_data_mutable = temp_bottom.mutable_cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + if (bottom_data[i] >= min_range_start && + bottom_data[i] <= min_range_end) { + temp_data_mutable[i] = bottom_data[i] - epsilon; + } else if (bottom_data[i] >= max_range_start && + bottom_data[i] <= max_range_end) { + temp_data_mutable[i] = bottom_data[i] + epsilon; + } else { + temp_data_mutable[i] = bottom_data[i]; + } + } + vector*> temp_bottom_vec; + temp_bottom_vec.push_back(&temp_bottom); + GradientChecker checker(epsilon, 1e-3); + checker.CheckGradientEltwise(&layer, temp_bottom_vec, this->blob_top_vec_); +} + TYPED_TEST(NeuronLayerTest, TestReLU) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; @@ -344,6 +406,84 @@ TYPED_TEST(NeuronLayerTest, TestSigmoidGradient) { this->blob_top_vec_); } +TYPED_TEST(NeuronLayerTest, TestSwish) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + SwishLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // Now, check values + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + const Dtype* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_FLOAT_EQ(top_data[i], bottom_data[i] / (1. + exp(-bottom_data[i]))); + } +} + +TYPED_TEST(NeuronLayerTest, TestSwishWithBeta) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CHECK(google::protobuf::TextFormat::ParseFromString( + "swish_param { beta: 1.5 }", &layer_param)); + SwishLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // Now, check values + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + const Dtype* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_FLOAT_EQ(top_data[i], bottom_data[i] / (1. + exp(-1.5 * + bottom_data[i]))); + } +} + +TYPED_TEST(NeuronLayerTest, TestSwishAsLinear) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CHECK(google::protobuf::TextFormat::ParseFromString( + "swish_param { beta: 0.0 }", &layer_param)); + SwishLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // Now, check values + const Dtype* bottom_data = this->blob_bottom_->cpu_data(); + const Dtype* top_data = this->blob_top_->cpu_data(); + for (int i = 0; i < this->blob_bottom_->count(); ++i) { + EXPECT_FLOAT_EQ(top_data[i], bottom_data[i] / 2.0); + } +} + +TYPED_TEST(NeuronLayerTest, TestSwishGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + SwishLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3, 1701, 0., 0.01); + checker.CheckGradientEltwise(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(NeuronLayerTest, TestSwishWithBetaGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CHECK(google::protobuf::TextFormat::ParseFromString( + "swish_param { beta: 1.5 }", &layer_param)); + SwishLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3, 1701, 0., 0.01); + checker.CheckGradientEltwise(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(NeuronLayerTest, TestSwishAsLinearGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CHECK(google::protobuf::TextFormat::ParseFromString( + "swish_param { beta: 0.0 }", &layer_param)); + SwishLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3, 1701, 0., 0.01); + checker.CheckGradientEltwise(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + TYPED_TEST(NeuronLayerTest, TestTanH) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; diff --git a/src/caffe/test/test_syncedmem.cpp b/src/caffe/test/test_syncedmem.cpp index 16dfb582..2ca9ca2f 100644 --- a/src/caffe/test/test_syncedmem.cpp +++ b/src/caffe/test/test_syncedmem.cpp @@ -80,7 +80,7 @@ TEST_F(SyncedMemoryTest, TestGPURead) { char* recovered_value = new char[10]; caffe_gpu_memcpy(10, gpu_data, recovered_value); for (int i = 0; i < mem.size(); ++i) { - EXPECT_EQ((static_cast(recovered_value))[i], 1); + EXPECT_EQ(recovered_value[i], 1); } // do another round cpu_data = mem.mutable_cpu_data(); @@ -94,7 +94,7 @@ TEST_F(SyncedMemoryTest, TestGPURead) { // check if values are the same caffe_gpu_memcpy(10, gpu_data, recovered_value); for (int i = 0; i < mem.size(); ++i) { - EXPECT_EQ((static_cast(recovered_value))[i], 2); + EXPECT_EQ(recovered_value[i], 2); } delete[] recovered_value; } diff --git a/src/caffe/test/test_upgrade_proto.cpp b/src/caffe/test/test_upgrade_proto.cpp index 9dcc2aa5..769112eb 100644 --- a/src/caffe/test/test_upgrade_proto.cpp +++ b/src/caffe/test/test_upgrade_proto.cpp @@ -2952,6 +2952,8 @@ TEST_F(SolverTypeUpgradeTest, TestSimple) { for (int i = 0; i < 6; ++i) { const string& input_proto = "net: 'examples/mnist/lenet_train_test.prototxt' " + "weights: 'examples/mnist/lenet_train_test1.caffemodel' " + "weights: 'examples/mnist/lenet_train_test2.caffemodel' " "test_iter: 100 " "test_interval: 500 " "base_lr: 0.01 " @@ -2968,6 +2970,8 @@ TEST_F(SolverTypeUpgradeTest, TestSimple) { "solver_type: " + std::string(old_type_vec[i]) + " "; const string& expected_output_proto = "net: 'examples/mnist/lenet_train_test.prototxt' " + "weights: 'examples/mnist/lenet_train_test1.caffemodel' " + "weights: 'examples/mnist/lenet_train_test2.caffemodel' " "test_iter: 100 " "test_interval: 500 " "base_lr: 0.01 " diff --git a/src/caffe/util/hdf5.cpp b/src/caffe/util/hdf5.cpp index ed737429..cefd853d 100644 --- a/src/caffe/util/hdf5.cpp +++ b/src/caffe/util/hdf5.cpp @@ -1,3 +1,4 @@ +#ifdef USE_HDF5 #include "caffe/util/hdf5.hpp" #include @@ -207,3 +208,4 @@ string hdf5_get_name_by_idx(hid_t loc_id, int idx) { } } // namespace caffe +#endif // USE_HDF5 diff --git a/src/caffe/util/io.cpp b/src/caffe/util/io.cpp index 29e952b3..58467c3e 100644 --- a/src/caffe/util/io.cpp +++ b/src/caffe/util/io.cpp @@ -122,7 +122,7 @@ cv::Mat ReadImageToCVMat(const string& filename) { static bool matchExt(const std::string & fn, std::string en) { size_t p = fn.rfind('.'); - std::string ext = p != fn.npos ? fn.substr(p) : fn; + std::string ext = p != fn.npos ? fn.substr(p+1) : fn; std::transform(ext.begin(), ext.end(), ext.begin(), ::tolower); std::transform(en.begin(), en.end(), en.begin(), ::tolower); if ( ext == en ) diff --git a/src/caffe/util/signal_handler.cpp b/src/caffe/util/signal_handler.cpp index 5d764ec5..9658fb39 100644 --- a/src/caffe/util/signal_handler.cpp +++ b/src/caffe/util/signal_handler.cpp @@ -48,7 +48,7 @@ namespace { void UnhookHandler() { if (already_hooked_up) { struct sigaction sa; - // Setup the sighub handler + // Setup the sighup handler sa.sa_handler = SIG_DFL; // Restart the system call, if at all possible sa.sa_flags = SA_RESTART; diff --git a/src/caffe/util/upgrade_proto.cpp b/src/caffe/util/upgrade_proto.cpp index 94771c8c..ad40b73d 100644 --- a/src/caffe/util/upgrade_proto.cpp +++ b/src/caffe/util/upgrade_proto.cpp @@ -2,6 +2,8 @@ #include #include +#include + #include #include @@ -1095,12 +1097,31 @@ bool UpgradeSolverAsNeeded(const string& param_file, SolverParameter* param) { return success; } +// Replaces snapshot_prefix of SolverParameter if it is not specified +// or is set to directory +void UpgradeSnapshotPrefixProperty(const string& param_file, + SolverParameter* param) { + using boost::filesystem::path; + using boost::filesystem::is_directory; + if (!param->has_snapshot_prefix()) { + param->set_snapshot_prefix(path(param_file).replace_extension().string()); + LOG(INFO) << "snapshot_prefix was not specified and is set to " + + param->snapshot_prefix(); + } else if (is_directory(param->snapshot_prefix())) { + param->set_snapshot_prefix((path(param->snapshot_prefix()) / + path(param_file).stem()).string()); + LOG(INFO) << "snapshot_prefix was a directory and is replaced to " + + param->snapshot_prefix(); + } +} + // Read parameters from a file into a SolverParameter proto message. void ReadSolverParamsFromTextFileOrDie(const string& param_file, SolverParameter* param) { CHECK(ReadProtoFromTextFile(param_file, param)) << "Failed to parse SolverParameter file: " << param_file; UpgradeSolverAsNeeded(param_file, param); + UpgradeSnapshotPrefixProperty(param_file, param); } } // namespace caffe diff --git a/tools/caffe.cpp b/tools/caffe.cpp index 3587d8aa..389cfb8a 100644 --- a/tools/caffe.cpp +++ b/tools/caffe.cpp @@ -146,20 +146,6 @@ int device_query() { } RegisterBrewFunction(device_query); -// Load the weights from the specified caffemodel(s) into the train and -// test nets. -void CopyLayers(caffe::Solver* solver, const std::string& model_list) { - std::vector model_names; - boost::split(model_names, model_list, boost::is_any_of(",") ); - for (int i = 0; i < model_names.size(); ++i) { - LOG(INFO) << "Finetuning from " << model_names[i]; - solver->net()->CopyTrainedLayersFrom(model_names[i]); - for (int j = 0; j < solver->test_nets().size(); ++j) { - solver->test_nets()[j]->CopyTrainedLayersFrom(model_names[i]); - } - } -} - // Translate the signal effect the user specified on the command-line to the // corresponding enumeration. caffe::SolverAction::Enum GetRequestedAction( @@ -233,6 +219,13 @@ int train() { GetRequestedAction(FLAGS_sigint_effect), GetRequestedAction(FLAGS_sighup_effect)); + if (FLAGS_snapshot.size()) { + solver_param.clear_weights(); + } else if (FLAGS_weights.size()) { + solver_param.clear_weights(); + solver_param.add_weights(FLAGS_weights); + } + shared_ptr > solver(caffe::SolverRegistry::CreateSolver(solver_param)); @@ -241,8 +234,6 @@ int train() { if (FLAGS_snapshot.size()) { LOG(INFO) << "Resuming from " << FLAGS_snapshot; solver->Restore(FLAGS_snapshot.c_str()); - } else if (FLAGS_weights.size()) { - CopyLayers(solver.get(), FLAGS_weights); } LOG(INFO) << "Starting Optimization"; diff --git a/tools/convert_imageset.cpp b/tools/convert_imageset.cpp index 90cdb15d..9c5d09f9 100644 --- a/tools/convert_imageset.cpp +++ b/tools/convert_imageset.cpp @@ -115,7 +115,7 @@ int main(int argc, char** argv) { size_t p = fn.rfind('.'); if ( p == fn.npos ) LOG(WARNING) << "Failed to guess the encoding of '" << fn << "'"; - enc = fn.substr(p); + enc = fn.substr(p+1); std::transform(enc.begin(), enc.end(), enc.begin(), ::tolower); } status = ReadImageToDatum(root_folder + lines[line_id].first, diff --git a/tools/device_query.cpp b/tools/device_query.cpp deleted file mode 100644 index 03799e52..00000000 --- a/tools/device_query.cpp +++ /dev/null @@ -1,7 +0,0 @@ -#include "caffe/common.hpp" - -int main(int argc, char** argv) { - LOG(FATAL) << "Deprecated. Use caffe device_query " - "[--device_id=0] instead."; - return 0; -} diff --git a/tools/finetune_net.cpp b/tools/finetune_net.cpp deleted file mode 100644 index 81c0c354..00000000 --- a/tools/finetune_net.cpp +++ /dev/null @@ -1,7 +0,0 @@ -#include "caffe/caffe.hpp" - -int main(int argc, char** argv) { - LOG(FATAL) << "Deprecated. Use caffe train --solver=... " - "[--weights=...] instead."; - return 0; -} diff --git a/tools/net_speed_benchmark.cpp b/tools/net_speed_benchmark.cpp deleted file mode 100644 index cd16e8d0..00000000 --- a/tools/net_speed_benchmark.cpp +++ /dev/null @@ -1,7 +0,0 @@ -#include "caffe/caffe.hpp" - -int main(int argc, char** argv) { - LOG(FATAL) << "Deprecated. Use caffe time --model=... " - "[--iterations=50] [--gpu] [--device_id=0]"; - return 0; -} diff --git a/tools/test_net.cpp b/tools/test_net.cpp deleted file mode 100644 index 92e14eee..00000000 --- a/tools/test_net.cpp +++ /dev/null @@ -1,7 +0,0 @@ -#include "caffe/caffe.hpp" - -int main(int argc, char** argv) { - LOG(FATAL) << "Deprecated. Use caffe test --model=... " - "--weights=... instead."; - return 0; -} diff --git a/tools/train_net.cpp b/tools/train_net.cpp deleted file mode 100644 index 622bca31..00000000 --- a/tools/train_net.cpp +++ /dev/null @@ -1,7 +0,0 @@ -#include "caffe/caffe.hpp" - -int main(int argc, char** argv) { - LOG(FATAL) << "Deprecated. Use caffe train --solver=... " - "[--snapshot=...] instead."; - return 0; -}