diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index c7b6a17aa69..d7eb59e333a 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -26,7 +26,7 @@ include(cmake/ProtoBuf.cmake) # ---[ HDF5 find_package(HDF5 COMPONENTS HL REQUIRED) include_directories(SYSTEM ${HDF5_INCLUDE_DIRS} ${HDF5_HL_INCLUDE_DIR}) -list(APPEND Caffe_LINKER_LIBS ${HDF5_LIBRARIES}) +list(APPEND Caffe_LINKER_LIBS ${HDF5_LIBRARIES} ${HDF5_HL_LIBRARIES}) # ---[ LMDB if(USE_LMDB) diff --git a/cmake/Modules/FindAtlas.cmake b/cmake/Modules/FindAtlas.cmake index 6e1564351c7..9c665a47bd5 100644 --- a/cmake/Modules/FindAtlas.cmake +++ b/cmake/Modules/FindAtlas.cmake @@ -26,9 +26,9 @@ set(Atlas_LIB_SEARCH_PATHS find_path(Atlas_CBLAS_INCLUDE_DIR NAMES cblas.h PATHS ${Atlas_INCLUDE_SEARCH_PATHS}) find_path(Atlas_CLAPACK_INCLUDE_DIR NAMES clapack.h PATHS ${Atlas_INCLUDE_SEARCH_PATHS}) -find_library(Atlas_CBLAS_LIBRARY NAMES ptcblas_r ptcblas cblas_r cblas PATHS ${Atlas_LIB_SEARCH_PATHS}) -find_library(Atlas_BLAS_LIBRARY NAMES atlas_r atlas PATHS ${Atlas_LIB_SEARCH_PATHS}) -find_library(Atlas_LAPACK_LIBRARY NAMES alapack_r alapack lapack_atlas PATHS ${Atlas_LIB_SEARCH_PATHS}) +find_library(Atlas_CBLAS_LIBRARY NAMES ptcblas_r ptcblas cblas_r cblas PATHS ${Atlas_LIB_SEARCH_PATHS}) +find_library(Atlas_BLAS_LIBRARY NAMES atlas_r atlas PATHS ${Atlas_LIB_SEARCH_PATHS}) +find_library(Atlas_LAPACK_LIBRARY NAMES lapack alapack_r alapack lapack_atlas PATHS ${Atlas_LIB_SEARCH_PATHS}) set(LOOKED_FOR Atlas_CBLAS_INCLUDE_DIR diff --git a/include/caffe/filler.hpp b/include/caffe/filler.hpp index dad9ad46b3b..2b25d761884 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,8 @@ 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); + int fan_out = blob->count() / blob->shape(1); Dtype n = fan_in; // default to fan_in if (this->filler_param_.variance_norm() == FillerParameter_VarianceNorm_AVERAGE) { @@ -189,8 +189,8 @@ 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); + int fan_out = blob->count() / blob->shape(1); Dtype n = fan_in; // default to fan_in if (this->filler_param_.variance_norm() == FillerParameter_VarianceNorm_AVERAGE) { diff --git a/include/caffe/layers/pooling_layer.hpp b/include/caffe/layers/pooling_layer.hpp index f4d6803ba8e..2b00b6793ef 100644 --- a/include/caffe/layers/pooling_layer.hpp +++ b/include/caffe/layers/pooling_layer.hpp @@ -44,12 +44,17 @@ class PoolingLayer : public Layer { virtual void Backward_gpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom); - int kernel_h_, kernel_w_; - int stride_h_, stride_w_; - int pad_h_, pad_w_; + /// @brief The spatial dimensions of a filter kernel. + std::vector kernel_shape_; + /// @brief The spatial dimensions of the stride. + std::vector stride_; + /// @brief The spatial dimensions of the padding. + std::vector pad_; + + int num_spatial_axes_; int channels_; - int height_, width_; - int pooled_height_, pooled_width_; + std::vector input_shape_; + std::vector pooled_shape_; bool global_pooling_; Blob rand_idx_; Blob max_idx_; diff --git a/include/caffe/util/cudnn.hpp b/include/caffe/util/cudnn.hpp index a7d8dbbad4c..0aaa9468c5a 100644 --- a/include/caffe/util/cudnn.hpp +++ b/include/caffe/util/cudnn.hpp @@ -4,6 +4,8 @@ #include +#include + #include "caffe/common.hpp" #include "caffe/proto/caffe.pb.h" @@ -68,6 +70,11 @@ inline void createTensor4dDesc(cudnnTensorDescriptor_t* desc) { CUDNN_CHECK(cudnnCreateTensorDescriptor(desc)); } +template +inline void createTensorDesc(cudnnTensorDescriptor_t* desc) { + CUDNN_CHECK(cudnnCreateTensorDescriptor(desc)); +} + template inline void setTensor4dDesc(cudnnTensorDescriptor_t* desc, int n, int c, int h, int w, @@ -76,6 +83,24 @@ inline void setTensor4dDesc(cudnnTensorDescriptor_t* desc, n, c, h, w, stride_n, stride_c, stride_h, stride_w)); } +template +inline void setTensorNdDesc(cudnnTensorDescriptor_t* desc, + std::vector shape, + std::vector stride) { + CHECK_EQ(shape.size(), stride.size()) + << "Dimensions of shape and stride don't match !"; + // fill shape with 1 to create tensors with at least 4 dimensions + // to prevent CUDNN_STATUS_BAD_PARAM error in CUDNN v4 + // TODO(christian.payer@gmx.net): check CUDNN doc, probably fixed + // in newer versions + for (int i = shape.size(); i < 4; ++i) { + shape.push_back(1); + stride.push_back(1); + } + CUDNN_CHECK(cudnnSetTensorNdDescriptor(*desc, dataType::type, + shape.size(), shape.data(), stride.data())); +} + template inline void setTensor4dDesc(cudnnTensorDescriptor_t* desc, int n, int c, int h, int w) { @@ -87,6 +112,17 @@ inline void setTensor4dDesc(cudnnTensorDescriptor_t* desc, stride_n, stride_c, stride_h, stride_w); } +template +inline void setTensorNdDesc(cudnnTensorDescriptor_t* desc, + std::vector shape) { + // set up stride + std::vector stride(shape.size(), 1); + for (int i = stride.size() - 2; i >= 0; --i) { + stride[i] = shape[i + 1] * stride[i + 1]; + } + setTensorNdDesc(desc, shape, stride); +} + template inline void createFilterDesc(cudnnFilterDescriptor_t* desc, int n, int c, int h, int w) { @@ -100,6 +136,19 @@ inline void createFilterDesc(cudnnFilterDescriptor_t* desc, #endif } +template +inline void createNdFilterDesc(cudnnFilterDescriptor_t* desc, + std::vector shape) { + CUDNN_CHECK(cudnnCreateFilterDescriptor(desc)); +#if CUDNN_VERSION_MIN(5, 0, 0) + CUDNN_CHECK(cudnnSetFilterNdDescriptor(*desc, dataType::type, + CUDNN_TENSOR_NCHW, shape.size(), shape.data())); +#else + CUDNN_CHECK(cudnnSetFilterNdDescriptor(*desc, dataType::type, + shape.size(), shape.data())); +#endif +} + template inline void createConvolutionDesc(cudnnConvolutionDescriptor_t* conv) { CUDNN_CHECK(cudnnCreateConvolutionDescriptor(conv)); @@ -113,6 +162,31 @@ inline void setConvolutionDesc(cudnnConvolutionDescriptor_t* conv, pad_h, pad_w, stride_h, stride_w, 1, 1, CUDNN_CROSS_CORRELATION)); } +template +inline void setNdConvolutionDesc(cudnnConvolutionDescriptor_t* conv, + cudnnTensorDescriptor_t bottom, cudnnFilterDescriptor_t filter, + std::vector pad, std::vector stride) { + int nbDims; + std::vector shape(pad.size() + 2); + cudnnDataType_t cudnn_type; +#if CUDNN_VERSION_MIN(5, 0, 0) + cudnnTensorFormat_t tensor_format; + cudnnGetFilterNdDescriptor(filter, + shape.size(), &cudnn_type, &tensor_format, &nbDims, shape.data()); +#else + cudnnGetFilterNdDescriptor(filter, + shape.size(), &cudnn_type, &nbDims, shape.data()); +#endif + CHECK_EQ(nbDims, pad.size() + 2) + << "Dimensions of filters and pad don't match !"; + CHECK_EQ(nbDims, stride.size() + 2) + << "Dimensions of filters and stride don't match !"; + std::vector upscale(pad.size(), 1); + CUDNN_CHECK(cudnnSetConvolutionNdDescriptor(*conv, + pad.size(), pad.data(), stride.data(), upscale.data(), + CUDNN_CROSS_CORRELATION, cudnn_type)); +} + template inline void createPoolingDesc(cudnnPoolingDescriptor_t* pool_desc, PoolingParameter_PoolMethod poolmethod, cudnnPoolingMode_t* mode, @@ -130,10 +204,10 @@ inline void createPoolingDesc(cudnnPoolingDescriptor_t* pool_desc, CUDNN_CHECK(cudnnCreatePoolingDescriptor(pool_desc)); #if CUDNN_VERSION_MIN(5, 0, 0) CUDNN_CHECK(cudnnSetPooling2dDescriptor(*pool_desc, *mode, - CUDNN_PROPAGATE_NAN, h, w, pad_h, pad_w, stride_h, stride_w)); + CUDNN_PROPAGATE_NAN, h, w, pad_h, pad_w, stride_h, stride_w)); #else - CUDNN_CHECK(cudnnSetPooling2dDescriptor_v4(*pool_desc, *mode, - CUDNN_PROPAGATE_NAN, h, w, pad_h, pad_w, stride_h, stride_w)); + CUDNN_CHECK(cudnnSetPooling2dDescriptor(*pool_desc, *mode, h, w, + pad_h, pad_w, stride_h, stride_w)); #endif } @@ -145,6 +219,36 @@ inline void createActivationDescriptor(cudnnActivationDescriptor_t* activ_desc, CUDNN_PROPAGATE_NAN, Dtype(0))); } +template +inline void createNdPoolingDesc(cudnnPoolingDescriptor_t* pool_desc, + PoolingParameter_PoolMethod poolmethod, cudnnPoolingMode_t* mode, + std::vector shape, std::vector pad, + std::vector stride) { + CHECK_EQ(shape.size(), pad.size()) + << "Dimensions of shape and pad don't match !"; + CHECK_EQ(shape.size(), stride.size()) + << "Dimensions of shape and stride don't match !"; + switch (poolmethod) { + case PoolingParameter_PoolMethod_MAX: + *mode = CUDNN_POOLING_MAX; + break; + case PoolingParameter_PoolMethod_AVE: + *mode = CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING; + break; + default: + LOG(FATAL) << "Unknown pooling method."; + } + CUDNN_CHECK(cudnnCreatePoolingDescriptor(pool_desc)); +#if CUDNN_VERSION_MIN(5, 0, 0) + CUDNN_CHECK(cudnnSetPoolingNdDescriptor(*pool_desc, *mode, + CUDNN_PROPAGATE_NAN, shape.size(), shape.data(), pad.data(), + stride.data())); +#else + CUDNN_CHECK(cudnnSetPoolingNdDescriptor(*pool_desc, *mode, shape.size(), + shape.data(), pad.data(), stride.data())); +#endif +} + } // namespace cudnn } // namespace caffe diff --git a/src/caffe/layers/cudnn_conv_layer.cpp b/src/caffe/layers/cudnn_conv_layer.cpp index 1987fb096b0..b32332ad8b8 100644 --- a/src/caffe/layers/cudnn_conv_layer.cpp +++ b/src/caffe/layers/cudnn_conv_layer.cpp @@ -59,20 +59,21 @@ void CuDNNConvolutionLayer::LayerSetUp( 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->num_output_ / this->group_, this->channels_ / this->group_, - kernel_h, kernel_w); + std::vector kernel_shape; + kernel_shape.push_back(this->num_output_ / this->group_); + kernel_shape.push_back(this->channels_ / this->group_); + for (unsigned int i = 0; i < this->num_spatial_axes_; ++i) + kernel_shape.push_back(this->kernel_shape_.cpu_data()[i]); + + cudnn::createNdFilterDesc(&filter_desc_, kernel_shape); // 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); + cudnn::createTensorDesc(&bottom_desc); bottom_descs_.push_back(bottom_desc); cudnnTensorDescriptor_t top_desc; - cudnn::createTensor4dDesc(&top_desc); + cudnn::createTensorDesc(&top_desc); top_descs_.push_back(top_desc); cudnnConvolutionDescriptor_t conv_desc; cudnn::createConvolutionDesc(&conv_desc); @@ -81,7 +82,7 @@ void CuDNNConvolutionLayer::LayerSetUp( // Tensor descriptor for bias. if (this->bias_term_) { - cudnn::createTensor4dDesc(&bias_desc_); + cudnn::createTensorDesc(&bias_desc_); } handles_setup_ = true; @@ -91,41 +92,42 @@ template void CuDNNConvolutionLayer::Reshape( const vector*>& bottom, const vector*>& top) { ConvolutionLayer::Reshape(bottom, top); - CHECK_EQ(2, this->num_spatial_axes_) - << "CuDNNConvolution 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]; + + std::vector bottom_tensor_shape(bottom[0]->shape()); + bottom_tensor_shape[1] /= this->group_; + std::vector bottom_tensor_stride(bottom[0]->shape().size(), 1); + for (int i = bottom[0]->shape().size() - 2; i >= 0; --i) { + bottom_tensor_stride[i] = + bottom[0]->shape(i + 1) * bottom_tensor_stride[i + 1]; + } + + std::vector top_tensor_shape(top[0]->shape()); + top_tensor_shape[1] /= this->group_; + std::vector top_tensor_stride(top[0]->shape().size(), 1); + for (int i = top[0]->shape().size() - 2; i >= 0; --i) { + top_tensor_stride[i] = top[0]->shape(i + 1) * top_tensor_stride[i + 1]; + } + + std::vector pad, stride; + for (unsigned int i = 0; i < this->num_spatial_axes_; ++i) { + pad.push_back(this->pad_.cpu_data()[i]); + stride.push_back(this->stride_.cpu_data()[i]); + } // 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_ * this->out_spatial_dim_, - this->out_spatial_dim_, width_out, 1); - cudnn::setConvolutionDesc(&conv_descs_[i], bottom_descs_[i], - filter_desc_, pad_h, pad_w, - stride_h, stride_w); + cudnn::setTensorNdDesc(&bottom_descs_[i], + bottom_tensor_shape, bottom_tensor_stride); + cudnn::setTensorNdDesc(&top_descs_[i], + top_tensor_shape, top_tensor_stride); + cudnn::setNdConvolutionDesc(&conv_descs_[i], bottom_descs_[i], + filter_desc_, pad, stride); // choose forward and backward algorithms + workspace(s) CUDNN_CHECK(cudnnGetConvolutionForwardAlgorithm(handle_[0], @@ -226,8 +228,9 @@ void CuDNNConvolutionLayer::Reshape( // Tensor descriptor for bias. if (this->bias_term_) { - cudnn::setTensor4dDesc(&bias_desc_, - 1, this->num_output_ / this->group_, 1, 1); + vector bias_shape(bottom[0]->shape().size(), 1); + bias_shape[1] = this->num_output_ / this->group_; + cudnn::setTensorNdDesc(&bias_desc_, bias_shape); } } diff --git a/src/caffe/layers/cudnn_pooling_layer.cpp b/src/caffe/layers/cudnn_pooling_layer.cpp index 24f14780b4f..a7d5ca4d53c 100644 --- a/src/caffe/layers/cudnn_pooling_layer.cpp +++ b/src/caffe/layers/cudnn_pooling_layer.cpp @@ -10,12 +10,11 @@ void CuDNNPoolingLayer::LayerSetUp(const vector*>& bottom, const vector*>& top) { PoolingLayer::LayerSetUp(bottom, top); CUDNN_CHECK(cudnnCreate(&handle_)); - cudnn::createTensor4dDesc(&bottom_desc_); - cudnn::createTensor4dDesc(&top_desc_); - cudnn::createPoolingDesc(&pooling_desc_, + cudnn::createTensorDesc(&bottom_desc_); + cudnn::createTensorDesc(&top_desc_); + cudnn::createNdPoolingDesc(&pooling_desc_, this->layer_param_.pooling_param().pool(), &mode_, - this->kernel_h_, this->kernel_w_, this->pad_h_, this->pad_w_, - this->stride_h_, this->stride_w_); + this->kernel_shape_, this->pad_, this->stride_); handles_setup_ = true; } @@ -23,10 +22,8 @@ template void CuDNNPoolingLayer::Reshape(const vector*>& bottom, const vector*>& top) { PoolingLayer::Reshape(bottom, top); - cudnn::setTensor4dDesc(&bottom_desc_, bottom[0]->num(), - this->channels_, this->height_, this->width_); - cudnn::setTensor4dDesc(&top_desc_, bottom[0]->num(), - this->channels_, this->pooled_height_, this->pooled_width_); + cudnn::setTensorNdDesc(&bottom_desc_, this->input_shape_); + cudnn::setTensorNdDesc(&top_desc_, this->pooled_shape_); } template diff --git a/src/caffe/layers/cudnn_relu_layer.cpp b/src/caffe/layers/cudnn_relu_layer.cpp index 795e0a9efb0..3c4eb3171c7 100644 --- a/src/caffe/layers/cudnn_relu_layer.cpp +++ b/src/caffe/layers/cudnn_relu_layer.cpp @@ -11,8 +11,8 @@ void CuDNNReLULayer::LayerSetUp(const vector*>& bottom, ReLULayer::LayerSetUp(bottom, top); // initialize cuDNN CUDNN_CHECK(cudnnCreate(&handle_)); - cudnn::createTensor4dDesc(&bottom_desc_); - cudnn::createTensor4dDesc(&top_desc_); + cudnn::createTensorDesc(&bottom_desc_); + cudnn::createTensorDesc(&top_desc_); cudnn::createActivationDescriptor(&activ_desc_, CUDNN_ACTIVATION_RELU); handles_setup_ = true; } @@ -21,12 +21,8 @@ template void CuDNNReLULayer::Reshape(const vector*>& bottom, const vector*>& top) { ReLULayer::Reshape(bottom, top); - const int N = bottom[0]->num(); - const int K = bottom[0]->channels(); - const int H = bottom[0]->height(); - const int W = bottom[0]->width(); - cudnn::setTensor4dDesc(&bottom_desc_, N, K, H, W); - cudnn::setTensor4dDesc(&top_desc_, N, K, H, W); + cudnn::setTensorNdDesc(&bottom_desc_, bottom[0]->shape()); + cudnn::setTensorNdDesc(&top_desc_, bottom[0]->shape()); } template diff --git a/src/caffe/layers/cudnn_sigmoid_layer.cpp b/src/caffe/layers/cudnn_sigmoid_layer.cpp index 3ce6aef1764..90f1ac6e0a3 100644 --- a/src/caffe/layers/cudnn_sigmoid_layer.cpp +++ b/src/caffe/layers/cudnn_sigmoid_layer.cpp @@ -11,8 +11,8 @@ void CuDNNSigmoidLayer::LayerSetUp(const vector*>& bottom, SigmoidLayer::LayerSetUp(bottom, top); // initialize cuDNN CUDNN_CHECK(cudnnCreate(&handle_)); - cudnn::createTensor4dDesc(&bottom_desc_); - cudnn::createTensor4dDesc(&top_desc_); + cudnn::createTensorDesc(&bottom_desc_); + cudnn::createTensorDesc(&top_desc_); cudnn::createActivationDescriptor(&activ_desc_, CUDNN_ACTIVATION_SIGMOID); handles_setup_ = true; @@ -22,12 +22,8 @@ template void CuDNNSigmoidLayer::Reshape(const vector*>& bottom, const vector*>& top) { SigmoidLayer::Reshape(bottom, top); - const int N = bottom[0]->num(); - const int K = bottom[0]->channels(); - const int H = bottom[0]->height(); - const int W = bottom[0]->width(); - cudnn::setTensor4dDesc(&bottom_desc_, N, K, H, W); - cudnn::setTensor4dDesc(&top_desc_, N, K, H, W); + cudnn::setTensorNdDesc(&bottom_desc_, bottom[0]->shape()); + cudnn::setTensorNdDesc(&top_desc_, bottom[0]->shape()); } template diff --git a/src/caffe/layers/cudnn_tanh_layer.cpp b/src/caffe/layers/cudnn_tanh_layer.cpp index e87dd9de0ab..c18d9b31c87 100644 --- a/src/caffe/layers/cudnn_tanh_layer.cpp +++ b/src/caffe/layers/cudnn_tanh_layer.cpp @@ -11,8 +11,8 @@ void CuDNNTanHLayer::LayerSetUp(const vector*>& bottom, TanHLayer::LayerSetUp(bottom, top); // initialize cuDNN CUDNN_CHECK(cudnnCreate(&handle_)); - cudnn::createTensor4dDesc(&bottom_desc_); - cudnn::createTensor4dDesc(&top_desc_); + cudnn::createTensorDesc(&bottom_desc_); + cudnn::createTensorDesc(&top_desc_); cudnn::createActivationDescriptor(&activ_desc_, CUDNN_ACTIVATION_TANH); handles_setup_ = true; } @@ -21,12 +21,8 @@ template void CuDNNTanHLayer::Reshape(const vector*>& bottom, const vector*>& top) { TanHLayer::Reshape(bottom, top); - const int N = bottom[0]->num(); - const int K = bottom[0]->channels(); - const int H = bottom[0]->height(); - const int W = bottom[0]->width(); - cudnn::setTensor4dDesc(&bottom_desc_, N, K, H, W); - cudnn::setTensor4dDesc(&top_desc_, N, K, H, W); + cudnn::setTensorNdDesc(&bottom_desc_, bottom[0]->shape()); + cudnn::setTensorNdDesc(&top_desc_, bottom[0]->shape()); } template diff --git a/src/caffe/layers/euclidean_loss_layer.cpp b/src/caffe/layers/euclidean_loss_layer.cpp index 300d991e765..c7dc52bbf97 100644 --- a/src/caffe/layers/euclidean_loss_layer.cpp +++ b/src/caffe/layers/euclidean_loss_layer.cpp @@ -24,7 +24,7 @@ void EuclideanLossLayer::Forward_cpu(const vector*>& bottom, bottom[1]->cpu_data(), diff_.mutable_cpu_data()); Dtype dot = caffe_cpu_dot(count, diff_.cpu_data(), diff_.cpu_data()); - Dtype loss = dot / bottom[0]->num() / Dtype(2); + Dtype loss = dot / bottom[0]->shape(0) / Dtype(2); top[0]->mutable_cpu_data()[0] = loss; } @@ -34,7 +34,7 @@ void EuclideanLossLayer::Backward_cpu(const vector*>& top, for (int i = 0; i < 2; ++i) { if (propagate_down[i]) { const Dtype sign = (i == 0) ? 1 : -1; - const Dtype alpha = sign * top[0]->cpu_diff()[0] / bottom[i]->num(); + const Dtype alpha = sign * top[0]->cpu_diff()[0] / bottom[i]->shape(0); caffe_cpu_axpby( bottom[i]->count(), // count alpha, // alpha diff --git a/src/caffe/layers/euclidean_loss_layer.cu b/src/caffe/layers/euclidean_loss_layer.cu index 4c221b64faf..7538324973b 100644 --- a/src/caffe/layers/euclidean_loss_layer.cu +++ b/src/caffe/layers/euclidean_loss_layer.cu @@ -16,7 +16,7 @@ void EuclideanLossLayer::Forward_gpu(const vector*>& bottom, diff_.mutable_gpu_data()); Dtype dot; caffe_gpu_dot(count, diff_.gpu_data(), diff_.gpu_data(), &dot); - Dtype loss = dot / bottom[0]->num() / Dtype(2); + Dtype loss = dot / bottom[0]->shape(0) / Dtype(2); top[0]->mutable_cpu_data()[0] = loss; } @@ -26,7 +26,7 @@ void EuclideanLossLayer::Backward_gpu(const vector*>& top, for (int i = 0; i < 2; ++i) { if (propagate_down[i]) { const Dtype sign = (i == 0) ? 1 : -1; - const Dtype alpha = sign * top[0]->cpu_diff()[0] / bottom[i]->num(); + const Dtype alpha = sign * top[0]->cpu_diff()[0] / bottom[i]->shape(0); caffe_gpu_axpby( bottom[i]->count(), // count alpha, // alpha diff --git a/src/caffe/layers/loss_layer.cpp b/src/caffe/layers/loss_layer.cpp index c0b7a862181..afb1ce94893 100644 --- a/src/caffe/layers/loss_layer.cpp +++ b/src/caffe/layers/loss_layer.cpp @@ -16,8 +16,8 @@ void LossLayer::LayerSetUp( template void LossLayer::Reshape( const vector*>& bottom, const vector*>& top) { - CHECK_EQ(bottom[0]->num(), bottom[1]->num()) - << "The data and label should have the same number."; + CHECK_EQ(bottom[0]->shape(0), bottom[1]->shape(0)) + << "The data and label should have the same first dimension."; vector loss_shape(0); // Loss layers output a scalar; 0 axes. top[0]->Reshape(loss_shape); } diff --git a/src/caffe/layers/lrn_layer.cpp b/src/caffe/layers/lrn_layer.cpp index 210525e20f3..8fcfdf9f502 100644 --- a/src/caffe/layers/lrn_layer.cpp +++ b/src/caffe/layers/lrn_layer.cpp @@ -38,8 +38,8 @@ void LRNLayer::LayerSetUp(const vector*>& bottom, LayerParameter pool_param; pool_param.mutable_pooling_param()->set_pool( PoolingParameter_PoolMethod_AVE); - pool_param.mutable_pooling_param()->set_pad(pre_pad_); - pool_param.mutable_pooling_param()->set_kernel_size(size_); + pool_param.mutable_pooling_param()->add_pad(pre_pad_); + pool_param.mutable_pooling_param()->add_kernel_size(size_); pool_layer_.reset(new PoolingLayer(pool_param)); pool_layer_->SetUp(square_top_vec_, pool_top_vec_); // Set up power_layer_ to compute (1 + alpha_/N^2 s)^-beta_, where s is diff --git a/src/caffe/layers/pooling_layer.cpp b/src/caffe/layers/pooling_layer.cpp index 90897db0f45..58b13f7d824 100644 --- a/src/caffe/layers/pooling_layer.cpp +++ b/src/caffe/layers/pooling_layer.cpp @@ -14,111 +14,147 @@ template void PoolingLayer::LayerSetUp(const vector*>& bottom, const vector*>& top) { PoolingParameter pool_param = this->layer_param_.pooling_param(); - if (pool_param.global_pooling()) { - CHECK(!(pool_param.has_kernel_size() || - pool_param.has_kernel_h() || pool_param.has_kernel_w())) - << "With Global_pooling: true Filter size cannot specified"; - } else { - CHECK(!pool_param.has_kernel_size() != - !(pool_param.has_kernel_h() && pool_param.has_kernel_w())) - << "Filter size is kernel_size OR kernel_h and kernel_w; not both"; - CHECK(pool_param.has_kernel_size() || - (pool_param.has_kernel_h() && pool_param.has_kernel_w())) - << "For non-square filters both kernel_h and kernel_w are required."; - } - CHECK((!pool_param.has_pad() && pool_param.has_pad_h() - && pool_param.has_pad_w()) - || (!pool_param.has_pad_h() && !pool_param.has_pad_w())) - << "pad is pad OR pad_h and pad_w are required."; - CHECK((!pool_param.has_stride() && pool_param.has_stride_h() - && pool_param.has_stride_w()) - || (!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(); + num_spatial_axes_ = bottom[0]->num_axes() - 2; + // Setup filter kernel dimensions (kernel_shape_). + kernel_shape_ = std::vector(num_spatial_axes_, 0); if (global_pooling_) { - kernel_h_ = bottom[0]->height(); - kernel_w_ = bottom[0]->width(); + CHECK(!((pool_param.kernel_size_size() > 0) || + pool_param.has_kernel_h() || pool_param.has_kernel_w())) + << "With Global_pooling: true Filter size cannot specified"; + for (int i = 0; i < num_spatial_axes_; ++i) + kernel_shape_[i] = bottom[0]->shape(i + 2); } else { - if (pool_param.has_kernel_size()) { - kernel_h_ = kernel_w_ = pool_param.kernel_size(); + if (pool_param.has_kernel_h() || pool_param.has_kernel_w()) { + CHECK_EQ(num_spatial_axes_, 2) + << "kernel_h & kernel_w can only be used for 2D pooling."; + CHECK_EQ(0, pool_param.kernel_size_size()) + << "Either kernel_size or kernel_h/w should be specified; not both."; + kernel_shape_[0] = pool_param.kernel_h(); + kernel_shape_[1] = pool_param.kernel_w(); } else { - kernel_h_ = pool_param.kernel_h(); - kernel_w_ = pool_param.kernel_w(); + const int num_kernel_dims = pool_param.kernel_size_size(); + CHECK(num_kernel_dims == 1 || num_kernel_dims == num_spatial_axes_) + << "kernel_size must be specified once, or once per spatial " + << "dimension (kernel_size specified " << num_kernel_dims + << " times; " << num_spatial_axes_ << " spatial dims)."; + for (int i = 0; i < num_spatial_axes_; ++i) { + kernel_shape_[i] = + pool_param.kernel_size((num_kernel_dims == 1) ? 0 : i); + } + } + for (int i = 0; i < num_spatial_axes_; ++i) { + CHECK_GT(kernel_shape_[i], 0) << "Filter dimensions must be nonzero."; } } - CHECK_GT(kernel_h_, 0) << "Filter dimensions cannot be zero."; - CHECK_GT(kernel_w_, 0) << "Filter dimensions cannot be zero."; - if (!pool_param.has_pad_h()) { - pad_h_ = pad_w_ = pool_param.pad(); + // Setup stride dimensions (stride_). + stride_ = std::vector(num_spatial_axes_, 0); + if (pool_param.has_stride_h() || pool_param.has_stride_w()) { + CHECK_EQ(num_spatial_axes_, 2) + << "stride_h & stride_w can only be used for 2D pooling."; + CHECK_EQ(0, pool_param.stride_size()) + << "Either stride or stride_h/w should be specified; not both."; + stride_[0] = pool_param.stride_h(); + stride_[1] = pool_param.stride_w(); } else { - pad_h_ = pool_param.pad_h(); - pad_w_ = pool_param.pad_w(); + const int num_stride_dims = pool_param.stride_size(); + CHECK(num_stride_dims == 0 || num_stride_dims == 1 || + num_stride_dims == num_spatial_axes_) + << "stride must be specified once, or once per spatial dimension " + << "(stride specified " << num_stride_dims << " times; " + << num_spatial_axes_ << " spatial dims)."; + const int kDefaultStride = 1; + for (int i = 0; i < num_spatial_axes_; ++i) { + stride_[i] = (num_stride_dims == 0) ? kDefaultStride : + pool_param.stride((num_stride_dims == 1) ? 0 : i); + CHECK_GT(stride_[i], 0) << "Stride dimensions must be nonzero."; + } } - if (!pool_param.has_stride_h()) { - stride_h_ = stride_w_ = pool_param.stride(); + // Setup pad dimensions (pad_). + pad_ = std::vector(num_spatial_axes_, 0); + if (pool_param.has_pad_h() || pool_param.has_pad_w()) { + CHECK_EQ(num_spatial_axes_, 2) + << "pad_h & pad_w can only be used for 2D pooling."; + CHECK_EQ(0, pool_param.pad_size()) + << "Either pad or pad_h/w should be specified; not both."; + pad_[0] = pool_param.pad_h(); + pad_[1] = pool_param.pad_w(); } else { - stride_h_ = pool_param.stride_h(); - stride_w_ = pool_param.stride_w(); - } - if (global_pooling_) { - CHECK(pad_h_ == 0 && pad_w_ == 0 && stride_h_ == 1 && stride_w_ == 1) - << "With Global_pooling: true; only pad = 0 and stride = 1"; + const int num_pad_dims = pool_param.pad_size(); + CHECK(num_pad_dims == 0 || num_pad_dims == 1 || + num_pad_dims == num_spatial_axes_) + << "pad must be specified once, or once per spatial dimension " + << "(pad specified " << num_pad_dims << " times; " + << num_spatial_axes_ << " spatial dims)."; + const int kDefaultPad = 0; + for (int i = 0; i < num_spatial_axes_; ++i) { + pad_[i] = (num_pad_dims == 0) ? kDefaultPad : + pool_param.pad((num_pad_dims == 1) ? 0 : i); + } } - if (pad_h_ != 0 || pad_w_ != 0) { - CHECK(this->layer_param_.pooling_param().pool() - == PoolingParameter_PoolMethod_AVE - || this->layer_param_.pooling_param().pool() - == PoolingParameter_PoolMethod_MAX) - << "Padding implemented only for average and max pooling."; - CHECK_LT(pad_h_, kernel_h_); - CHECK_LT(pad_w_, kernel_w_); + // remaining pooling sanity checks + for (int i = 0; i < num_spatial_axes_; ++i) { + if (global_pooling_) { + CHECK(pad_[i] == 0 && stride_[i] == 1) + << "With Global_pooling: true; only pad = 0 and stride = 1"; + } + if (pad_[i] != 0) { + CHECK(this->layer_param_.pooling_param().pool() + == PoolingParameter_PoolMethod_AVE + || this->layer_param_.pooling_param().pool() + == PoolingParameter_PoolMethod_MAX) + << "Padding implemented only for average and max pooling."; + } + CHECK_LT(pad_[i], kernel_shape_[i]); } } template void PoolingLayer::Reshape(const vector*>& bottom, const vector*>& top) { - CHECK_EQ(4, bottom[0]->num_axes()) << "Input must have 4 axes, " - << "corresponding to (num, channels, height, width)"; - channels_ = bottom[0]->channels(); - height_ = bottom[0]->height(); - width_ = bottom[0]->width(); + CHECK_EQ(bottom[0]->num_axes() - 2, num_spatial_axes_) + << "bottom num_axes may not change."; + channels_ = bottom[0]->shape(1); + input_shape_ = bottom[0]->shape(); if (global_pooling_) { - kernel_h_ = bottom[0]->height(); - kernel_w_ = bottom[0]->width(); + for (int i = 0; i < num_spatial_axes_; ++i) + kernel_shape_[i] = input_shape_[i + 2]; } - 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; - 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. - if ((pooled_height_ - 1) * stride_h_ >= height_ + pad_h_) { - --pooled_height_; - } - if ((pooled_width_ - 1) * stride_w_ >= width_ + pad_w_) { - --pooled_width_; + // setup pooled shape + pooled_shape_ = std::vector(input_shape_.size()); + pooled_shape_[0] = input_shape_[0]; + pooled_shape_[1] = input_shape_[1]; + for (unsigned int i = 0; i < num_spatial_axes_; ++i) { + pooled_shape_[i + 2] = static_cast(std::ceil(static_cast( + input_shape_[i + 2] + 2 * pad_[i] - kernel_shape_[i]) / + stride_[i])) + 1; + } + for (unsigned int i = 0; i < num_spatial_axes_; ++i) { + if (pad_[i]) { + // If we have padding, ensure that the last pooling starts strictly + // inside the image (instead of at the padding); otherwise clip the last. + if ((pooled_shape_[i + 2] - 1) * stride_[i] >= + input_shape_[i + 2] + pad_[i]) { + --pooled_shape_[i + 2]; + } + CHECK_LT((pooled_shape_[i + 2] - 1) * stride_[i], + input_shape_[i + 2] + pad_[i]); } - CHECK_LT((pooled_height_ - 1) * stride_h_, height_ + pad_h_); - CHECK_LT((pooled_width_ - 1) * stride_w_, width_ + pad_w_); } - top[0]->Reshape(bottom[0]->num(), channels_, pooled_height_, - pooled_width_); + // reshape outputs + top[0]->Reshape(pooled_shape_); if (top.size() > 1) { top[1]->ReshapeLike(*top[0]); } // If max pooling, we will initialize the vector index part. if (this->layer_param_.pooling_param().pool() == PoolingParameter_PoolMethod_MAX && top.size() == 1) { - max_idx_.Reshape(bottom[0]->num(), channels_, pooled_height_, - pooled_width_); + max_idx_.Reshape(pooled_shape_); } // If stochastic pooling, we will initialize the random index part. if (this->layer_param_.pooling_param().pool() == PoolingParameter_PoolMethod_STOCHASTIC) { - rand_idx_.Reshape(bottom[0]->num(), channels_, pooled_height_, - pooled_width_); + rand_idx_.Reshape(pooled_shape_); } } @@ -134,6 +170,16 @@ void PoolingLayer::Forward_cpu(const vector*>& bottom, const bool use_top_mask = top.size() > 1; int* mask = NULL; // suppress warnings about uninitalized variables Dtype* top_mask = NULL; + const int height_ = input_shape_[2]; + const int width_ = input_shape_[3]; + const int pooled_height_ = pooled_shape_[2]; + const int pooled_width_ = pooled_shape_[3]; + const int kernel_h_ = kernel_shape_[0]; + const int kernel_w_ = kernel_shape_[1]; + const int pad_h_ = pad_[0]; + const int pad_w_ = pad_[1]; + const int stride_h_ = stride_[0]; + const int stride_w_ = stride_[1]; // Different pooling methods. We explicitly do the switch outside the for // loop to save time, although this results in more code. switch (this->layer_param_.pooling_param().pool()) { @@ -241,6 +287,16 @@ void PoolingLayer::Backward_cpu(const vector*>& top, const bool use_top_mask = top.size() > 1; const int* mask = NULL; // suppress warnings about uninitialized variables const Dtype* top_mask = NULL; + const int height_ = input_shape_[2]; + const int width_ = input_shape_[3]; + const int pooled_height_ = pooled_shape_[2]; + const int pooled_width_ = pooled_shape_[3]; + const int kernel_h_ = kernel_shape_[0]; + const int kernel_w_ = kernel_shape_[1]; + const int pad_h_ = pad_[0]; + const int pad_w_ = pad_[1]; + const int stride_h_ = stride_[0]; + const int stride_w_ = stride_[1]; switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: // The main loop diff --git a/src/caffe/layers/pooling_layer.cu b/src/caffe/layers/pooling_layer.cu index 1ea46cc81b1..8d616e0f411 100644 --- a/src/caffe/layers/pooling_layer.cu +++ b/src/caffe/layers/pooling_layer.cu @@ -164,6 +164,16 @@ void PoolingLayer::Forward_gpu(const vector*>& bottom, const bool use_top_mask = top.size() > 1; int* mask = NULL; Dtype* top_mask = NULL; + const int height_ = input_shape_[2]; + const int width_ = input_shape_[3]; + const int pooled_height_ = pooled_shape_[2]; + const int pooled_width_ = pooled_shape_[3]; + const int kernel_h_ = kernel_shape_[0]; + const int kernel_w_ = kernel_shape_[1]; + const int pad_h_ = pad_[0]; + const int pad_w_ = pad_[1]; + const int stride_h_ = stride_[0]; + const int stride_w_ = stride_[1]; switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: if (use_top_mask) { @@ -344,6 +354,16 @@ void PoolingLayer::Backward_gpu(const vector*>& top, const bool use_top_mask = top.size() > 1; const int* mask = NULL; const Dtype* top_mask = NULL; + const int height_ = input_shape_[2]; + const int width_ = input_shape_[3]; + const int pooled_height_ = pooled_shape_[2]; + const int pooled_width_ = pooled_shape_[3]; + const int kernel_h_ = kernel_shape_[0]; + const int kernel_w_ = kernel_shape_[1]; + const int pad_h_ = pad_[0]; + const int pad_w_ = pad_[1]; + const int stride_h_ = stride_[0]; + const int stride_w_ = stride_[1]; switch (this->layer_param_.pooling_param().pool()) { case PoolingParameter_PoolMethod_MAX: if (use_top_mask) { diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 1556781cbc2..0df91704e52 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -887,24 +887,28 @@ message PoolingParameter { } optional PoolMethod pool = 1 [default = MAX]; // The pooling method // Pad, kernel size, and stride are all given as a single value for equal - // dimensions in height and width or as Y, X pairs. - optional uint32 pad = 4 [default = 0]; // The padding size (equal in Y, X) - optional uint32 pad_h = 9 [default = 0]; // The padding height - optional uint32 pad_w = 10 [default = 0]; // The padding width - optional uint32 kernel_size = 2; // The kernel size (square) - optional uint32 kernel_h = 5; // The kernel height - optional uint32 kernel_w = 6; // The kernel width - optional uint32 stride = 3 [default = 1]; // The stride (equal in Y, X) - optional uint32 stride_h = 7; // The stride height - optional uint32 stride_w = 8; // The stride width + // dimensions in all spatial dimensions, or once per spatial dimension. + repeated uint32 pad = 4; // The padding size; defaults to 0 + repeated uint32 kernel_size = 2; // The kernel size + repeated uint32 stride = 3; // The stride; defaults to 1 + + // For 2D padding only, the *_h and *_w versions may also be used to + // specify both spatial dimensions. + optional uint32 pad_h = 9 [default = 0]; // The padding height (2D only) + optional uint32 pad_w = 10 [default = 0]; // The padding width (2D only) + optional uint32 kernel_h = 5; // The kernel height (2D only) + optional uint32 kernel_w = 6; // The kernel width (2D only) + optional uint32 stride_h = 7; // The stride height (2D only) + optional uint32 stride_w = 8; // The stride width (2D only) + enum Engine { DEFAULT = 0; CAFFE = 1; CUDNN = 2; } optional Engine engine = 11 [default = DEFAULT]; - // If global_pooling then it will pool over the size of the bottom by doing - // kernel_h = bottom->height and kernel_w = bottom->width + // If global_pooling then it will pool over the size of the bottom by setting + // the size of the kernel to the size of the input image optional bool global_pooling = 12 [default = false]; } diff --git a/src/caffe/test/test_convolution_layer.cpp b/src/caffe/test/test_convolution_layer.cpp index 9bb19d13592..93f090903d1 100644 --- a/src/caffe/test/test_convolution_layer.cpp +++ b/src/caffe/test/test_convolution_layer.cpp @@ -1081,6 +1081,80 @@ TYPED_TEST(CuDNNConvolutionLayerTest, TestGradientGroupCuDNN) { this->blob_top_vec_); } +TYPED_TEST(CuDNNConvolutionLayerTest, TestSimple3DConvolutionCuDNN) { + typedef TypeParam Dtype; + this->blob_bottom_vec_.push_back(this->blob_bottom_2_); + this->blob_top_vec_.push_back(this->blob_top_2_); + vector bottom_shape(5); + bottom_shape[0] = this->blob_bottom_vec_[0]->shape(0); + bottom_shape[1] = this->blob_bottom_vec_[0]->shape(1); + bottom_shape[2] = 5; + bottom_shape[3] = this->blob_bottom_vec_[0]->shape(2); + bottom_shape[4] = this->blob_bottom_vec_[0]->shape(3); + 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->add_kernel_size(3); + convolution_param->add_stride(2); + convolution_param->set_num_output(4); + convolution_param->mutable_weight_filler()->set_type("gaussian"); + convolution_param->mutable_bias_filler()->set_type("gaussian"); + CuDNNConvolutionLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + // Check against reference convolution. + const Dtype* top_data; + const Dtype* ref_top_data; + caffe_conv(this->blob_bottom_, convolution_param, layer.blobs(), + this->MakeReferenceTop(this->blob_top_)); + top_data = this->blob_top_->cpu_data(); + ref_top_data = this->ref_blob_top_->cpu_data(); + for (int i = 0; i < this->blob_top_->count(); ++i) { + EXPECT_NEAR(top_data[i], ref_top_data[i], 1e-4); + } + caffe_conv(this->blob_bottom_2_, convolution_param, layer.blobs(), + this->MakeReferenceTop(this->blob_top_2_)); + top_data = this->blob_top_2_->cpu_data(); + ref_top_data = this->ref_blob_top_->cpu_data(); + for (int i = 0; i < this->blob_top_->count(); ++i) { + EXPECT_NEAR(top_data[i], ref_top_data[i], 1e-4); + } +} + +TYPED_TEST(CuDNNConvolutionLayerTest, TestGradient3DCuDNN) { + typedef TypeParam Dtype; + LayerParameter layer_param; + ConvolutionParameter* convolution_param = + layer_param.mutable_convolution_param(); + vector bottom_shape(5); + bottom_shape[0] = this->blob_bottom_vec_[0]->shape(0); + bottom_shape[1] = this->blob_bottom_vec_[0]->shape(1); + bottom_shape[2] = 5; + bottom_shape[3] = this->blob_bottom_vec_[0]->shape(2); + bottom_shape[4] = this->blob_bottom_vec_[0]->shape(3); + 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]); + } + convolution_param->add_kernel_size(3); + convolution_param->add_stride(2); + convolution_param->set_num_output(2); + convolution_param->mutable_weight_filler()->set_type("gaussian"); + convolution_param->mutable_bias_filler()->set_type("gaussian"); + CuDNNConvolutionLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + #endif } // namespace caffe diff --git a/src/caffe/test/test_maxpool_dropout_layers.cpp b/src/caffe/test/test_maxpool_dropout_layers.cpp index 4f0e20ac3a7..0fd5d33409a 100644 --- a/src/caffe/test/test_maxpool_dropout_layers.cpp +++ b/src/caffe/test/test_maxpool_dropout_layers.cpp @@ -44,8 +44,8 @@ TYPED_TEST(MaxPoolingDropoutTest, TestSetup) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); PoolingLayer max_layer(layer_param); max_layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); DropoutLayer dropout_layer(layer_param); @@ -61,8 +61,8 @@ TYPED_TEST(MaxPoolingDropoutTest, TestForward) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); @@ -91,8 +91,8 @@ TYPED_TEST(MaxPoolingDropoutTest, TestBackward) { LayerParameter layer_param; layer_param.set_phase(TRAIN); PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); diff --git a/src/caffe/test/test_pooling_layer.cpp b/src/caffe/test/test_pooling_layer.cpp index bb95cae032d..0f6c970c1f1 100644 --- a/src/caffe/test/test_pooling_layer.cpp +++ b/src/caffe/test/test_pooling_layer.cpp @@ -49,7 +49,7 @@ class PoolingLayerTest : public MultiDeviceTest { void TestForwardSquare() { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(2); + pooling_param->add_kernel_size(2); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); const int num = 2; const int channels = 2; @@ -377,8 +377,8 @@ TYPED_TEST(PoolingLayerTest, TestSetup) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); EXPECT_EQ(this->blob_top_->num(), this->blob_bottom_->num()); @@ -391,9 +391,9 @@ TYPED_TEST(PoolingLayerTest, TestSetupPadded) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); - pooling_param->set_pad(1); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(1); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); @@ -421,7 +421,7 @@ TYPED_TEST(PoolingLayerTest, TestSetupGlobalPooling) { TYPED_TEST(PoolingLayerTest, PrintBackward) { LayerParameter layer_param; layer_param.set_kernelsize(3); - layer_param.set_stride(2); + layer_param.add_stride(2); layer_param.set_pool(LayerParameter_PoolMethod_MAX); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); @@ -464,8 +464,8 @@ TYPED_TEST(PoolingLayerTest, TestGradientMax) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); - pooling_param->set_pad(1); + pooling_param->add_stride(2); + pooling_param->add_pad(1); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); PoolingLayer layer(layer_param); GradientChecker checker(1e-4, 1e-2); @@ -479,9 +479,9 @@ TYPED_TEST(PoolingLayerTest, TestForwardMaxPadded) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); - pooling_param->set_pad(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(2); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); this->blob_bottom_->Reshape(1, 1, 3, 3); // Input: @@ -528,7 +528,7 @@ TYPED_TEST(PoolingLayerTest, TestGradientMaxTopMask) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); this->blob_top_vec_.push_back(this->blob_top_mask_); PoolingLayer layer(layer_param); @@ -544,9 +544,9 @@ TYPED_TEST(PoolingLayerTest, TestForwardAve) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(1); - pooling_param->set_pad(1); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(1); + pooling_param->add_pad(1); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); this->blob_bottom_->Reshape(1, 1, 3, 3); FillerParameter filler_param; @@ -580,7 +580,7 @@ TYPED_TEST(PoolingLayerTest, TestGradientAve) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); PoolingLayer layer(layer_param); GradientChecker checker(1e-2, 1e-2); @@ -598,8 +598,8 @@ TYPED_TEST(PoolingLayerTest, TestGradientAvePadded) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); - pooling_param->set_pad(2); + pooling_param->add_stride(2); + pooling_param->add_pad(2); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); PoolingLayer layer(layer_param); GradientChecker checker(1e-2, 1e-2); @@ -641,7 +641,7 @@ class CuDNNPoolingLayerTest : public GPUDeviceTest { void TestForwardSquare() { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(2); + pooling_param->add_kernel_size(2); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); const int num = 2; const int channels = 2; @@ -968,8 +968,8 @@ TYPED_TEST_CASE(CuDNNPoolingLayerTest, TestDtypes); TYPED_TEST(CuDNNPoolingLayerTest, TestSetupCuDNN) { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); CuDNNPoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); EXPECT_EQ(this->blob_top_->num(), this->blob_bottom_->num()); @@ -981,9 +981,9 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestSetupCuDNN) { TYPED_TEST(CuDNNPoolingLayerTest, TestSetupPaddedCuDNN) { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); - pooling_param->set_pad(1); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(1); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); CuDNNPoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); @@ -997,7 +997,7 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestSetupPaddedCuDNN) { TYPED_TEST(CuDNNPoolingLayerTest, PrintBackwardCuDNN) { LayerParameter layer_param; layer_param.set_kernelsize(3); - layer_param.set_stride(2); + layer_param.add_stride(2); layer_param.set_pool(LayerParameter_PoolMethod_MAX); CuDNNPoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); @@ -1043,9 +1043,9 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientMaxCuDNN) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); + pooling_param->add_stride(2); // currenty, cuDNN pooling does not support padding - pooling_param->set_pad(0); + pooling_param->add_pad(0); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); CuDNNPoolingLayer layer(layer_param); GradientChecker checker(1e-4, 1e-2); @@ -1058,9 +1058,9 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientMaxCuDNN) { TYPED_TEST(CuDNNPoolingLayerTest, TestForwardMaxPaddedCuDNN) { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); - pooling_param->set_pad(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(2); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); this->blob_bottom_->Reshape(1, 1, 3, 3); // Input: @@ -1107,7 +1107,7 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientMaxTopMaskCuDNN) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); this->blob_top_vec_.push_back(this->blob_top_mask_); CuDNNPoolingLayer layer(layer_param); @@ -1123,11 +1123,11 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientMaxTopMaskCuDNN) { TYPED_TEST(CuDNNPoolingLayerTest, TestForwardAveCuDNN) { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(1); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(1); // Currently, cuDNN pooling does not support padding, so we use // a simplified version of this test. - pooling_param->set_pad(0); + pooling_param->add_pad(0); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); this->blob_bottom_->Reshape(1, 1, 3, 3); FillerParameter filler_param; @@ -1152,7 +1152,7 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientAveCuDNN) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); CuDNNPoolingLayer layer(layer_param); GradientChecker checker(1e-2, 1e-2); @@ -1169,8 +1169,8 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientAvePaddedCuDNN) { PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); pooling_param->set_kernel_h(kernel_h); pooling_param->set_kernel_w(kernel_w); - pooling_param->set_stride(2); - pooling_param->set_pad(2); + pooling_param->add_stride(2); + pooling_param->add_pad(2); pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); CuDNNPoolingLayer layer(layer_param); GradientChecker checker(1e-2, 1e-2); @@ -1180,6 +1180,493 @@ TYPED_TEST(CuDNNPoolingLayerTest, TestGradientAvePaddedCuDNN) { } } +// 3D cuDNN pooling tests +template +class CuDNNPoolingLayerTest3D : public GPUDeviceTest { + protected: + CuDNNPoolingLayerTest3D() + : blob_bottom_(new Blob()), + blob_top_(new Blob()), + blob_top_mask_(new Blob()) {} + virtual void SetUp() { + blob_bottom_vec_.push_back(blob_bottom_); + blob_top_vec_.push_back(blob_top_); + } + virtual ~CuDNNPoolingLayerTest3D() { + delete blob_bottom_; + delete blob_top_; + delete blob_top_mask_; + } + Blob* const blob_bottom_; + Blob* const blob_top_; + Blob* const blob_top_mask_; + std::vector*> blob_bottom_vec_; + std::vector*> blob_top_vec_; + + // set up random blob for gradient tests + void SetUp3DRandomBottomBlob() { + Caffe::set_random_seed(1701); + std::vector shape; + shape.push_back(2); + shape.push_back(3); + shape.push_back(6); + shape.push_back(5); + shape.push_back(4); + this->blob_bottom_->Reshape(shape); + // fill the values + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_); + } + + // set up fixed blob for forward tests + void SetUp3DTestBottomBlob(const int num, const int channels) { + std::vector shape; + shape.push_back(num); + shape.push_back(channels); + shape.push_back(4); + shape.push_back(3); + shape.push_back(6); + blob_bottom_->Reshape(shape); + // generated with matlab + // reshape(randperm(72, 72), [6 3 4]); + const int input[] = { 23, 41, 17, 36, 9, 31, + 56, 18, 55, 69, 70, 29, + 39, 61, 52, 20, 63, 26, + // + 71, 30, 59, 48, 32, 47, + 21, 19, 38, 27, 57, 13, + 37, 49, 64, 44, 33, 60, + // + 53, 68, 51, 16, 35, 15, + 8, 10, 6, 65, 62, 3, + 7, 45, 40, 22, 1, 50, + // + 46, 58, 2, 4, 66, 54, + 72, 25, 43, 12, 67, 28, + 5, 24, 34, 11, 42, 14 }; + const unsigned int num_elements = sizeof(input) / sizeof(*input); + for (int i = 0; i < num_elements * num * channels; i += num_elements) { + for (int j = 0; j < num_elements; ++j) { + blob_bottom_->mutable_cpu_data()[i + j] = input[j]; + } + } + } + + // test for 2x2x2 pooling + void TestForwardCube() { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + const int num = 2; + const int channels = 2; + SetUp3DTestBottomBlob(num, channels); + CuDNNPoolingLayer layer(layer_param); + layer.SetUp(blob_bottom_vec_, blob_top_vec_); + EXPECT_EQ(blob_top_->shape(0), num); + EXPECT_EQ(blob_top_->shape(1), channels); + EXPECT_EQ(blob_top_->shape(2), 3); + EXPECT_EQ(blob_top_->shape(3), 2); + EXPECT_EQ(blob_top_->shape(4), 5); + layer.Forward(blob_bottom_vec_, blob_top_vec_); + // expected output + const int output[] = { 71, 59, 69, 70, 70, + 61, 64, 69, 70, 70, + // + 71, 68, 65, 65, 62, + 49, 64, 65, 65, 62, + // + 72, 68, 65, 67, 67, + 72, 45, 65, 67, 67 }; + const unsigned int num_elements = sizeof(output) / sizeof(*output); + for (int i = 0; i < num_elements * num * channels; i += num_elements) { + for (int j = 0; j < num_elements; ++j) { + EXPECT_EQ(blob_top_->cpu_data()[i + j], output[j]); + } + } + } + + // test for 2x2x3 pooling + void TestForwardCuboidX() { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(2); + pooling_param->add_kernel_size(2); + pooling_param->add_kernel_size(3); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + const int num = 2; + const int channels = 2; + SetUp3DTestBottomBlob(num, channels); + CuDNNPoolingLayer layer(layer_param); + layer.SetUp(blob_bottom_vec_, blob_top_vec_); + EXPECT_EQ(blob_top_->shape(0), num); + EXPECT_EQ(blob_top_->shape(1), channels); + EXPECT_EQ(blob_top_->shape(2), 3); + EXPECT_EQ(blob_top_->shape(3), 2); + EXPECT_EQ(blob_top_->shape(4), 4); + layer.Forward(blob_bottom_vec_, blob_top_vec_); + // expected output + const int output[] = { 71, 69, 70, 70, + 64, 69, 70, 70, + // + 71, 68, 65, 65, + 64, 65, 65, 65, + // + 72, 68, 67, 67, + 72, 65, 67, 67}; + const unsigned int num_elements = sizeof(output) / sizeof(*output); + for (int i = 0; i < num_elements * num * channels; i += num_elements) { + for (int j = 0; j < num_elements; ++j) { + EXPECT_EQ(blob_top_->cpu_data()[i + j], output[j]); + } + } + } + + // test for 2x3x2 pooling + void TestForwardCuboidY() { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(2); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + const int num = 2; + const int channels = 2; + SetUp3DTestBottomBlob(num, channels); + CuDNNPoolingLayer layer(layer_param); + layer.SetUp(blob_bottom_vec_, blob_top_vec_); + EXPECT_EQ(blob_top_->shape(0), num); + EXPECT_EQ(blob_top_->shape(1), channels); + EXPECT_EQ(blob_top_->shape(2), 3); + EXPECT_EQ(blob_top_->shape(3), 1); + EXPECT_EQ(blob_top_->shape(4), 5); + layer.Forward(blob_bottom_vec_, blob_top_vec_); + // expected output + const int output[] = { 71, 64, 69, 70, 70, + // + 71, 68, 65, 65, 62, + // + 72, 68, 65, 67, 67}; + const unsigned int num_elements = sizeof(output) / sizeof(*output); + for (int i = 0; i < num_elements * num * channels; i += num_elements) { + for (int j = 0; j < num_elements; ++j) { + EXPECT_EQ(blob_top_->cpu_data()[i + j], output[j]); + } + } + } + + // test for 3x2x2 pooling + void TestForwardCuboidZ() { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(2); + pooling_param->add_kernel_size(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + const int num = 2; + const int channels = 2; + SetUp3DTestBottomBlob(num, channels); + CuDNNPoolingLayer layer(layer_param); + layer.SetUp(blob_bottom_vec_, blob_top_vec_); + EXPECT_EQ(blob_top_->shape(0), num); + EXPECT_EQ(blob_top_->shape(1), channels); + EXPECT_EQ(blob_top_->shape(2), 2); + EXPECT_EQ(blob_top_->shape(3), 2); + EXPECT_EQ(blob_top_->shape(4), 5); + layer.Forward(blob_bottom_vec_, blob_top_vec_); + // expected output + const int output[] = { 71, 68, 69, 70, 70, + 61, 64, 69, 70, 70, + // + 72, 68, 65, 67, 67, + 72, 64, 65, 67, 67}; + const unsigned int num_elements = sizeof(output) / sizeof(*output); + for (int i = 0; i < num_elements * num * channels; i += num_elements) { + for (int j = 0; j < num_elements; ++j) { + EXPECT_EQ(blob_top_->cpu_data()[i + j], output[j]); + } + } + } +}; + +TYPED_TEST_CASE(CuDNNPoolingLayerTest3D, TestDtypes); + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestSetup3DCuDNN) { + typedef TypeParam Dtype; + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + CuDNNPoolingLayer layer(layer_param); + // input shape {2, 3, 6, 5, 4} + this->SetUp3DRandomBottomBlob(); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->shape(0), this->blob_bottom_->shape(0)); + EXPECT_EQ(this->blob_top_->shape(1), this->blob_bottom_->shape(1)); + EXPECT_EQ(this->blob_top_->shape(2), 3); + EXPECT_EQ(this->blob_top_->shape(3), 2); + EXPECT_EQ(this->blob_top_->shape(4), 2); +} + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestSetupPadded3DCuDNN) { + typedef TypeParam Dtype; + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(1); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + CuDNNPoolingLayer layer(layer_param); + // input shape {2, 3, 6, 5, 4} + this->SetUp3DRandomBottomBlob(); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->shape(0), this->blob_bottom_->shape(0)); + EXPECT_EQ(this->blob_top_->shape(1), this->blob_bottom_->shape(1)); + EXPECT_EQ(this->blob_top_->shape(2), 4); + EXPECT_EQ(this->blob_top_->shape(3), 3); + EXPECT_EQ(this->blob_top_->shape(4), 3); +} + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestSetupGlobalPooling3DCuDNN) { + typedef TypeParam Dtype; + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->set_global_pooling(true); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + CuDNNPoolingLayer layer(layer_param); + // input shape {2, 3, 6, 5, 4} + this->SetUp3DRandomBottomBlob(); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->shape(0), this->blob_bottom_->shape(0)); + EXPECT_EQ(this->blob_top_->shape(1), this->blob_bottom_->shape(1)); + EXPECT_EQ(this->blob_top_->shape(2), 1); + EXPECT_EQ(this->blob_top_->shape(3), 1); + EXPECT_EQ(this->blob_top_->shape(4), 1); +} + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestForwardMax3DCuDNN) { + this->TestForwardCube(); + this->TestForwardCuboidX(); + this->TestForwardCuboidY(); + this->TestForwardCuboidZ(); +} + +// Currently, cuDNN does not support a top mask, so we comment this and +// the corresponding backward test. +/* +TYPED_TEST(CuDNNPoolingLayerTest, TestForwardMaxTopMask3DCuDNN) { + this->blob_top_vec_.push_back(this->blob_top_mask_); + this->TestForwardCube(); + this->TestForwardCuboidX(); + this->TestForwardCuboidY(); + this->TestForwardCuboidZ(); +} +*/ + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestForwardMaxPadded3DCuDNN) { + typedef TypeParam Dtype; + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + std::vector shape; + shape.push_back(1); + shape.push_back(1); + shape.push_back(3); + shape.push_back(3); + shape.push_back(3); + this->blob_bottom_->Reshape(shape); + // input + const int input[] = { 23, 12, 17, + 6, 9, 1, + 4, 18, 5, + // + 10, 11, 2, + 16, 14, 22, + 20, 8, 26, + // + 15, 13, 27, + 3, 7, 25, + 21, 19, 24 }; + for (int i = 0; i < sizeof(input) / sizeof(*input); ++i) { + this->blob_bottom_->mutable_cpu_data()[i] = input[i]; + } + CuDNNPoolingLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->shape(0), 1); + EXPECT_EQ(this->blob_top_->shape(1), 1); + EXPECT_EQ(this->blob_top_->shape(2), 3); + EXPECT_EQ(this->blob_top_->shape(3), 3); + EXPECT_EQ(this->blob_top_->shape(4), 3); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + Dtype epsilon = 1e-8; + // expected output + const int output[] = { 23, 23, 17, + 23, 23, 17, + 4, 18, 5, + // + 23, 27, 27, + 23, 27, 27, + 21, 26, 26, + // + 15, 27, 27, + 21, 27, 27, + 21, 24, 24}; + for (int i = 0; i < sizeof(output) / sizeof(*output); ++i) { + EXPECT_NEAR(this->blob_top_->cpu_data()[i], output[i], epsilon); + } +} +// Currently, cuDNN does not support a top mask, so we comment this and +// the corresponding backward test. +/* +TYPED_TEST(PoolingLayerTest, TestGradientMaxTopMask3DCuDNN) { + typedef TypeParam Dtype; + std::vector shape; + shape.push_back(2); + shape.push_back(3); + shape.push_back(6); + shape.push_back(5); + shape.push_back(4); + this->blob_bottom_->Reshape(shape); + for (int kernel_d = 3; kernel_d <= 4; kernel_d++) { + for (int kernel_h = 3; kernel_h <= 4; kernel_h++) { + for (int kernel_w = 3; kernel_w <= 4; kernel_w++) { + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(kernel_d); + pooling_param->add_kernel_size(kernel_h); + pooling_param->add_kernel_size(kernel_w); + pooling_param->add_stride(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + this->blob_top_vec_.push_back(this->blob_top_mask_); + PoolingLayer layer(layer_param); + GradientChecker checker(1e-4, 1e-2); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); + this->blob_top_vec_.pop_back(); + } + } + } +} +*/ + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestForwardAve3DCuDNN) { + typedef TypeParam Dtype; + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(1); + pooling_param->add_pad(1); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + std::vector shape; + shape.push_back(1); + shape.push_back(1); + shape.push_back(3); + shape.push_back(3); + shape.push_back(3); + this->blob_bottom_->Reshape(shape); + // input + const int input[] = { 23, 12, 17, + 6, 9, 1, + 4, 18, 5, + // + 10, 11, 2, + 16, 14, 22, + 20, 8, 26, + // + 15, 13, 27, + 3, 7, 25, + 21, 19, 24 }; + for (int i = 0; i < sizeof(input) / sizeof(*input); ++i) { + this->blob_bottom_->mutable_cpu_data()[i] = input[i]; + } + CuDNNPoolingLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + EXPECT_EQ(this->blob_top_->shape(0), 1); + EXPECT_EQ(this->blob_top_->shape(1), 1); + EXPECT_EQ(this->blob_top_->shape(2), 3); + EXPECT_EQ(this->blob_top_->shape(3), 3); + EXPECT_EQ(this->blob_top_->shape(4), 3); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + Dtype epsilon = 1e-4; + // expected output + const Dtype output[] = { 3.7407, 5.2963, 3.2593, + 5.5926, 8.2963, 5.3704, + 3.5185, 5.5185, 3.8148, + // + 5.1481, 8.6296, 5.9259, + 8.4815, 14.0000, 9.6296, + 5.3704, 9.1852, 6.5926, + // + 3.2963, 6.1111, 4.4815, + 5.8148, 10.4815, 7.3333, + 4.0000, 7.5926, 5.3704}; + for (int i = 0; i < sizeof(output) / sizeof(*output); ++i) { + EXPECT_NEAR(this->blob_top_->cpu_data()[i], output[i], epsilon); + } +} + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestGradientMax3DCuDNN) { + typedef TypeParam Dtype; + this->SetUp3DRandomBottomBlob(); + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(1); + pooling_param->set_pool(PoolingParameter_PoolMethod_MAX); + CuDNNPoolingLayer layer(layer_param); + GradientChecker checker(1e-4, 1e-2); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestGradientAve3DCuDNN) { + typedef TypeParam Dtype; + this->SetUp3DRandomBottomBlob(); + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + CuDNNPoolingLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(CuDNNPoolingLayerTest3D, TestGradientAvePadded3DCuDNN) { + typedef TypeParam Dtype; + this->SetUp3DRandomBottomBlob(); + LayerParameter layer_param; + PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); + pooling_param->clear_kernel_size(); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(3); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); + pooling_param->add_pad(2); + pooling_param->set_pool(PoolingParameter_PoolMethod_AVE); + CuDNNPoolingLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-2); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + #endif } // namespace caffe diff --git a/src/caffe/test/test_stochastic_pooling.cpp b/src/caffe/test/test_stochastic_pooling.cpp index cd5db8383ab..8b8a8ce1d4f 100644 --- a/src/caffe/test/test_stochastic_pooling.cpp +++ b/src/caffe/test/test_stochastic_pooling.cpp @@ -56,8 +56,8 @@ TYPED_TEST_CASE(CPUStochasticPoolingLayerTest, TestDtypes); TYPED_TEST(CPUStochasticPoolingLayerTest, TestSetup) { LayerParameter layer_param; PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); EXPECT_EQ(this->blob_top_->num(), this->blob_bottom_->num()); @@ -79,8 +79,8 @@ TYPED_TEST(GPUStochasticPoolingLayerTest, TestStochastic) { LayerParameter layer_param; layer_param.set_phase(TRAIN); PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_STOCHASTIC); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); @@ -122,8 +122,8 @@ TYPED_TEST(GPUStochasticPoolingLayerTest, TestStochasticTestPhase) { LayerParameter layer_param; layer_param.set_phase(TEST); PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_STOCHASTIC); PoolingLayer layer(layer_param); layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); @@ -159,8 +159,8 @@ TYPED_TEST(GPUStochasticPoolingLayerTest, TestGradient) { LayerParameter layer_param; layer_param.set_phase(TRAIN); PoolingParameter* pooling_param = layer_param.mutable_pooling_param(); - pooling_param->set_kernel_size(3); - pooling_param->set_stride(2); + pooling_param->add_kernel_size(3); + pooling_param->add_stride(2); pooling_param->set_pool(PoolingParameter_PoolMethod_STOCHASTIC); PoolingLayer layer(layer_param); GradientChecker checker(1e-4, 1e-2); diff --git a/src/caffe/util/upgrade_proto.cpp b/src/caffe/util/upgrade_proto.cpp index 9e186915b43..b9746618157 100644 --- a/src/caffe/util/upgrade_proto.cpp +++ b/src/caffe/util/upgrade_proto.cpp @@ -267,7 +267,7 @@ bool UpgradeV0LayerParameter(const V1LayerParameter& v0_layer_connection, if (type == "conv") { layer_param->mutable_convolution_param()->add_pad(v0_layer_param.pad()); } else if (type == "pool") { - layer_param->mutable_pooling_param()->set_pad(v0_layer_param.pad()); + layer_param->mutable_pooling_param()->add_pad(v0_layer_param.pad()); } else { LOG(ERROR) << "Unknown parameter pad for layer type " << type; is_fully_compatible = false; @@ -278,7 +278,7 @@ bool UpgradeV0LayerParameter(const V1LayerParameter& v0_layer_connection, layer_param->mutable_convolution_param()->add_kernel_size( v0_layer_param.kernelsize()); } else if (type == "pool") { - layer_param->mutable_pooling_param()->set_kernel_size( + layer_param->mutable_pooling_param()->add_kernel_size( v0_layer_param.kernelsize()); } else { LOG(ERROR) << "Unknown parameter kernelsize for layer type " << type; @@ -299,7 +299,7 @@ bool UpgradeV0LayerParameter(const V1LayerParameter& v0_layer_connection, layer_param->mutable_convolution_param()->add_stride( v0_layer_param.stride()); } else if (type == "pool") { - layer_param->mutable_pooling_param()->set_stride( + layer_param->mutable_pooling_param()->add_stride( v0_layer_param.stride()); } else { LOG(ERROR) << "Unknown parameter stride for layer type " << type;