From 7a4660130668e51cdfe60233c23ed50422f84b75 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 20 Dec 2018 11:36:10 +0100 Subject: [PATCH] yolo3 layer 15 --- include/Layer.h | 2 +- src/Shortcut.cpp | 6 ++-- tests/yolo3_berkeley/yolo3_berkeley.cpp | 46 ++++++++++++++++++++----- 3 files changed, 42 insertions(+), 12 deletions(-) diff --git a/include/Layer.h b/include/Layer.h index 6d59a03..5b9ace5 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -293,7 +293,7 @@ public: class Shortcut : public Layer { public: - Shortcut(Network *net, Layer *backLayer, int layers_n); + Shortcut(Network *net, Layer *backLayer); virtual ~Shortcut(); virtual layerType_t getLayerType() { return LAYER_SHORTCUT; }; diff --git a/src/Shortcut.cpp b/src/Shortcut.cpp index 3c40a8e..1d817fe 100644 --- a/src/Shortcut.cpp +++ b/src/Shortcut.cpp @@ -5,7 +5,7 @@ namespace tk { namespace dnn { -Shortcut::Shortcut(Network *net, Layer *backLayer, int layers_n) : Layer(net) { +Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) { this->backLayer = backLayer; checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); @@ -18,8 +18,10 @@ Shortcut::~Shortcut() { dnnType* Shortcut::infer(dataDim_t &dim, dnnType* srcData) { + dataDim_t bdim = this->backLayer->output_dim; - + checkCuda(cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + shortcutForward(this->backLayer->dstData, dstData, dim.n, dim.c, dim.h, dim.w, 1, bdim.n, bdim.c, bdim.h, bdim.w, 1); //update data dimensions dim = output_dim; diff --git a/tests/yolo3_berkeley/yolo3_berkeley.cpp b/tests/yolo3_berkeley/yolo3_berkeley.cpp index 2cdadfc..718551e 100644 --- a/tests/yolo3_berkeley/yolo3_berkeley.cpp +++ b/tests/yolo3_berkeley/yolo3_berkeley.cpp @@ -6,7 +6,15 @@ const char *c0_bin = "../tests/yolo3_berkeley/layers/c0.bin"; const char *c1_bin = "../tests/yolo3_berkeley/layers/c1.bin"; const char *c2_bin = "../tests/yolo3_berkeley/layers/c2.bin"; const char *c3_bin = "../tests/yolo3_berkeley/layers/c3.bin"; -const char *output_bin = "../tests/yolo3_berkeley/debug/layer3_out.bin"; +const char *c5_bin = "../tests/yolo3_berkeley/layers/c5.bin"; +const char *c6_bin = "../tests/yolo3_berkeley/layers/c6.bin"; +const char *c7_bin = "../tests/yolo3_berkeley/layers/c7.bin"; +const char *c9_bin = "../tests/yolo3_berkeley/layers/c9.bin"; +const char *c10_bin = "../tests/yolo3_berkeley/layers/c10.bin"; +const char *c12_bin = "../tests/yolo3_berkeley/layers/c12.bin"; +const char *c13_bin = "../tests/yolo3_berkeley/layers/c13.bin"; +const char *c14_bin = "../tests/yolo3_berkeley/layers/c14.bin"; +const char *output_bin = "../tests/yolo3_berkeley/debug/layer15_out.bin"; int main() { @@ -14,14 +22,34 @@ int main() { tk::dnn::dataDim_t dim(1, 3, 320, 544, 1); tk::dnn::Network net(dim); - tk::dnn::Conv2d c0 (&net, 32, 3, 3, 1, 1, 1, 1, c0_bin, true); - tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c1 (&net, 64, 3, 3, 2, 2, 1, 1, c1_bin, true); - tk::dnn::Activation a1 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c2 (&net, 32, 1, 1, 1, 1, 0, 0, c2_bin, true); - tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c3 (&net, 64, 3, 3, 1, 1, 1, 1, c3_bin, true); - tk::dnn::Activation a3 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c0 (&net, 32, 3, 3, 1, 1, 1, 1, c0_bin, true); + tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c1 (&net, 64, 3, 3, 2, 2, 1, 1, c1_bin, true); + tk::dnn::Activation a1 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c2 (&net, 32, 1, 1, 1, 1, 0, 0, c2_bin, true); + tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c3 (&net, 64, 3, 3, 1, 1, 1, 1, c3_bin, true); + tk::dnn::Activation a3 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Shortcut s4 (&net, &a1); + tk::dnn::Conv2d c5 (&net, 128, 3, 3, 2, 2, 1, 1, c5_bin, true); + tk::dnn::Activation a5 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c6 (&net, 64, 1, 1, 1, 1, 0, 0, c6_bin, true); + tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c7 (&net, 128, 3, 3, 1, 1, 1, 1, c7_bin, true); + tk::dnn::Activation a7 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Shortcut s8 (&net, &a5); + tk::dnn::Conv2d c9 (&net, 64, 1, 1, 1, 1, 0, 0, c9_bin, true); + tk::dnn::Activation a9 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c10 (&net, 128, 3, 3, 1, 1, 1, 1, c10_bin, true); + tk::dnn::Activation a10 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Shortcut s11 (&net, &s8); + tk::dnn::Conv2d c12 (&net, 256, 3, 3, 2, 2, 1, 1, c12_bin, true); + tk::dnn::Activation a12 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c13 (&net, 128, 1, 1, 1, 1, 0, 0, c13_bin, true); + tk::dnn::Activation a13 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c14 (&net, 256, 3, 3, 1, 1, 1, 1, c14_bin, true); + tk::dnn::Activation a14 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Shortcut s15 (&net, &a12); // Load input dnnType *data;