From 64e78bdc76b8cabdf4282506438ab2d7f321adf3 Mon Sep 17 00:00:00 2001 From: Jonathan L Long Date: Sat, 27 Dec 2014 01:44:36 -0800 Subject: [PATCH 1/4] add CropLayer: crop blob to another blob's dimensions with offsets configure offset(s) through proto definition. --- include/caffe/layers/crop_layer.hpp | 49 +++++++++++++++++++++++ src/caffe/layers/crop_layer.cpp | 78 +++++++++++++++++++++++++++++++++++++ src/caffe/layers/crop_layer.cu | 60 ++++++++++++++++++++++++++++ src/caffe/proto/caffe.proto | 10 ++++- 4 files changed, 196 insertions(+), 1 deletion(-) create mode 100644 include/caffe/layers/crop_layer.hpp create mode 100644 src/caffe/layers/crop_layer.cpp create mode 100644 src/caffe/layers/crop_layer.cu diff --git a/include/caffe/layers/crop_layer.hpp b/include/caffe/layers/crop_layer.hpp new file mode 100644 index 00000000000..bab290718e0 --- /dev/null +++ b/include/caffe/layers/crop_layer.hpp @@ -0,0 +1,49 @@ +#ifndef CAFFE_CROP_LAYER_HPP_ +#define CAFFE_CROP_LAYER_HPP_ + +#include +#include + +#include "caffe/blob.hpp" +#include "caffe/layer.hpp" +#include "caffe/proto/caffe.pb.h" + +namespace caffe { + +/** + * @brief Takes a Blob and crop it along either the width or height dimension, + * outputting a cropped Blob. + * + * TODO(dox): thorough documentation for Forward, Backward, and proto params. + */ + +template +class CropLayer : public Layer { + public: + explicit CropLayer(const LayerParameter& param) + : Layer(param) {} + 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 "Crop"; } + virtual inline int ExactNumBottomBlobs() const { return 2; } + virtual inline int ExactNumTopBlobs() const { return 1; } + + protected: + virtual void Forward_cpu(const vector*>& bottom, + const vector*>& top); + virtual void Backward_cpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + virtual void Forward_gpu(const vector*>& bottom, + const vector*>& top); + virtual void Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom); + + int crop_h_, crop_w_; +}; + +} // namespace caffe + +#endif // CAFFE_CROP_LAYER_HPP_ diff --git a/src/caffe/layers/crop_layer.cpp b/src/caffe/layers/crop_layer.cpp new file mode 100644 index 00000000000..76409bd7408 --- /dev/null +++ b/src/caffe/layers/crop_layer.cpp @@ -0,0 +1,78 @@ +#include +#include +#include +#include + +#include "caffe/layer.hpp" +#include "caffe/layers/crop_layer.hpp" +#include "caffe/net.hpp" + + +namespace caffe { + +template +void CropLayer::LayerSetUp(const vector*>& bottom, + const vector*>& top) { + const CropParameter& param = this->layer_param_.crop_param(); + CHECK_EQ(bottom.size(), 2) << "Wrong number of bottom blobs."; + CHECK_EQ(bottom[0]->num_axes(), 4) << "Only works with 4D blobs."; + CHECK_EQ(bottom[1]->num_axes(), 4) << "Only works with 4D blobs."; + crop_h_ = param.offset_height(); + crop_w_ = param.offset_width(); +} + +template +void CropLayer::Reshape(const vector*>& bottom, + const vector*>& top) { + // Check that the image we are cropping minus the margin is bigger than the + // destination image. + CHECK_GT(bottom[0]->height()-crop_h_, bottom[1]->height()) + << "invalid offset"; + CHECK_GT(bottom[0]->width()-crop_w_, bottom[1]->width()) << "invalid offset"; + top[0]->Reshape(bottom[0]->num(), bottom[0]->channels(), bottom[1]->height(), + bottom[1]->width()); +} + +template +void CropLayer::Forward_cpu(const vector*>& bottom, + const vector*>& top) { + const Dtype* bottom_data = bottom[0]->cpu_data(); + Dtype* top_data = top[0]->mutable_cpu_data(); + for (int n = 0; n < top[0]->num(); ++n) { + for (int c = 0; c < top[0]->channels(); ++c) { + for (int h = 0; h < top[0]->height(); ++h) { + caffe_copy(top[0]->width(), + bottom_data + bottom[0]->offset(n, c, crop_h_ + h, crop_w_), + top_data + top[0]->offset(n, c, h)); + } + } + } +} + +template +void CropLayer::Backward_cpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom) { + const Dtype* top_diff = top[0]->cpu_diff(); + Dtype* bottom_diff = bottom[0]->mutable_cpu_diff(); + if (propagate_down[0]) { + caffe_set(bottom[0]->count(), static_cast(0), bottom_diff); + for (int n = 0; n < top[0]->num(); ++n) { + for (int c = 0; c < top[0]->channels(); ++c) { + for (int h = 0; h < top[0]->height(); ++h) { + caffe_copy(top[0]->width(), + top_diff + top[0]->offset(n, c, h), + bottom_diff + bottom[0]->offset(n, c, crop_h_ + h, crop_w_)); + } + } + } + } +} + +#ifdef CPU_ONLY +STUB_GPU(CropLayer); +#endif + +INSTANTIATE_CLASS(CropLayer); +REGISTER_LAYER_CLASS(Crop); + +} // namespace caffe diff --git a/src/caffe/layers/crop_layer.cu b/src/caffe/layers/crop_layer.cu new file mode 100644 index 00000000000..262f5fa8483 --- /dev/null +++ b/src/caffe/layers/crop_layer.cu @@ -0,0 +1,60 @@ +#include + +#include "caffe/layers/crop_layer.hpp" + +namespace caffe { + +// Copy (one line per thread) from one array to another, with arbitrary +// strides in the last two dimensions. +template +__global__ void copy_kernel(const int n, const int height, const int width, + const int src_outer_stride, const int src_inner_stride, + const int dest_outer_stride, const int dest_inner_stride, + const Dtype* src, Dtype* dest) { + CUDA_KERNEL_LOOP(index, n) { + int src_start = index / height * src_outer_stride + + index % height * src_inner_stride; + int dest_start = index / height * dest_outer_stride + + index % height * dest_inner_stride; + for (int i = 0; i < width; ++i) { + dest[dest_start + i] = src[src_start + i]; + } + } +} + +template +void CropLayer::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 lines = top[0]->count() / top[0]->width(); + + // NOLINT_NEXT_LINE(whitespace/operators) + copy_kernel<<>>( + lines, top[0]->height(), top[0]->width(), + bottom[0]->height() * bottom[0]->width(), bottom[0]->width(), + top[0]->height() * top[0]->width(), top[0]->width(), + bottom_data + bottom[0]->offset(0, 0, crop_h_, crop_w_), top_data); +} + +template +void CropLayer::Backward_gpu(const vector*>& top, + const vector& propagate_down, const vector*>& bottom) { + const Dtype* top_diff = top[0]->gpu_diff(); + Dtype* bottom_diff = bottom[0]->mutable_gpu_diff(); + const int lines = top[0]->count() / top[0]->width(); + + if (propagate_down[0]) { + caffe_gpu_set(bottom[0]->count(), static_cast(0), bottom_diff); + // NOLINT_NEXT_LINE(whitespace/operators) + copy_kernel<<>>( + lines, top[0]->height(), top[0]->width(), + top[0]->height() * top[0]->width(), top[0]->width(), + bottom[0]->height() * bottom[0]->width(), bottom[0]->width(), + top_diff, bottom_diff + bottom[0]->offset(0, 0, crop_h_, crop_w_)); + } +} + +INSTANTIATE_LAYER_GPU_FUNCS(CropLayer); + +} // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 3b27bbd94d2..ace1a418592 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -306,7 +306,7 @@ message ParamSpec { // NOTE // Update the next available ID when you add a new LayerParameter field. // -// LayerParameter next available layer-specific ID: 144 (last added: input_param) +// LayerParameter next available layer-specific ID: 145 (last added: crop_param) message LayerParameter { optional string name = 1; // the layer name optional string type = 2; // the layer type @@ -360,6 +360,7 @@ message LayerParameter { optional ConcatParameter concat_param = 104; optional ContrastiveLossParameter contrastive_loss_param = 105; optional ConvolutionParameter convolution_param = 106; + optional CropParameter crop_param = 144; optional DataParameter data_param = 107; optional DropoutParameter dropout_param = 108; optional DummyDataParameter dummy_data_param = 109; @@ -598,6 +599,13 @@ message ConvolutionParameter { optional bool force_nd_im2col = 17 [default = false]; } +message CropParameter { + // Assumes standard dimensions: ( N,C,H,W ) + // This could possibly be extended to use "optional BlobShape offsets" + optional uint32 offset_height = 1[default = 0]; + optional uint32 offset_width = 2[default = 0]; +} + message DataParameter { enum DB { LEVELDB = 0; From 952fd17e52b24736f2644c32c249538241b34474 Mon Sep 17 00:00:00 2001 From: max argus Date: Tue, 19 Jan 2016 18:35:04 +0000 Subject: [PATCH 2/4] Extend Crop to N-D, changed CropParameter. --- include/caffe/layers/crop_layer.hpp | 22 ++++++- src/caffe/layers/crop_layer.cpp | 124 +++++++++++++++++++++++++++--------- src/caffe/layers/crop_layer.cu | 97 +++++++++++++++++++++++----- src/caffe/proto/caffe.proto | 19 ++++-- 4 files changed, 210 insertions(+), 52 deletions(-) diff --git a/include/caffe/layers/crop_layer.hpp b/include/caffe/layers/crop_layer.hpp index bab290718e0..f84bab1ec49 100644 --- a/include/caffe/layers/crop_layer.hpp +++ b/include/caffe/layers/crop_layer.hpp @@ -41,9 +41,27 @@ class CropLayer : public Layer { virtual void Backward_gpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom); - int crop_h_, crop_w_; -}; + vector offsets; + + private: + void crop_copy(const vector*>& bottom, + const vector*>& top, + const vector& offsets, + vector indices, + int cur_dim, + const Dtype* src_data, + Dtype* dest_data, + bool is_forward); + void crop_copy_gpu(const vector*>& bottom, + const vector*>& top, + const vector& offsets, + vector indices, + int cur_dim, + const Dtype* src_data, + Dtype* dest_data, + bool is_forward); +}; } // namespace caffe #endif // CAFFE_CROP_LAYER_HPP_ diff --git a/src/caffe/layers/crop_layer.cpp b/src/caffe/layers/crop_layer.cpp index 76409bd7408..82729f17c9f 100644 --- a/src/caffe/layers/crop_layer.cpp +++ b/src/caffe/layers/crop_layer.cpp @@ -1,8 +1,10 @@ #include +#include #include #include #include + #include "caffe/layer.hpp" #include "caffe/layers/crop_layer.hpp" #include "caffe/net.hpp" @@ -13,40 +15,108 @@ namespace caffe { template void CropLayer::LayerSetUp(const vector*>& bottom, const vector*>& top) { - const CropParameter& param = this->layer_param_.crop_param(); CHECK_EQ(bottom.size(), 2) << "Wrong number of bottom blobs."; - CHECK_EQ(bottom[0]->num_axes(), 4) << "Only works with 4D blobs."; - CHECK_EQ(bottom[1]->num_axes(), 4) << "Only works with 4D blobs."; - crop_h_ = param.offset_height(); - crop_w_ = param.offset_width(); + // parameter setup moved to Reshape because it depends on size. } template void CropLayer::Reshape(const vector*>& bottom, const vector*>& top) { - // Check that the image we are cropping minus the margin is bigger than the - // destination image. - CHECK_GT(bottom[0]->height()-crop_h_, bottom[1]->height()) - << "invalid offset"; - CHECK_GT(bottom[0]->width()-crop_w_, bottom[1]->width()) << "invalid offset"; - top[0]->Reshape(bottom[0]->num(), bottom[0]->channels(), bottom[1]->height(), - bottom[1]->width()); + const CropParameter& param = this->layer_param_.crop_param(); + // bottom[0] supplies the data + // bottom[1] supplies the size + int input_dim = bottom[0]->num_axes(); + CHECK_LT(param.axis(), input_dim) << "crop axis bigger than input dim"; + // initialize all offsets to 0 + offsets = vector(input_dim, 0); + // initialize new shape to bottom[0] + vector new_shape(bottom[0]->shape()); + + if (param.offset_size() > 1) { + // the number of crop values specified must be equal to the number + // of dimensions following axis + CHECK_EQ(param.axis() + param.offset_size(), input_dim) + << "number of crop values specified must be equal to the number of " + << "dimensions following axis."; + } + // apply crops + for (int i = 0; i < input_dim; ++i) { + int crop_offset = 0; + int new_size = bottom[0]->shape(i); + if (i >= param.axis() && param.offset_size() == 1) { + // if only one crop value is supplied, crop all dimensions after axis + // by this crop value + crop_offset = param.offset(0); + new_size = bottom[1]->shape(i); + } else if (i >= param.axis() && param.offset_size() > 1) { + // crop values specified must be equal to the number of dimensions + // following axis + crop_offset = param.offset(i - param.axis()); + new_size = bottom[1]->shape(i); + } + // Check that the image we are cropping minus the margin is bigger + // than the destination image. + CHECK_GE(bottom[0]->shape(i) - crop_offset, + bottom[1]->shape(i)) + << "invalid crop parameters in dimension: " << i; + // Now set new size and offsets + new_shape[i] = new_size; + offsets[i] = crop_offset; + } + top[0]->Reshape(new_shape); +} + +// recursive copy function +template +void CropLayer::crop_copy(const vector*>& bottom, + const vector*>& top, + const vector& offsets, + vector indices, + int cur_dim, + const Dtype* src_data, + Dtype* dest_data, + bool is_forward) { + if (cur_dim + 1 < top[0]->num_axes()) { + // We are not yet at the final dimension, call copy recursivley + for (int i = 0; i < top[0]->shape(cur_dim); ++i) { + indices[cur_dim] = i; + crop_copy(bottom, top, offsets, indices, cur_dim+1, + src_data, dest_data, is_forward); + } + } else { + // We are at the last dimensions, which is stored continously in memory + for (int i = 0; i < top[0]->shape(cur_dim); ++i) { + // prepare index vector reduced(red) and with offsets(off) + std::vector ind_red(cur_dim, 0); + std::vector ind_off(cur_dim+1, 0); + for (int j = 0; j < cur_dim; ++j) { + ind_red[j] = indices[j]; + ind_off[j] = indices[j] + offsets[j]; + } + ind_off[cur_dim] = offsets[cur_dim]; + // do the copy + if (is_forward) { + caffe_copy(top[0]->shape(cur_dim), + src_data + bottom[0]->offset(ind_off), + dest_data + top[0]->offset(ind_red)); + } else { + // in the backwards pass the src_data is top_diff + // and the dest_data is bottom_diff + caffe_copy(top[0]->shape(cur_dim), + src_data + top[0]->offset(ind_red), + dest_data + bottom[0]->offset(ind_off)); + } + } + } } template void CropLayer::Forward_cpu(const vector*>& bottom, const vector*>& top) { + std::vector indices(top[0]->num_axes(), 0); const Dtype* bottom_data = bottom[0]->cpu_data(); Dtype* top_data = top[0]->mutable_cpu_data(); - for (int n = 0; n < top[0]->num(); ++n) { - for (int c = 0; c < top[0]->channels(); ++c) { - for (int h = 0; h < top[0]->height(); ++h) { - caffe_copy(top[0]->width(), - bottom_data + bottom[0]->offset(n, c, crop_h_ + h, crop_w_), - top_data + top[0]->offset(n, c, h)); - } - } - } + crop_copy(bottom, top, offsets, indices, 0, bottom_data, top_data, true); } template @@ -54,17 +124,11 @@ void CropLayer::Backward_cpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom) { const Dtype* top_diff = top[0]->cpu_diff(); Dtype* bottom_diff = bottom[0]->mutable_cpu_diff(); + if (propagate_down[0]) { caffe_set(bottom[0]->count(), static_cast(0), bottom_diff); - for (int n = 0; n < top[0]->num(); ++n) { - for (int c = 0; c < top[0]->channels(); ++c) { - for (int h = 0; h < top[0]->height(); ++h) { - caffe_copy(top[0]->width(), - top_diff + top[0]->offset(n, c, h), - bottom_diff + bottom[0]->offset(n, c, crop_h_ + h, crop_w_)); - } - } - } + std::vector indices(top[0]->num_axes(), 0); + crop_copy(bottom, top, offsets, indices, 0, top_diff, bottom_diff, false); } } diff --git a/src/caffe/layers/crop_layer.cu b/src/caffe/layers/crop_layer.cu index 262f5fa8483..7b832c0a0dc 100644 --- a/src/caffe/layers/crop_layer.cu +++ b/src/caffe/layers/crop_layer.cu @@ -22,19 +22,90 @@ __global__ void copy_kernel(const int n, const int height, const int width, } } +// recursive copy function, this function is similar to crop_copy but loops +// over all but the last two dimensions. It is implemented this way to allow +// for ND cropping while still relying on a CUDA kernel for the innermost +// two dimensions for performance reasons. +// An alternative way to implement ND cropping relying more on the kernel +// would require passing offsets to the kernel, which is a bit problematic +// because it is of variable length. Since in the standard (N,C,W,H) case +// N,C are usually not cropped a speedup could be achieved by not looping +// the application of the copy_kernel around these dimensions. +template +void CropLayer::crop_copy_gpu(const vector*>& bottom, + const vector*>& top, + const vector& offsets, + vector indices, + int cur_dim, + const Dtype* src_data, + Dtype* dest_data, + bool is_forward) { + if (cur_dim + 2 < top[0]->num_axes()) { + // We are not yet at the final dimension, call copy recursivley + for (int i = 0; i < top[0]->shape(cur_dim); ++i) { + indices[cur_dim] = i; + crop_copy_gpu(bottom, top, offsets, indices, cur_dim+1, + src_data, dest_data, is_forward); + } + } else { + // We are at the last two dimensions, which are stored continously in memory + // With (N,C,H,W) + // (0,1,2,3) cur_dim -> H + // cur_dim+1 -> W + const int lines = top[0]->shape(cur_dim); + const int height = top[0]->shape(cur_dim); + const int width = top[0]->shape(cur_dim+1); + std::vector ind_off(cur_dim+2, 0); + for (int j = 0; j < cur_dim; ++j) { + ind_off[j] = indices[j] + offsets[j]; + } + ind_off[cur_dim] = offsets[cur_dim]; + ind_off[cur_dim+1] = offsets[cur_dim+1]; + // Compute copy strides + const int src_outer_stride = + bottom[0]->shape(cur_dim)*bottom[0]->shape(cur_dim+1); + const int src_inner_stride = bottom[0]->shape(cur_dim+1); + const int dest_outer_stride = + top[0]->shape(cur_dim)*top[0]->shape(cur_dim+1); + const int dest_inner_stride = top[0]->shape(cur_dim+1); + + if (is_forward) { + const Dtype* bottom_data = bottom[0]->gpu_data() + + bottom[0]->offset(ind_off); + Dtype* top_data = top[0]->mutable_gpu_data() + + top[0]->offset(indices); + // NOLINT_NEXT_LINE(whitespace/operators) + copy_kernel<<>>( + lines, height, width, + src_outer_stride, src_inner_stride, + dest_outer_stride, dest_inner_stride, + bottom_data, top_data); + + } else { + const Dtype* top_diff = top[0]->gpu_diff() + + top[0]->offset(indices); + Dtype* bottom_diff = bottom[0]->mutable_gpu_diff() + + bottom[0]->offset(ind_off); + // NOLINT_NEXT_LINE(whitespace/operators) + copy_kernel<<>>( + lines, height, width, + dest_outer_stride, dest_inner_stride, + src_outer_stride, src_inner_stride, + top_diff, bottom_diff); + } + } +} + template void CropLayer::Forward_gpu(const vector*>& bottom, const vector*>& top) { + std::vector indices(top[0]->num_axes(), 0); + // This works because crop_copy uses caffe_copy which calls cudaMemcpy. + // My intuition is that calling this thousands of times is probably less + // efficient than writing a custom kernel. const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = top[0]->mutable_gpu_data(); - const int lines = top[0]->count() / top[0]->width(); - - // NOLINT_NEXT_LINE(whitespace/operators) - copy_kernel<<>>( - lines, top[0]->height(), top[0]->width(), - bottom[0]->height() * bottom[0]->width(), bottom[0]->width(), - top[0]->height() * top[0]->width(), top[0]->width(), - bottom_data + bottom[0]->offset(0, 0, crop_h_, crop_w_), top_data); + crop_copy_gpu(bottom, top, offsets, indices, 0, bottom_data, top_data, true); } template @@ -42,16 +113,12 @@ void CropLayer::Backward_gpu(const vector*>& top, const vector& propagate_down, const vector*>& bottom) { const Dtype* top_diff = top[0]->gpu_diff(); Dtype* bottom_diff = bottom[0]->mutable_gpu_diff(); - const int lines = top[0]->count() / top[0]->width(); if (propagate_down[0]) { caffe_gpu_set(bottom[0]->count(), static_cast(0), bottom_diff); - // NOLINT_NEXT_LINE(whitespace/operators) - copy_kernel<<>>( - lines, top[0]->height(), top[0]->width(), - top[0]->height() * top[0]->width(), top[0]->width(), - bottom[0]->height() * bottom[0]->width(), bottom[0]->width(), - top_diff, bottom_diff + bottom[0]->offset(0, 0, crop_h_, crop_w_)); + std::vector indices(top[0]->num_axes(), 0); + crop_copy_gpu(bottom, top, offsets, indices, 0, top_diff, bottom_diff, + false); } } diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index ace1a418592..60d387a7de1 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -600,10 +600,19 @@ message ConvolutionParameter { } message CropParameter { - // Assumes standard dimensions: ( N,C,H,W ) - // This could possibly be extended to use "optional BlobShape offsets" - optional uint32 offset_height = 1[default = 0]; - optional uint32 offset_width = 2[default = 0]; + // To crop, elements of the first bottom are selected to fit the dimensions + // of the second, reference bottom. The crop is configured by + // - the crop `axis` to pick the dimensions for cropping + // - the crop `offset` to set the shift for all/each dimension + // to align the cropped bottom with the reference bottom. + // All dimensions up to but excluding `axis` are preserved, while + // the dimensions including and trailing `axis` are cropped. + // If only one `offset` is set, then all dimensions are offset by this amount. + // Otherwise, the number of offsets must equal the number of cropped axes to + // shift the crop in each dimension accordingly. + // Note: standard dimensions are N,C,H,W so the default is a spatial crop. + optional uint32 axis = 1 [default = 2]; + repeated uint32 offset = 2; } message DataParameter { @@ -680,7 +689,7 @@ message EltwiseParameter { // Message that stores parameters used by ELULayer message ELUParameter { // Described in: - // Clevert, D.-A., Unterthiner, T., & Hochreiter, S. (2015). Fast and Accurate + // Clevert, D.-A., Unterthiner, T., & Hochreiter, S. (2015). Fast and Accurate // Deep Network Learning by Exponential Linear Units (ELUs). arXiv optional float alpha = 1 [default = 1]; } From ca9fa498fb76a83befe6cb299a71f7f779aeb5d9 Mon Sep 17 00:00:00 2001 From: max argus Date: Mon, 29 Feb 2016 11:24:25 +0000 Subject: [PATCH 3/4] Crop: fixes, tests and negative axis indexing. --- include/caffe/layers/crop_layer.hpp | 4 +- src/caffe/layers/crop_layer.cpp | 50 ++++---- src/caffe/layers/crop_layer.cu | 3 - src/caffe/proto/caffe.proto | 6 +- src/caffe/test/test_crop_layer.cpp | 228 ++++++++++++++++++++++++++++++++++++ 5 files changed, 263 insertions(+), 28 deletions(-) create mode 100644 src/caffe/test/test_crop_layer.cpp diff --git a/include/caffe/layers/crop_layer.hpp b/include/caffe/layers/crop_layer.hpp index f84bab1ec49..5c605b2ae9e 100644 --- a/include/caffe/layers/crop_layer.hpp +++ b/include/caffe/layers/crop_layer.hpp @@ -11,8 +11,8 @@ namespace caffe { /** - * @brief Takes a Blob and crop it along either the width or height dimension, - * outputting a cropped Blob. + * @brief Takes a Blob and crop it, to the shape specified by the second input + * Blob, across all dimensions after the specified axis. * * TODO(dox): thorough documentation for Forward, Backward, and proto params. */ diff --git a/src/caffe/layers/crop_layer.cpp b/src/caffe/layers/crop_layer.cpp index 82729f17c9f..e81bdd732f3 100644 --- a/src/caffe/layers/crop_layer.cpp +++ b/src/caffe/layers/crop_layer.cpp @@ -15,44 +15,52 @@ namespace caffe { template void CropLayer::LayerSetUp(const vector*>& bottom, const vector*>& top) { + // All logic that depends only on the number of dimensions is here, + // the rest is in Reshape because it depends on Blob size. + // bottom[0] supplies the data + // bottom[1] supplies the size + const CropParameter& param = this->layer_param_.crop_param(); CHECK_EQ(bottom.size(), 2) << "Wrong number of bottom blobs."; - // parameter setup moved to Reshape because it depends on size. + int input_dim = bottom[0]->num_axes(); + const int start_axis = bottom[0]->CanonicalAxisIndex(param.axis()); + CHECK_LT(start_axis, input_dim) << "crop axis bigger than input dim"; + if (param.offset_size() > 1) { + // the number of crop values specified must be equal to the number + // of dimensions following axis + CHECK_EQ(start_axis + param.offset_size(), input_dim) + << "number of offset values specified must be equal to the number of " + << "dimensions following axis."; + } } template void CropLayer::Reshape(const vector*>& bottom, const vector*>& top) { const CropParameter& param = this->layer_param_.crop_param(); - // bottom[0] supplies the data - // bottom[1] supplies the size int input_dim = bottom[0]->num_axes(); - CHECK_LT(param.axis(), input_dim) << "crop axis bigger than input dim"; + const int start_axis = bottom[0]->CanonicalAxisIndex(param.axis()); + // initialize all offsets to 0 offsets = vector(input_dim, 0); // initialize new shape to bottom[0] vector new_shape(bottom[0]->shape()); - if (param.offset_size() > 1) { - // the number of crop values specified must be equal to the number - // of dimensions following axis - CHECK_EQ(param.axis() + param.offset_size(), input_dim) - << "number of crop values specified must be equal to the number of " - << "dimensions following axis."; - } // apply crops for (int i = 0; i < input_dim; ++i) { int crop_offset = 0; int new_size = bottom[0]->shape(i); - if (i >= param.axis() && param.offset_size() == 1) { - // if only one crop value is supplied, crop all dimensions after axis - // by this crop value - crop_offset = param.offset(0); - new_size = bottom[1]->shape(i); - } else if (i >= param.axis() && param.offset_size() > 1) { - // crop values specified must be equal to the number of dimensions - // following axis - crop_offset = param.offset(i - param.axis()); + if (i >= start_axis) { new_size = bottom[1]->shape(i); + + if (param.offset_size() == 1) { + // if only one crop value is supplied, crop all dimensions after axis + // by this crop value + crop_offset = param.offset(0); + } else if (param.offset_size() > 1) { + // crop values specified must be equal to the number of dimensions + // following axis + crop_offset = param.offset(i - start_axis); + } } // Check that the image we are cropping minus the margin is bigger // than the destination image. @@ -77,7 +85,7 @@ void CropLayer::crop_copy(const vector*>& bottom, Dtype* dest_data, bool is_forward) { if (cur_dim + 1 < top[0]->num_axes()) { - // We are not yet at the final dimension, call copy recursivley + // We are not yet at the final dimension, call copy recursively for (int i = 0; i < top[0]->shape(cur_dim); ++i) { indices[cur_dim] = i; crop_copy(bottom, top, offsets, indices, cur_dim+1, diff --git a/src/caffe/layers/crop_layer.cu b/src/caffe/layers/crop_layer.cu index 7b832c0a0dc..9ed8f7cce57 100644 --- a/src/caffe/layers/crop_layer.cu +++ b/src/caffe/layers/crop_layer.cu @@ -100,9 +100,6 @@ template void CropLayer::Forward_gpu(const vector*>& bottom, const vector*>& top) { std::vector indices(top[0]->num_axes(), 0); - // This works because crop_copy uses caffe_copy which calls cudaMemcpy. - // My intuition is that calling this thousands of times is probably less - // efficient than writing a custom kernel. const Dtype* bottom_data = bottom[0]->gpu_data(); Dtype* top_data = top[0]->mutable_gpu_data(); crop_copy_gpu(bottom, top, offsets, indices, 0, bottom_data, top_data, true); diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 60d387a7de1..6900bb71482 100644 --- a/src/caffe/proto/caffe.proto +++ b/src/caffe/proto/caffe.proto @@ -610,8 +610,10 @@ message CropParameter { // If only one `offset` is set, then all dimensions are offset by this amount. // Otherwise, the number of offsets must equal the number of cropped axes to // shift the crop in each dimension accordingly. - // Note: standard dimensions are N,C,H,W so the default is a spatial crop. - optional uint32 axis = 1 [default = 2]; + // Note: standard dimensions are N,C,H,W so the default is a spatial crop, + // and `axis` may be negative to index from the end (e.g., -1 for the last + // axis). + optional int32 axis = 1 [default = 2]; repeated uint32 offset = 2; } diff --git a/src/caffe/test/test_crop_layer.cpp b/src/caffe/test/test_crop_layer.cpp new file mode 100644 index 00000000000..ba962e4c6f1 --- /dev/null +++ b/src/caffe/test/test_crop_layer.cpp @@ -0,0 +1,228 @@ +#include + +#include "gtest/gtest.h" + +#include "caffe/blob.hpp" +#include "caffe/common.hpp" +#include "caffe/filler.hpp" +#include "caffe/layers/crop_layer.hpp" + +#include "caffe/test/test_caffe_main.hpp" +#include "caffe/test/test_gradient_check_util.hpp" + +namespace caffe { + +template +class CropLayerTest : public MultiDeviceTest { + typedef typename TypeParam::Dtype Dtype; + + protected: + CropLayerTest() + : blob_bottom_0_(new Blob(2, 5, 6, 5)), + blob_bottom_1_(new Blob(2, 4, 5, 3)), + blob_top_(new Blob()) {} + virtual void SetUp() { + // fill the values + for (int i = 0; i < this->blob_bottom_0_->count(); ++i) { + this->blob_bottom_0_->mutable_cpu_data()[i] = i; + } + + + blob_bottom_vec_0_.push_back(blob_bottom_0_); + blob_bottom_vec_0_.push_back(blob_bottom_1_); + blob_top_vec_.push_back(blob_top_); + } + + virtual ~CropLayerTest() { + delete blob_bottom_0_; delete blob_bottom_1_; + delete blob_top_; + } + + Blob* const blob_bottom_0_; + Blob* const blob_bottom_1_; + Blob* const blob_top_; + vector*> blob_bottom_vec_0_; + vector*> blob_top_vec_; +}; + + +TYPED_TEST_CASE(CropLayerTest, TestDtypesAndDevices); + +TYPED_TEST(CropLayerTest, TestSetupShapeAll) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + + // Crop all dimensions + layer_param.mutable_crop_param()->set_axis(0); + CropLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + for (int i = 0; i < this->blob_top_->num_axes(); ++i) { + EXPECT_EQ(this->blob_bottom_1_->shape(i), this->blob_top_->shape(i)); + } +} + +TYPED_TEST(CropLayerTest, TestSetupShapeDefault) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + // Crop last two dimensions, axis is 2 by default + CropLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + for (int i = 0; i < this->blob_top_->num_axes(); ++i) { + if (i < 2) { + EXPECT_EQ(this->blob_bottom_0_->shape(i), this->blob_top_->shape(i)); + } else { + EXPECT_EQ(this->blob_bottom_1_->shape(i), this->blob_top_->shape(i)); + } + } +} + +TYPED_TEST(CropLayerTest, TestSetupShapeNegativeIndexing) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + // Crop last dimension by negative indexing + layer_param.mutable_crop_param()->set_axis(-1); + CropLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + for (int i = 0; i < this->blob_top_->num_axes(); ++i) { + if (i < 3) { + EXPECT_EQ(this->blob_bottom_0_->shape(i), this->blob_top_->shape(i)); + } else { + EXPECT_EQ(this->blob_bottom_1_->shape(i), this->blob_top_->shape(i)); + } + } +} + + +TYPED_TEST(CropLayerTest, TestForwardNum) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(0); + + CropLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); + for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { + for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { + for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { + for (int w = 0; w < this->blob_bottom_0_->width(); ++w) { + if ( n < this->blob_top_->shape(0) && + c < this->blob_top_->shape(1) && + h < this->blob_top_->shape(2) && + w < this->blob_top_->shape(3) ) { + EXPECT_EQ(this->blob_top_->data_at(n, c, h, w), + this->blob_bottom_0_->data_at(n, c, h, w)); + } + } + } + } + } +} + +TYPED_TEST(CropLayerTest, TestForwardNumOffsets) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(0); + layer_param.mutable_crop_param()->add_offset(0); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(2); + CropLayer layer(layer_param); + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); + for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { + for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { + for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { + for (int w = 0; w < this->blob_bottom_0_->width(); ++w) { + if ( n < this->blob_top_->shape(0) && + c < this->blob_top_->shape(1) && + h < this->blob_top_->shape(2) && + w < this->blob_top_->shape(3) ) { + EXPECT_EQ(this->blob_top_->data_at(n, c, h, w), + this->blob_bottom_0_->data_at(n+0, c+1, h+1, w+2)); + } + } + } + } + } +} + +TYPED_TEST(CropLayerTest, TestGradientNum) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + CropLayer layer(layer_param); + + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); + + // Copy top data into diff + caffe_copy(this->blob_top_->count(), this->blob_top_->cpu_data(), + this->blob_top_->mutable_cpu_diff()); + + // Do backward pass + vector propagate_down(2, true); + layer.Backward(this->blob_top_vec_, propagate_down, this->blob_bottom_vec_0_); + + + // Check results + for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { + for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { + for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { + for (int w = 0; w < this->blob_bottom_0_->width(); ++w) { + if ( n < this->blob_top_->shape(0) && + c < this->blob_top_->shape(1) && + h < this->blob_top_->shape(2) && + w < this->blob_top_->shape(3) ) { + EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), + this->blob_bottom_0_->data_at(n, c, h, w)); + } else { + EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), 0); + } + } + } + } + } +} + +TYPED_TEST(CropLayerTest, TestGradientNumOffset) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(0); + layer_param.mutable_crop_param()->add_offset(0); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(2); + CropLayer layer(layer_param); + + layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); + + // Copy top data into diff + caffe_copy(this->blob_top_->count(), this->blob_top_->cpu_data(), + this->blob_top_->mutable_cpu_diff()); + + // Do backward pass + vector propagate_down(2, true); + layer.Backward(this->blob_top_vec_, propagate_down, this->blob_bottom_vec_0_); + + + // Check results + for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { + for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { + for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { + for (int w = 0; w < this->blob_bottom_0_->width(); ++w) { + if ( 0 <= n && n < 0 + this->blob_top_->shape(0) && + 1 <= c && c < 1 + this->blob_top_->shape(1) && + 1 <= h && h < 1 + this->blob_top_->shape(2) && + 2 <= w && w < 2 + this->blob_top_->shape(3) ) { + EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), + this->blob_bottom_0_->data_at(n, c, h, w)); + } else { + EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), 0); + } + } + } + } + } +} + +} // namespace caffe From e03a2873a0fdfd871b11a86d7feed45fb71013b0 Mon Sep 17 00:00:00 2001 From: Evan Shelhamer Date: Fri, 4 Mar 2016 02:20:13 -0800 Subject: [PATCH 4/4] Crop: more tests and test tuning. Changes are: reduce test blob dims for speed use standard Gaussian filler, polish formatting and rename tests, test HW crop and 5D crop, standard gradient checks --- src/caffe/test/test_crop_layer.cpp | 183 ++++++++++++++++++++++--------------- 1 file changed, 110 insertions(+), 73 deletions(-) diff --git a/src/caffe/test/test_crop_layer.cpp b/src/caffe/test/test_crop_layer.cpp index ba962e4c6f1..45f24e2ee8d 100644 --- a/src/caffe/test/test_crop_layer.cpp +++ b/src/caffe/test/test_crop_layer.cpp @@ -18,18 +18,18 @@ class CropLayerTest : public MultiDeviceTest { protected: CropLayerTest() - : blob_bottom_0_(new Blob(2, 5, 6, 5)), - blob_bottom_1_(new Blob(2, 4, 5, 3)), + : blob_bottom_0_(new Blob(2, 4, 5, 4)), + blob_bottom_1_(new Blob(2, 3, 4, 2)), blob_top_(new Blob()) {} virtual void SetUp() { // fill the values - for (int i = 0; i < this->blob_bottom_0_->count(); ++i) { - this->blob_bottom_0_->mutable_cpu_data()[i] = i; - } - + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_0_); + filler.Fill(this->blob_bottom_1_); - blob_bottom_vec_0_.push_back(blob_bottom_0_); - blob_bottom_vec_0_.push_back(blob_bottom_1_); + blob_bottom_vec_.push_back(blob_bottom_0_); + blob_bottom_vec_.push_back(blob_bottom_1_); blob_top_vec_.push_back(blob_top_); } @@ -41,7 +41,7 @@ class CropLayerTest : public MultiDeviceTest { Blob* const blob_bottom_0_; Blob* const blob_bottom_1_; Blob* const blob_top_; - vector*> blob_bottom_vec_0_; + vector*> blob_bottom_vec_; vector*> blob_top_vec_; }; @@ -51,11 +51,10 @@ TYPED_TEST_CASE(CropLayerTest, TestDtypesAndDevices); TYPED_TEST(CropLayerTest, TestSetupShapeAll) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; - // Crop all dimensions layer_param.mutable_crop_param()->set_axis(0); CropLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); for (int i = 0; i < this->blob_top_->num_axes(); ++i) { EXPECT_EQ(this->blob_bottom_1_->shape(i), this->blob_top_->shape(i)); } @@ -66,7 +65,7 @@ TYPED_TEST(CropLayerTest, TestSetupShapeDefault) { LayerParameter layer_param; // Crop last two dimensions, axis is 2 by default CropLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); for (int i = 0; i < this->blob_top_->num_axes(); ++i) { if (i < 2) { EXPECT_EQ(this->blob_bottom_0_->shape(i), this->blob_top_->shape(i)); @@ -82,7 +81,7 @@ TYPED_TEST(CropLayerTest, TestSetupShapeNegativeIndexing) { // Crop last dimension by negative indexing layer_param.mutable_crop_param()->set_axis(-1); CropLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); for (int i = 0; i < this->blob_top_->num_axes(); ++i) { if (i < 3) { EXPECT_EQ(this->blob_bottom_0_->shape(i), this->blob_top_->shape(i)); @@ -92,15 +91,13 @@ TYPED_TEST(CropLayerTest, TestSetupShapeNegativeIndexing) { } } - -TYPED_TEST(CropLayerTest, TestForwardNum) { +TYPED_TEST(CropLayerTest, TestCropAll) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; layer_param.mutable_crop_param()->set_axis(0); - CropLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { @@ -118,7 +115,7 @@ TYPED_TEST(CropLayerTest, TestForwardNum) { } } -TYPED_TEST(CropLayerTest, TestForwardNumOffsets) { +TYPED_TEST(CropLayerTest, TestCropAllOffset) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; layer_param.mutable_crop_param()->set_axis(0); @@ -127,8 +124,8 @@ TYPED_TEST(CropLayerTest, TestForwardNumOffsets) { layer_param.mutable_crop_param()->add_offset(1); layer_param.mutable_crop_param()->add_offset(2); CropLayer layer(layer_param); - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { @@ -138,7 +135,7 @@ TYPED_TEST(CropLayerTest, TestForwardNumOffsets) { h < this->blob_top_->shape(2) && w < this->blob_top_->shape(3) ) { EXPECT_EQ(this->blob_top_->data_at(n, c, h, w), - this->blob_bottom_0_->data_at(n+0, c+1, h+1, w+2)); + this->blob_bottom_0_->data_at(n, c+1, h+1, w+2)); } } } @@ -146,36 +143,25 @@ TYPED_TEST(CropLayerTest, TestForwardNumOffsets) { } } -TYPED_TEST(CropLayerTest, TestGradientNum) { +TYPED_TEST(CropLayerTest, TestCropHW) { typedef typename TypeParam::Dtype Dtype; LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(2); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(2); CropLayer layer(layer_param); - - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); - - // Copy top data into diff - caffe_copy(this->blob_top_->count(), this->blob_top_->cpu_data(), - this->blob_top_->mutable_cpu_diff()); - - // Do backward pass - vector propagate_down(2, true); - layer.Backward(this->blob_top_vec_, propagate_down, this->blob_bottom_vec_0_); - - - // Check results + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { for (int w = 0; w < this->blob_bottom_0_->width(); ++w) { - if ( n < this->blob_top_->shape(0) && + if (n < this->blob_top_->shape(0) && c < this->blob_top_->shape(1) && h < this->blob_top_->shape(2) && - w < this->blob_top_->shape(3) ) { - EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), - this->blob_bottom_0_->data_at(n, c, h, w)); - } else { - EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), 0); + w < this->blob_top_->shape(3)) { + EXPECT_EQ(this->blob_top_->data_at(n, c, h, w), + this->blob_bottom_0_->data_at(n, c, h+1, w+2)); } } } @@ -183,41 +169,50 @@ TYPED_TEST(CropLayerTest, TestGradientNum) { } } -TYPED_TEST(CropLayerTest, TestGradientNumOffset) { +TYPED_TEST(CropLayerTest, TestCrop5D) { typedef typename TypeParam::Dtype Dtype; + // Add dimension to each bottom for >4D check + vector bottom_0_shape = this->blob_bottom_0_->shape(); + vector bottom_1_shape = this->blob_bottom_1_->shape(); + bottom_0_shape.push_back(2); + bottom_1_shape.push_back(1); + this->blob_bottom_0_->Reshape(bottom_0_shape); + this->blob_bottom_1_->Reshape(bottom_1_shape); + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_0_); + filler.Fill(this->blob_bottom_1_); + // Make layer LayerParameter layer_param; - layer_param.mutable_crop_param()->set_axis(0); - layer_param.mutable_crop_param()->add_offset(0); - layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->set_axis(2); layer_param.mutable_crop_param()->add_offset(1); layer_param.mutable_crop_param()->add_offset(2); + layer_param.mutable_crop_param()->add_offset(0); CropLayer layer(layer_param); - - layer.SetUp(this->blob_bottom_vec_0_, this->blob_top_vec_); - layer.Forward(this->blob_bottom_vec_0_, this->blob_top_vec_); - - // Copy top data into diff - caffe_copy(this->blob_top_->count(), this->blob_top_->cpu_data(), - this->blob_top_->mutable_cpu_diff()); - - // Do backward pass - vector propagate_down(2, true); - layer.Backward(this->blob_top_vec_, propagate_down, this->blob_bottom_vec_0_); - - - // Check results - for (int n = 0; n < this->blob_bottom_0_->num(); ++n) { - for (int c = 0; c < this->blob_bottom_0_->channels(); ++c) { - for (int h = 0; h < this->blob_bottom_0_->height(); ++h) { - for (int w = 0; w < this->blob_bottom_0_->width(); ++w) { - if ( 0 <= n && n < 0 + this->blob_top_->shape(0) && - 1 <= c && c < 1 + this->blob_top_->shape(1) && - 1 <= h && h < 1 + this->blob_top_->shape(2) && - 2 <= w && w < 2 + this->blob_top_->shape(3) ) { - EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), - this->blob_bottom_0_->data_at(n, c, h, w)); - } else { - EXPECT_EQ(this->blob_bottom_0_->diff_at(n, c, h, w), 0); + layer.SetUp(this->blob_bottom_vec_, this->blob_top_vec_); + layer.Forward(this->blob_bottom_vec_, this->blob_top_vec_); + vector bottom_idx = vector(5, 0); + vector top_idx = vector(5, 0); + for (int n = 0; n < this->blob_bottom_0_->shape(0); ++n) { + for (int c = 0; c < this->blob_bottom_0_->shape(1); ++c) { + for (int z = 0; z < this->blob_bottom_0_->shape(2); ++z) { + for (int h = 0; h < this->blob_bottom_0_->shape(3); ++h) { + for (int w = 0; w < this->blob_bottom_0_->shape(4); ++w) { + if (n < this->blob_top_->shape(0) && + c < this->blob_top_->shape(1) && + z < this->blob_top_->shape(2) && + h < this->blob_top_->shape(3) && + w < this->blob_top_->shape(4)) { + bottom_idx[0] = top_idx[0] = n; + bottom_idx[1] = top_idx[1] = c; + bottom_idx[2] = z; + bottom_idx[3] = h; + bottom_idx[4] = top_idx[4] = w; + top_idx[2] = z+1; + top_idx[3] = h+2; + EXPECT_EQ(this->blob_top_->data_at(bottom_idx), + this->blob_bottom_0_->data_at(top_idx)); + } } } } @@ -225,4 +220,46 @@ TYPED_TEST(CropLayerTest, TestGradientNumOffset) { } } +TYPED_TEST(CropLayerTest, TestCropAllGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(0); + CropLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(CropLayerTest, TestCropHWGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(2); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(2); + CropLayer layer(layer_param); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + +TYPED_TEST(CropLayerTest, TestCrop5DGradient) { + typedef typename TypeParam::Dtype Dtype; + LayerParameter layer_param; + layer_param.mutable_crop_param()->set_axis(2); + layer_param.mutable_crop_param()->add_offset(1); + layer_param.mutable_crop_param()->add_offset(2); + layer_param.mutable_crop_param()->add_offset(0); + CropLayer layer(layer_param); + // Add dimension to each bottom for >4D check + vector bottom_0_shape = this->blob_bottom_0_->shape(); + vector bottom_1_shape = this->blob_bottom_1_->shape(); + bottom_0_shape.push_back(2); + bottom_1_shape.push_back(1); + this->blob_bottom_0_->Reshape(bottom_0_shape); + this->blob_bottom_1_->Reshape(bottom_1_shape); + GradientChecker checker(1e-2, 1e-3); + checker.CheckGradientExhaustive(&layer, this->blob_bottom_vec_, + this->blob_top_vec_); +} + } // namespace caffe