From ab6d2d1766ea6761e79715a4c3c2eb60a19dc399 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Fri, 12 Jun 2020 11:39:46 +0200 Subject: [PATCH 01/14] Update README Signed-off-by: Micaela Verucchi --- README.md | 30 ++++++++++++++++++++++++++++-- 1 file changed, 28 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 1175465..19b98a2 100644 --- a/README.md +++ b/README.md @@ -6,14 +6,40 @@ The main goal of this project is to exploit NVIDIA boards as much as possible to If you use tkDNN in your research, please cite one of the following papers. For use in commercial solutions, write at gattifrancesco@hotmail.it or refer to https://hipert.unimore.it/ . ``` -Accepted paper @ IRC 2020, will soon been published. +Accepted paper @ IRC 2020, will soon be published. M. Verucchi, L. Bartoli, F. Bagni, F. Gatti, P. Burgio and M. Bertogna, "Real-Time clustering and LiDAR-camera fusion on embedded platforms for self-driving cars", in proceedings in IEEE Robotic Computing (2020) -Accepted paper @ ETFA 2020, will soon been published. +Accepted paper @ ETFA 2020, will soon be published. M. Verucchi, G. Brilli, D. Sapienza, M. Verasani, M. Arena, F. Gatti, A. Capotondi, R. Cavicchioli, M. Bertogna, M. Solieri "A Systematic Assessment of Embedded Neural Networks for Object Detection", in IEEE International Conference on Emerging Technologies and Factory Automation (2020) ``` +## Results +Inference FPS of yolov4 with tkDNN, average of 1200 images with the same dimesion as the input size, on + * RTX 2080Ti (CUDA 10.2, TensorRT 7.0.0, Cudnn 7.6.5); + * Xavier AGX, Jetpack 4.3 (CUDA 10.0, CUDNN 7.6.3, tensorrt 6.0.1 ); + * Tx2, Jetpack 4.2 (CUDA 10.0, CUDNN 7.3.1, tensorrt 5.0.6 ); + * Jetson Nano, Jetpack 4.4 (CUDA 10.2, CUDNN 8.0.0, tensorrt 7.1.0 ). + +| Platform | Network | FP32, B=1 | FP32, B=4 | FP16, B=1 | FP16, B=4 | INT8, B=1 | INT8, B=4 | +| :------: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | :-----: | +| RTX 2080Ti | yolo4 320 | 118,59 |237,31 | 207,81 | 443,32 | 262,37 | 530,93 | +| RTX 2080Ti | yolo4 416 | 104,81 |162,86 | 169,06 | 293,78 | 206,93 | 353,26 | +| RTX 2080Ti | yolo4 512 | 92,98 |132,43 | 140,36 | 215,17 | 165,35 | 254,96 | +| RTX 2080Ti | yolo4 608 | 63,77 |81,53 | 111,39 | 152,89 | 127,79 | 184,72 | +| AGX Xavier | yolo4 320 | 26,78 |32,05 | 57,14 | 79,05 | 73,15 | 97,56 | +| AGX Xavier | yolo4 416 | 19,96 |21,52 | 41,01 | 49,00 | 50,81 | 60,61 | +| AGX Xavier | yolo4 512 | 16,58 |16,98 | 31,12 | 33,84 | 37,82 | 41,28 | +| AGX Xavier | yolo4 608 | 9,45 |10,13 | 21,92 | 23,36 | 27,05 | 28,93 | +| Tx2 | yolo4 320 | 11,18 | 12,07 | 15,32 | 16,31 | - | - | +| Tx2 | yolo4 416 | 7,30 | 7,58 | 9,45 | 9,90 | - | - | +| Tx2 | yolo4 512 | 5,96 | 5,95 | 7,22 | 7,23 | - | - | +| Tx2 | yolo4 608 | 3,63 | 3,65 | 4,67 | 4,70 | - | - | +| Nano | yolo4 320 | 4,23 | 4,55 | 6,14 | 6,53 | - | - | +| Nano | yolo4 416 | 2,88 | 3,00 | 3,90 | 4,04 | - | - | +| Nano | yolo4 512 | 2,32 | 2,34 | 3,02 | 3,04 | - | - | +| Nano | yolo4 608 | 1,40 | 1,41 | 1,92 | 1,93 | - | - | + ## Index - [tkDNN](#tkdnn) - [Index](#index) From 567dc0f75d5951eb677f95a14e8d3d9216d90634 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Sun, 14 Jun 2020 12:57:29 +0200 Subject: [PATCH 02/14] cmake cudnn fix --- CMakeLists.txt | 2 + cmake/FindCUDNN.cmake | 91 +++++++++++++++++++++++++++++-------------- 2 files changed, 64 insertions(+), 29 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 376c175..4a64372 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -20,6 +20,8 @@ SET(CUDA_SEPARABLE_COMPILATION ON) set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32) find_package(CUDNN REQUIRED) +include_directories(${CUDNN_INCLUDE_DIR}) + # compile file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu") diff --git a/cmake/FindCUDNN.cmake b/cmake/FindCUDNN.cmake index f240fcb..583b4a6 100644 --- a/cmake/FindCUDNN.cmake +++ b/cmake/FindCUDNN.cmake @@ -1,33 +1,66 @@ -# Find the header files +# find the library +if(CUDA_FOUND) + find_cuda_helper_libs(cudnn) + set(CUDNN_LIBRARY ${CUDA_cudnn_LIBRARY} CACHE FILEPATH "location of the cuDNN library") + unset(CUDA_cudnn_LIBRARY CACHE) -find_path(CUDNN_INCLUDE_DIR - ${CMAKE_SYSROOT}/usr/local/include - ${CMAKE_SYSROOT}/usr/include - /usr/local/nvidia/tensorrt/include/ - NO_DEFAULT_PATH -) + find_cuda_helper_libs(nvinfer) + set(NVINFER_LIBRARY ${CUDA_nvinfer_LIBRARY} CACHE FILEPATH "location of the nvinfer library") + unset(CUDA_nvinfer_LIBRARY CACHE) +endif() -set(OLD_ROOT ${CMAKE_FIND_ROOT_PATH}) -list(APPEND CMAKE_FIND_ROOT_PATH /) -list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.7) -list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.5) -find_library(CUDNN_LIB - NAMES cudnn - PATHS - /usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib - /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/ +# find the include +if(CUDNN_LIBRARY) + find_path(CUDNN_INCLUDE_DIR + cudnn.h + PATHS ${CUDA_TOOLKIT_INCLUDE} + DOC "location of cudnn.h" NO_DEFAULT_PATH -) -find_library(CUDNN_NVLIB - NAMES "nvinfer" - PATHS - /usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib - /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/ - NO_DEFAULT_PATH -) -set(CMAKE_FIND_ROOT_PATH ${OLD_ROOT}) + ) -set(CUDNN_LIBRARIES ${CUDNN_LIB} ${CUDNN_NVLIB}) -message("-- Found CUDNN: " ${CUDNN_LIB}) -message("-- Found NVINFER: " ${CUDNN_NVLIB}) -set(CUDNN_FOUND true) + if(NOT CUDNN_INCLUDE_DIR) + find_path(CUDNN_INCLUDE_DIR + cudnn.h + DOC "location of cudnn.h" + ) + endif() + + message("-- Found CUDNN: " ${CUDNN_LIBRARY}) + message("-- Found CUDNN include: " ${CUDNN_INCLUDE_DIR}) +endif() + +if(NVINFER_LIBRARY) + find_path(NVINFER_INCLUDE_DIR + NvInfer.h + PATHS ${CUDA_TOOLKIT_INCLUDE} + DOC "location of NvInfer.h" + NO_DEFAULT_PATH + ) + + if(NOT NVINFER_INCLUDE_DIR) + find_path(NVINFER_INCLUDE_DIR + NvInfer.h + DOC "location of NvInfer.h" + ) + endif() + + message("-- Found NVINFER: " ${NVINFER_LIBRARY}) + message("-- Found NVINFER include: " ${NVINFER_INCLUDE_DIR}) +endif() + + +include(FindPackageHandleStandardArgs) +find_package_handle_standard_args(CUDNN + FOUND_VAR CUDNN_FOUND + REQUIRED_VARS + CUDNN_LIBRARY + CUDNN_INCLUDE_DIR + VERSION_VAR CUDNN_VERSION +) + +if(CUDNN_FOUND) + set(CUDNN_LIBRARIES ${CUDNN_LIBRARY} ${NVINFER_LIBRARY}) + set(CUDNN_INCLUDE_DIRS ${CUDNN_INCLUDE_DIR} ${NVINFER_INCLUDE_DIR}) +endif() + +set(CUDNN_FOUND true) \ No newline at end of file From cbfc8ea4f2d8d52763b36e2c8ad2db4d4aa9d621 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Sun, 14 Jun 2020 13:01:45 +0200 Subject: [PATCH 03/14] serialize fix --- src/NetworkRT.cpp | 2 +- tests/darknet/yolo3.cpp | 14 +++++++------- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 9f53b06..6e86de1 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -595,7 +595,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { bool NetworkRT::serialize(const char *filename) { - std::ofstream p(filename); + std::ofstream p(filename, std::ios::binary); if (!p) { FatalError("could not open plan output file"); return false; diff --git a/tests/darknet/yolo3.cpp b/tests/darknet/yolo3.cpp index ea53b84..464c3cf 100644 --- a/tests/darknet/yolo3.cpp +++ b/tests/darknet/yolo3.cpp @@ -5,19 +5,19 @@ #include "DarknetParser.h" int main() { - std::string bin_path = "yolo3"; + std::string bin_path = "yolov3-sppx"; std::vector input_bins = { bin_path + "/layers/input.bin" }; std::vector output_bins = { - bin_path + "/debug/layer82_out.bin", - bin_path + "/debug/layer94_out.bin", - bin_path + "/debug/layer106_out.bin" + bin_path + "/debug/layer89_out.bin", + bin_path + "/debug/layer101_out.bin", + bin_path + "/debug/layer113_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; - downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); + std::string cfg_path = "../tests/darknet/cfg/yolov3-sppx.cfg"; + std::string name_path = "../tests/darknet/names/coco4.names"; + //downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); // parse darknet network tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); From 6dff675db7a9f58798535dfce8029b2dddac174b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Sun, 14 Jun 2020 13:03:48 +0200 Subject: [PATCH 04/14] fix wrong commit --- tests/darknet/yolo3.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/tests/darknet/yolo3.cpp b/tests/darknet/yolo3.cpp index 464c3cf..e73d226 100644 --- a/tests/darknet/yolo3.cpp +++ b/tests/darknet/yolo3.cpp @@ -5,19 +5,19 @@ #include "DarknetParser.h" int main() { - std::string bin_path = "yolov3-sppx"; + std::string bin_path = "yolo3"; std::vector input_bins = { bin_path + "/layers/input.bin" }; std::vector output_bins = { - bin_path + "/debug/layer89_out.bin", - bin_path + "/debug/layer101_out.bin", - bin_path + "/debug/layer113_out.bin" + bin_path + "/debug/layer82_out.bin", + bin_path + "/debug/layer94_out.bin", + bin_path + "/debug/layer106_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolov3-sppx.cfg"; - std::string name_path = "../tests/darknet/names/coco4.names"; - //downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); + std::string cfg_path = "../tests/darknet/cfg/yolo3.cfg"; + std::string name_path = "../tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); // parse darknet network tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); @@ -31,4 +31,4 @@ int main() { delete net; delete netRT; return ret; -} +} \ No newline at end of file From 285c77d6dd22fba207ec9c5fb5a43b3551e1e62a Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Mon, 15 Jun 2020 17:15:37 +0200 Subject: [PATCH 05/14] Imptove relative paths Signed-off-by: Micaela Verucchi Francesco Gatti --- CMakeLists.txt | 1 + tests/darknet/csresnext50-panet-spp.cpp | 4 ++-- tests/darknet/csresnext50-panet-spp_berkeley.cpp | 4 ++-- tests/darknet/yolo2.cpp | 4 ++-- tests/darknet/yolo2_voc.cpp | 4 ++-- tests/darknet/yolo2tiny.cpp | 4 ++-- tests/darknet/yolo3.cpp | 4 ++-- tests/darknet/yolo3_512.cpp | 4 ++-- tests/darknet/yolo3_berkeley.cpp | 4 ++-- tests/darknet/yolo3_coco4.cpp | 4 ++-- tests/darknet/yolo3_flir.cpp | 4 ++-- tests/darknet/yolo3tiny.cpp | 4 ++-- tests/darknet/yolo3tiny_512.cpp | 4 ++-- tests/darknet/yolo4.cpp | 4 ++-- tests/darknet/yolo4_berkeley.cpp | 4 ++-- 15 files changed, 29 insertions(+), 28 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 4a64372..8c8619d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -10,6 +10,7 @@ if(DEBUG) add_definitions(-DDEBUG) endif() +add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}") #------------------------------------------------------------------------------- # CUDA diff --git a/tests/darknet/csresnext50-panet-spp.cpp b/tests/darknet/csresnext50-panet-spp.cpp index 1da95b2..a366e14 100644 --- a/tests/darknet/csresnext50-panet-spp.cpp +++ b/tests/darknet/csresnext50-panet-spp.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer137_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/csresnext50-panet-spp.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/csresnext50-panet-spp.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/Kcs4xBozwY4wFx8/download"); // parse darknet network diff --git a/tests/darknet/csresnext50-panet-spp_berkeley.cpp b/tests/darknet/csresnext50-panet-spp_berkeley.cpp index 47bbfd4..3cd0d52 100644 --- a/tests/darknet/csresnext50-panet-spp_berkeley.cpp +++ b/tests/darknet/csresnext50-panet-spp_berkeley.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer137_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/csresnext50-panet-spp_berkeley.cfg"; - std::string name_path = "../tests/darknet/names/berkeley.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/csresnext50-panet-spp_berkeley.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/berkeley.names"; // FIXME: wrong weights // downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s//download"); diff --git a/tests/darknet/yolo2.cpp b/tests/darknet/yolo2.cpp index 7a46c31..978c137 100644 --- a/tests/darknet/yolo2.cpp +++ b/tests/darknet/yolo2.cpp @@ -13,8 +13,8 @@ int main() { bin_path + "/layers/output.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo2.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo2.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/nf4PJ3k8bxBETwL/download"); // parse darknet network diff --git a/tests/darknet/yolo2_voc.cpp b/tests/darknet/yolo2_voc.cpp index eab215b..94111e6 100644 --- a/tests/darknet/yolo2_voc.cpp +++ b/tests/darknet/yolo2_voc.cpp @@ -13,8 +13,8 @@ int main() { bin_path + "/layers/output.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo2_voc.cfg"; - std::string name_path = "../tests/darknet/names/voc.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo2_voc.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/voc.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/DJC5Fi2pEjfNDP9/download"); // parse darknet network diff --git a/tests/darknet/yolo2tiny.cpp b/tests/darknet/yolo2tiny.cpp index 64faa36..cc12109 100644 --- a/tests/darknet/yolo2tiny.cpp +++ b/tests/darknet/yolo2tiny.cpp @@ -13,8 +13,8 @@ int main() { bin_path + "/layers/output.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo2tiny.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo2tiny.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; // FIXME: wrong weights //downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s//download"); diff --git a/tests/darknet/yolo3.cpp b/tests/darknet/yolo3.cpp index e73d226..d9a684b 100644 --- a/tests/darknet/yolo3.cpp +++ b/tests/darknet/yolo3.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer106_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); // parse darknet network diff --git a/tests/darknet/yolo3_512.cpp b/tests/darknet/yolo3_512.cpp index a67c3a7..6a5c20e 100644 --- a/tests/darknet/yolo3_512.cpp +++ b/tests/darknet/yolo3_512.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer106_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3_512.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_512.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/RGecMeGLD4cXEWL/download"); // parse darknet network diff --git a/tests/darknet/yolo3_berkeley.cpp b/tests/darknet/yolo3_berkeley.cpp index a71fe83..016a8a2 100644 --- a/tests/darknet/yolo3_berkeley.cpp +++ b/tests/darknet/yolo3_berkeley.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer106_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3_berkeley.cfg"; - std::string name_path = "../tests/darknet/names/berkeley.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_berkeley.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/berkeley.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/o5cHa4AjTKS64oD/download"); // parse darknet network diff --git a/tests/darknet/yolo3_coco4.cpp b/tests/darknet/yolo3_coco4.cpp index a651430..eaf9bd8 100644 --- a/tests/darknet/yolo3_coco4.cpp +++ b/tests/darknet/yolo3_coco4.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer106_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3_coco4.cfg"; - std::string name_path = "../tests/darknet/names/coco4.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_coco4.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco4.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/o27NDzSAartbyc4/download"); // parse darknet network diff --git a/tests/darknet/yolo3_flir.cpp b/tests/darknet/yolo3_flir.cpp index 678f10f..24aac7f 100644 --- a/tests/darknet/yolo3_flir.cpp +++ b/tests/darknet/yolo3_flir.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer106_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3_flir.cfg"; - std::string name_path = "../tests/darknet/names/flir.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_flir.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/flir.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/62DECncmF6bMMiH/download"); // parse darknet network diff --git a/tests/darknet/yolo3tiny.cpp b/tests/darknet/yolo3tiny.cpp index 01fc6f9..c33f7a8 100644 --- a/tests/darknet/yolo3tiny.cpp +++ b/tests/darknet/yolo3tiny.cpp @@ -14,8 +14,8 @@ int main() { bin_path + "/debug/layer23_out.bin", }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3tiny.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3tiny.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/LMcSHtWaLeps8yN/download"); // parse darknet network diff --git a/tests/darknet/yolo3tiny_512.cpp b/tests/darknet/yolo3tiny_512.cpp index 8153b0d..ce4ce86 100644 --- a/tests/darknet/yolo3tiny_512.cpp +++ b/tests/darknet/yolo3tiny_512.cpp @@ -14,8 +14,8 @@ int main() { bin_path + "/debug/layer23_out.bin", }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo3tiny_512.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3tiny_512.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/8Zt6bHwHADqP4JC/download"); // parse darknet network diff --git a/tests/darknet/yolo4.cpp b/tests/darknet/yolo4.cpp index 80b12ce..65ac6ee 100644 --- a/tests/darknet/yolo4.cpp +++ b/tests/darknet/yolo4.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer161_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo4.cfg"; - std::string name_path = "../tests/darknet/names/coco.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download"); // parse darknet network diff --git a/tests/darknet/yolo4_berkeley.cpp b/tests/darknet/yolo4_berkeley.cpp index 8642eb1..89e9f04 100644 --- a/tests/darknet/yolo4_berkeley.cpp +++ b/tests/darknet/yolo4_berkeley.cpp @@ -15,8 +15,8 @@ int main() { bin_path + "/debug/layer161_out.bin" }; std::string wgs_path = bin_path + "/layers"; - std::string cfg_path = "../tests/darknet/cfg/yolo4_berkeley.cfg"; - std::string name_path = "../tests/darknet/names/berkeley.names"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4_berkeley.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/berkeley.names"; downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/nkWFa5fgb4NTdnB/download"); // parse darknet network From 3d3a2427c9d17e6bc7ac83b56486568b1adba180 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 16 Jun 2020 12:36:45 +0200 Subject: [PATCH 06/14] viz yolo3 --- include/tkDNN/DarknetParser.h | 2 +- include/tkDNN/NetworkViz.h | 12 ++++++ include/tkDNN/utils.h | 4 ++ src/NetworkViz.cpp | 69 +++++++++++++++++++++++++++++++++++ tests/darknet/viz_yolo3.cpp | 54 +++++++++++++++++++++++++++ 5 files changed, 140 insertions(+), 1 deletion(-) create mode 100644 include/tkDNN/NetworkViz.h create mode 100644 src/NetworkViz.cpp create mode 100644 tests/darknet/viz_yolo3.cpp diff --git a/include/tkDNN/DarknetParser.h b/include/tkDNN/DarknetParser.h index cbeed48..cd7fa48 100644 --- a/include/tkDNN/DarknetParser.h +++ b/include/tkDNN/DarknetParser.h @@ -1,6 +1,6 @@ #pragma once #include -#include "tkdnn.h" +#include "tkDNN/tkdnn.h" namespace tk { namespace dnn { diff --git a/include/tkDNN/NetworkViz.h b/include/tkDNN/NetworkViz.h new file mode 100644 index 0000000..c8b1bea --- /dev/null +++ b/include/tkDNN/NetworkViz.h @@ -0,0 +1,12 @@ +#pragma once +#include +#include +#include "tkdnn.h" + +namespace tk { namespace dnn { + +cv::Mat vizFloat2colorMap(cv::Mat map); +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim); +cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim = 1000); + +}} diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index aa73e9e..538a3f3 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -118,4 +118,8 @@ void matrixMulAdd( cublasHandle_t handle, dnnType* srcData, dnnType* dstData, void getMemUsage(double& vm_usage_kb, double& resident_set_kb); void printCudaMemUsage(); void removePathAndExtension(const std::string &full_string, std::string &name); +static inline bool isCudaPointer(void *data) { + cudaPointerAttributes attr; + return cudaPointerGetAttributes(&attr, data) == 0; +} #endif //UTILS_H diff --git a/src/NetworkViz.cpp b/src/NetworkViz.cpp new file mode 100644 index 0000000..6ac274c --- /dev/null +++ b/src/NetworkViz.cpp @@ -0,0 +1,69 @@ +#include +#include +#include +#include +#include "tkDNN/NetworkViz.h" + +namespace tk { namespace dnn { + +cv::Mat vizFloat2colorMap(cv::Mat map) { + + double min; + double max; + cv::minMaxIdx(map, &min, &max); + cv::Mat adjMap; + // expand your range to 0..255. Similar to histEq(); + map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); + //return adjMap; + + + cv::Mat falseColorsMap; + applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_HOT); + return falseColorsMap; +} + +cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) { + dnnType *data = nullptr; + + // copy to CPU + if(isCudaPointer(dataInput)) { + data = new dnnType[dim.tot()]; + checkCuda( cudaMemcpy(data, dataInput, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) ); + } else { + data = dataInput; + } + + int gridDim = ceil(sqrt(dim.c)); + cv::Size gridSize(dim.w*gridDim, dim.h*gridDim); + cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0)); + + for(int i=0; i= net->num_layers) + FatalError("Could not viz layer\n"); + return vizData2Mat(net->layers[layer]->dstData, net->layers[layer]->output_dim, imgdim); + + //cv::imwrite("viz/layer" + std::to_string(layer) + ".png", viz); + //cv::imshow("layer", viz); + //cv::waitKey(0); +} + +}} \ No newline at end of file diff --git a/tests/darknet/viz_yolo3.cpp b/tests/darknet/viz_yolo3.cpp new file mode 100644 index 0000000..374f7dd --- /dev/null +++ b/tests/darknet/viz_yolo3.cpp @@ -0,0 +1,54 @@ +#include +#include +#include + +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" +#include "NetworkViz.h" + +int main() { + std::string bin_path = "yolo3"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + // Load input and infer + dnnType *input_d; + dnnType *input_h; + readBinaryFile(input_bins[0], net->input_dim.tot(), &input_h, &input_d); + tk::dnn::dataDim_t dim = net->input_dim; + + dim.print(); + std::cout<<"infer\n"; + net->infer(dim, input_d); + + // output directory + std::string output_viz = "viz/"; + system( (std::string("mkdir -p ") + output_viz).c_str() ); + + for(int i=0; inum_layers; i++) { + std::string output_png = output_viz + "/layer" + std::to_string(i) + ".png"; + std::cout<<"saving "<releaseLayers(); + delete net; + return 0; +} + + \ No newline at end of file From 1b8f45703f0f42a5a9bd1ca45fdda556345984c2 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 16 Jun 2020 12:41:48 +0200 Subject: [PATCH 07/14] darknet parser cpp --- include/tkDNN/DarknetParser.h | 272 ++-------------------------------- src/DarknetParser.cpp | 261 ++++++++++++++++++++++++++++++++ 2 files changed, 274 insertions(+), 259 deletions(-) create mode 100644 src/DarknetParser.cpp diff --git a/include/tkDNN/DarknetParser.h b/include/tkDNN/DarknetParser.h index cd7fa48..f36469b 100644 --- a/include/tkDNN/DarknetParser.h +++ b/include/tkDNN/DarknetParser.h @@ -27,267 +27,21 @@ namespace tk { namespace dnn { std::vector layers; std::string activation = "linear"; + friend std::ostream& operator<<(std::ostream& os, const darknetFields_t& f){ + os << f.width << " " << f.height << " " << f.channels << " " << f.batch_normalize<< " " << f.filters << " " << f.activation<< " " << f.scale_xy; + return os; + } }; - std::ostream& operator<<(std::ostream& os, const darknetFields_t& f){ - os << f.width << " " << f.height << " " << f.channels << " " << f.batch_normalize<< " " << f.filters << " " << f.activation<< " " << f.scale_xy; - return os; - } - - std::string darknetParseType(const std::string& line){ - size_t start = line.find("["); - size_t end = line.find("]"); - if( start == std::string::npos || end == std::string::npos) - return ""; - start++; - std::string type = line.substr(start, end-start); - return type; - } - - bool divideNameAndValue(const std::string& line, std::string&name, std::string& value){ - size_t sep = line.find("="); - if(sep == std::string::npos) - return false; - - name = line.substr(0, sep); - value = line.substr(sep+1, line.size() - (sep+1)); - return true; - } - - std::vector fromStringToIntVec(const std::string& line, const char delimiter){ - std::stringstream linestream(line); - std::string value; - std::vector values; - - while(getline(linestream,value,delimiter)) - values.push_back(std::stoi(value)); - return values; - } - - bool darknetParseFields(const std::string& line, darknetFields_t& fields){ - - std::string name,value; - if(!divideNameAndValue(line, name, value)) - return false; - if(name.find("width") != std::string::npos) - fields.width = std::stoi(value); - else if(name.find("height") != std::string::npos) - fields.height = std::stoi(value); - else if(name.find("channels") != std::string::npos) - fields.channels = std::stoi(value); - else if(name.find("batch_normalize") != std::string::npos) - fields.batch_normalize = std::stoi(value); - else if(name.find("filters") != std::string::npos) - fields.filters = std::stoi(value); - else if(name.find("activation") != std::string::npos) - fields.activation = value; - else if(name.find("size") != std::string::npos){ - fields.size_x = std::stoi(value); - fields.size_y = std::stoi(value); - } - else if(name.find("size_x") != std::string::npos) - fields.size_x = std::stoi(value); - else if(name.find("size_y") != std::string::npos) - fields.size_y = std::stoi(value); - else if(name.find("stride") != std::string::npos){ - fields.stride_x = std::stoi(value); - fields.stride_y = std::stoi(value); - } - else if(name.find("stride_x") != std::string::npos) - fields.stride_x = std::stoi(value); - else if(name.find("stride_y") != std::string::npos) - fields.stride_y = std::stoi(value); - else if(name.find("pad") != std::string::npos) - fields.pad = std::stoi(value); - else if(name.find("classes") != std::string::npos) - fields.classes = std::stoi(value); - else if(name.find("num") != std::string::npos) - fields.num = std::stoi(value); - else if(name.find("coords") != std::string::npos) - fields.coords = std::stoi(value); - else if(name.find("groups") != std::string::npos) - fields.groups = std::stoi(value); - else if(name.find("scale_x_y") != std::string::npos) - fields.scale_xy = std::stof(value); - else if(name.find("from") != std::string::npos) - fields.layers.push_back(std::stof(value)); - else if(name.find("mask") != std::string::npos){ - auto vec = fromStringToIntVec(value, ','); - fields.n_mask = vec.size(); - } - else if(name.find("layers") != std::string::npos) - fields.layers = fromStringToIntVec(value, ','); - - else - std::cout<<"Not supported field: "< &netLayers, const std::vector& names) { - if(net == nullptr) - FatalError("Cant add a layer without a Net\n"); - - // padding compute - if(f.pad == 1) { - f.padding_x = f.padding_y = f.size_x /2; - } - //std::cout<<"Add layer: "<= netLayers.size()) FatalError("impossible to shortcut\n"); - //std::cout<<"shortcut to "<getLayerName()<<"\n"; - netLayers.push_back(new tk::dnn::Shortcut(net, netLayers[layerIdx])); - - } else if(f.type == "upsample") { - netLayers.push_back(new tk::dnn::Upsample(net, f.stride_x)); - - } else if(f.type == "route") { - if(f.layers.size() == 0) FatalError("no layers to Route\n"); - std::vector layers; - for(int i=0; i= netLayers.size()) FatalError("impossible to route\n"); - //std::cout<<"Route to "<getLayerName()<<"\n"; - layers.push_back(netLayers[layerIdx]); - } - netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size())); - - } else if(f.type == "reorg") { - netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x)); - - } else if(f.type == "region") { - netLayers.push_back(new tk::dnn::Region(net, f.classes, f.coords, f.num)); - - } else if(f.type == "yolo") { - std::string wgs = wgs_path + "/g" + std::to_string(netLayers.size()) + ".bin"; - //printf("%d %d %s %d %f\n", f.classes, f.num/f.n_mask, wgs.c_str(), f.n_mask, f.scale_xy); - tk::dnn::Yolo *l = new tk::dnn::Yolo(net, f.classes, f.num/f.n_mask, wgs, f.n_mask, f.scale_xy); - if(names.size() != f.classes) - FatalError("Mismatch between number of classes and names"); - l->classesNames = names; - netLayers.push_back(l); - - } else{ - FatalError("layer not supported: " + f.type); - } - - // add activation - if(netLayers.size() > 0 && f.activation != "linear") { - tkdnnActivationMode_t act; - if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); - else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; - else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; - else { FatalError("activation not supported: " + f.activation); } - netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); - }; - } - - std::vector darknetReadNames(const std::string& names_file){ - std::ifstream if_names(names_file); - if(!if_names.is_open()) - FatalError("cloud not open names file: " + names_file); - - std::vector names; - std::string line; - while(std::getline(if_names, line)) - if(line != "") - names.push_back(line); - - if_names.close(); - return names; - } - - tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file) { - - tk::dnn::Network *net = nullptr; - - // layers without activations to retrive correct id number - std::vector netLayers; - - std::ifstream if_cfg(cfg_file); - if(!if_cfg.is_open()) - FatalError("cloud not open cfg file: " + cfg_file); - - std::vector names = darknetReadNames(names_file); - - darknetFields_t fields; // will be filled with layers fields - std::string line; - while(std::getline(if_cfg, line)) { - // remove comments - std::size_t found = line.find("#"); - if ( found != std::string::npos ) { - line = line.substr(0, found); - } - - // skip empty lines - if(line.size() == 0) - continue; - - std::string type = darknetParseType(line); - if(type.size() > 0) { - // end of filled type - if(fields.type != "") { - if(fields.type == "net") - net = darknetAddNet(fields); - else - darknetAddLayer(net, fields, wgs_path, netLayers, names); - } - - // new type - //std::cout<<"type: "< fromStringToIntVec(const std::string& line, const char delimiter); + bool darknetParseFields(const std::string& line, darknetFields_t& fields); + tk::dnn::Network *darknetAddNet(darknetFields_t &fields); + void darknetAddLayer(tk::dnn::Network *net, darknetFields_t &f, std::string wgs_path, + std::vector &netLayers, const std::vector& names); + std::vector darknetReadNames(const std::string& names_file); + tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file); }} diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp new file mode 100644 index 0000000..5092595 --- /dev/null +++ b/src/DarknetParser.cpp @@ -0,0 +1,261 @@ +#include "tkDNN/DarknetParser.h" + +namespace tk { namespace dnn { + + std::string darknetParseType(const std::string& line){ + size_t start = line.find("["); + size_t end = line.find("]"); + if( start == std::string::npos || end == std::string::npos) + return ""; + start++; + std::string type = line.substr(start, end-start); + return type; + } + + bool divideNameAndValue(const std::string& line, std::string&name, std::string& value){ + size_t sep = line.find("="); + if(sep == std::string::npos) + return false; + + name = line.substr(0, sep); + value = line.substr(sep+1, line.size() - (sep+1)); + return true; + } + + std::vector fromStringToIntVec(const std::string& line, const char delimiter){ + std::stringstream linestream(line); + std::string value; + std::vector values; + + while(getline(linestream,value,delimiter)) + values.push_back(std::stoi(value)); + return values; + } + + bool darknetParseFields(const std::string& line, darknetFields_t& fields){ + + std::string name,value; + if(!divideNameAndValue(line, name, value)) + return false; + if(name.find("width") != std::string::npos) + fields.width = std::stoi(value); + else if(name.find("height") != std::string::npos) + fields.height = std::stoi(value); + else if(name.find("channels") != std::string::npos) + fields.channels = std::stoi(value); + else if(name.find("batch_normalize") != std::string::npos) + fields.batch_normalize = std::stoi(value); + else if(name.find("filters") != std::string::npos) + fields.filters = std::stoi(value); + else if(name.find("activation") != std::string::npos) + fields.activation = value; + else if(name.find("size") != std::string::npos){ + fields.size_x = std::stoi(value); + fields.size_y = std::stoi(value); + } + else if(name.find("size_x") != std::string::npos) + fields.size_x = std::stoi(value); + else if(name.find("size_y") != std::string::npos) + fields.size_y = std::stoi(value); + else if(name.find("stride") != std::string::npos){ + fields.stride_x = std::stoi(value); + fields.stride_y = std::stoi(value); + } + else if(name.find("stride_x") != std::string::npos) + fields.stride_x = std::stoi(value); + else if(name.find("stride_y") != std::string::npos) + fields.stride_y = std::stoi(value); + else if(name.find("pad") != std::string::npos) + fields.pad = std::stoi(value); + else if(name.find("classes") != std::string::npos) + fields.classes = std::stoi(value); + else if(name.find("num") != std::string::npos) + fields.num = std::stoi(value); + else if(name.find("coords") != std::string::npos) + fields.coords = std::stoi(value); + else if(name.find("groups") != std::string::npos) + fields.groups = std::stoi(value); + else if(name.find("scale_x_y") != std::string::npos) + fields.scale_xy = std::stof(value); + else if(name.find("from") != std::string::npos) + fields.layers.push_back(std::stof(value)); + else if(name.find("mask") != std::string::npos){ + auto vec = fromStringToIntVec(value, ','); + fields.n_mask = vec.size(); + } + else if(name.find("layers") != std::string::npos) + fields.layers = fromStringToIntVec(value, ','); + + else + std::cout<<"Not supported field: "< &netLayers, const std::vector& names) { + if(net == nullptr) + FatalError("Cant add a layer without a Net\n"); + + // padding compute + if(f.pad == 1) { + f.padding_x = f.padding_y = f.size_x /2; + } + //std::cout<<"Add layer: "<= netLayers.size()) FatalError("impossible to shortcut\n"); + //std::cout<<"shortcut to "<getLayerName()<<"\n"; + netLayers.push_back(new tk::dnn::Shortcut(net, netLayers[layerIdx])); + + } else if(f.type == "upsample") { + netLayers.push_back(new tk::dnn::Upsample(net, f.stride_x)); + + } else if(f.type == "route") { + if(f.layers.size() == 0) FatalError("no layers to Route\n"); + std::vector layers; + for(int i=0; i= netLayers.size()) FatalError("impossible to route\n"); + //std::cout<<"Route to "<getLayerName()<<"\n"; + layers.push_back(netLayers[layerIdx]); + } + netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size())); + + } else if(f.type == "reorg") { + netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x)); + + } else if(f.type == "region") { + netLayers.push_back(new tk::dnn::Region(net, f.classes, f.coords, f.num)); + + } else if(f.type == "yolo") { + std::string wgs = wgs_path + "/g" + std::to_string(netLayers.size()) + ".bin"; + //printf("%d %d %s %d %f\n", f.classes, f.num/f.n_mask, wgs.c_str(), f.n_mask, f.scale_xy); + tk::dnn::Yolo *l = new tk::dnn::Yolo(net, f.classes, f.num/f.n_mask, wgs, f.n_mask, f.scale_xy); + if(names.size() != f.classes) + FatalError("Mismatch between number of classes and names"); + l->classesNames = names; + netLayers.push_back(l); + + } else{ + FatalError("layer not supported: " + f.type); + } + + // add activation + if(netLayers.size() > 0 && f.activation != "linear") { + tkdnnActivationMode_t act; + if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); + else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; + else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; + else { FatalError("activation not supported: " + f.activation); } + netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); + }; + } + + std::vector darknetReadNames(const std::string& names_file){ + std::ifstream if_names(names_file); + if(!if_names.is_open()) + FatalError("cloud not open names file: " + names_file); + + std::vector names; + std::string line; + while(std::getline(if_names, line)) + if(line != "") + names.push_back(line); + + if_names.close(); + return names; + } + + tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file) { + + tk::dnn::Network *net = nullptr; + + // layers without activations to retrive correct id number + std::vector netLayers; + + std::ifstream if_cfg(cfg_file); + if(!if_cfg.is_open()) + FatalError("cloud not open cfg file: " + cfg_file); + + std::vector names = darknetReadNames(names_file); + + darknetFields_t fields; // will be filled with layers fields + std::string line; + while(std::getline(if_cfg, line)) { + // remove comments + std::size_t found = line.find("#"); + if ( found != std::string::npos ) { + line = line.substr(0, found); + } + + // skip empty lines + if(line.size() == 0) + continue; + + std::string type = darknetParseType(line); + if(type.size() > 0) { + // end of filled type + if(fields.type != "") { + if(fields.type == "net") + net = darknetAddNet(fields); + else + darknetAddLayer(net, fields, wgs_path, netLayers, names); + } + + // new type + //std::cout<<"type: "< Date: Tue, 16 Jun 2020 12:44:12 +0200 Subject: [PATCH 08/14] dealloc in test.h fix #36 --- include/tkDNN/test.h | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/include/tkDNN/test.h b/include/tkDNN/test.h index be3d891..4e238ff 100644 --- a/include/tkDNN/test.h +++ b/include/tkDNN/test.h @@ -1,7 +1,7 @@ #include int testInference(std::vector input_bins, std::vector output_bins, - tk::dnn::Network *net, tk::dnn::NetworkRT *netRT = nullptr) { + tk::dnn::Network *net, tk::dnn::NetworkRT *netRT = nullptr) { std::vector outputs; for(int i=0; inum_layers; i++) { @@ -67,7 +67,11 @@ int testInference(std::vector input_bins, std::vector std::cout<<"CUDNN vs TRT "; ret_cudnn_tensorrt |= checkResult(odim, cudnn_out[i], rt_out[i]) == 0 ? 0 : ERROR_CUDNNvsTENSORRT; } - } - return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt; - } \ No newline at end of file + delete [] out_h; + checkCuda( cudaFree(out) ); + } + delete [] input_h; + checkCuda( cudaFree(data) ); + return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt; +} \ No newline at end of file From 1dfc69ba8916f913282c1fb68a117cdc72468f78 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 16 Jun 2020 13:19:30 +0200 Subject: [PATCH 09/14] viz yolo3 preprocess --- tests/darknet/viz_yolo3.cpp | 36 ++++++++++++++++++++++++++---------- 1 file changed, 26 insertions(+), 10 deletions(-) diff --git a/tests/darknet/viz_yolo3.cpp b/tests/darknet/viz_yolo3.cpp index 374f7dd..9e53116 100644 --- a/tests/darknet/viz_yolo3.cpp +++ b/tests/darknet/viz_yolo3.cpp @@ -1,32 +1,49 @@ #include #include #include +#include #include "tkdnn.h" #include "test.h" #include "DarknetParser.h" #include "NetworkViz.h" -int main() { +int main(int argc, char *argv[]) { + if(argc <2) + FatalError("you must provide an input image"); + std::string input_image = argv[1]; std::string bin_path = "yolo3"; - std::vector input_bins = { - bin_path + "/layers/input.bin" - }; std::string wgs_path = bin_path + "/layers"; std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3.cfg"; std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; - downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); + downloadWeightsifDoNotExist(wgs_path, bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download"); // parse darknet network tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); net->print(); - // Load input and infer + // input data dnnType *input_d; - dnnType *input_h; - readBinaryFile(input_bins[0], net->input_dim.tot(), &input_h, &input_d); - tk::dnn::dataDim_t dim = net->input_dim; + checkCuda( cudaMalloc(&input_d, sizeof(dnnType)*net->input_dim.tot())); + + // load image + cv::Mat frame, frameFloat; + frame = cv::imread(input_image); + cv::resize(frame, frame, cv::Size(net->input_dim.w, net->input_dim.h)); + frame.convertTo(frameFloat, CV_32FC3, 1/255.0); + + //split channels + cv::Mat bgr[3]; + cv::split(frameFloat,bgr);//split source + + //write channels + for(int i=0; iinput_dim.c; i++) { + int idx = i*frameFloat.rows*frameFloat.cols; + int ch = net->input_dim.c-1 -i; + checkCuda( cudaMemcpy(input_d + idx, (void*)bgr[ch].data, frameFloat.rows*frameFloat.cols*sizeof(dnnType), cudaMemcpyHostToDevice)); + } + tk::dnn::dataDim_t dim = net->input_dim; dim.print(); std::cout<<"infer\n"; net->infer(dim, input_d); @@ -44,7 +61,6 @@ int main() { //cv::waitKey(0); } - delete [] input_h; checkCuda(cudaFree(input_d)); net->releaseLayers(); delete net; From 2817ade782cc13b8f42d9616e4b3a935400b5179 Mon Sep 17 00:00:00 2001 From: Omar Alvarez Date: Mon, 22 Jun 2020 13:40:15 +0200 Subject: [PATCH 10/14] Fix parsing label files with unexpected chars --- src/Int8BatchStream.cpp | 19 ++++++------------- 1 file changed, 6 insertions(+), 13 deletions(-) diff --git a/src/Int8BatchStream.cpp b/src/Int8BatchStream.cpp index dd4399f..fdc1db2 100644 --- a/src/Int8BatchStream.cpp +++ b/src/Int8BatchStream.cpp @@ -132,21 +132,14 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector& res void BatchStream::readLabels(std::string inputFileName, std::vector& ris) { std::ifstream is(inputFileName.c_str()); - //read only the first number: the image sub-portion class - while (true) { + + std::string line; + while (std::getline(is, line)) + { + std::istringstream iss(line); float val; - is >> val; - if (!is) { - break; - } - // insert the first number and skip all others + if(!(iss >> val)) { break; } // error ris.push_back(val); - while( true ) { - char c; - is >> c; - if (is.peek() == '\n') //detect "\n" - break; - } } } From 61aa24c6b7716b833c8a0840dc0c47a00ca7e8ff Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Tue, 30 Jun 2020 15:19:03 +0200 Subject: [PATCH 11/14] yolov4tiny works on CUDNN Signed-off-by: Micaela Verucchi --- include/tkDNN/DarknetParser.h | 1 + include/tkDNN/Layer.h | 4 +- src/DarknetParser.cpp | 4 +- src/Route.cpp | 10 +- tests/darknet/cfg/yolo4tiny.cfg | 281 ++++++++++++++++++++++++++++++++ tests/darknet/yolo4tiny.cpp | 33 ++++ 6 files changed, 328 insertions(+), 5 deletions(-) create mode 100644 tests/darknet/cfg/yolo4tiny.cfg create mode 100644 tests/darknet/yolo4tiny.cpp diff --git a/include/tkDNN/DarknetParser.h b/include/tkDNN/DarknetParser.h index f36469b..29d1e8e 100644 --- a/include/tkDNN/DarknetParser.h +++ b/include/tkDNN/DarknetParser.h @@ -11,6 +11,7 @@ namespace tk { namespace dnn { int channels = 3; int batch_normalize=0; int groups = 1; + int group_id = 0; int filters=1; int size_x=1; int size_y=1; diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 9bd8432..bd544b2 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -509,7 +509,7 @@ public: class Route : public Layer { public: - Route(Network *net, Layer **layers, int layers_n); + Route(Network *net, Layer **layers, int layers_n, int groups = 1, int group_id = 0); virtual ~Route(); virtual layerType_t getLayerType() { return LAYER_ROUTE; }; @@ -519,6 +519,8 @@ public: static const int MAX_LAYERS = 32; Layer *layers[MAX_LAYERS]; //ids of layers to be merged int layers_n; //number of layers + int groups; + int group_id; }; diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 5092595..7bbc7ba 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -75,6 +75,8 @@ namespace tk { namespace dnn { fields.coords = std::stoi(value); else if(name.find("groups") != std::string::npos) fields.groups = std::stoi(value); + else if(name.find("group_id") != std::string::npos) + fields.group_id = std::stoi(value); else if(name.find("scale_x_y") != std::string::npos) fields.scale_xy = std::stof(value); else if(name.find("from") != std::string::npos) @@ -148,7 +150,7 @@ namespace tk { namespace dnn { //std::cout<<"Route to "<getLayerName()<<"\n"; layers.push_back(netLayers[layerIdx]); } - netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size())); + netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size(), f.groups, f.group_id)); } else if(f.type == "reorg") { netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x)); diff --git a/src/Route.cpp b/src/Route.cpp index 39bb14e..816566e 100644 --- a/src/Route.cpp +++ b/src/Route.cpp @@ -5,7 +5,7 @@ namespace tk { namespace dnn { -Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) { +Route::Route(Network *net, Layer **layers, int layers_n, int groups, int group_id) : Layer(net) { // copy input layers if(layers_n > MAX_LAYERS) { @@ -15,6 +15,8 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) { this->layers[i] = layers[i]; } this->layers_n = layers_n; + this->groups = groups; + this->group_id = group_id; //get dims output_dim.l = 1; @@ -32,6 +34,7 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) { output_dim.c += layers[i]->output_dim.c; } + output_dim.c /= this->groups; input_dim = output_dim; checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); @@ -49,8 +52,9 @@ dnnType* Route::infer(dataDim_t &dim, dnnType* srcData) { for(int i=0; idstData; int in_dim = layers[i]->output_dim.tot(); - checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - offset += in_dim; + int part_in_dim = in_dim / this->groups; + checkCuda( cudaMemcpy(dstData + offset, input + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + offset += part_in_dim; } //update data dimensions diff --git a/tests/darknet/cfg/yolo4tiny.cfg b/tests/darknet/cfg/yolo4tiny.cfg new file mode 100644 index 0000000..dc6f5bf --- /dev/null +++ b/tests/darknet/cfg/yolo4tiny.cfg @@ -0,0 +1,281 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +subdivisions=1 +width=416 +height=416 +channels=3 +momentum=0.9 +decay=0.0005 +angle=0 +saturation = 1.5 +exposure = 1.5 +hue=.1 + +learning_rate=0.00261 +burn_in=1000 +max_batches = 500200 +policy=steps +steps=400000,450000 +scales=.1,.1 + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=2 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers=-1 +groups=2 +group_id=1 + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -6,-1 + +[maxpool] +size=2 +stride=2 + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers=-1 +groups=2 +group_id=1 + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -6,-1 + +[maxpool] +size=2 +stride=2 + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers=-1 +groups=2 +group_id=1 + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -1,-2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[route] +layers = -6,-1 + +[maxpool] +size=2 +stride=2 + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=leaky + +################################## + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=leaky + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + + + +[yolo] +mask = 3,4,5 +anchors = 10,14, 23,27, 37,58, 81,82, 135,169, 344,319 +classes=80 +num=6 +jitter=.3 +scale_x_y = 1.05 +cls_normalizer=1.0 +iou_normalizer=0.07 +iou_loss=ciou +ignore_thresh = .7 +truth_thresh = 1 +random=0 +resize=1.5 +nms_kind=greedynms +beta_nms=0.6 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[upsample] +stride=2 + +[route] +layers = -1, 23 + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + +[yolo] +mask = 1,2,3 +anchors = 10,14, 23,27, 37,58, 81,82, 135,169, 344,319 +classes=80 +num=6 +jitter=.3 +scale_x_y = 1.05 +cls_normalizer=1.0 +iou_normalizer=0.07 +iou_loss=ciou +ignore_thresh = .7 +truth_thresh = 1 +random=0 +resize=1.5 +nms_kind=greedynms +beta_nms=0.6 diff --git a/tests/darknet/yolo4tiny.cpp b/tests/darknet/yolo4tiny.cpp new file mode 100644 index 0000000..d9011a8 --- /dev/null +++ b/tests/darknet/yolo4tiny.cpp @@ -0,0 +1,33 @@ +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "yolo4tiny"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector output_bins = { + bin_path + "/debug/layer30_out.bin", + bin_path + "/debug/layer37_out.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4tiny.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download"); + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + //convert network to tensorRT + // tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + + int ret = testInference(input_bins, output_bins, net, nullptr); + net->releaseLayers(); + delete net; + // delete netRT; + return ret; +} From fe2e4eae92324d4c03482e20b5746ebc6aed75c5 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Tue, 30 Jun 2020 15:52:58 +0200 Subject: [PATCH 12/14] yolo4tiny works on tensorRT :dolphin: :dolphin: :dolphin: :dolphin: Signed-off-by: Micaela Verucchi --- include/tkDNN/NetworkRT.h | 2 +- include/tkDNN/pluginsRT/RouteRT.h | 17 ++++++++++++----- src/NetworkRT.cpp | 19 +++++++++++-------- tests/darknet/yolo4tiny.cpp | 6 +++--- 4 files changed, 27 insertions(+), 17 deletions(-) diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index ee1f728..66b4f3d 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -28,7 +28,7 @@ using namespace nvinfer1; #include "pluginsRT/ActivationMishRT.h" #include "pluginsRT/ReorgRT.h" #include "pluginsRT/RegionRT.h" -//#include "pluginsRT/RouteRT.h" +#include "pluginsRT/RouteRT.h" #include "pluginsRT/ShortcutRT.h" #include "pluginsRT/YoloRT.h" #include "pluginsRT/UpsampleRT.h" diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 0e94a97..263893f 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -8,7 +8,9 @@ class RouteRT : public IPlugin { */ public: - RouteRT() { + RouteRT(int groups, int group_id) { + this->groups = groups; + this->group_id = group_id; } ~RouteRT(){ @@ -22,7 +24,7 @@ public: Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { int out_c = 0; for(int i=0; i(inputs[i]); int in_dim = c_in[i]*h*w; - checkCuda( cudaMemcpyAsync(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); - offset += in_dim; + int part_in_dim = in_dim / this->groups; + checkCuda( cudaMemcpyAsync(dstData + offset, input + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); + offset += part_in_dim; } return 0; @@ -65,11 +69,13 @@ public: virtual size_t getSerializationSize() override { - return (4+MAX_INPUTS)*sizeof(int); + return (6+MAX_INPUTS)*sizeof(int); } virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, groups); + tk::dnn::writeBUF(buf, group_id); tk::dnn::writeBUF(buf, in); for(int i=0; iaddConcatenation(tens, l->layers_n); - //IPlugin *plugin = new RouteRT(); - //IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin); - checkNULL(lRT); + if(l->groups > 1){ + IPlugin *plugin = new RouteRT(l->groups, l->group_id); + IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin); + checkNULL(lRT); + return lRT; + } + IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n); + checkNULL(lRT); return lRT; } @@ -766,9 +769,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); return r; } -/* + if(name.find("Route") == 0) { - RouteRT *r = new RouteRT(); + RouteRT *r = new RouteRT(readBUF(buf),readBUF(buf)); r->in = readBUF(buf); for(int i=0; ic_in[i] = readBUF(buf); @@ -777,7 +780,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); return r; } -*/ + if(name.find("Deformable") == 0) { DeformableConvRT *r = new DeformableConvRT(readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), diff --git a/tests/darknet/yolo4tiny.cpp b/tests/darknet/yolo4tiny.cpp index d9011a8..44fbac8 100644 --- a/tests/darknet/yolo4tiny.cpp +++ b/tests/darknet/yolo4tiny.cpp @@ -23,11 +23,11 @@ int main() { net->print(); //convert network to tensorRT - // tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); - int ret = testInference(input_bins, output_bins, net, nullptr); + int ret = testInference(input_bins, output_bins, net, netRT); net->releaseLayers(); delete net; - // delete netRT; + delete netRT; return ret; } From 04602f395236d993d79cff7d1be5ced3ae738b0b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 30 Jun 2020 19:37:16 +0200 Subject: [PATCH 13/14] Yolo4-tiny batched fix #59 --- README.md | 2 ++ include/tkDNN/pluginsRT/RouteRT.h | 18 ++++++++++-------- scripts/test_all_tests.sh | 1 + 3 files changed, 13 insertions(+), 8 deletions(-) diff --git a/README.md b/README.md index 19b98a2..3ff6fe7 100644 --- a/README.md +++ b/README.md @@ -317,6 +317,8 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing | resnet101_cnet | Centernet (Resnet101 backend)4 | [COCO 2017](http://cocodataset.org/) | 80 | 512x512 | [weights](https://cloud.hipert.unimore.it/s/5BTjHMWBcJk8g3i/download) | | csresnext50-panet-spp | Cross Stage Partial Network 7 | [COCO 2014](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/Kcs4xBozwY4wFx8/download) | | yolo4 | Yolov4 8 | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download) | +| yolo4_berkeley | Yolov4 8 | [BDD100K ](https://bair.berkeley.edu/blog/2018/05/30/bdd/) | 10 | 540x320 | [weights](https://cloud.hipert.unimore.it/s/nkWFa5fgb4NTdnB/download) | +| yolo4tiny | Yolov4 tiny | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download) | ## References diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 263893f..23f30b7 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -52,16 +52,18 @@ public: } virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - + dnnType *dstData = reinterpret_cast(outputs[0]); - int offset = 0; - for(int i=0; i(inputs[i]); - int in_dim = c_in[i]*h*w; - int part_in_dim = in_dim / this->groups; - checkCuda( cudaMemcpyAsync(dstData + offset, input + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); - offset += part_in_dim; + for(int b=0; b(inputs[i]); + int in_dim = c_in[i]*h*w; + int part_in_dim = in_dim / this->groups; + checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) ); + offset += part_in_dim; + } } return 0; diff --git a/scripts/test_all_tests.sh b/scripts/test_all_tests.sh index 193a90c..770aa22 100644 --- a/scripts/test_all_tests.sh +++ b/scripts/test_all_tests.sh @@ -74,6 +74,7 @@ do test_net yolo4 test_net yolo4_berkeley + test_net yolo4tiny test_net yolo3 test_net yolo3_berkeley test_net yolo3_coco4 From 7c2155decfc2d225f523350d55fb6d773d3b3b6c Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Wed, 1 Jul 2020 11:56:46 +0200 Subject: [PATCH 14/14] Add weights download link for csresnext50-panet-spp_berkeley ( fix #63 ) Signed-off-by: Micaela Verucchi --- tests/darknet/csresnext50-panet-spp_berkeley.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/darknet/csresnext50-panet-spp_berkeley.cpp b/tests/darknet/csresnext50-panet-spp_berkeley.cpp index 3cd0d52..a8ba59f 100644 --- a/tests/darknet/csresnext50-panet-spp_berkeley.cpp +++ b/tests/darknet/csresnext50-panet-spp_berkeley.cpp @@ -17,8 +17,7 @@ int main() { std::string wgs_path = bin_path + "/layers"; std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/csresnext50-panet-spp_berkeley.cfg"; std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/berkeley.names"; - // FIXME: wrong weights - // downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s//download"); + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/q82qHAtqpoaFYo5/download"); // parse darknet network tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);