Conv 在BM1682上的实现

saber/funcs/impl/cuda/base/cuda_c/saber_unpool.cu

@@ -0,0 +1,57 @@
+#include "saber/funcs/impl/cuda/saber_unpool.h"
+#include "cuda_fp16.h"
+
+namespace anakin {
+namespace saber {
+
+template <typename Dtype>
+__global__ void ker_unpool_max_fwd(Dtype * out_data, \
+                    const Dtype* in_data,
+                    const Dtype* in_max_index,
+                    const int in_n_stride,
+                    const int in_c_stride,
+                    const int out_n_stride,
+                    const int out_c_stride,
+                    const int in_n, 
+                    const int in_c,
+                    const int num_threads)
+{
+    CUDA_KERNEL_LOOP(tid, num_threads){
+        int n = tid / in_n_stride;
+        int c = (tid / in_c_stride) % in_c;
+        int out_offset = n * out_n_stride + c * out_c_stride; 
+        int index = in_max_index[tid];
+        out_data[out_offset + index] = in_data[tid];
+    }
+}
+
+template <DataType OpDtype>
+SaberStatus SaberUnpool<NV, OpDtype>::dispatch(\
+    const std::vector<Tensor<NV> *>& inputs, \
+    std::vector<Tensor<NV> *>& outputs,\
+    PoolingParam<NV>& param) {
+
+    const InDataType* in_data = (const InDataType*)inputs[0]->data();
+    const OutDataType* in_max_index = (const OutDataType*)inputs[1]->data();
+    OutDataType* out_data = (OutDataType*)outputs[0]->mutable_data();
+    cudaStream_t cuda_stream = this->_ctx->get_compute_stream();
+    int count = inputs[0]->valid_size();
+    int in_n = inputs[0]->num();
+    int in_c = inputs[0]->channel();
+    if (inputs[0]->is_continue_mem() && outputs[0]->is_continue_mem()) {
+        cudaMemsetAsync(out_data, 0, sizeof(InDataType) * outputs[0]->valid_size(), cuda_stream);
+        ker_unpool_max_fwd<InDataType>\
+                 <<<CUDA_GET_BLOCKS(count), CUDA_NUM_THREADS, 0, cuda_stream>>>(\
+                 out_data,
+                 in_data, in_max_index,\
+                 _in_n_stride, _in_c_stride,\
+                 _out_n_stride, _out_c_stride,\
+                 in_n, in_c, count);
+    }
+
+    return SaberSuccess;
+}
+
+} //namespace saber
+
+} //namespace anakin

saber/funcs/impl/cuda/saber_unpool.h

@@ -0,0 +1,85 @@
+/* Copyright (c) 2018 Anakin Authors, Inc. All Rights Reserved.
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+   
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License. 
+*/
+
+#ifndef ANAKIN_SABER_FUNCS_IMPL_CUDA_SABER_UNPOOL_H
+#define ANAKIN_SABER_FUNCS_IMPL_CUDA_SABER_UNPOOL_H
+
+#include "saber/funcs/impl/impl_unpool.h"
+
+namespace anakin{
+
+namespace saber{
+
+template <DataType OpDtype>
+class SaberUnpool<NV, OpDtype>:\
+    public ImplBase<NV, OpDtype, PoolingParam<NV>> {
+
+public:
+    typedef Tensor<NV> DataTensor_in;
+    typedef Tensor<NV> DataTensor_out;
+    typedef Tensor<NV> OpTensor;
+
+    typedef typename DataTrait<NV, OpDtype>::Dtype InDataType;
+    typedef typename DataTrait<NV, OpDtype>::Dtype OutDataType;
+    typedef typename DataTrait<NV, OpDtype>::Dtype OpDataType;
+
+    SaberUnpool()
+    {}
+
+    ~SaberUnpool() {
+
+    }
+
+    virtual SaberStatus init(const std::vector<DataTensor_in*>& inputs,
+                             std::vector<DataTensor_out*>& outputs,
+                             PoolingParam<NV> &param,
+                             Context<NV> &ctx)  {
+        this->_ctx = &ctx;
+        return create(inputs, outputs, param, ctx);
+    }
+
+    virtual SaberStatus create(const std::vector<DataTensor_in*>& inputs,
+                               std::vector<DataTensor_out*>& outputs,
+                               PoolingParam<NV> &param,
+                               Context<NV> &ctx)  {
+        Shape out_stride = outputs[0]->get_stride();
+        Shape in_stride = inputs[0]->get_stride();
+        int in_n_index = inputs[0]->num_index();
+        int in_c_index = inputs[0]->channel_index();
+        int out_n_index = outputs[0]->num_index();
+        int out_c_index = outputs[0]->channel_index();
+        _in_n_stride = in_stride[in_n_index];
+        _in_c_stride = in_stride[in_c_index];
+        _out_n_stride = out_stride[out_n_index];
+        _out_c_stride = out_stride[out_c_index];
+        return SaberSuccess;
+    }
+
+    virtual SaberStatus dispatch(const std::vector<DataTensor_in*>& inputs,
+                                 std::vector<DataTensor_out*>& outputs,
+                                 PoolingParam<NV> &param);
+
+private:
+    int _in_n_stride;
+    int _in_c_stride;
+    int _out_n_stride;
+    int _out_c_stride;
+};
+template class SaberUnpool<NV, AK_FLOAT>;
+}
+
+}
+
+#endif //ANAKIN_SABER_FUNCS_SABER_CONV2D_H

saber/funcs/unpool.h

@@ -19,7 +19,7 @@
 #include "saber/funcs/impl/impl_base.h"
 #include "saber/funcs/impl/impl_unpool.h"
 #ifdef NVIDIA_GPU
-//#include "saber/funcs/impl/cuda/saber_unpool.h"
+#include "saber/funcs/impl/cuda/saber_unpool.h"
 #endif
 
 namespace anakin {
@@ -95,4 +95,4 @@ class Unpool : public BaseFunc<
 
 }
 
-#endif //ANAKIN_SABER_FUNCS_UNPOOL_H
+#endif //ANAKIN_SABER_FUNCS_UNPOOL_H

test/saber/test_saber_base.h

@@ -176,7 +176,7 @@ class TestSaberBase{
         clear_datas();
         std::vector<Shape> shape_v;
         for (int i=0; i<_op_input_num; ++i){
-            shape_v.push_back(input[0] -> valid_shape());
+            shape_v.push_back(input[i] -> valid_shape());
         }
         add_inputs_shape(shape_v);
         for(int i = 0; i < _op_input_num; ++i)

test/saber/test_saber_unpool.cpp

@@ -0,0 +1,96 @@
+
+#include "saber/core/context.h"
+#include "saber/funcs/unpool.h"
+#include "test_saber_func.h"
+#include "test_saber_base.h"
+#include "saber/core/tensor_op.h"
+#include "saber/saber_types.h"
+#include <vector>
+
+#include <chrono>
+
+
+using namespace anakin::saber;
+
+template <typename dtype,typename TargetType_D,typename TargetType_H>
+void unpool_cpu(const std::vector<Tensor<TargetType_H>*>& input,std::vector<Tensor<TargetType_H>*>& output,\
+                PoolingParam<TargetType_D>& param){
+    const dtype* in_data = (const dtype*)input[0]->data();
+    const dtype* max_data = (const dtype*)input[1]->data();
+    dtype* out_data = (dtype*)output[0]->mutable_data();
+    memset(out_data, 0, output[0]->valid_size() * sizeof(dtype));
+    int num_in = input[0]->num();
+    int c_in = input[0]->channel();
+    int num_stride_in = input[0]->get_stride()[input[0]->num_index()];
+    int c_stride_in = input[0]->get_stride()[input[0]->channel_index()];
+    int num_stride_out = output[0]->get_stride()[output[0]->num_index()];
+    int c_stride_out = output[0]->get_stride()[output[0]->channel_index()];
+    int total_count = input[0]->valid_size();
+    for(int i = 0; i < total_count; ++i){
+        int num_out = i / num_stride_in;
+        int c_out = (i / c_stride_in) % c_in;
+        int out_index = num_out * num_stride_out + c_out * c_stride_out;
+        int max_index = max_data[i];
+        out_data[out_index + max_index] = in_data[i];
+    }
+
+}
+template<typename TargetType>
+void create_pooling_index(float* pool_index, int n, int c, int h, int w, PoolingParam<TargetType> param){
+    int out_w = (w - 1) * param.stride_w + param.window_w - 2 * param.pad_w;
+    int count  = 0;
+    for (int n_id = 0; n_id < n; n_id++) {
+        for (int c_id = 0; c_id < c; c_id++) {
+            for (int h_id = 0; h_id < h; h_id++) {
+                for (int w_id = 0; w_id < w; w_id++) {
+                    pool_index[count] = (h_id * param.stride_h + rand() % param.window_h) * out_w +
+                                   (w_id * param.stride_w + rand() % param.window_w);
+                    count++;
+                }
+            }
+        }
+    }
+}
+
+TEST(TestSaberFunc, test_func_resize){
+
+#ifdef USE_CUDA
+    //Init the test_base
+    TestSaberBase<NV,NVHX86,AK_FLOAT,Unpool, PoolingParam> testbase(2);
+    for(int num_in:{1,3,32}){
+        for(int c_in:{1,3,128}){
+            for(int h_in:{2,3,256}){
+                for(int w_in:{2,3,64}){
+
+                    PoolingParam<NV> param(3, 3, 0, 0, 3, 3,Pooling_max, false);
+                    Tensor<NV> td, td_pool_index;
+                    Tensor<X86> th_pool_index;
+                    Shape sh_in({num_in, c_in, h_in, w_in});
+                    td.re_alloc(sh_in, AK_FLOAT);
+                    th_pool_index.re_alloc(sh_in, AK_FLOAT);
+                    td_pool_index.re_alloc(sh_in, AK_FLOAT);
+                    fill_tensor_rand(td, -1.0, 1.0);
+                    float* pool_index_data = (float*)th_pool_index.mutable_data();
+                    create_pooling_index(pool_index_data, num_in, c_in, h_in, w_in, param);
+                    td_pool_index.copy_from(th_pool_index);
+                    std::vector<Tensor<NV>*> input;
+                    input.push_back(&td);
+                    input.push_back(&td_pool_index);
+                    testbase.add_custom_input(input);
+                    testbase.set_param(param);
+                    testbase.run_test(unpool_cpu<float, NV, NVHX86>);
+                }
+            }
+        }
+    }
+
+
+#endif
+}
+int main(int argc, const char** argv) {
+    // initial logger
+    //logger::init(argv[0]);
+    InitTest();
+    RUN_ALL_TESTS(argv[0]);
+    return 0;
+}