diff --git a/include/caffe/layers/crop_layer.hpp b/include/caffe/layers/crop_layer.hpp new file mode 100644 index 00000000000..5c605b2ae9e --- /dev/null +++ b/include/caffe/layers/crop_layer.hpp @@ -0,0 +1,67 @@ +#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, 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. + */ + +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); + + 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 new file mode 100644 index 00000000000..e81bdd732f3 --- /dev/null +++ b/src/caffe/layers/crop_layer.cpp @@ -0,0 +1,150 @@ +#include +#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) { + // 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."; + 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(); + int input_dim = bottom[0]->num_axes(); + 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()); + + // apply crops + for (int i = 0; i < input_dim; ++i) { + int crop_offset = 0; + int new_size = bottom[0]->shape(i); + 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. + 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 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, + 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(); + crop_copy(bottom, top, offsets, indices, 0, bottom_data, top_data, true); +} + +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); + std::vector indices(top[0]->num_axes(), 0); + crop_copy(bottom, top, offsets, indices, 0, top_diff, bottom_diff, false); + } +} + +#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..9ed8f7cce57 --- /dev/null +++ b/src/caffe/layers/crop_layer.cu @@ -0,0 +1,124 @@ +#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]; + } + } +} + +// 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); + 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); +} + +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(); + + if (propagate_down[0]) { + caffe_gpu_set(bottom[0]->count(), static_cast(0), bottom_diff); + std::vector indices(top[0]->num_axes(), 0); + crop_copy_gpu(bottom, top, offsets, indices, 0, top_diff, bottom_diff, + false); + } +} + +INSTANTIATE_LAYER_GPU_FUNCS(CropLayer); + +} // namespace caffe diff --git a/src/caffe/proto/caffe.proto b/src/caffe/proto/caffe.proto index 3b27bbd94d2..6900bb71482 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,24 @@ message ConvolutionParameter { optional bool force_nd_im2col = 17 [default = false]; } +message CropParameter { + // 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, + // 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; +} + message DataParameter { enum DB { LEVELDB = 0; @@ -672,7 +691,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]; } diff --git a/src/caffe/test/test_crop_layer.cpp b/src/caffe/test/test_crop_layer.cpp new file mode 100644 index 00000000000..45f24e2ee8d --- /dev/null +++ b/src/caffe/test/test_crop_layer.cpp @@ -0,0 +1,265 @@ +#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, 4, 5, 4)), + blob_bottom_1_(new Blob(2, 3, 4, 2)), + blob_top_(new Blob()) {} + virtual void SetUp() { + // fill the values + FillerParameter filler_param; + GaussianFiller filler(filler_param); + filler.Fill(this->blob_bottom_0_); + filler.Fill(this->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_); + } + + 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_; + 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_, 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_, 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_, 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, 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_, 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) && + 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, TestCropAllOffset) { + 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_, 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) && + 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+1, h+1, w+2)); + } + } + } + } + } +} + +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_, 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) && + 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+1, w+2)); + } + } + } + } + } +} + +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(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_, 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)); + } + } + } + } + } + } +} + +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