提交 9401f737 编写于 作者: Z Zhi

add cpu reverse.

上级 c941fcd8
#include "caffe/common_layers.hpp"
namespace caffe {
template <typename Dtype>
void reverse_cpu(const int count, const Dtype* from_data, Dtype* to_data,
const int* counts, const int axis_count, const int axis) {
for(int index=0; index<count; index++) {
int ind=(index/counts[axis])%axis_count;
int to_index=counts[axis]*(axis_count-2*ind-1)+index;
*(to_data+to_index)=*(from_data+index);
}
}
template <typename Dtype>
void ReverseLayer<Dtype>::LayerSetUp(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
......@@ -33,13 +41,20 @@ void ReverseLayer<Dtype>::Reshape(const vector<Blob<Dtype>*>& bottom,
template <typename Dtype>
void ReverseLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
NOT_IMPLEMENTED;
reverse_cpu<Dtype>(bottom[0]->count(), bottom[0]->cpu_data(),
top[0]->mutable_cpu_data(), bottom_counts_.cpu_data(),
bottom[0]->shape(axis_), axis_);
}
template <typename Dtype>
void ReverseLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
NOT_IMPLEMENTED;
if (!propagate_down[0]) {
return;
}
reverse_cpu<Dtype>(bottom[0]->count(), top[0]->cpu_diff(),
bottom[0]->mutable_cpu_diff(), bottom_counts_.cpu_data(),
bottom[0]->shape(axis_), axis_);
}
INSTANTIATE_CLASS(ReverseLayer);
......
......@@ -3,7 +3,7 @@
namespace caffe {
template <typename Dtype>
__global__ void Reverse(const int nthreads, const Dtype* from_data, Dtype* to_data,
__global__ void reverse_gpu(const int nthreads, const Dtype* from_data, Dtype* to_data,
const int* counts, const int axis_count, const int axis) {
CUDA_KERNEL_LOOP(index, nthreads) {
int ind=(index/counts[axis])%axis_count;
......@@ -16,7 +16,7 @@ template <typename Dtype>
void ReverseLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
const vector<Blob<Dtype>*>& top) {
const int nthreads=bottom[0]->count();
Reverse<Dtype> // NOLINT_NEXT_LINE(whitespace/operators)
reverse_gpu<Dtype> // NOLINT_NEXT_LINE(whitespace/operators)
<<<CAFFE_GET_BLOCKS(nthreads), CAFFE_CUDA_NUM_THREADS>>>(
nthreads, bottom[0]->gpu_data(), top[0]->mutable_gpu_data(),
bottom_counts_.gpu_data(), bottom[0]->shape(axis_), axis_);
......@@ -29,7 +29,7 @@ void ReverseLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
return;
}
const int nthreads=bottom[0]->count();
Reverse<Dtype> // NOLINT_NEXT_LINE(whitespace/operators)
reverse_gpu<Dtype> // NOLINT_NEXT_LINE(whitespace/operators)
<<<CAFFE_GET_BLOCKS(nthreads), CAFFE_CUDA_NUM_THREADS>>>(
nthreads, top[0]->gpu_diff(), bottom[0]->mutable_gpu_diff(),
bottom_counts_.gpu_data(), bottom[0]->shape(axis_), axis_);
......
......@@ -13,8 +13,6 @@
#include "caffe/test/test_gradient_check_util.hpp"
namespace caffe {
typedef ::testing::Types<GPUDevice<float>, GPUDevice<double> > TestDtypesGPU;
template <typename TypeParam>
class ReverseLayerTest : public MultiDeviceTest<TypeParam> {
typedef typename TypeParam::Dtype Dtype;
......@@ -36,7 +34,7 @@ namespace caffe {
vector<Blob<Dtype>*> blob_top_vec_;
};
TYPED_TEST_CASE(ReverseLayerTest, TestDtypesGPU);
TYPED_TEST_CASE(ReverseLayerTest, TestDtypesAndDevices);
TYPED_TEST(ReverseLayerTest, TestGradient) {
typedef typename TypeParam::Dtype Dtype;
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册