From 4ebbb6af2b9ccbebc957e79ca6e8a66b3a307587 Mon Sep 17 00:00:00 2001 From: Autochaffeur <228594@studenti.unimore.it> Date: Mon, 13 May 2019 17:32:18 +0200 Subject: [PATCH 01/30] README update --- README.md | 1 + 1 file changed, 1 insertion(+) diff --git a/README.md b/README.md index 1201cfb..9441e0f 100644 --- a/README.md +++ b/README.md @@ -20,6 +20,7 @@ Build with cmake mkdir build cd build cmake .. +# use -DBUILD_DEPS=False to skip dataset download make ``` during the cmake configuration it will be dowloaded the weights needed for running From 6656c3d0e863e1c8fed68a4de8d5bb3186acb819 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Fri, 28 Jun 2019 17:44:32 +0200 Subject: [PATCH 02/30] fix cmake --- CMakeLists.txt | 6 +++--- README.md | 2 +- tkDNNConfig.cmake | 24 ++++++++++++++++++++++-- 3 files changed, 26 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index fef13a0..34130fd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,15 +1,15 @@ cmake_minimum_required(VERSION 2.8) project (tkDNN) -set(BUILD_DEPS true CACHE BOOL "If true download deps") +set(TEST_DATA true CACHE BOOL "If true download deps") -if( ${BUILD_DEPS} ) +if( ${TEST_DATA} ) message("Launching pre-build dependency installer script...") execute_process (COMMAND bash -c "bash build_models.sh download" WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tests) - set(BUILD_DEPS false CACHE BOOL "If true download deps" FORCE) + set(TEST_DATA false CACHE BOOL "If true download deps" FORCE) message("Finished dowloading test weights") endif() diff --git a/README.md b/README.md index 9441e0f..b166f63 100644 --- a/README.md +++ b/README.md @@ -20,7 +20,7 @@ Build with cmake mkdir build cd build cmake .. -# use -DBUILD_DEPS=False to skip dataset download +# use -DTEST_DATA=False to skip dataset download make ``` during the cmake configuration it will be dowloaded the weights needed for running diff --git a/tkDNNConfig.cmake b/tkDNNConfig.cmake index 8665415..fd06286 100644 --- a/tkDNNConfig.cmake +++ b/tkDNNConfig.cmake @@ -1,3 +1,7 @@ +message("-- Found tkDNN") +set(CMAKE_MODULE_PATH ${CMAKE_CURRENT_LIST_DIR}) +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++11 -fPIC") + find_package(CUDA REQUIRED) find_package(OpenCV REQUIRED) find_library(NVINFER NAMES nvinfer) @@ -6,5 +10,21 @@ if(NVINFER STREQUAL "NVINFER-NOTFOUND") link_directories(/usr/local/nvidia/tensorrt/targets/x86_64-linux-gnu/lib/ /usr/local/cuda/targets/x86_64-linux/lib/) endif() -set(tkDNN_INCLUDE_DIRS ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES}) -set(tkDNN_LIBRARIES tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS}) + +set(tkDNN_INCLUDE_DIRS + ${CUDA_INCLUDE_DIRS} + ${OPENCV_INCLUDE_DIRS} + ${NVINFER_INCLUDES} +) + +set(tkDNN_LIBRARIES + tkDNN + kernels + ${CUDA_LIBRARIES} + ${CUDA_CUBLAS_LIBRARIES} + -lcudnn + -lnvinfer + ${OpenCV_LIBS} +) + +set(tkDNN_FOUND true) From f50aa4ad1a0cad5c497b7d8196fa04da4947c3a8 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Fri, 28 Jun 2019 18:51:01 +0200 Subject: [PATCH 03/30] fix cmake --- CMakeLists.txt | 62 ++++++++++++++++++++++++++++++++++++-------------- 1 file changed, 45 insertions(+), 17 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 34130fd..0e33f1e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,25 +1,23 @@ -cmake_minimum_required(VERSION 2.8) +cmake_minimum_required(VERSION 3.5) + project (tkDNN) +set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) +set(CMAKE_BUILD_TYPE "Release") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC") +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) -set(TEST_DATA true CACHE BOOL "If true download deps") - -if( ${TEST_DATA} ) - message("Launching pre-build dependency installer script...") - - execute_process (COMMAND bash -c "bash build_models.sh download" - WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tests) - - set(TEST_DATA false CACHE BOOL "If true download deps" FORCE) - message("Finished dowloading test weights") -endif() - +# project specific flags if(DEBUG) add_definitions(-DDEBUG) endif() -find_package(CUDA REQUIRED) -find_package(OpenCV REQUIRED) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") + +#------------------------------------------------------------------------------- +# CUDA +#------------------------------------------------------------------------------- +SET(CUDA_SEPARABLE_COMPILATION ON) +find_package(CUDA 9.0 REQUIRED) +#set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'") # compile Discovery only if TensorRT is installed find_library(NVINFER NAMES nvinfer) @@ -29,10 +27,22 @@ if(NVINFER STREQUAL "NVINFER-NOTFOUND") /usr/local/cuda/targets/x86_64-linux/lib/) endif() +# compile file(GLOB tkdnn_CUSRC "src/kernels/*.cu") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${NVINFER_INCLUDES}) cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) + +#------------------------------------------------------------------------------- +# External Libraries +#------------------------------------------------------------------------------- +find_package(OpenCV REQUIRED) +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") + + +#------------------------------------------------------------------------------- +# Build Libraries +#------------------------------------------------------------------------------- file(GLOB tkdnn_SRC "src/*.cpp") set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS}) @@ -88,7 +98,10 @@ target_link_libraries(test_rtinference tkDNN) add_executable(yolo3_demo demo/demo/demo.cpp) target_link_libraries(yolo3_demo tkDNN) -#install + +#------------------------------------------------------------------------------- +# Install +#------------------------------------------------------------------------------- #if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) # set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install" # CACHE PATH "default install path" FORCE) @@ -100,3 +113,18 @@ install(TARGETS tkDNN kernels DESTINATION lib) install(FILES "${CMAKE_SOURCE_DIR}/${CMAKE_PROJECT_NAME}Config.cmake" # source directory DESTINATION "share/${CMAKE_PROJECT_NAME}/cmake/" # target directory ) + +#------------------------------------------------------------------------------- +# Prepare for test +#------------------------------------------------------------------------------- +set(TEST_DATA true CACHE BOOL "If true download deps") +if( ${TEST_DATA} ) + message("Launching pre-build dependency installer script...") + + execute_process (COMMAND bash -c "bash build_models.sh download" + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tests) + + set(TEST_DATA false CACHE BOOL "If true download deps" FORCE) + message("Finished dowloading test weights") +endif() + From 041968f38a830def13c0aad2070297817340cda2 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Sat, 29 Jun 2019 11:08:30 +0200 Subject: [PATCH 04/30] cmake fix --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 0e33f1e..acf3613 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -15,8 +15,8 @@ endif() #------------------------------------------------------------------------------- # CUDA #------------------------------------------------------------------------------- -SET(CUDA_SEPARABLE_COMPILATION ON) find_package(CUDA 9.0 REQUIRED) +SET(CUDA_SEPARABLE_COMPILATION ON) #set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'") # compile Discovery only if TensorRT is installed From 8c629ebe7b5a845cc8b1cc6f1661d043ebe7b4a8 Mon Sep 17 00:00:00 2001 From: mbosi <205839@studenti.unimore.it> Date: Sat, 14 Sep 2019 19:03:13 +0200 Subject: [PATCH 05/30] string input and flir test --- include/Layer.h | 10 +- include/models/Yolo3.h | 289 +++++++++ include/utils.h | 2 +- src/Conv2d.cpp | 2 +- src/Dense.cpp | 2 +- src/LayerWgs.cpp | 2 +- src/Region.cpp | 2 +- src/Yolo.cpp | 4 +- src/Yolo3Detection.cpp | 2 +- src/utils.cpp | 2 +- tests/yolo3_berkeley/yolo3_berkeley.cpp | 308 +--------- tests/yolo3_coco4/yolo3_coco4.cpp | 308 +--------- tests/yolo3_flir/yolo3_flir.cfg | 785 ++++++++++++++++++++++++ tests/yolo3_flir/yolo3_flir.cpp | 88 +++ 14 files changed, 1198 insertions(+), 608 deletions(-) create mode 100644 include/models/Yolo3.h create mode 100644 tests/yolo3_flir/yolo3_flir.cfg create mode 100644 tests/yolo3_flir/yolo3_flir.cpp diff --git a/include/Layer.h b/include/Layer.h index c7871b5..55729ce 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -75,7 +75,7 @@ class LayerWgs : public Layer { public: LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, - const char* fname_weights, bool batchnorm = false); + std::string fname_weights, bool batchnorm = false); virtual ~LayerWgs(); int inputs, outputs; @@ -108,7 +108,7 @@ public: class Dense : public LayerWgs { public: - Dense(Network *net, int out_ch, const char* fname_weights); + Dense(Network *net, int out_ch, std::string fname_weights); virtual ~Dense(); virtual layerType_t getLayerType() { return LAYER_DENSE; }; @@ -151,7 +151,7 @@ class Conv2d : public LayerWgs { public: Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - const char* fname_weights, bool batchnorm = false); + std::string fname_weights, bool batchnorm = false); virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; @@ -352,7 +352,7 @@ public: int sort_class; }; - Yolo(Network *net, int classes, int num, const char* fname_weights); + Yolo(Network *net, int classes, int num, std::string fname_weights); virtual ~Yolo(); virtual layerType_t getLayerType() { return LAYER_YOLO; }; @@ -389,7 +389,7 @@ class RegionInterpret { public: RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, - int classes, int coords, int num, float thresh, const char* fname_weights); + int classes, int coords, int num, float thresh, std::string fname_weights); ~RegionInterpret(); dataDim_t input_dim, output_dim; diff --git a/include/models/Yolo3.h b/include/models/Yolo3.h new file mode 100644 index 0000000..cd69b32 --- /dev/null +++ b/include/models/Yolo3.h @@ -0,0 +1,289 @@ +int preYoloFilters = (classes+5)*3; + +std::string input_bin = 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" +}; +std::string c0_bin = bin_path + "/layers/c0.bin"; +std::string c1_bin = bin_path + "/layers/c1.bin"; +std::string c2_bin = bin_path + "/layers/c2.bin"; +std::string c3_bin = bin_path + "/layers/c3.bin"; +std::string c5_bin = bin_path + "/layers/c5.bin"; +std::string c6_bin = bin_path + "/layers/c6.bin"; +std::string c7_bin = bin_path + "/layers/c7.bin"; +std::string c9_bin = bin_path + "/layers/c9.bin"; +std::string c10_bin = bin_path + "/layers/c10.bin"; +std::string c12_bin = bin_path + "/layers/c12.bin"; +std::string c13_bin = bin_path + "/layers/c13.bin"; +std::string c14_bin = bin_path + "/layers/c14.bin"; +std::string c16_bin = bin_path + "/layers/c16.bin"; +std::string c17_bin = bin_path + "/layers/c17.bin"; +std::string c19_bin = bin_path + "/layers/c19.bin"; +std::string c20_bin = bin_path + "/layers/c20.bin"; +std::string c22_bin = bin_path + "/layers/c22.bin"; +std::string c23_bin = bin_path + "/layers/c23.bin"; +std::string c25_bin = bin_path + "/layers/c25.bin"; +std::string c26_bin = bin_path + "/layers/c26.bin"; +std::string c28_bin = bin_path + "/layers/c28.bin"; +std::string c29_bin = bin_path + "/layers/c29.bin"; +std::string c31_bin = bin_path + "/layers/c31.bin"; +std::string c32_bin = bin_path + "/layers/c32.bin"; +std::string c34_bin = bin_path + "/layers/c34.bin"; +std::string c35_bin = bin_path + "/layers/c35.bin"; +std::string c37_bin = bin_path + "/layers/c37.bin"; +std::string c38_bin = bin_path + "/layers/c38.bin"; +std::string c39_bin = bin_path + "/layers/c39.bin"; +std::string c41_bin = bin_path + "/layers/c41.bin"; +std::string c42_bin = bin_path + "/layers/c42.bin"; +std::string c44_bin = bin_path + "/layers/c44.bin"; +std::string c45_bin = bin_path + "/layers/c45.bin"; +std::string c47_bin = bin_path + "/layers/c47.bin"; +std::string c48_bin = bin_path + "/layers/c48.bin"; +std::string c50_bin = bin_path + "/layers/c50.bin"; +std::string c51_bin = bin_path + "/layers/c51.bin"; +std::string c53_bin = bin_path + "/layers/c53.bin"; +std::string c54_bin = bin_path + "/layers/c54.bin"; +std::string c56_bin = bin_path + "/layers/c56.bin"; +std::string c57_bin = bin_path + "/layers/c57.bin"; +std::string c59_bin = bin_path + "/layers/c59.bin"; +std::string c60_bin = bin_path + "/layers/c60.bin"; +std::string c62_bin = bin_path + "/layers/c62.bin"; +std::string c63_bin = bin_path + "/layers/c63.bin"; +std::string c64_bin = bin_path + "/layers/c64.bin"; +std::string c66_bin = bin_path + "/layers/c66.bin"; +std::string c67_bin = bin_path + "/layers/c67.bin"; +std::string c69_bin = bin_path + "/layers/c69.bin"; +std::string c70_bin = bin_path + "/layers/c70.bin"; +std::string c72_bin = bin_path + "/layers/c72.bin"; +std::string c73_bin = bin_path + "/layers/c73.bin"; +std::string c75_bin = bin_path + "/layers/c75.bin"; +std::string c76_bin = bin_path + "/layers/c76.bin"; +std::string c77_bin = bin_path + "/layers/c77.bin"; +std::string c78_bin = bin_path + "/layers/c78.bin"; +std::string c79_bin = bin_path + "/layers/c79.bin"; +std::string c80_bin = bin_path + "/layers/c80.bin"; +std::string c81_bin = bin_path + "/layers/c81.bin"; +std::string g82_bin = bin_path + "/layers/g82.bin"; +std::string c84_bin = bin_path + "/layers/c84.bin"; +std::string c87_bin = bin_path + "/layers/c87.bin"; +std::string c88_bin = bin_path + "/layers/c88.bin"; +std::string c89_bin = bin_path + "/layers/c89.bin"; +std::string c90_bin = bin_path + "/layers/c90.bin"; +std::string c91_bin = bin_path + "/layers/c91.bin"; +std::string c92_bin = bin_path + "/layers/c92.bin"; +std::string c93_bin = bin_path + "/layers/c93.bin"; +std::string g94_bin = bin_path + "/layers/g94.bin"; +std::string c96_bin = bin_path + "/layers/c96.bin"; +std::string c99_bin = bin_path + "/layers/c99.bin"; +std::string c100_bin = bin_path + "/layers/c100.bin"; +std::string c101_bin = bin_path + "/layers/c101.bin"; +std::string c102_bin = bin_path + "/layers/c102.bin"; +std::string c103_bin = bin_path + "/layers/c103.bin"; +std::string c104_bin = bin_path + "/layers/c104.bin"; +std::string c105_bin = bin_path + "/layers/c105.bin"; +std::string g106_bin = bin_path + "/layers/g106.bin"; + +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); + +tk::dnn::Conv2d c16 (&net, 128, 1, 1, 1, 1, 0, 0, c16_bin, true); +tk::dnn::Activation a16 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c17 (&net, 256, 3, 3, 1, 1, 1, 1, c17_bin, true); +tk::dnn::Activation a17 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s18 (&net, &s15); +tk::dnn::Conv2d c19 (&net, 128, 1, 1, 1, 1, 0, 0, c19_bin, true); +tk::dnn::Activation a19 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c20 (&net, 256, 3, 3, 1, 1, 1, 1, c20_bin, true); +tk::dnn::Activation a20 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s21 (&net, &s18); +tk::dnn::Conv2d c22 (&net, 128, 1, 1, 1, 1, 0, 0, c22_bin, true); +tk::dnn::Activation a22 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c23 (&net, 256, 3, 3, 1, 1, 1, 1, c23_bin, true); +tk::dnn::Activation a23 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s24 (&net, &s21); +tk::dnn::Conv2d c25 (&net, 128, 1, 1, 1, 1, 0, 0, c25_bin, true); +tk::dnn::Activation a25 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c26 (&net, 256, 3, 3, 1, 1, 1, 1, c26_bin, true); +tk::dnn::Activation a26 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s27 (&net, &s24); +tk::dnn::Conv2d c28 (&net, 128, 1, 1, 1, 1, 0, 0, c28_bin, true); +tk::dnn::Activation a28 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c29 (&net, 256, 3, 3, 1, 1, 1, 1, c29_bin, true); +tk::dnn::Activation a29 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s30 (&net, &s27); +tk::dnn::Conv2d c31 (&net, 128, 1, 1, 1, 1, 0, 0, c31_bin, true); +tk::dnn::Activation a31 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c32 (&net, 256, 3, 3, 1, 1, 1, 1, c32_bin, true); +tk::dnn::Activation a32 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s33 (&net, &s30); +tk::dnn::Conv2d c34 (&net, 128, 1, 1, 1, 1, 0, 0, c34_bin, true); +tk::dnn::Activation a34 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c35 (&net, 256, 3, 3, 1, 1, 1, 1, c35_bin, true); +tk::dnn::Activation a35 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s36 (&net, &s33); + +tk::dnn::Conv2d c37 (&net, 512, 3, 3, 2, 2, 1, 1, c37_bin, true); +tk::dnn::Activation a37 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c38 (&net, 256, 1, 1, 1, 1, 0, 0, c38_bin, true); +tk::dnn::Activation a38 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c39 (&net, 512, 3, 3, 1, 1, 1, 1, c39_bin, true); +tk::dnn::Activation a39 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s40 (&net, &a37); + +tk::dnn::Conv2d c41 (&net, 256, 1, 1, 1, 1, 0, 0, c41_bin, true); +tk::dnn::Activation a41 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c42 (&net, 512, 3, 3, 1, 1, 1, 1, c42_bin, true); +tk::dnn::Activation a42 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s43 (&net, &s40); +tk::dnn::Conv2d c44 (&net, 256, 1, 1, 1, 1, 0, 0, c44_bin, true); +tk::dnn::Activation a44 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c45 (&net, 512, 3, 3, 1, 1, 1, 1, c45_bin, true); +tk::dnn::Activation a45 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s46 (&net, &s43); +tk::dnn::Conv2d c47 (&net, 256, 1, 1, 1, 1, 0, 0, c47_bin, true); +tk::dnn::Activation a47 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c48 (&net, 512, 3, 3, 1, 1, 1, 1, c48_bin, true); +tk::dnn::Activation a48 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s49 (&net, &s46); +tk::dnn::Conv2d c50 (&net, 256, 1, 1, 1, 1, 0, 0, c50_bin, true); +tk::dnn::Activation a50 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c51 (&net, 512, 3, 3, 1, 1, 1, 1, c51_bin, true); +tk::dnn::Activation a51 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s52 (&net, &s49); +tk::dnn::Conv2d c53 (&net, 256, 1, 1, 1, 1, 0, 0, c53_bin, true); +tk::dnn::Activation a53 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c54 (&net, 512, 3, 3, 1, 1, 1, 1, c54_bin, true); +tk::dnn::Activation a54 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s55 (&net, &s52); +tk::dnn::Conv2d c56 (&net, 256, 1, 1, 1, 1, 0, 0, c56_bin, true); +tk::dnn::Activation a56 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c57 (&net, 512, 3, 3, 1, 1, 1, 1, c57_bin, true); +tk::dnn::Activation a57 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s58 (&net, &s55); +tk::dnn::Conv2d c59 (&net, 256, 1, 1, 1, 1, 0, 0, c59_bin, true); +tk::dnn::Activation a59 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c60 (&net, 512, 3, 3, 1, 1, 1, 1, c60_bin, true); +tk::dnn::Activation a60 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s61 (&net, &s58); + +tk::dnn::Conv2d c62 (&net,1024, 3, 3, 2, 2, 1, 1, c62_bin, true); +tk::dnn::Activation a62 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c63 (&net, 512, 1, 1, 1, 1, 0, 0, c63_bin, true); +tk::dnn::Activation a63 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c64 (&net,1024, 3, 3, 1, 1, 1, 1, c64_bin, true); +tk::dnn::Activation a64 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s65 (&net, &a62); + +tk::dnn::Conv2d c66 (&net, 512, 1, 1, 1, 1, 0, 0, c66_bin, true); +tk::dnn::Activation a66 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c67 (&net,1024, 3, 3, 1, 1, 1, 1, c67_bin, true); +tk::dnn::Activation a67 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s68 (&net, &s65); + +tk::dnn::Conv2d c69 (&net, 512, 1, 1, 1, 1, 0, 0, c69_bin, true); +tk::dnn::Activation a69 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c70 (&net,1024, 3, 3, 1, 1, 1, 1, c70_bin, true); +tk::dnn::Activation a70 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s71 (&net, &s68); + +tk::dnn::Conv2d c72 (&net, 512, 1, 1, 1, 1, 0, 0, c72_bin, true); +tk::dnn::Activation a72 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c73 (&net,1024, 3, 3, 1, 1, 1, 1, c73_bin, true); +tk::dnn::Activation a73 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Shortcut s74 (&net, &s71); + +tk::dnn::Conv2d c75 (&net, 512, 1, 1, 1, 1, 0, 0, c75_bin, true); +tk::dnn::Activation a75 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c76 (&net,1024, 3, 3, 1, 1, 1, 1, c76_bin, true); +tk::dnn::Activation a76 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c77 (&net, 512, 1, 1, 1, 1, 0, 0, c77_bin, true); +tk::dnn::Activation a77 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c78 (&net,1024, 3, 3, 1, 1, 1, 1, c78_bin, true); +tk::dnn::Activation a78 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c79 (&net, 512, 1, 1, 1, 1, 0, 0, c79_bin, true); +tk::dnn::Activation a79 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c80 (&net,1024, 3, 3, 1, 1, 1, 1, c80_bin, true); +tk::dnn::Activation a80 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c81 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c81_bin, false); +tk::dnn::Yolo yolo0 (&net, classes, 3, g82_bin); + +tk::dnn::Layer *m83_layers[1] = { &a79 }; +tk::dnn::Route m83 (&net, m83_layers, 1); +tk::dnn::Conv2d c84 (&net, 256, 1, 1, 1, 1, 0, 0, c84_bin, true); +tk::dnn::Activation a84 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Upsample u85 (&net, 2); + +tk::dnn::Layer *m86_layers[2] = { &u85, &s61 }; +tk::dnn::Route m86 (&net, m86_layers, 2); +tk::dnn::Conv2d c87 (&net, 256, 1, 1, 1, 1, 0, 0, c87_bin, true); +tk::dnn::Activation a87 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c88 (&net, 512, 3, 3, 1, 1, 1, 1, c88_bin, true); +tk::dnn::Activation a88 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c89 (&net, 256, 1, 1, 1, 1, 0, 0, c89_bin, true); +tk::dnn::Activation a89 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c90 (&net, 512, 3, 3, 1, 1, 1, 1, c90_bin, true); +tk::dnn::Activation a90 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c91 (&net, 256, 1, 1, 1, 1, 0, 0, c91_bin, true); +tk::dnn::Activation a91 (&net, tk::dnn::ACTIVATION_LEAKY); + +tk::dnn::Conv2d c92 (&net, 512, 3, 3, 1, 1, 1, 1, c92_bin, true); +tk::dnn::Activation a92 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c93 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c93_bin, false); +tk::dnn::Yolo yolo1 (&net, classes, 3, g94_bin); + +tk::dnn::Layer *m95_layers[1] = { &a91 }; +tk::dnn::Route m95 (&net, m95_layers, 1); +tk::dnn::Conv2d c96 (&net, 128, 1, 1, 1, 1, 0, 0, c96_bin, true); +tk::dnn::Activation a96 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Upsample u97 (&net, 2); + +tk::dnn::Layer *m98_layers[2] = { &u97, &s36 }; +tk::dnn::Route m98 (&net, m98_layers, 2); +tk::dnn::Conv2d c99 (&net, 128, 1, 1, 1, 1, 0, 0, c99_bin, true); +tk::dnn::Activation a99 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c100 (&net, 256, 3, 3, 1, 1, 1, 1, c100_bin, true); +tk::dnn::Activation a100 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c101 (&net, 128, 1, 1, 1, 1, 0, 0, c101_bin, true); +tk::dnn::Activation a101 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c102 (&net, 256, 3, 3, 1, 1, 1, 1, c102_bin, true); +tk::dnn::Activation a102 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c103 (&net, 128, 1, 1, 1, 1, 0, 0, c103_bin, true); +tk::dnn::Activation a103 (&net, tk::dnn::ACTIVATION_LEAKY); + +tk::dnn::Conv2d c104 (&net, 256, 3, 3, 1, 1, 1, 1, c104_bin, true); +tk::dnn::Activation a104 (&net, tk::dnn::ACTIVATION_LEAKY); +tk::dnn::Conv2d c105 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c105_bin, false); +tk::dnn::Yolo yolo2 (&net, classes, 3, g106_bin); + +yolo[0] = &yolo0; +yolo[1] = &yolo1; +yolo[2] = &yolo2; \ No newline at end of file diff --git a/include/utils.h b/include/utils.h index a46b2c4..dc34a31 100644 --- a/include/utils.h +++ b/include/utils.h @@ -90,7 +90,7 @@ void printCenteredTitle(const char *title, char fill, int dim); bool fileExist(const char *fname); -void readBinaryFile(const char* fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0); +void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0); int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true); void printDeviceVector(int size, dnnType* vec_d, bool device = true); void resize(int size, dnnType **data); diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index bfb8316..3dfdce6 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -6,7 +6,7 @@ namespace tk { namespace dnn { Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - const char* fname_weights, bool batchnorm) : + std::string fname_weights, bool batchnorm) : LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, fname_weights, batchnorm) { diff --git a/src/Dense.cpp b/src/Dense.cpp index f86aa4a..b6a9af2 100644 --- a/src/Dense.cpp +++ b/src/Dense.cpp @@ -4,7 +4,7 @@ namespace tk { namespace dnn { -Dense::Dense(Network *net, int out_ch, const char* fname_weights) : +Dense::Dense(Network *net, int out_ch, std::string fname_weights) : LayerWgs(net, net->getOutputDim().tot(), out_ch, 1, 1, 1, fname_weights) { output_dim.n = 1; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index d8bb881..f27da0f 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -8,7 +8,7 @@ namespace tk { namespace dnn { LayerWgs::LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kl, - const char* fname_weights, bool batchnorm) : Layer(net) { + std::string fname_weights, bool batchnorm) : Layer(net) { this->inputs = inputs; this->outputs = outputs; diff --git a/src/Region.cpp b/src/Region.cpp index 56010dd..1b52375 100644 --- a/src/Region.cpp +++ b/src/Region.cpp @@ -66,7 +66,7 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) { /* Intepret class */ RegionInterpret::RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, - int classes, int coords, int num, float thresh, const char* fname_weights) { + int classes, int coords, int num, float thresh, std::string fname_weights) { this->input_dim = input_dim; this->output_dim = output_dim; diff --git a/src/Yolo.cpp b/src/Yolo.cpp index 0dcea7e..f2ec53c 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -11,14 +11,14 @@ namespace tk { namespace dnn { -Yolo::Yolo(Network *net, int classes, int num, const char* fname_weights) : +Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights) : Layer(net) { this->classes = classes; this->num = num; // load anchors - if(fname_weights != nullptr) { + if(fname_weights != "") { int seek = 0; readBinaryFile(fname_weights, num, &mask_h, &mask_d, seek); seek += num; diff --git a/src/Yolo3Detection.cpp b/src/Yolo3Detection.cpp index e121503..ce9fd6a 100644 --- a/src/Yolo3Detection.cpp +++ b/src/Yolo3Detection.cpp @@ -32,7 +32,7 @@ bool Yolo3Detection::init(std::string tensor_path) { num = yRT->num; // make a yolo layer for interpret predictions - yolo[i] = new tk::dnn::Yolo(nullptr, classes, num, nullptr); // yolo without input and bias + yolo[i] = new tk::dnn::Yolo(nullptr, classes, num, ""); // yolo without input and bias yolo[i]->mask_h = new dnnType[num]; yolo[i]->bias_h = new dnnType[num*3*2]; memcpy(yolo[i]->mask_h, yRT->mask, sizeof(dnnType)*num); diff --git a/src/utils.cpp b/src/utils.cpp index 7955d15..e6f71f8 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -21,7 +21,7 @@ bool fileExist(const char *fname) { } -void readBinaryFile(const char* fname, int size, dnnType** data_h, dnnType** data_d, int seek) +void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek) { std::ifstream dataFile (fname, std::ios::in | std::ios::binary); std::stringstream error_s; diff --git a/tests/yolo3_berkeley/yolo3_berkeley.cpp b/tests/yolo3_berkeley/yolo3_berkeley.cpp index 3a69e32..43f5da5 100644 --- a/tests/yolo3_berkeley/yolo3_berkeley.cpp +++ b/tests/yolo3_berkeley/yolo3_berkeley.cpp @@ -1,295 +1,18 @@ #include +#include #include "tkdnn.h" -const char *input_bin = "../tests/yolo3_berkeley/layers/input.bin"; -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 *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 *c16_bin = "../tests/yolo3_berkeley/layers/c16.bin"; -const char *c17_bin = "../tests/yolo3_berkeley/layers/c17.bin"; -const char *c19_bin = "../tests/yolo3_berkeley/layers/c19.bin"; -const char *c20_bin = "../tests/yolo3_berkeley/layers/c20.bin"; -const char *c22_bin = "../tests/yolo3_berkeley/layers/c22.bin"; -const char *c23_bin = "../tests/yolo3_berkeley/layers/c23.bin"; -const char *c25_bin = "../tests/yolo3_berkeley/layers/c25.bin"; -const char *c26_bin = "../tests/yolo3_berkeley/layers/c26.bin"; -const char *c28_bin = "../tests/yolo3_berkeley/layers/c28.bin"; -const char *c29_bin = "../tests/yolo3_berkeley/layers/c29.bin"; -const char *c31_bin = "../tests/yolo3_berkeley/layers/c31.bin"; -const char *c32_bin = "../tests/yolo3_berkeley/layers/c32.bin"; -const char *c34_bin = "../tests/yolo3_berkeley/layers/c34.bin"; -const char *c35_bin = "../tests/yolo3_berkeley/layers/c35.bin"; -const char *c37_bin = "../tests/yolo3_berkeley/layers/c37.bin"; -const char *c38_bin = "../tests/yolo3_berkeley/layers/c38.bin"; -const char *c39_bin = "../tests/yolo3_berkeley/layers/c39.bin"; -const char *c41_bin = "../tests/yolo3_berkeley/layers/c41.bin"; -const char *c42_bin = "../tests/yolo3_berkeley/layers/c42.bin"; -const char *c44_bin = "../tests/yolo3_berkeley/layers/c44.bin"; -const char *c45_bin = "../tests/yolo3_berkeley/layers/c45.bin"; -const char *c47_bin = "../tests/yolo3_berkeley/layers/c47.bin"; -const char *c48_bin = "../tests/yolo3_berkeley/layers/c48.bin"; -const char *c50_bin = "../tests/yolo3_berkeley/layers/c50.bin"; -const char *c51_bin = "../tests/yolo3_berkeley/layers/c51.bin"; -const char *c53_bin = "../tests/yolo3_berkeley/layers/c53.bin"; -const char *c54_bin = "../tests/yolo3_berkeley/layers/c54.bin"; -const char *c56_bin = "../tests/yolo3_berkeley/layers/c56.bin"; -const char *c57_bin = "../tests/yolo3_berkeley/layers/c57.bin"; -const char *c59_bin = "../tests/yolo3_berkeley/layers/c59.bin"; -const char *c60_bin = "../tests/yolo3_berkeley/layers/c60.bin"; -const char *c62_bin = "../tests/yolo3_berkeley/layers/c62.bin"; -const char *c63_bin = "../tests/yolo3_berkeley/layers/c63.bin"; -const char *c64_bin = "../tests/yolo3_berkeley/layers/c64.bin"; -const char *c66_bin = "../tests/yolo3_berkeley/layers/c66.bin"; -const char *c67_bin = "../tests/yolo3_berkeley/layers/c67.bin"; -const char *c69_bin = "../tests/yolo3_berkeley/layers/c69.bin"; -const char *c70_bin = "../tests/yolo3_berkeley/layers/c70.bin"; -const char *c72_bin = "../tests/yolo3_berkeley/layers/c72.bin"; -const char *c73_bin = "../tests/yolo3_berkeley/layers/c73.bin"; -const char *c75_bin = "../tests/yolo3_berkeley/layers/c75.bin"; -const char *c76_bin = "../tests/yolo3_berkeley/layers/c76.bin"; -const char *c77_bin = "../tests/yolo3_berkeley/layers/c77.bin"; -const char *c78_bin = "../tests/yolo3_berkeley/layers/c78.bin"; -const char *c79_bin = "../tests/yolo3_berkeley/layers/c79.bin"; -const char *c80_bin = "../tests/yolo3_berkeley/layers/c80.bin"; -const char *c81_bin = "../tests/yolo3_berkeley/layers/c81.bin"; -const char *g82_bin = "../tests/yolo3_berkeley/layers/g82.bin"; -const char *c84_bin = "../tests/yolo3_berkeley/layers/c84.bin"; -const char *c87_bin = "../tests/yolo3_berkeley/layers/c87.bin"; -const char *c88_bin = "../tests/yolo3_berkeley/layers/c88.bin"; -const char *c89_bin = "../tests/yolo3_berkeley/layers/c89.bin"; -const char *c90_bin = "../tests/yolo3_berkeley/layers/c90.bin"; -const char *c91_bin = "../tests/yolo3_berkeley/layers/c91.bin"; -const char *c92_bin = "../tests/yolo3_berkeley/layers/c92.bin"; -const char *c93_bin = "../tests/yolo3_berkeley/layers/c93.bin"; -const char *g94_bin = "../tests/yolo3_berkeley/layers/g94.bin"; -const char *c96_bin = "../tests/yolo3_berkeley/layers/c96.bin"; -const char *c99_bin = "../tests/yolo3_berkeley/layers/c99.bin"; -const char *c100_bin = "../tests/yolo3_berkeley/layers/c100.bin"; -const char *c101_bin = "../tests/yolo3_berkeley/layers/c101.bin"; -const char *c102_bin = "../tests/yolo3_berkeley/layers/c102.bin"; -const char *c103_bin = "../tests/yolo3_berkeley/layers/c103.bin"; -const char *c104_bin = "../tests/yolo3_berkeley/layers/c104.bin"; -const char *c105_bin = "../tests/yolo3_berkeley/layers/c105.bin"; -const char *g106_bin = "../tests/yolo3_berkeley/layers/g106.bin"; -const char *output_bins[3] = { - "../tests/yolo3_berkeley/debug/layer82_out.bin", - "../tests/yolo3_berkeley/debug/layer94_out.bin", - "../tests/yolo3_berkeley/debug/layer106_out.bin" -}; - int main() { // Network layout 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::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); - - tk::dnn::Conv2d c16 (&net, 128, 1, 1, 1, 1, 0, 0, c16_bin, true); - tk::dnn::Activation a16 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c17 (&net, 256, 3, 3, 1, 1, 1, 1, c17_bin, true); - tk::dnn::Activation a17 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s18 (&net, &s15); - tk::dnn::Conv2d c19 (&net, 128, 1, 1, 1, 1, 0, 0, c19_bin, true); - tk::dnn::Activation a19 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c20 (&net, 256, 3, 3, 1, 1, 1, 1, c20_bin, true); - tk::dnn::Activation a20 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s21 (&net, &s18); - tk::dnn::Conv2d c22 (&net, 128, 1, 1, 1, 1, 0, 0, c22_bin, true); - tk::dnn::Activation a22 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c23 (&net, 256, 3, 3, 1, 1, 1, 1, c23_bin, true); - tk::dnn::Activation a23 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s24 (&net, &s21); - tk::dnn::Conv2d c25 (&net, 128, 1, 1, 1, 1, 0, 0, c25_bin, true); - tk::dnn::Activation a25 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c26 (&net, 256, 3, 3, 1, 1, 1, 1, c26_bin, true); - tk::dnn::Activation a26 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s27 (&net, &s24); - tk::dnn::Conv2d c28 (&net, 128, 1, 1, 1, 1, 0, 0, c28_bin, true); - tk::dnn::Activation a28 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c29 (&net, 256, 3, 3, 1, 1, 1, 1, c29_bin, true); - tk::dnn::Activation a29 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s30 (&net, &s27); - tk::dnn::Conv2d c31 (&net, 128, 1, 1, 1, 1, 0, 0, c31_bin, true); - tk::dnn::Activation a31 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c32 (&net, 256, 3, 3, 1, 1, 1, 1, c32_bin, true); - tk::dnn::Activation a32 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s33 (&net, &s30); - tk::dnn::Conv2d c34 (&net, 128, 1, 1, 1, 1, 0, 0, c34_bin, true); - tk::dnn::Activation a34 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c35 (&net, 256, 3, 3, 1, 1, 1, 1, c35_bin, true); - tk::dnn::Activation a35 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s36 (&net, &s33); - - tk::dnn::Conv2d c37 (&net, 512, 3, 3, 2, 2, 1, 1, c37_bin, true); - tk::dnn::Activation a37 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c38 (&net, 256, 1, 1, 1, 1, 0, 0, c38_bin, true); - tk::dnn::Activation a38 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c39 (&net, 512, 3, 3, 1, 1, 1, 1, c39_bin, true); - tk::dnn::Activation a39 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s40 (&net, &a37); - - tk::dnn::Conv2d c41 (&net, 256, 1, 1, 1, 1, 0, 0, c41_bin, true); - tk::dnn::Activation a41 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c42 (&net, 512, 3, 3, 1, 1, 1, 1, c42_bin, true); - tk::dnn::Activation a42 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s43 (&net, &s40); - tk::dnn::Conv2d c44 (&net, 256, 1, 1, 1, 1, 0, 0, c44_bin, true); - tk::dnn::Activation a44 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c45 (&net, 512, 3, 3, 1, 1, 1, 1, c45_bin, true); - tk::dnn::Activation a45 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s46 (&net, &s43); - tk::dnn::Conv2d c47 (&net, 256, 1, 1, 1, 1, 0, 0, c47_bin, true); - tk::dnn::Activation a47 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c48 (&net, 512, 3, 3, 1, 1, 1, 1, c48_bin, true); - tk::dnn::Activation a48 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s49 (&net, &s46); - tk::dnn::Conv2d c50 (&net, 256, 1, 1, 1, 1, 0, 0, c50_bin, true); - tk::dnn::Activation a50 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c51 (&net, 512, 3, 3, 1, 1, 1, 1, c51_bin, true); - tk::dnn::Activation a51 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s52 (&net, &s49); - tk::dnn::Conv2d c53 (&net, 256, 1, 1, 1, 1, 0, 0, c53_bin, true); - tk::dnn::Activation a53 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c54 (&net, 512, 3, 3, 1, 1, 1, 1, c54_bin, true); - tk::dnn::Activation a54 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s55 (&net, &s52); - tk::dnn::Conv2d c56 (&net, 256, 1, 1, 1, 1, 0, 0, c56_bin, true); - tk::dnn::Activation a56 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c57 (&net, 512, 3, 3, 1, 1, 1, 1, c57_bin, true); - tk::dnn::Activation a57 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s58 (&net, &s55); - tk::dnn::Conv2d c59 (&net, 256, 1, 1, 1, 1, 0, 0, c59_bin, true); - tk::dnn::Activation a59 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c60 (&net, 512, 3, 3, 1, 1, 1, 1, c60_bin, true); - tk::dnn::Activation a60 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s61 (&net, &s58); - - tk::dnn::Conv2d c62 (&net,1024, 3, 3, 2, 2, 1, 1, c62_bin, true); - tk::dnn::Activation a62 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c63 (&net, 512, 1, 1, 1, 1, 0, 0, c63_bin, true); - tk::dnn::Activation a63 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c64 (&net,1024, 3, 3, 1, 1, 1, 1, c64_bin, true); - tk::dnn::Activation a64 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s65 (&net, &a62); - - tk::dnn::Conv2d c66 (&net, 512, 1, 1, 1, 1, 0, 0, c66_bin, true); - tk::dnn::Activation a66 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c67 (&net,1024, 3, 3, 1, 1, 1, 1, c67_bin, true); - tk::dnn::Activation a67 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s68 (&net, &s65); - - tk::dnn::Conv2d c69 (&net, 512, 1, 1, 1, 1, 0, 0, c69_bin, true); - tk::dnn::Activation a69 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c70 (&net,1024, 3, 3, 1, 1, 1, 1, c70_bin, true); - tk::dnn::Activation a70 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s71 (&net, &s68); - - tk::dnn::Conv2d c72 (&net, 512, 1, 1, 1, 1, 0, 0, c72_bin, true); - tk::dnn::Activation a72 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c73 (&net,1024, 3, 3, 1, 1, 1, 1, c73_bin, true); - tk::dnn::Activation a73 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Shortcut s74 (&net, &s71); - - tk::dnn::Conv2d c75 (&net, 512, 1, 1, 1, 1, 0, 0, c75_bin, true); - tk::dnn::Activation a75 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c76 (&net,1024, 3, 3, 1, 1, 1, 1, c76_bin, true); - tk::dnn::Activation a76 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c77 (&net, 512, 1, 1, 1, 1, 0, 0, c77_bin, true); - tk::dnn::Activation a77 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c78 (&net,1024, 3, 3, 1, 1, 1, 1, c78_bin, true); - tk::dnn::Activation a78 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c79 (&net, 512, 1, 1, 1, 1, 0, 0, c79_bin, true); - tk::dnn::Activation a79 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c80 (&net,1024, 3, 3, 1, 1, 1, 1, c80_bin, true); - tk::dnn::Activation a80 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c81 (&net, 45, 1, 1, 1, 1, 0, 0, c81_bin, false); - tk::dnn::Yolo yolo0 (&net, 10, 3, g82_bin); - - tk::dnn::Layer *m83_layers[1] = { &a79 }; - tk::dnn::Route m83 (&net, m83_layers, 1); - tk::dnn::Conv2d c84 (&net, 256, 1, 1, 1, 1, 0, 0, c84_bin, true); - tk::dnn::Activation a84 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Upsample u85 (&net, 2); - - tk::dnn::Layer *m86_layers[2] = { &u85, &s61 }; - tk::dnn::Route m86 (&net, m86_layers, 2); - tk::dnn::Conv2d c87 (&net, 256, 1, 1, 1, 1, 0, 0, c87_bin, true); - tk::dnn::Activation a87 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c88 (&net, 512, 3, 3, 1, 1, 1, 1, c88_bin, true); - tk::dnn::Activation a88 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c89 (&net, 256, 1, 1, 1, 1, 0, 0, c89_bin, true); - tk::dnn::Activation a89 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c90 (&net, 512, 3, 3, 1, 1, 1, 1, c90_bin, true); - tk::dnn::Activation a90 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c91 (&net, 256, 1, 1, 1, 1, 0, 0, c91_bin, true); - tk::dnn::Activation a91 (&net, tk::dnn::ACTIVATION_LEAKY); - - tk::dnn::Conv2d c92 (&net, 512, 3, 3, 1, 1, 1, 1, c92_bin, true); - tk::dnn::Activation a92 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c93 (&net, 45, 1, 1, 1, 1, 0, 0, c93_bin, false); - tk::dnn::Yolo yolo1 (&net, 10, 3, g94_bin); - - tk::dnn::Layer *m95_layers[1] = { &a91 }; - tk::dnn::Route m95 (&net, m95_layers, 1); - tk::dnn::Conv2d c96 (&net, 128, 1, 1, 1, 1, 0, 0, c96_bin, true); - tk::dnn::Activation a96 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Upsample u97 (&net, 2); - - tk::dnn::Layer *m98_layers[2] = { &u97, &s36 }; - tk::dnn::Route m98 (&net, m98_layers, 2); - tk::dnn::Conv2d c99 (&net, 128, 1, 1, 1, 1, 0, 0, c99_bin, true); - tk::dnn::Activation a99 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c100 (&net, 256, 3, 3, 1, 1, 1, 1, c100_bin, true); - tk::dnn::Activation a100 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c101 (&net, 128, 1, 1, 1, 1, 0, 0, c101_bin, true); - tk::dnn::Activation a101 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c102 (&net, 256, 3, 3, 1, 1, 1, 1, c102_bin, true); - tk::dnn::Activation a102 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c103 (&net, 128, 1, 1, 1, 1, 0, 0, c103_bin, true); - tk::dnn::Activation a103 (&net, tk::dnn::ACTIVATION_LEAKY); - - tk::dnn::Conv2d c104 (&net, 256, 3, 3, 1, 1, 1, 1, c104_bin, true); - tk::dnn::Activation a104 (&net, tk::dnn::ACTIVATION_LEAKY); - tk::dnn::Conv2d c105 (&net, 45, 1, 1, 1, 1, 0, 0, c105_bin, false); - tk::dnn::Yolo yolo2 (&net, 10, 3, g106_bin); + // create yolo3 model + std::string bin_path = "../tests/yolo3_berkeley"; + int classes = 10; + tk::dnn::Yolo *yolo [3]; + #include "models/Yolo3.h" // Load input dnnType *data; @@ -304,9 +27,7 @@ int main() { // the network have 3 outputs tk::dnn::dataDim_t out_dim[3]; - out_dim[0] = yolo0.output_dim; - out_dim[1] = yolo1.output_dim; - out_dim[2] = yolo2.output_dim; + for(int i=0; i<3; i++) out_dim[i] = yolo[i]->output_dim; dnnType *cudnn_out[3], *rt_out[3]; tk::dnn::dataDim_t dim1 = dim; //input dim @@ -317,18 +38,13 @@ int main() { TIMER_STOP dim1.print(); } - cudnn_out[0] = yolo0.dstData; - cudnn_out[1] = yolo1.dstData; - cudnn_out[2] = yolo2.dstData; - + for(int i=0; i<3; i++) cudnn_out[i] = yolo[i]->dstData; + printCenteredTitle(" compute detections ", '=', 30); TIMER_START int ndets = 0; - int classes = yolo0.classes; tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); - yolo0.computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); - yolo1.computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); - yolo2.computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); + for(int i=0; i<3; i++) yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); tk::dnn::Yolo::mergeDetections(dets, ndets, classes); for(int j=0; jcomputeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); tk::dnn::Yolo::mergeDetections(dets, ndets, classes); for(int j=0; jdstData; + + printCenteredTitle(" compute detections ", '=', 30); + TIMER_START + int ndets = 0; + tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); + for(int i=0; i<3; i++) yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); + tk::dnn::Yolo::mergeDetections(dets, ndets, classes); + + for(int j=0; j 0) + cl = c; + } + std::cout< Date: Sun, 15 Sep 2019 16:19:30 +0200 Subject: [PATCH 06/30] yolo3 flir ok --- CMakeLists.txt | 3 +++ demo/demo/demo.cpp | 4 +++- src/Yolo3Detection.cpp | 13 +++++++------ tests/yolo3_flir/yolo3_flir.cfg | 14 +++++++------- tests/yolo3_flir/yolo3_flir.cpp | 2 +- 5 files changed, 21 insertions(+), 15 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index acf3613..f94ccbb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -89,6 +89,9 @@ target_link_libraries(test_yolo3_coco4 tkDNN) add_executable(test_yolo3_berkeley tests/yolo3_berkeley/yolo3_berkeley.cpp) target_link_libraries(test_yolo3_berkeley tkDNN) + +add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp) +target_link_libraries(test_yolo3_flir tkDNN) ################################################################################ diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index e57a36e..82f724b 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -12,7 +12,9 @@ #include "Yolo3Detection.h" bool gRun; -std::string obj_class[10] {"person", "car", "truck", "bus", "motor", "bike", "rider", "traffic light", "traffic sign", "train"}; +//std::string obj_class[10] {"person", "car", "truck", "bus", "motor", "bike", "rider", "traffic light", "traffic sign", "train"}; +//std::string obj_class[3] {"person", "bike", "car"}; +std::string obj_class[10] {"0", "1", "2", "3", "4", "5", "6", "7", "8", "9"}; void sig_handler(int signo) { diff --git a/src/Yolo3Detection.cpp b/src/Yolo3Detection.cpp index ce9fd6a..244aa36 100644 --- a/src/Yolo3Detection.cpp +++ b/src/Yolo3Detection.cpp @@ -68,18 +68,19 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { float yRatio = float(imageORIG.rows) / float(netRT->input_dim.h); resize(imageORIG, imageORIG, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); + imageORIG.convertTo(imageF, CV_32FC3, 1/255.0); //split channels cv::split(imageF,bgr);//split source //write channels - int idx = 0; - memcpy((void*)&input[idx], (void*)bgr[2].data, imageF.rows*imageF.cols*sizeof(dnnType)); - idx = imageF.rows*imageF.cols; - memcpy((void*)&input[idx], (void*)bgr[1].data, imageF.rows*imageF.cols*sizeof(dnnType)); - idx *= 2; - memcpy((void*)&input[idx], (void*)bgr[0].data, imageF.rows*imageF.cols*sizeof(dnnType)); + for(int i=0; iinput_dim.c; i++) { + int idx = i*imageF.rows*imageF.cols; + int ch = netRT->input_dim.c-1 -i; + memcpy((void*)&input[idx], (void*)bgr[ch].data, imageF.rows*imageF.cols*sizeof(dnnType)); + } + //DO INFERENCE dnnType *rt_out[3]; diff --git a/tests/yolo3_flir/yolo3_flir.cfg b/tests/yolo3_flir/yolo3_flir.cfg index 3a39bd8..1bbf5c1 100644 --- a/tests/yolo3_flir/yolo3_flir.cfg +++ b/tests/yolo3_flir/yolo3_flir.cfg @@ -17,7 +17,7 @@ hue=.1 learning_rate=0.001 burn_in=1000 -max_batches = 10000 +max_batches = 20000 policy=steps steps=8000,9000 scales=.1,.1 @@ -602,13 +602,13 @@ activation=leaky size=1 stride=1 pad=1 -filters=30 +filters=24 activation=linear [yolo] mask = 6,7,8 anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648 -classes=5 +classes=3 num=9 jitter=.3 ignore_thresh = .5 @@ -686,13 +686,13 @@ activation=leaky size=1 stride=1 pad=1 -filters=30 +filters=24 activation=linear [yolo] mask = 3,4,5 anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648 -classes=5 +classes=3 num=9 jitter=.3 ignore_thresh = .5 @@ -770,13 +770,13 @@ activation=leaky size=1 stride=1 pad=1 -filters=30 +filters=24 activation=linear [yolo] mask = 0,1,2 anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648 -classes=5 +classes=3 num=9 jitter=.3 ignore_thresh = .5 diff --git a/tests/yolo3_flir/yolo3_flir.cpp b/tests/yolo3_flir/yolo3_flir.cpp index 62c1889..2bc45e4 100644 --- a/tests/yolo3_flir/yolo3_flir.cpp +++ b/tests/yolo3_flir/yolo3_flir.cpp @@ -11,7 +11,7 @@ int main() { // create yolo3 model std::string bin_path = "../tests/yolo3_flir"; - int classes = 5; + int classes = 3; tk::dnn::Yolo *yolo [3]; #include "models/Yolo3.h" From 77f031c0f488c0ffb36f63a5215e39f5b3c81625 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Mon, 16 Sep 2019 10:35:29 +0200 Subject: [PATCH 07/30] save video result --- demo/demo/demo.cpp | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 82f724b..7ed80da 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -16,6 +16,7 @@ bool gRun; //std::string obj_class[3] {"person", "bike", "car"}; std::string obj_class[10] {"0", "1", "2", "3", "4", "5", "6", "7", "8", "9"}; +bool SAVE_RESULT = false; void sig_handler(int signo) { std::cout<<"request gateway stop\n"; @@ -46,6 +47,14 @@ int main(int argc, char *argv[]) { else std::cout<<"camera started\n"; + + cv::VideoWriter resultVideo; + if(SAVE_RESULT) { + int w = cap.get(CV_CAP_PROP_FRAME_WIDTH); + int h = cap.get(CV_CAP_PROP_FRAME_HEIGHT); + resultVideo.open("result.mp4", CV_FOURCC('M','P','4','V'), 30, cv::Size(w, h)); + } + cv::Mat frame; cv::Mat dnn_input; cv::namedWindow("detection", cv::WINDOW_NORMAL); @@ -53,7 +62,7 @@ int main(int argc, char *argv[]) { while(gRun) { cap >> frame; if(!frame.data) { - continue; + break; } // this will be resized to the net format @@ -86,6 +95,8 @@ int main(int argc, char *argv[]) { cv::imshow("detection", frame); cv::waitKey(1); + if(SAVE_RESULT) + resultVideo << frame; } std::cout<<"detection end\n"; From ec02c7292fa7113140b923c25535aca17618e670 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Mon, 16 Sep 2019 19:41:59 +0200 Subject: [PATCH 08/30] save layer names in rt file --- demo/demo/demo.cpp | 6 +----- include/Layer.h | 2 ++ include/Yolo3Detection.h | 7 +++++++ include/pluginsRT/YoloRT.h | 15 ++++++++++++++- src/NetworkRT.cpp | 9 +++++++++ src/Yolo.cpp | 6 ++++++ src/Yolo3Detection.cpp | 1 + tests/yolo3_berkeley/yolo3_berkeley.cpp | 5 +++++ tests/yolo3_flir/yolo3_flir.cpp | 5 +++++ 9 files changed, 50 insertions(+), 6 deletions(-) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 7ed80da..ac6c63d 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -12,10 +12,6 @@ #include "Yolo3Detection.h" bool gRun; -//std::string obj_class[10] {"person", "car", "truck", "bus", "motor", "bike", "rider", "traffic light", "traffic sign", "train"}; -//std::string obj_class[3] {"person", "bike", "car"}; -std::string obj_class[10] {"0", "1", "2", "3", "4", "5", "6", "7", "8", "9"}; - bool SAVE_RESULT = false; void sig_handler(int signo) { @@ -77,7 +73,7 @@ int main(int argc, char *argv[]) { int x1 = b.x + b.w; int y0 = b.y; int y1 = b.y + b.h; - std::string det_class = obj_class[b.cl]; + std::string det_class = yolo.getYoloLayer()->classesNames[b.cl]; float prob = b.prob; std::cout< +#include #include "utils.h" #include "Network.h" @@ -359,6 +360,7 @@ public: int classes, num; dnnType *mask_h, *mask_d; //anchors dnnType *bias_h, *bias_d; //anchors + std::vector classesNames; virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); int computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh); diff --git a/include/Yolo3Detection.h b/include/Yolo3Detection.h index 0ff4b1e..0a7ba25 100644 --- a/include/Yolo3Detection.h +++ b/include/Yolo3Detection.h @@ -52,6 +52,13 @@ class Yolo3Detection { void update(cv::Mat &frame); + tk::dnn::Yolo* getYoloLayer(int n=0) { + if(n<3) + return yolo[n]; + else + return nullptr; + } + }; }} diff --git a/include/pluginsRT/YoloRT.h b/include/pluginsRT/YoloRT.h index 2e52142..dab7c18 100644 --- a/include/pluginsRT/YoloRT.h +++ b/include/pluginsRT/YoloRT.h @@ -1,6 +1,8 @@ #include #include "../kernels.h" +#define YOLORT_CLASSNAME_W 256 + class YoloRT : public IPlugin { @@ -16,6 +18,7 @@ public: if(yolo != nullptr) { memcpy(mask, yolo->mask_h, sizeof(dnnType)*num); memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*3*2); + classesNames = yolo->classesNames; } } @@ -72,7 +75,7 @@ public: virtual size_t getSerializationSize() override { - return 5*sizeof(int) + num*sizeof(dnnType) + num*3*2*sizeof(dnnType); + return 5*sizeof(int) + num*sizeof(dnnType) + num*3*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char); } virtual void serialize(void* buffer) override { @@ -86,10 +89,20 @@ public: tk::dnn::writeBUF(buf, mask[i]); for(int i=0; i<3*2*num; i++) tk::dnn::writeBUF(buf, bias[i]); + + // save classes names + for(int i=0; i classesNames; dnnType *mask; dnnType *bias; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 2a73efa..7430b6f 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -461,6 +461,15 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa for(int i=0; i<3*2*r->num; i++) r->bias[i] = readBUF(buf); + // save classes names + r->classesNames.resize(r->classes); + for(int i=0; iclasses; i++) { + char tmp[YOLORT_CLASSNAME_W]; + for(int j=0; j(buf); + r->classesNames[i] = std::string(tmp); + } + yolos[n_yolos++] = r; return r; } diff --git a/src/Yolo.cpp b/src/Yolo.cpp index f2ec53c..babef38 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -25,6 +25,12 @@ Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights) : readBinaryFile(fname_weights, 3*num*2, &bias_h, &bias_d, seek); } + // init default classes name + classesNames.clear(); + for(int i=0; imask_h, yRT->mask, sizeof(dnnType)*num); memcpy(yolo[i]->bias_h, yRT->bias, sizeof(dnnType)*num*3*2); yolo[i]->input_dim = yolo[i]->output_dim = tk::dnn::dataDim_t(1, yRT->c, yRT->h, yRT->w); + yolo[i]->classesNames = yRT->classesNames; } dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); diff --git a/tests/yolo3_berkeley/yolo3_berkeley.cpp b/tests/yolo3_berkeley/yolo3_berkeley.cpp index 43f5da5..70e8b60 100644 --- a/tests/yolo3_berkeley/yolo3_berkeley.cpp +++ b/tests/yolo3_berkeley/yolo3_berkeley.cpp @@ -14,6 +14,11 @@ int main() { tk::dnn::Yolo *yolo [3]; #include "models/Yolo3.h" + // fill classes names + for(int i=0; i<3; i++) { + yolo[i]->classesNames = {"person", "car", "truck", "bus", "motor", "bike", "rider", "traffic light", "traffic sign", "train"}; + } + // Load input dnnType *data; dnnType *input_h; diff --git a/tests/yolo3_flir/yolo3_flir.cpp b/tests/yolo3_flir/yolo3_flir.cpp index 2bc45e4..83b53af 100644 --- a/tests/yolo3_flir/yolo3_flir.cpp +++ b/tests/yolo3_flir/yolo3_flir.cpp @@ -15,6 +15,11 @@ int main() { tk::dnn::Yolo *yolo [3]; #include "models/Yolo3.h" + // fill classes names + for(int i=0; i<3; i++) { + yolo[i]->classesNames = {"person", "bike", "car"}; + } + // Load input dnnType *data; dnnType *input_h; From ca62784f57d177e3fc9e5ec360fc13d44586a2a0 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 17 Sep 2019 15:22:39 +0200 Subject: [PATCH 09/30] include dir fix, cmake dir --- CMakeLists.txt | 5 ++--- tkDNNConfig.cmake => cmake/tkDNNConfig.cmake | 0 include/{ => tkDNN}/Layer.h | 0 include/{ => tkDNN}/Network.h | 0 include/{ => tkDNN}/NetworkRT.h | 0 include/{ => tkDNN}/Yolo3Detection.h | 0 include/{ => tkDNN}/kernels.h | 0 include/{ => tkDNN}/models/Yolo3.h | 0 include/{ => tkDNN}/pluginsRT/ActivationLeakyRT.h | 0 include/{ => tkDNN}/pluginsRT/Int8Calibrator.h | 0 include/{ => tkDNN}/pluginsRT/RegionRT.h | 0 include/{ => tkDNN}/pluginsRT/ReorgRT.h | 0 include/{ => tkDNN}/pluginsRT/RouteRT.h | 0 include/{ => tkDNN}/pluginsRT/ShortcutRT.h | 0 include/{ => tkDNN}/pluginsRT/UpsampleRT.h | 0 include/{ => tkDNN}/pluginsRT/YoloRT.h | 0 include/{ => tkDNN}/tkdnn.h | 0 include/{ => tkDNN}/utils.h | 0 18 files changed, 2 insertions(+), 3 deletions(-) rename tkDNNConfig.cmake => cmake/tkDNNConfig.cmake (100%) rename include/{ => tkDNN}/Layer.h (100%) rename include/{ => tkDNN}/Network.h (100%) rename include/{ => tkDNN}/NetworkRT.h (100%) rename include/{ => tkDNN}/Yolo3Detection.h (100%) rename include/{ => tkDNN}/kernels.h (100%) rename include/{ => tkDNN}/models/Yolo3.h (100%) rename include/{ => tkDNN}/pluginsRT/ActivationLeakyRT.h (100%) rename include/{ => tkDNN}/pluginsRT/Int8Calibrator.h (100%) rename include/{ => tkDNN}/pluginsRT/RegionRT.h (100%) rename include/{ => tkDNN}/pluginsRT/ReorgRT.h (100%) rename include/{ => tkDNN}/pluginsRT/RouteRT.h (100%) rename include/{ => tkDNN}/pluginsRT/ShortcutRT.h (100%) rename include/{ => tkDNN}/pluginsRT/UpsampleRT.h (100%) rename include/{ => tkDNN}/pluginsRT/YoloRT.h (100%) rename include/{ => tkDNN}/tkdnn.h (100%) rename include/{ => tkDNN}/utils.h (100%) diff --git a/CMakeLists.txt b/CMakeLists.txt index f94ccbb..a35264c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -4,7 +4,7 @@ project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) set(CMAKE_BUILD_TYPE "Release") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC") -include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) # project specific flags if(DEBUG) @@ -110,8 +110,7 @@ target_link_libraries(yolo3_demo tkDNN) # CACHE PATH "default install path" FORCE) #endif() message("install dir:" ${CMAKE_INSTALL_PREFIX}) -install(DIRECTORY include/ DESTINATION include/${CMAKE_PROJECT_NAME} -FILES_MATCHING PATTERN "*.h") +install(DIRECTORY include/ DESTINATION include/) install(TARGETS tkDNN kernels DESTINATION lib) install(FILES "${CMAKE_SOURCE_DIR}/${CMAKE_PROJECT_NAME}Config.cmake" # source directory DESTINATION "share/${CMAKE_PROJECT_NAME}/cmake/" # target directory diff --git a/tkDNNConfig.cmake b/cmake/tkDNNConfig.cmake similarity index 100% rename from tkDNNConfig.cmake rename to cmake/tkDNNConfig.cmake diff --git a/include/Layer.h b/include/tkDNN/Layer.h similarity index 100% rename from include/Layer.h rename to include/tkDNN/Layer.h diff --git a/include/Network.h b/include/tkDNN/Network.h similarity index 100% rename from include/Network.h rename to include/tkDNN/Network.h diff --git a/include/NetworkRT.h b/include/tkDNN/NetworkRT.h similarity index 100% rename from include/NetworkRT.h rename to include/tkDNN/NetworkRT.h diff --git a/include/Yolo3Detection.h b/include/tkDNN/Yolo3Detection.h similarity index 100% rename from include/Yolo3Detection.h rename to include/tkDNN/Yolo3Detection.h diff --git a/include/kernels.h b/include/tkDNN/kernels.h similarity index 100% rename from include/kernels.h rename to include/tkDNN/kernels.h diff --git a/include/models/Yolo3.h b/include/tkDNN/models/Yolo3.h similarity index 100% rename from include/models/Yolo3.h rename to include/tkDNN/models/Yolo3.h diff --git a/include/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h similarity index 100% rename from include/pluginsRT/ActivationLeakyRT.h rename to include/tkDNN/pluginsRT/ActivationLeakyRT.h diff --git a/include/pluginsRT/Int8Calibrator.h b/include/tkDNN/pluginsRT/Int8Calibrator.h similarity index 100% rename from include/pluginsRT/Int8Calibrator.h rename to include/tkDNN/pluginsRT/Int8Calibrator.h diff --git a/include/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h similarity index 100% rename from include/pluginsRT/RegionRT.h rename to include/tkDNN/pluginsRT/RegionRT.h diff --git a/include/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h similarity index 100% rename from include/pluginsRT/ReorgRT.h rename to include/tkDNN/pluginsRT/ReorgRT.h diff --git a/include/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h similarity index 100% rename from include/pluginsRT/RouteRT.h rename to include/tkDNN/pluginsRT/RouteRT.h diff --git a/include/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h similarity index 100% rename from include/pluginsRT/ShortcutRT.h rename to include/tkDNN/pluginsRT/ShortcutRT.h diff --git a/include/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h similarity index 100% rename from include/pluginsRT/UpsampleRT.h rename to include/tkDNN/pluginsRT/UpsampleRT.h diff --git a/include/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h similarity index 100% rename from include/pluginsRT/YoloRT.h rename to include/tkDNN/pluginsRT/YoloRT.h diff --git a/include/tkdnn.h b/include/tkDNN/tkdnn.h similarity index 100% rename from include/tkdnn.h rename to include/tkDNN/tkdnn.h diff --git a/include/utils.h b/include/tkDNN/utils.h similarity index 100% rename from include/utils.h rename to include/tkDNN/utils.h From de8b02fe5091e99c33aa5b688b11bbf4dc10fb66 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 17 Sep 2019 16:12:59 +0200 Subject: [PATCH 10/30] install fix --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index a35264c..20fa8d5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -112,7 +112,7 @@ target_link_libraries(yolo3_demo tkDNN) message("install dir:" ${CMAKE_INSTALL_PREFIX}) install(DIRECTORY include/ DESTINATION include/) install(TARGETS tkDNN kernels DESTINATION lib) -install(FILES "${CMAKE_SOURCE_DIR}/${CMAKE_PROJECT_NAME}Config.cmake" # source directory +install(FILES "${CMAKE_SOURCE_DIR}/cmake/${CMAKE_PROJECT_NAME}Config.cmake" # source directory DESTINATION "share/${CMAKE_PROJECT_NAME}/cmake/" # target directory ) From bbc4dda63512486b40646d309d569fec46b2e2e9 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 17 Sep 2019 17:16:46 +0200 Subject: [PATCH 11/30] removed buildtype --- CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 20fa8d5..e305d1a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,7 +2,6 @@ cmake_minimum_required(VERSION 3.5) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) -set(CMAKE_BUILD_TYPE "Release") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC") include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) From 92f3d1c548210978d1f14a6178a59fba93691e7a Mon Sep 17 00:00:00 2001 From: mbosi <205839@studenti.unimore.it> Date: Tue, 1 Oct 2019 18:47:31 +0200 Subject: [PATCH 12/30] fixed install cmake --- CMakeLists.txt | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e305d1a..467e823 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -111,10 +111,11 @@ target_link_libraries(yolo3_demo tkDNN) message("install dir:" ${CMAKE_INSTALL_PREFIX}) install(DIRECTORY include/ DESTINATION include/) install(TARGETS tkDNN kernels DESTINATION lib) -install(FILES "${CMAKE_SOURCE_DIR}/cmake/${CMAKE_PROJECT_NAME}Config.cmake" # source directory - DESTINATION "share/${CMAKE_PROJECT_NAME}/cmake/" # target directory +install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/cmake/" # source directory + DESTINATION "share/tkDNN/cmake/" # target directory ) + #------------------------------------------------------------------------------- # Prepare for test #------------------------------------------------------------------------------- From aa5927d8a168b44f32a13dfa0551d78594e9fe33 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Wed, 6 Nov 2019 14:04:23 +0100 Subject: [PATCH 13/30] findCUDNN --- CMakeLists.txt | 12 +++--------- cmake/FindCUDNN.cmake | 33 +++++++++++++++++++++++++++++++++ cmake/tkDNNConfig.cmake | 16 +++++----------- 3 files changed, 41 insertions(+), 20 deletions(-) create mode 100644 cmake/FindCUDNN.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 467e823..b4b3c38 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -18,17 +18,11 @@ find_package(CUDA 9.0 REQUIRED) SET(CUDA_SEPARABLE_COMPILATION ON) #set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'") -# compile Discovery only if TensorRT is installed -find_library(NVINFER NAMES nvinfer) -if(NVINFER STREQUAL "NVINFER-NOTFOUND") - set(NVINFER_INCLUDES "/usr/local/nvidia/tensorrt/include/") - link_directories(/usr/local/nvidia/tensorrt/targets/x86_64-linux-gnu/lib/ - /usr/local/cuda/targets/x86_64-linux/lib/) -endif() +find_package(CUDNN REQUIRED) # compile file(GLOB tkdnn_CUSRC "src/kernels/*.cu") -cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${NVINFER_INCLUDES}) +cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) @@ -43,7 +37,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") # Build Libraries #------------------------------------------------------------------------------- file(GLOB tkdnn_SRC "src/*.cpp") -set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS}) +set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS}) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11") include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES}) diff --git a/cmake/FindCUDNN.cmake b/cmake/FindCUDNN.cmake new file mode 100644 index 0000000..0e99052 --- /dev/null +++ b/cmake/FindCUDNN.cmake @@ -0,0 +1,33 @@ +# Find the header files + +find_path(CUDNN_INCLUDE_DIR + ${CMAKE_SYSROOT}/usr/local/include + ${CMAKE_SYSROOT}/usr/include + /usr/local/nvidia/tensorrt/include/ + NO_DEFAULT_PATH +) + +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-2.0/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib + /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/ + NO_DEFAULT_PATH +) +find_library(CUDNN_NVLIB + NAMES "nvinfer" + PATHS + /usr/local/driveworks-2.0/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) diff --git a/cmake/tkDNNConfig.cmake b/cmake/tkDNNConfig.cmake index fd06286..4dbaf21 100644 --- a/cmake/tkDNNConfig.cmake +++ b/cmake/tkDNNConfig.cmake @@ -4,27 +4,21 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++11 -fPIC") find_package(CUDA REQUIRED) find_package(OpenCV REQUIRED) -find_library(NVINFER NAMES nvinfer) -if(NVINFER STREQUAL "NVINFER-NOTFOUND") - set(NVINFER_INCLUDES "/usr/local/nvidia/tensorrt/include/") - link_directories(/usr/local/nvidia/tensorrt/targets/x86_64-linux-gnu/lib/ - /usr/local/cuda/targets/x86_64-linux/lib/) -endif() +find_package(CUDNN REQUIRED) set(tkDNN_INCLUDE_DIRS ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} - ${NVINFER_INCLUDES} + ${CUDNN_INCLUDE_DIRS} ) set(tkDNN_LIBRARIES tkDNN kernels ${CUDA_LIBRARIES} - ${CUDA_CUBLAS_LIBRARIES} - -lcudnn - -lnvinfer - ${OpenCV_LIBS} + ${CUDA_CUBLAS_LIBRARIES} + ${CUDNN_LIBRARIES} + ${OpenCV_LIBS} ) set(tkDNN_FOUND true) From b218b18a02b19c4e84ef85fd312c8085b46bbdce Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Mon, 2 Dec 2019 20:24:12 +0100 Subject: [PATCH 14/30] readme update --- README.md | 2 -- 1 file changed, 2 deletions(-) diff --git a/README.md b/README.md index b166f63..93a7574 100644 --- a/README.md +++ b/README.md @@ -48,5 +48,3 @@ this will genereate a yolo3_berkeley.rt file that can be used for live detection ./yolo3_demo # launch detection on a demo video ./yolo3_demo yolo3_berkeley.rt /dev/video0 # launch detection on device 0 ``` - - From 6bf9179acc1f8430b5b4b3cb3f96c69f6332f84e Mon Sep 17 00:00:00 2001 From: mbosi <205839@studenti.unimore.it> Date: Thu, 12 Dec 2019 12:30:24 +0100 Subject: [PATCH 15/30] fix to drivework global path --- cmake/FindCUDNN.cmake | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cmake/FindCUDNN.cmake b/cmake/FindCUDNN.cmake index 0e99052..f240fcb 100644 --- a/cmake/FindCUDNN.cmake +++ b/cmake/FindCUDNN.cmake @@ -14,14 +14,14 @@ list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.5) find_library(CUDNN_LIB NAMES cudnn PATHS - /usr/local/driveworks-2.0/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib + /usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/ NO_DEFAULT_PATH ) find_library(CUDNN_NVLIB NAMES "nvinfer" PATHS - /usr/local/driveworks-2.0/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib + /usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/ NO_DEFAULT_PATH ) From 57d7743f7e950a9d72870ea811758cd156775403 Mon Sep 17 00:00:00 2001 From: xavier Date: Wed, 15 Jan 2020 09:55:10 +0100 Subject: [PATCH 16/30] Change opencv funcion call (due to OpenCV 4) Signed-off-by: xavier --- demo/demo/demo.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index ac6c63d..3a9ef6e 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -7,6 +7,7 @@ #include #include +#include #include #include "Yolo3Detection.h" @@ -46,9 +47,9 @@ int main(int argc, char *argv[]) { cv::VideoWriter resultVideo; if(SAVE_RESULT) { - int w = cap.get(CV_CAP_PROP_FRAME_WIDTH); - int h = cap.get(CV_CAP_PROP_FRAME_HEIGHT); - resultVideo.open("result.mp4", CV_FOURCC('M','P','4','V'), 30, cv::Size(w, h)); + int w = cap.get(cv::CAP_PROP_FRAME_WIDTH); + int h = cap.get(cv::CAP_PROP_FRAME_HEIGHT); + resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h)); } cv::Mat frame; From c32a0be257acb29cbe4f4ef20be1bc80027a3475 Mon Sep 17 00:00:00 2001 From: xavier Date: Wed, 15 Jan 2020 18:06:02 +0100 Subject: [PATCH 17/30] Batchnorm eps fix, works on jetpack 4.3 --- README.md | 7 ++++--- include/tkDNN/Layer.h | 2 ++ src/Conv2d.cpp | 2 +- src/LayerWgs.cpp | 2 +- src/Network.cpp | 9 +++++---- 5 files changed, 13 insertions(+), 9 deletions(-) diff --git a/README.md b/README.md index 93a7574..3b16ec3 100644 --- a/README.md +++ b/README.md @@ -3,9 +3,10 @@ tkDNN is a Deep Neural Network library built with cuDNN primitives specifically The main scope is to do high performance inference on already trained models. this branch actually work on every NVIDIA GPU that support the dependencies: -* CUDA 9 -* CUDNN 7.105 -* TENSORRT 4.02 +* CUDA 10.0 +* CUDNN 7.603 +* TENSORRT 6.01 +* OPENCV 4.1 ## Workflow The recommended workflow follow these step: diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index c6bee42..ee73109 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -24,6 +24,8 @@ enum layerType_t { LAYER_YOLO }; +#define TKDNN_BN_MIN_EPSILON 1e-5 + /** Simple layer Father class */ diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 3dfdce6..b275198 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -117,7 +117,7 @@ dnnType* Conv2d::infer(dataDim_t &dim, dnnType* srcData) { dstTensorDesc, dstData, dstTensorDesc, dstData, biasTensorDesc, //same tensor descriptor as bias scales_d, bias_d, mean_d, variance_d, - CUDNN_BN_MIN_EPSILON); + TKDNN_BN_MIN_EPSILON); } //update data dimensions dim = output_dim; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index f27da0f..7851164 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -29,7 +29,7 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs, seek += outputs; readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek); - float eps = CUDNN_BN_MIN_EPSILON; + float eps = TKDNN_BN_MIN_EPSILON; power_h = new dnnType[outputs]; for(int i=0; i Date: Wed, 15 Jan 2020 18:07:40 +0100 Subject: [PATCH 18/30] support clion --- .gitignore | 1 + 1 file changed, 1 insertion(+) diff --git a/.gitignore b/.gitignore index 446d490..02f2a8e 100644 --- a/.gitignore +++ b/.gitignore @@ -8,3 +8,4 @@ build/ *.h5 *.tar.gz *.weights +.idea/ \ No newline at end of file From da4f24615741b1ab1a44dd550826e8da3ec7648e Mon Sep 17 00:00:00 2001 From: xavier Date: Wed, 15 Jan 2020 21:48:18 +0100 Subject: [PATCH 19/30] add DLA, plugin for shortcut and leaky. new verison 0.4 --- demo/demo/demo.cpp | 7 +++++++ include/tkDNN/Yolo3Detection.h | 3 +++ include/tkDNN/tkdnn.h | 2 +- src/NetworkRT.cpp | 22 ++++++++++++++++++---- src/Yolo3Detection.cpp | 2 ++ 5 files changed, 31 insertions(+), 5 deletions(-) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 3a9ef6e..ba08caa 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -97,6 +97,13 @@ int main(int argc, char *argv[]) { } std::cout<<"detection end\n"; + + + std::cout< detected; + // keep track of inference times (ms) + std::vector stats; + Yolo3Detection() {} virtual ~Yolo3Detection() {} diff --git a/include/tkDNN/tkdnn.h b/include/tkDNN/tkdnn.h index dde3126..554daa1 100644 --- a/include/tkDNN/tkdnn.h +++ b/include/tkDNN/tkdnn.h @@ -5,4 +5,4 @@ #include "Layer.h" #include "NetworkRT.h" -#define TKDNN_VERSION 300 +#define TKDNN_VERSION 400 diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 7430b6f..d4fdf90 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -35,7 +35,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { builderRT = createInferBuilder(loggerRT); std::cout<<"Float16 support: "<platformHasFastFp16()<<"\n"; std::cout<<"Int8 support: "<platformHasFastInt8()<<"\n"; - //std::cout<<"DLAs: "<getNbDLACores()<<"\n"; + std::cout<<"DLAs: "<getNbDLACores()<<"\n"; networkRT = builderRT->createNetwork(); if(!fileExist(name)) { @@ -51,7 +51,6 @@ NetworkRT::NetworkRT(Network *net, const char *name) { dtRT = DataType::kHALF; builderRT->setHalf2Mode(true); } - /* if(net->dla && builderRT->getNbDLACores() > 0) { dtRT = DataType::kHALF; builderRT->setFp16Mode(true); @@ -59,7 +58,6 @@ NetworkRT::NetworkRT(Network *net, const char *name) { builderRT->setDefaultDeviceType(DeviceType::kDLA); builderRT->setDLACore(0); } - */ //add input layer ITensor *input = networkRT->addInput("data", DataType::kFLOAT, @@ -276,10 +274,19 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { if(l->act_mode == ACTIVATION_LEAKY) { //std::cout<<"New plugin LEAKY\n"; + + /* + // plugin version IPlugin *plugin = new ActivationLeakyRT(); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); checkNULL(lRT); return lRT; + */ + + IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU); + lRT->setAlpha(0.1); + checkNULL(lRT); + return lRT; } else if(l->act_mode == CUDNN_ACTIVATION_RELU) { IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU); @@ -340,14 +347,21 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) { //std::cout<<"convert Shortcut\n"; //std::cout<<"New plugin Shortcut\n"; + ITensor *back_tens = tensors[l->backLayer]; + /* + // plugin version IPlugin *plugin = new ShortcutRT(); - ITensor **inputs = new ITensor*[2]; inputs[0] = input; inputs[1] = back_tens; IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); checkNULL(lRT); + */ + + IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); + checkNULL(lRT); + return lRT; } diff --git a/src/Yolo3Detection.cpp b/src/Yolo3Detection.cpp index ed392f7..869c792 100644 --- a/src/Yolo3Detection.cpp +++ b/src/Yolo3Detection.cpp @@ -94,6 +94,8 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { netRT->infer(dim, input_d); TIMER_STOP dim.print(); + + stats.push_back(t_ns); } TIMER_START From 2f57ba122201a256e1853382b4aa7543c6faaa7d Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 16 Jan 2020 18:21:34 +0100 Subject: [PATCH 20/30] Update README.md --- README.md | 1 + 1 file changed, 1 insertion(+) diff --git a/README.md b/README.md index 3b16ec3..6a0a7ff 100644 --- a/README.md +++ b/README.md @@ -49,3 +49,4 @@ this will genereate a yolo3_berkeley.rt file that can be used for live detection ./yolo3_demo # launch detection on a demo video ./yolo3_demo yolo3_berkeley.rt /dev/video0 # launch detection on device 0 ``` +![Demo preview](https://github.com/ceccocats/tkDNN/releases/download/v0.4/demo.gif) From 146e1442495919de5627b7e4f3ad2adbacd8943c Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 16 Jan 2020 18:24:35 +0100 Subject: [PATCH 21/30] Update README.md --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 6a0a7ff..3d956a5 100644 --- a/README.md +++ b/README.md @@ -49,4 +49,4 @@ this will genereate a yolo3_berkeley.rt file that can be used for live detection ./yolo3_demo # launch detection on a demo video ./yolo3_demo yolo3_berkeley.rt /dev/video0 # launch detection on device 0 ``` -![Demo preview](https://github.com/ceccocats/tkDNN/releases/download/v0.4/demo.gif) +![demo](https://user-images.githubusercontent.com/11562617/72547657-540e7800-388d-11ea-83c6-49dfea2a0607.gif) From 8a4d1cac17f14f14c7db6bf3402c10bc83faeccb Mon Sep 17 00:00:00 2001 From: luca <228618@studenti.unimore.it> Date: Mon, 20 Jan 2020 14:51:38 +0100 Subject: [PATCH 22/30] compile with tensorrt 5 --- src/NetworkRT.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index d4fdf90..f27f70d 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -275,18 +275,18 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { if(l->act_mode == ACTIVATION_LEAKY) { //std::cout<<"New plugin LEAKY\n"; - /* +#if NV_TENSORRT_MAJOR < 6 // plugin version IPlugin *plugin = new ActivationLeakyRT(); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); checkNULL(lRT); return lRT; - */ - +#else IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU); lRT->setAlpha(0.1); checkNULL(lRT); return lRT; +#endif } else if(l->act_mode == CUDNN_ACTIVATION_RELU) { IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU); From a9c0db0bf6abc3f21611b1c13f55aaba383159e5 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 19:27:18 +0100 Subject: [PATCH 23/30] LSTM cudnn test --- .gitignore | 3 +- CMakeLists.txt | 3 + include/tkDNN/Layer.h | 76 +++++++++++++++++ src/LSTM.cpp | 142 ++++++++++++++++++++++++++++++++ src/utils.cpp | 3 +- tests/imuodom/imuodom.cpp | 70 ++++++++++++++++ tests/imuodom/infer.py | 74 +++++++++++++++++ tests/simple/test_model.py | 60 +++++++------- tests/simple/test_simple.cpp | 14 +--- tests/weights_exporter.py | 153 ++++++++++++++--------------------- 10 files changed, 464 insertions(+), 134 deletions(-) create mode 100644 src/LSTM.cpp create mode 100644 tests/imuodom/imuodom.cpp create mode 100644 tests/imuodom/infer.py diff --git a/.gitignore b/.gitignore index 02f2a8e..de78410 100644 --- a/.gitignore +++ b/.gitignore @@ -8,4 +8,5 @@ build/ *.h5 *.tar.gz *.weights -.idea/ \ No newline at end of file +.idea/ +*.hdf5 \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index b4b3c38..e3d8607 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -85,6 +85,9 @@ target_link_libraries(test_yolo3_berkeley tkDNN) add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp) target_link_libraries(test_yolo3_flir tkDNN) + +add_executable(test_imuodom tests/imuodom/imuodom.cpp) +target_link_libraries(test_imuodom tkDNN) ################################################################################ diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index ee73109..f15c781 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -9,8 +9,10 @@ namespace tk { namespace dnn { enum layerType_t { + LAYER_INPUT, LAYER_DENSE, LAYER_CONV2D, + LAYER_LSTM, LAYER_ACTIVATION, LAYER_FLATTEN, LAYER_MULADD, @@ -47,8 +49,10 @@ public: std::string getLayerName() { layerType_t type = getLayerType(); switch(type) { + case LAYER_INPUT: return "Input"; case LAYER_DENSE: return "Dense"; case LAYER_CONV2D: return "Conv2d"; + case LAYER_LSTM: return "LSTM"; case LAYER_ACTIVATION: return "Activation"; case LAYER_FLATTEN: return "Flatten"; case LAYER_MULADD: return "MulAdd"; @@ -105,6 +109,27 @@ public: }; +/** + Input layer (it doesnt need weigths) +*/ +class Input : public Layer { + +public: + + Input(Network *net, dataDim_t &dim, dnnType* srcData) : Layer(net) { + input_dim = dim; + output_dim = dim; + dstData = srcData; + } + virtual ~Input() {} + virtual layerType_t getLayerType() { return LAYER_INPUT; }; + + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData) { + return dstData; + } +}; + + /** Dense (full interconnection) layer */ @@ -148,6 +173,14 @@ protected: /** Convolutional 2D layer + + WEIGHTS shape: OUTCH, INCH, KH, KW ... + BIAS shape: OUTCH + + with BATCHNORM: + scales: OUTCH + means: OUTCH + variance: OUTCH */ class Conv2d : public LayerWgs { @@ -172,6 +205,49 @@ protected: size_t ws_sizeInBytes; }; +/** + Bidirectional LSTM layer + https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp + + numlayers = 1 # hardcoded as 1 + + PARAMS (numlayers*2): + layer0: + ( INCH, ? ) ??? + ( HIDDEN, ? ) ??? + ( HIDDEN * 8 ) ??? + layer2: + ( INCH, ? ) ??? + ( HIDDEN, ? ) ??? + ( HIDDEN * 8 ) ??? + + output shape: ( 2*HIDDEN, INH, INW ) +*/ +class LSTM : public Layer { + +public: + LSTM(Network *net, int hiddensize, std::string fname_weights); + virtual ~LSTM(); + virtual layerType_t getLayerType() { return LAYER_LSTM; }; + + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); + + int kernelH, kernelW, strideH, strideW, paddingH, paddingW; + +protected: + cudnnFilterDescriptor_t paramDesc; + cudnnTensorDescriptor_t hiddenStateTensorDesc, cellStateTensorDesc; + cudnnRNNDescriptor_t rnnDesc; + cudnnRNNDataDescriptor_t rnnDataDesc; + cudnnDropoutDescriptor_t dropDesc; + cudnnRNNAlgo_t algo; + + dnnType *hiddenStateData, *cellStateData; + dnnType *paramsSpace; + void* workSpace; + size_t ws_sizeInBytes; +}; + /** Flatten layer diff --git a/src/LSTM.cpp b/src/LSTM.cpp new file mode 100644 index 0000000..3176819 --- /dev/null +++ b/src/LSTM.cpp @@ -0,0 +1,142 @@ +#include + +#include "Layer.h" + +namespace tk { namespace dnn { + +LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : + Layer(net) { + + checkCUDNN( cudnnCreateFilterDescriptor(¶mDesc)); + checkCUDNN( cudnnCreateRNNDescriptor(&rnnDesc) ); + checkCUDNN( cudnnCreateRNNDataDescriptor(&rnnDataDesc) ); + checkCUDNN( cudnnCreateDropoutDescriptor(&dropDesc)); + + int n = input_dim.n; + int c = input_dim.c; + int h = input_dim.h; + int w = input_dim.w; + checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, + net->tensorFormat, net->dataType, n, 1, h, w) ); + + int numlayers = 1; + checkCUDNN( cudnnSetRNNDescriptor(net->cudnnHandle, rnnDesc, hiddensize, numlayers, dropDesc, + cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, + cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, cudnnRNNMode_t::CUDNN_LSTM, + cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType) ); + + // find dimension of params + size_t params_size = 0; + checkCUDNN( cudnnGetRNNParamsSize(net->cudnnHandle, rnnDesc, srcTensorDesc, ¶ms_size, net->dataType) ); + std::cout<<"Params size bytes: "<dataType, net->tensorFormat, 3, dimW)); + checkCuda( cudaMalloc(¶msSpace, params_size) ); + + + int numlinearlayers = 8; + + for(int i=0; icudnnHandle, rnnDesc, + i, srcTensorDesc, paramDesc, paramsSpace, + j, linLayerMatDesc, (void **)&linLayerMat)); + + if(linLayerMat == nullptr) { + FatalError("LSTM No weights in hidden layer"); + } + + cudnnDataType_t dataType; + cudnnTensorFormat_t format; + int nbDims; + int filterDimA[3]; + checkCUDNN(cudnnGetFilterNdDescriptor(linLayerMatDesc, 3, &dataType, + &format, &nbDims, filterDimA)); + std::cout<<"Wgs Dims: "<cudnnHandle, rnnDesc, + i, srcTensorDesc, paramDesc, paramsSpace, + j, linLayerBiasDesc, (void **)&linLayerBias)); + + if(linLayerMat == nullptr) { + FatalError("LSTM No bias in hidden layer"); + } + + checkCUDNN(cudnnGetFilterNdDescriptor(linLayerBiasDesc, 3, &dataType, + &format, &nbDims, filterDimA)); + std::cout<<"bias Dims: "<tensorFormat, net->dataType, 2*n, c, h, w) ); + checkCuda( cudaMalloc(&hiddenStateData, 2*input_dim.tot()*sizeof(dnnType)) ); + checkCUDNN( cudnnCreateTensorDescriptor(&cellStateTensorDesc)); + checkCUDNN( cudnnSetTensor4dDescriptor(cellStateTensorDesc, + net->tensorFormat, net->dataType, 2*n, c, h, w) ); + checkCuda( cudaMalloc(&cellStateData, 2*input_dim.tot()*sizeof(dnnType)) ); + + + output_dim = input_dim; + output_dim.c = hiddensize*2; + checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, + net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w) ); + + + + + //allocate data for infer result + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); +} + +LSTM::~LSTM() { + + checkCuda( cudaFree(dstData) ); +} + +dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { + + checkCUDNN(cudnnRNNForwardInference( + net->cudnnHandle, rnnDesc, 1, + &srcTensorDesc, srcData, + hiddenStateTensorDesc, hiddenStateData, + cellStateTensorDesc, cellStateData, + paramDesc, paramsSpace, + &dstTensorDesc, dstData, + hiddenStateTensorDesc, hiddenStateData, + cellStateTensorDesc, cellStateData, + workSpace, ws_sizeInBytes + )); + + return dstData; +} + +}} diff --git a/src/utils.cpp b/src/utils.cpp index e6f71f8..444318d 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -39,7 +39,8 @@ void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** dat *data_h = new dnnType[size]; if (!dataFile.read ((char*) *data_h, size_b)) { - error_s << "Error reading file " << fname; + error_s << "Error reading file " << fname << " with n of float: "< +#include "tkdnn.h" + +const char *i0_bin = "../tests/imuodom/layers/input0.bin"; +const char *i1_bin = "../tests/imuodom/layers/input1.bin"; +const char *i2_bin = "../tests/imuodom/layers/input2.bin"; +const char *o0_bin = "../tests/imuodom/layers/output0.bin"; +const char *o1_bin = "../tests/imuodom/layers/output1.bin"; +const char *output_bin = "../tests/imuodom/layers/output.bin"; + +const char *c0_bin = "../tests/imuodom/layers/conv1d_7.bin"; +const char *c1_bin = "../tests/imuodom/layers/conv1d_8.bin"; +const char *c2_bin = "../tests/imuodom/layers/conv1d_9.bin"; +const char *c3_bin = "../tests/imuodom/layers/conv1d_10.bin"; +const char *c4_bin = "../tests/imuodom/layers/conv1d_11.bin"; +const char *c5_bin = "../tests/imuodom/layers/conv1d_12.bin"; + +int main() { + + // Network layout + tk::dnn::dataDim_t dim0(1, 4, 1, 100); + tk::dnn::dataDim_t dim1(1, 3, 1, 100); + tk::dnn::dataDim_t dim2(1, 3, 1, 100); + + // Load input + dnnType *i0_d, *i1_d, *i2_d; + dnnType *i0_h, *i1_h, *i2_h; + readBinaryFile(i0_bin, dim0.tot(), &i0_h, &i0_d); + readBinaryFile(i1_bin, dim1.tot(), &i1_h, &i1_d); + readBinaryFile(i2_bin, dim2.tot(), &i2_h, &i2_d); + + tk::dnn::Network net(dim0); + tk::dnn::Input x0 (&net, dim0, i0_d); + tk::dnn::Conv2d x0_0(&net, 128, 1, 11, 1, 1, 0, 0, c0_bin); + tk::dnn::Conv2d x0_1(&net, 128, 1, 11, 1, 1, 0, 0, c1_bin); + tk::dnn::Pooling x0_2(&net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX); + + tk::dnn::Input x1 (&net, dim1, i1_d); + tk::dnn::Conv2d x1_0(&net, 128, 1, 11, 1, 1, 0, 0, c2_bin); + tk::dnn::Conv2d x1_1(&net, 128, 1, 11, 1, 1, 0, 0, c3_bin); + tk::dnn::Pooling x1_2(&net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX); + + tk::dnn::Input x2 (&net, dim2, i2_d); + tk::dnn::Conv2d x2_0(&net, 128, 1, 11, 1, 1, 0, 0, c4_bin); + tk::dnn::Conv2d x2_1(&net, 128, 1, 11, 1, 1, 0, 0, c5_bin); + tk::dnn::Pooling x2_2(&net, 1, 3, 1, 3, tk::dnn::tkdnnPoolingMode_t::POOLING_MAX); + + tk::dnn::Layer *concat_l[3] = { &x0_2, &x1_2, &x2_2 }; + tk::dnn::Route concat (&net, concat_l, 3); + + tk::dnn::LSTM lstm0(&net, 128, "ciao"); + + net.print(); + + dnnType *data; + tk::dnn::dataDim_t dim; + + TIMER_START + // Inference + data = net.infer(dim, data); dim.print(); + TIMER_STOP + + // Print real test + std::cout<<"\n==== CHECK RESULT ====\n"; + dnnType *out; + dnnType *out_h; + readBinaryFile(output_bin, dim.tot(), &out_h, &out); + checkResult(dim.tot(), data, out); + return 0; +} diff --git a/tests/imuodom/infer.py b/tests/imuodom/infer.py new file mode 100644 index 0000000..9d9929d --- /dev/null +++ b/tests/imuodom/infer.py @@ -0,0 +1,74 @@ +import keras +from keras.models import load_model +import keras.backend.tensorflow_backend as KTF +import numpy as np +import argparse +import tensorflow as tf +import os +import random +import struct +from keras.models import Sequential, Model + +def bin_write(f, data): + data = data.flatten() + fmt = 'f'*len(data) + bin = struct.pack(fmt, *data) + f.write(bin) + + +if __name__ == '__main__': + + + print("DATA FORMAT: ", keras.backend.image_data_format()) + + print("Load model: ", "ferrariS1.hdf5") + model = load_model("ferrariS1.hdf5") + model.summary() + + weights = model.get_weights() + + x_angle = np.random.rand(1,100,4) + x_gyro = np.random.rand(1,100,3) + x_acc = np.random.rand(1,100,3) + + [yhat_delta_p, yhat_delta_q] = model.predict([x_angle, x_gyro, x_acc], batch_size=1, verbose=1) + + layer_name = 'bidirectional_3' + intermediate_layer_model = Model(inputs=model.input, + outputs=model.get_layer(layer_name).output) + intermediate_output = intermediate_layer_model.predict([x_angle, x_gyro, x_acc]) + + x_angle = np.array([x_angle]) + x_gyro = np.array([x_gyro]) + x_acc = np.array([x_acc]) + intermediate_output = np.array([intermediate_output]) + + x_angle = x_angle.transpose(0, 3, 1, 2) + x_gyro = x_gyro.transpose(0, 3, 1, 2) + x_acc = x_acc.transpose(0, 3, 1, 2) + intermediate_output = intermediate_output.transpose(0, 3, 1, 2) + + print("x0: ", np.shape(x_angle)) + print("out: ",np.shape(intermediate_output)) + + x_angle = np.array(x_angle.flatten(), dtype=np.float32) + x_gyro = np.array(x_gyro.flatten(), dtype=np.float32) + x_acc = np.array(x_acc.flatten(), dtype=np.float32) + yhat_delta_p = np.array(yhat_delta_p.flatten(), dtype=np.float32) + yhat_delta_q = np.array(yhat_delta_q.flatten(), dtype=np.float32) + intermediate_output = np.array(intermediate_output.flatten(), dtype=np.float32) + + + f = open("layers/input0.bin", mode='wb') + bin_write(f, x_angle) + f = open("layers/input1.bin", mode='wb') + bin_write(f, x_gyro) + f = open("layers/input2.bin", mode='wb') + bin_write(f, x_acc) + f = open("layers/output0.bin", mode='wb') + bin_write(f, yhat_delta_p) + f = open("layers/output1.bin", mode='wb') + bin_write(f, yhat_delta_q) + f = open("layers/output.bin", mode='wb') + bin_write(f, intermediate_output) + diff --git a/tests/simple/test_model.py b/tests/simple/test_model.py index 60f17d8..3b1d053 100644 --- a/tests/simple/test_model.py +++ b/tests/simple/test_model.py @@ -1,43 +1,49 @@ import keras import numpy as np from keras.models import Sequential -from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape, Lambda +from keras.layers import Input, Dense, Activation, Flatten, Dropout, ELU, Reshape, Lambda, Conv1D from keras.layers.convolutional import Convolution2D, Convolution3D from keras.layers.pooling import MaxPooling2D, MaxPooling3D, AveragePooling3D from keras.models import Sequential, Model from keras.layers import Cropping2D import keras.backend.tensorflow_backend as KTF +import struct +from keras.models import Sequential, Model -def dense_model(): - model = Sequential() +def bin_write(f, data): + data = data.flatten() + fmt = 'f'*len(data) + bin = struct.pack(fmt, *data) + f.write(bin) + +def create_model(): + x1 = Input((6, 16), name='x1') + conv = Conv1D(4, 2)(x1) + model = Model([x1], [conv]) + model.summary() - model.add(Reshape((10, 10, 1), input_shape=(10, 10))) - model.add(Convolution2D(2, (4, 4), subsample=(2, 2), - bias_initializer='random_uniform', activation="relu")) - model.add(Convolution2D(4, (2, 2), subsample=(1, 1), - bias_initializer='random_uniform', activation="relu")) - model.add(Flatten()) - model.add(Dense(4, bias_initializer='random_uniform', activation="relu")) - sgd = keras.optimizers.Adam(lr=1e-4, decay=1e-8) - model.compile(optimizer=sgd, loss="mse") return model - if __name__ == '__main__': - print "DATA FORMAT: ", keras.backend.image_data_format() + print ("DATA FORMAT: ", keras.backend.image_data_format()) - model = dense_model() - model.save("net.h5") + model = create_model() + model.save("net.hdf5") - grid = np.random.rand(10,10) - X = grid[None,:,:] - i = np.array(grid.flatten(), dtype=np.float32) - print i - i.tofile("input.bin", format="f") - print "Input: ", X + x = np.random.rand(1,1,6,16) + r = model.predict( x[0], batch_size=1) + r = np.array([r]) + + x = x.transpose(0, 3, 1, 2) + r = r.transpose(0, 3, 1, 2) + print("in: ", np.shape(x)) + print("out: ", np.shape(r)) + + x = np.array(x.flatten(), dtype=np.float32) + f = open("input.bin", mode='wb') + bin_write(f, x) + + r = np.array(r.flatten(), dtype=np.float32) + f = open("output.bin", mode='wb') + bin_write(f, r) - r = model.predict( X, batch_size=1) - print np.shape(r) - print "Result: ", r - print "Result shape: ", np.shape(r) - r.tofile("output.bin", format="f") diff --git a/tests/simple/test_simple.cpp b/tests/simple/test_simple.cpp index a7628ea..3427d6d 100644 --- a/tests/simple/test_simple.cpp +++ b/tests/simple/test_simple.cpp @@ -2,23 +2,15 @@ #include "tkdnn.h" const char *input_bin = "../tests/simple/input.bin"; -const char *c0_bin = "../tests/simple/layers/c0.bin"; -const char *c1_bin = "../tests/simple/layers/c1.bin"; -const char *d2_bin = "../tests/simple/layers/d2.bin"; +const char *c0_bin = "../tests/simple/layers/conv1d_1.bin"; const char *output_bin = "../tests/simple/output.bin"; int main() { // Network layout - tk::dnn::dataDim_t dim(1, 1, 10, 10, 1); + tk::dnn::dataDim_t dim(1, 16, 1, 6); tk::dnn::Network net(dim); - tk::dnn::Conv2d l0(&net, 2, 4, 4, 2, 2, 0, 0, c0_bin); - tk::dnn::Activation l1(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Conv2d l2(&net, 4, 2, 2, 1, 1, 0, 0, c1_bin); - tk::dnn::Activation l3(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Flatten l4(&net); - tk::dnn::Dense l5(&net, 4, d2_bin); - tk::dnn::Activation l6(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d l0(&net, 4, 1, 2, 1, 1, 0, 0, c0_bin); // Load input dnnType *data; diff --git a/tests/weights_exporter.py b/tests/weights_exporter.py index df8bab3..912ef82 100644 --- a/tests/weights_exporter.py +++ b/tests/weights_exporter.py @@ -5,96 +5,51 @@ import numpy as np import argparse import tensorflow as tf import os -import msgpack -import lmdb import random +import struct +from keras.models import Sequential, Model -def export_dense(name, weights, bias): - print "######## EXPORT", name, "LAYER ########" - print "Original weighs:" - print weights - print bias, "\n" +def bin_write(f, data): + data = data.flatten() + fmt = 'f'*len(data) + bin = struct.pack(fmt, *data) + f.write(bin) - #input, filters - I, C = np.shape(weights) - B = np.shape(bias) - print "w shape: ", I, C - print "b shape: ", B +def export_layer(name, weights, bias): + print ("######## EXPORT", name, "LAYER ########") - wgs = [ [ j[i] for j in weights ] for i in xrange(C) ] - wgs = np.array(wgs, dtype=np.float32) - - print "REPOSITIONED WEIGHTS:" - print wgs + print("wgs pretranpose: ", np.shape(weights)) + # convert NHWC to NCHW + if(weights.ndim == 4): + weights = weights.transpose(3,2,0,1) + elif(weights.ndim == 3): + weights = weights.transpose(2,1,0) + else: + print("Ndim", weights.ndim) + raise("not implemented with dim" ) + print("weights: ", np.shape(weights)) + print("bias: ", np.shape(bias)) + + weights = np.array(weights.flatten(), dtype=np.float32) bias = np.array(bias, dtype=np.float32) - wgs.tofile(name + ".bin", format="f") - bias.tofile(name + ".bias.bin", format="f") - print "WEIGHTS saved\n" + print(len(weights) + len(bias)) -def export_conv2d(name, weights, bias): - print "######## EXPORT", name, "LAYER ########" - print "Original weighs:" - print weights - print bias, "\n" - - # height, width, input, filters - H, W, N, C = np.shape(weights) - B = np.shape(bias) - print "w shape: ", N, C, H, W - print "b shape: ", B + f = open(name + ".bin", mode='wb') + bin_write(f, weights) + bin_write(f, bias) + print ("WEIGHTS saved\n") - wgs = weights.transpose() - wgs = wgs.transpose(0, 1, 3, 2) - print "Final shape:", np.shape(wgs) - wgs = np.array(wgs.flatten(), dtype=np.float32) - - print "REPOSITIONED WEIGHTS:" - print wgs - - bias = np.array(bias, dtype=np.float32) - - wgs.tofile(name + ".bin", format="f") - bias.tofile(name + ".bias.bin", format="f") - print "WEIGHTS saved\n" - -def export_conv3d(name, weights, bias): - print "######## EXPORT", name, "LAYER ########" - print "Original weighs:" - print weights - print bias, "\n" - - print np.shape(weights) - # height, width, input, thickness, filters - H, W, T, N, C = np.shape(weights) - B = np.shape(bias) - print "w shape: ", T, C, H, W #thickness is number of images for cudnn - print "b shape: ", B - - wgs = weights.transpose() - wgs = wgs.transpose(0, 1, 4, 3, 2) - print "Final shape:", np.shape(wgs) - wgs = np.array(wgs.flatten(), dtype=np.float32) - - print "REPOSITIONED WEIGHTS:" - print wgs - - bias = np.array(bias, dtype=np.float32) - - wgs.tofile(name + ".bin", format="f") - bias.tofile(name + ".bias.bin", format="f") - print "WEIGHTS saved\n" - - -def get_session(gpu_fraction=0.5): - gpu_options = tf.GPUOptions(allow_growth=True) - #per_process_gpu_memory_fraction=gpu_fraction) - return tf.Session(config=tf.ConfigProto(gpu_options=gpu_options)) +def export_bidir(name, weights): + print ("######## EXPORT", name, "LAYER ########") + + for w in weights: + print(np.shape(w)) #https://github.com/fchollet/keras/wiki/Converting-convolution-kernels-from-Theano-to-TensorFlow-and-vice-versa if __name__ == '__main__': - KTF.set_session(get_session()) + print("DATA FORMAT: ", keras.backend.image_data_format()) parser = argparse.ArgumentParser(description='KERAS WEIGHTS EXPORTER TO CUDNN') parser.add_argument('model',type=str, @@ -103,31 +58,41 @@ if __name__ == '__main__': args = parser.parse_args() - print "DATA FORMAT: ", keras.backend.image_data_format() + print("DATA FORMAT: ", keras.backend.image_data_format()) - print "Load model: ", args.model + print("Load model: ", args.model) model = load_model(args.model) + model.summary() + weights = model.get_weights() ws = np.shape(weights) - print "Weights shape:", ws + print("Weights shape:", ws) if not os.path.exists(args.output): os.makedirs(args.output) - num = 0 + name_num = 0 for l in model.layers: - name = l.name - if name.startswith("conv3d"): - export_conv3d(args.output + "/conv" + str(name_num), weights[num], weights[num+1]) - elif name.startswith("conv2d"): - export_conv2d(args.output + "/conv" + str(name_num), weights[num], weights[num+1]) - elif name.startswith("dense"): - export_dense(args.output + "/dense" + str(name_num), weights[num], weights[num+1]) - else: - print "skip:", name, "has no weights" - continue - name_num += 1 - num += 2 + print("\n\nNAME: ", l.name) + print("input: ", l.input_shape, " output: ", l.output_shape) + wgs = l.get_weights() + print("wgs num: ", len(wgs)) + + name = l.name + if name.startswith("conv3d"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("conv2d"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("conv1d"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("dense"): + export_layer(args.output + "/" + name, wgs[0], wgs[1]) + elif name.startswith("bidirectional"): + export_bidir(args.output + "/" + name, wgs) + else: + print ("skip:", name, "has no weights") + continue + From 03d39d991c5932aa4c94f7b86fc8c64e9d48238b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 23:04:29 +0100 Subject: [PATCH 24/30] LSTM to be tested --- include/tkDNN/Layer.h | 40 ++++--- src/LSTM.cpp | 246 ++++++++++++++++++++++++------------------ 2 files changed, 170 insertions(+), 116 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index f15c781..c538fc9 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -207,9 +207,11 @@ protected: /** Bidirectional LSTM layer + + implementation info: https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp - - numlayers = 1 # hardcoded as 1 + https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h + https://stackoverflow.com/a/38737941 PARAMS (numlayers*2): layer0: @@ -221,7 +223,9 @@ protected: ( HIDDEN, ? ) ??? ( HIDDEN * 8 ) ??? - output shape: ( 2*HIDDEN, INH, INW ) + OUTPUT shape: + (N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=True) ---> (N, 2*HIDDEN, 1, W) # W is seqLength + (N, C, 1, W) ---> LSTM(HIDDEN, returnSeq=False) ---> (N, 2*HIDDEN, 1, 1) */ class LSTM : public Layer { @@ -232,20 +236,28 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - int kernelH, kernelW, strideH, strideW, paddingH, paddingW; + const bool bidirectional = 1; /**> is the net bidir */ + int stateSize = 0; /**> number of hidden states */ + int seqLen = 0; /**> number of timesteps */ + int numLayers = 1; /**> number of internal layers */ protected: - cudnnFilterDescriptor_t paramDesc; - cudnnTensorDescriptor_t hiddenStateTensorDesc, cellStateTensorDesc; - cudnnRNNDescriptor_t rnnDesc; - cudnnRNNDataDescriptor_t rnnDataDesc; - cudnnDropoutDescriptor_t dropDesc; - cudnnRNNAlgo_t algo; + cudnnRNNDescriptor_t rnnDesc; + cudnnDropoutDescriptor_t dropoutDesc; + dnnType *dropout_states_, *work_space_; - dnnType *hiddenStateData, *cellStateData; - dnnType *paramsSpace; - void* workSpace; - size_t ws_sizeInBytes; + size_t workspace_byte_, reserve_space_byte_, dropout_byte_; + int workspace_size_, dropout_size_; + + std::vector x_desc_vec_, y_desc_vec_, dx_desc_vec_, dy_desc_vec_; + cudnnTensorDescriptor_t hx_desc_, cx_desc_; + cudnnTensorDescriptor_t hy_desc_, cy_desc_; + cudnnTensorDescriptor_t dhx_desc_, dcx_desc_; + cudnnTensorDescriptor_t dhy_desc_, dcy_desc_; + dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + + cudnnFilterDescriptor_t w_desc_, dw_desc_; + dnnType *w_ptr, *dw_ptr; }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index 3176819..be12cec 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -7,111 +7,143 @@ namespace tk { namespace dnn { LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : Layer(net) { - checkCUDNN( cudnnCreateFilterDescriptor(¶mDesc)); - checkCUDNN( cudnnCreateRNNDescriptor(&rnnDesc) ); - checkCUDNN( cudnnCreateRNNDataDescriptor(&rnnDataDesc) ); - checkCUDNN( cudnnCreateDropoutDescriptor(&dropDesc)); + int batchSize = input_dim.n; + int inputSize = input_dim.c; + seqLen = input_dim.w; + stateSize = hiddensize; - int n = input_dim.n; - int c = input_dim.c; - int h = input_dim.h; - int w = input_dim.w; - checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc, - net->tensorFormat, net->dataType, n, 1, h, w) ); + std::cout<<"LSTM seqLen: "<cudnnHandle, rnnDesc, hiddensize, numlayers, dropDesc, - cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, - cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, cudnnRNNMode_t::CUDNN_LSTM, - cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType) ); + // init Tensor Descriptors + std::vector x_vec(seqLen); + std::vector y_vec(seqLen); + std::vector dx_vec(seqLen); + std::vector dy_vec(seqLen); - // find dimension of params - size_t params_size = 0; - checkCUDNN( cudnnGetRNNParamsSize(net->cudnnHandle, rnnDesc, srcTensorDesc, ¶ms_size, net->dataType) ); - std::cout<<"Params size bytes: "<dataType, net->tensorFormat, 3, dimW)); - checkCuda( cudaMalloc(¶msSpace, params_size) ); + checkCUDNN(cudnnSetTensorNdDescriptor(x_vec[i], + net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dx_vec[i], + net->dataType, 3, dimA, strideA)); + dimA[0] = batchSize; + dimA[1] = bidirectional ? stateSize*2 : stateSize; + dimA[2] = 1; + strideA[0] = dimA[2] * dimA[1]; + strideA[1] = dimA[2]; + strideA[2] = 1; - - int numlinearlayers = 8; - - for(int i=0; icudnnHandle, rnnDesc, - i, srcTensorDesc, paramDesc, paramsSpace, - j, linLayerMatDesc, (void **)&linLayerMat)); - - if(linLayerMat == nullptr) { - FatalError("LSTM No weights in hidden layer"); - } - - cudnnDataType_t dataType; - cudnnTensorFormat_t format; - int nbDims; - int filterDimA[3]; - checkCUDNN(cudnnGetFilterNdDescriptor(linLayerMatDesc, 3, &dataType, - &format, &nbDims, filterDimA)); - std::cout<<"Wgs Dims: "<cudnnHandle, rnnDesc, - i, srcTensorDesc, paramDesc, paramsSpace, - j, linLayerBiasDesc, (void **)&linLayerBias)); - - if(linLayerMat == nullptr) { - FatalError("LSTM No bias in hidden layer"); - } - - checkCUDNN(cudnnGetFilterNdDescriptor(linLayerBiasDesc, 3, &dataType, - &format, &nbDims, filterDimA)); - std::cout<<"bias Dims: "<dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dy_vec[i], + net->dataType, 3, dimA, strideA)); } + // apply tensordesc + x_desc_vec_ = x_vec; + y_desc_vec_ = y_vec; + dx_desc_vec_ = dx_vec; + dy_desc_vec_ = dy_vec; + // set the state tensors + dimA[0] = numLayers * (bidirectional ? 2 : 1); + dimA[1] = batchSize; + dimA[2] = stateSize; + strideA[0] = dimA[2] * dimA[1]; + strideA[1] = dimA[2]; + strideA[2] = 1; + checkCUDNN(cudnnCreateTensorDescriptor(&hx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&cx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&hy_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&cy_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dhx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dcx_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dhy_desc_)); + checkCUDNN(cudnnCreateTensorDescriptor(&dcy_desc_)); + checkCUDNN(cudnnSetTensorNdDescriptor(hx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(cx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dhx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dcx_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dhy_desc_, net->dataType, 3, dimA, strideA)); + checkCUDNN(cudnnSetTensorNdDescriptor(dcy_desc_, net->dataType, 3, dimA, strideA)); + // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + checkCuda( cudaMalloc(&hx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&hy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCUDNN( cudnnCreateTensorDescriptor(&hiddenStateTensorDesc)); - checkCUDNN( cudnnSetTensor4dDescriptor(hiddenStateTensorDesc, - net->tensorFormat, net->dataType, 2*n, c, h, w) ); - checkCuda( cudaMalloc(&hiddenStateData, 2*input_dim.tot()*sizeof(dnnType)) ); - checkCUDNN( cudnnCreateTensorDescriptor(&cellStateTensorDesc)); - checkCUDNN( cudnnSetTensor4dDescriptor(cellStateTensorDesc, - net->tensorFormat, net->dataType, 2*n, c, h, w) ); - checkCuda( cudaMalloc(&cellStateData, 2*input_dim.tot()*sizeof(dnnType)) ); + + // Create Dropout descriptors // TODO: ??? IS IT NECESSARY ??? + float dropoutprob = 0.1f; // random val ???? + checkCUDNN(cudnnCreateDropoutDescriptor(&dropoutDesc)); + checkCUDNN(cudnnDropoutGetStatesSize(net->cudnnHandle, &dropout_byte_)); + dropout_size_ = dropout_byte_ / sizeof(dnnType); + checkCuda( cudaMalloc(&dropout_states_, dropout_byte_) ); + uint64_t seed_ = 17 + rand() % 4096; // NOLINT(runtime/threadsafe_fn) + checkCUDNN(cudnnSetDropoutDescriptor(dropoutDesc, + net->cudnnHandle, dropoutprob, dropout_states_, dropout_byte_, seed_)); + + + // RNN descriptors + checkCUDNN(cudnnCreateRNNDescriptor(&rnnDesc)); + + checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, + rnnDesc, stateSize, numLayers, dropoutDesc, + cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, + cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, + cudnnRNNMode_t::CUDNN_LSTM, + cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, + net->dataType)); + + + // Get temp space sizes + checkCUDNN(cudnnGetRNNWorkspaceSize(net->cudnnHandle, + rnnDesc, seqLen, x_desc_vec_.data(), &workspace_byte_)); + workspace_size_ = workspace_byte_ / sizeof(dnnType); + checkCuda( cudaMalloc(&work_space_, workspace_byte_) ); + + + // Check that number of params are correct + size_t cudnn_param_size; + checkCUDNN(cudnnGetRNNParamsSize(net->cudnnHandle, + rnnDesc,x_desc_vec_[0], &cudnn_param_size, net->dataType)); + int cudnn_params = cudnn_param_size/sizeof(dnnType); + std::cout<<"LSTM params size: "<dataType, net->tensorFormat, 3, dim_w)); + checkCUDNN(cudnnSetFilterNdDescriptor(dw_desc_, + net->dataType, net->tensorFormat, 3, dim_w)); + // allocate params dnnType *w_ptr, *dw_ptr; + checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&dw_ptr, cudnn_params*sizeof(dnnType)) ); output_dim = input_dim; - output_dim.c = hiddensize*2; - checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, - net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w) ); - - - + output_dim.c = stateSize*2; //allocate data for infer result checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); @@ -123,19 +155,29 @@ LSTM::~LSTM() { } dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { + std::cout<<"LSTM infer\n"; - checkCUDNN(cudnnRNNForwardInference( - net->cudnnHandle, rnnDesc, 1, - &srcTensorDesc, srcData, - hiddenStateTensorDesc, hiddenStateData, - cellStateTensorDesc, cellStateData, - paramDesc, paramsSpace, - &dstTensorDesc, dstData, - hiddenStateTensorDesc, hiddenStateData, - cellStateTensorDesc, cellStateData, - workSpace, ws_sizeInBytes - )); + checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle, + rnnDesc, + seqLen, + x_desc_vec_.data(), // input array of desc + srcData, // input pointer + hx_desc_, // initial hidden state desc + hx_ptr, // initial hidden state pointer + cx_desc_, // initial cell state desc + cx_ptr, // initial cell state pointer + w_desc_, // weights desc + w_ptr, // weights pointer + y_desc_vec_.data(), // output desc + dstData, // output pointer + hy_desc_, // final hidden state desc + hy_ptr, // final hidden state pointer + cy_desc_, // final cell state desc + cy_ptr, // final cell state pointer + work_space_, // workspace pointer + workspace_byte_)); // workspace size + dim = output_dim; return dstData; } From c1c2173e4d9ea4f59d95ee069fa9e0fc1f58f31b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 23:10:48 +0100 Subject: [PATCH 25/30] removed unused var --- include/tkDNN/Layer.h | 10 ++++------ src/LSTM.cpp | 42 +++++++++++++----------------------------- 2 files changed, 17 insertions(+), 35 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index c538fc9..596fbec 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -246,18 +246,16 @@ protected: cudnnDropoutDescriptor_t dropoutDesc; dnnType *dropout_states_, *work_space_; - size_t workspace_byte_, reserve_space_byte_, dropout_byte_; + size_t workspace_byte_, dropout_byte_; int workspace_size_, dropout_size_; - std::vector x_desc_vec_, y_desc_vec_, dx_desc_vec_, dy_desc_vec_; + std::vector x_desc_vec_, y_desc_vec_; cudnnTensorDescriptor_t hx_desc_, cx_desc_; cudnnTensorDescriptor_t hy_desc_, cy_desc_; - cudnnTensorDescriptor_t dhx_desc_, dcx_desc_; - cudnnTensorDescriptor_t dhy_desc_, dcy_desc_; dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; - cudnnFilterDescriptor_t w_desc_, dw_desc_; - dnnType *w_ptr, *dw_ptr; + cudnnFilterDescriptor_t w_desc_; + dnnType *w_ptr; }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index be12cec..46a0593 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -17,16 +17,12 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : // init Tensor Descriptors std::vector x_vec(seqLen); std::vector y_vec(seqLen); - std::vector dx_vec(seqLen); - std::vector dy_vec(seqLen); int dimA[3]; int strideA[3]; for (int i = 0; i < seqLen; i++) { checkCUDNN(cudnnCreateTensorDescriptor(&x_vec[i])); checkCUDNN(cudnnCreateTensorDescriptor(&y_vec[i])); - checkCUDNN(cudnnCreateTensorDescriptor(&dx_vec[i])); - checkCUDNN(cudnnCreateTensorDescriptor(&dy_vec[i])); dimA[0] = batchSize; dimA[1] = inputSize; @@ -36,29 +32,21 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : strideA[0] = dimA[2] * dimA[1]; strideA[1] = dimA[2]; strideA[2] = 1; - checkCUDNN(cudnnSetTensorNdDescriptor(x_vec[i], net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dx_vec[i], - net->dataType, 3, dimA, strideA)); + dimA[0] = batchSize; dimA[1] = bidirectional ? stateSize*2 : stateSize; dimA[2] = 1; strideA[0] = dimA[2] * dimA[1]; strideA[1] = dimA[2]; strideA[2] = 1; - checkCUDNN(cudnnSetTensorNdDescriptor(y_vec[i], net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dy_vec[i], - net->dataType, 3, dimA, strideA)); } - // apply tensordesc x_desc_vec_ = x_vec; y_desc_vec_ = y_vec; - dx_desc_vec_ = dx_vec; - dy_desc_vec_ = dy_vec; // set the state tensors @@ -72,18 +60,10 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : checkCUDNN(cudnnCreateTensorDescriptor(&cx_desc_)); checkCUDNN(cudnnCreateTensorDescriptor(&hy_desc_)); checkCUDNN(cudnnCreateTensorDescriptor(&cy_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dhx_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dcx_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dhy_desc_)); - checkCUDNN(cudnnCreateTensorDescriptor(&dcy_desc_)); checkCUDNN(cudnnSetTensorNdDescriptor(hx_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(cx_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dhx_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dcx_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dhy_desc_, net->dataType, 3, dimA, strideA)); - checkCUDNN(cudnnSetTensorNdDescriptor(dcy_desc_, net->dataType, 3, dimA, strideA)); // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; checkCuda( cudaMalloc(&hx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); checkCuda( cudaMalloc(&cx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); @@ -130,28 +110,32 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : // Set param descriptors checkCUDNN(cudnnCreateFilterDescriptor(&w_desc_)); - checkCUDNN(cudnnCreateFilterDescriptor(&dw_desc_)); int dim_w[3] = {1, 1, 1}; dim_w[0] = cudnn_params; checkCUDNN(cudnnSetFilterNdDescriptor(w_desc_, net->dataType, net->tensorFormat, 3, dim_w)); - checkCUDNN(cudnnSetFilterNdDescriptor(dw_desc_, - net->dataType, net->tensorFormat, 3, dim_w)); - // allocate params dnnType *w_ptr, *dw_ptr; + // allocate params dnnType *w_ptr; checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&dw_ptr, cudnn_params*sizeof(dnnType)) ); - + // set output dim output_dim = input_dim; output_dim.c = stateSize*2; - + //allocate data for infer result checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); } LSTM::~LSTM() { + checkCuda(cudaFree(hx_ptr)); + checkCuda(cudaFree(cx_ptr)); + checkCuda(cudaFree(hy_ptr)); + checkCuda(cudaFree(cy_ptr)); + checkCuda(cudaFree(w_ptr )); - checkCuda( cudaFree(dstData) ); + checkCuda(cudaFree(work_space_ )); + checkCuda(cudaFree(dropout_states_)); + + checkCuda(cudaFree(dstData)); } dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { From 4fa5d2c231900e8ee9bfbab908b2ac97add974db Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 13 Feb 2020 23:21:28 +0100 Subject: [PATCH 26/30] lstm return seq --- include/tkDNN/Layer.h | 3 ++- src/LSTM.cpp | 16 ++++++++++++---- tests/imuodom/imuodom.cpp | 3 ++- 3 files changed, 16 insertions(+), 6 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 596fbec..240d9a2 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -230,13 +230,14 @@ protected: class LSTM : public Layer { public: - LSTM(Network *net, int hiddensize, std::string fname_weights); + LSTM(Network *net, int hiddensize, bool returnSeq, std::string fname_weights); virtual ~LSTM(); virtual layerType_t getLayerType() { return LAYER_LSTM; }; virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); const bool bidirectional = 1; /**> is the net bidir */ + bool returnSeq = false; /**> if false return only the result of last timestep */ int stateSize = 0; /**> number of hidden states */ int seqLen = 0; /**> number of timesteps */ int numLayers = 1; /**> number of internal layers */ diff --git a/src/LSTM.cpp b/src/LSTM.cpp index 46a0593..c31bccf 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -4,9 +4,10 @@ namespace tk { namespace dnn { -LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : +LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weights) : Layer(net) { + this->returnSeq = returnSeq; int batchSize = input_dim.n; int inputSize = input_dim.c; seqLen = input_dim.w; @@ -117,12 +118,19 @@ LSTM::LSTM( Network *net, int hiddensize, std::string fname_weights) : // allocate params dnnType *w_ptr; checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); + + + //allocate data for infer result + int dstDim = input_dim.n * stateSize*2 * input_dim.h * input_dim.w; + checkCuda( cudaMalloc(&dstData, dstDim*sizeof(dnnType)) ); + // set output dim output_dim = input_dim; output_dim.c = stateSize*2; - - //allocate data for infer result - checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); + if(!returnSeq) { + output_dim.h = 1; + output_dim.w = 1; + } } LSTM::~LSTM() { diff --git a/tests/imuodom/imuodom.cpp b/tests/imuodom/imuodom.cpp index 14146b2..6fc3fba 100644 --- a/tests/imuodom/imuodom.cpp +++ b/tests/imuodom/imuodom.cpp @@ -48,7 +48,8 @@ int main() { tk::dnn::Layer *concat_l[3] = { &x0_2, &x1_2, &x2_2 }; tk::dnn::Route concat (&net, concat_l, 3); - tk::dnn::LSTM lstm0(&net, 128, "ciao"); + tk::dnn::LSTM lstm0(&net, 128, true, "ciao"); + tk::dnn::LSTM lstm1(&net, 128, false, "ciao"); net.print(); From 4746121d438c72d0287a6ff7edb193a2da68029d Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Sat, 15 Feb 2020 20:37:08 +0100 Subject: [PATCH 27/30] LSTM params --- include/tkDNN/Layer.h | 5 ++- src/LSTM.cpp | 94 +++++++++++++++++++++++++++++++++------ tests/imuodom/imuodom.cpp | 20 ++++++--- tests/imuodom/infer.py | 1 + tests/weights_exporter.py | 16 +++++-- 5 files changed, 111 insertions(+), 25 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 240d9a2..319df6a 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -212,6 +212,7 @@ protected: https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h https://stackoverflow.com/a/38737941 + https://colah.github.io/posts/2015-08-Understanding-LSTMs/ PARAMS (numlayers*2): layer0: @@ -236,7 +237,7 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - const bool bidirectional = 1; /**> is the net bidir */ + const bool bidirectional = false; /**> is the net bidir */ bool returnSeq = false; /**> if false return only the result of last timestep */ int stateSize = 0; /**> number of hidden states */ int seqLen = 0; /**> number of timesteps */ @@ -254,9 +255,11 @@ protected: cudnnTensorDescriptor_t hx_desc_, cx_desc_; cudnnTensorDescriptor_t hy_desc_, cy_desc_; dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + int stateDataDim; cudnnFilterDescriptor_t w_desc_; dnnType *w_ptr; + dnnType *w_h; }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index c31bccf..ce2a9f5 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -66,10 +66,12 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnSetTensorNdDescriptor(hy_desc_, net->dataType, 3, dimA, strideA)); checkCUDNN(cudnnSetTensorNdDescriptor(cy_desc_, net->dataType, 3, dimA, strideA)); // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; - checkCuda( cudaMalloc(&hx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&cx_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&hy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); - checkCuda( cudaMalloc(&cy_ptr, dimA[0]*dimA[1]*dimA[2]*sizeof(dnnType)) ); + stateDataDim = dimA[0]*dimA[1]*dimA[2]; + checkCuda( cudaMalloc(&hx_ptr, stateDataDim*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cx_ptr, stateDataDim*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&hy_ptr, stateDataDim*sizeof(dnnType)) ); + checkCuda( cudaMalloc(&cy_ptr, stateDataDim*sizeof(dnnType)) ); + // Create Dropout descriptors // TODO: ??? IS IT NECESSARY ??? @@ -89,7 +91,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, rnnDesc, stateSize, numLayers, dropoutDesc, cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, - cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL, + (bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL), cudnnRNNMode_t::CUDNN_LSTM, cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType)); @@ -115,22 +117,81 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig dim_w[0] = cudnn_params; checkCUDNN(cudnnSetFilterNdDescriptor(w_desc_, net->dataType, net->tensorFormat, 3, dim_w)); - // allocate params dnnType *w_ptr; - checkCuda( cudaMalloc(&w_ptr, cudnn_params*sizeof(dnnType)) ); - + // load params + readBinaryFile(fname_weights, cudnn_params, &w_h, &w_ptr); //allocate data for infer result - int dstDim = input_dim.n * stateSize*2 * input_dim.h * input_dim.w; + int dstDim = input_dim.n * stateSize*(bidirectional ? 2 : 1) * input_dim.h * input_dim.w; checkCuda( cudaMalloc(&dstData, dstDim*sizeof(dnnType)) ); // set output dim output_dim = input_dim; - output_dim.c = stateSize*2; + output_dim.c = stateSize*(bidirectional ? 2 : 1); if(!returnSeq) { output_dim.h = 1; output_dim.w = 1; } + + + + + // Query weight layout + cudnnFilterDescriptor_t m_desc; + checkCUDNN(cudnnCreateFilterDescriptor(&m_desc)); + dnnType *p; + int n = 8; // lstm layers + + printCenteredTitle("WEIGHTS", '=', 20); + for (int i = 0; i < numLayers*(bidirectional?2:1); ++i) { + for (int j = 0; j < n; ++j) { + + checkCUDNN(cudnnGetRNNLinLayerMatrixParams(net->cudnnHandle, rnnDesc, + i, x_desc_vec_[0], w_desc_, 0, j, m_desc, (void**)&p)); + + std::cout << "ptr: " << ((int64_t)(p - NULL))/sizeof(dnnType)<<"\n"; + + cudnnDataType_t t; + cudnnTensorFormat_t f; + int ndim = 5; + int dims[5] = {0, 0, 0, 0, 0}; + checkCUDNN(cudnnGetFilterNdDescriptor(m_desc, ndim, &t, &f, &ndim, &dims[0])); + std::cout << "(layer, linlayer): " << i << " " << j << "\n"; + + int tot = 1; + for (int i = 0; i < ndim; ++i) { + std::cout << dims[i] << " "; + tot *= dims[i]; + } + std::cout<<"\t-> "< Date: Sun, 16 Feb 2020 16:28:39 +0100 Subject: [PATCH 28/30] works but it need cleaning --- include/tkDNN/Layer.h | 3 +- include/tkDNN/utils.h | 2 +- src/LSTM.cpp | 163 ++++++++++++++++++++++++++++------- src/utils.cpp | 4 +- tests/imuodom/imuodom.cpp | 26 +++--- tests/imuodom/infer.py | 25 +++--- tests/simple/test_model.py | 15 ++-- tests/simple/test_simple.cpp | 8 +- tests/weights_exporter.py | 44 ++++++++-- 9 files changed, 219 insertions(+), 71 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 319df6a..f12027f 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -237,7 +237,7 @@ public: virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); - const bool bidirectional = false; /**> is the net bidir */ + const bool bidirectional = true; /**> is the net bidir */ bool returnSeq = false; /**> if false return only the result of last timestep */ int stateSize = 0; /**> number of hidden states */ int seqLen = 0; /**> number of timesteps */ @@ -260,6 +260,7 @@ protected: cudnnFilterDescriptor_t w_desc_; dnnType *w_ptr; dnnType *w_h; + dnnType *wf_ptr, *wb_ptr; // params pointer forward and backward layer }; diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index dc34a31..3fa9d34 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -91,7 +91,7 @@ void printCenteredTitle(const char *title, char fill, int dim); bool fileExist(const char *fname); void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0); -int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true); +int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true, int limit = 10); void printDeviceVector(int size, dnnType* vec_d, bool device = true); void resize(int size, dnnType **data); diff --git a/src/LSTM.cpp b/src/LSTM.cpp index ce2a9f5..00d8f76 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -37,7 +37,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig net->dataType, 3, dimA, strideA)); dimA[0] = batchSize; - dimA[1] = bidirectional ? stateSize*2 : stateSize; + dimA[1] = stateSize; dimA[2] = 1; strideA[0] = dimA[2] * dimA[1]; strideA[1] = dimA[2]; @@ -51,7 +51,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig // set the state tensors - dimA[0] = numLayers * (bidirectional ? 2 : 1); + dimA[0] = numLayers; dimA[1] = batchSize; dimA[2] = stateSize; strideA[0] = dimA[2] * dimA[1]; @@ -91,7 +91,8 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, rnnDesc, stateSize, numLayers, dropoutDesc, cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT, - (bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL), + //(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL), + cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL, cudnnRNNMode_t::CUDNN_LSTM, cudnnRNNAlgo_t::CUDNN_RNN_ALGO_STANDARD, net->dataType)); @@ -119,23 +120,26 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig net->dataType, net->tensorFormat, 3, dim_w)); // load params - readBinaryFile(fname_weights, cudnn_params, &w_h, &w_ptr); - - //allocate data for infer result - int dstDim = input_dim.n * stateSize*(bidirectional ? 2 : 1) * input_dim.h * input_dim.w; - checkCuda( cudaMalloc(&dstData, dstDim*sizeof(dnnType)) ); + readBinaryFile(fname_weights, cudnn_params*2, &w_h, &w_ptr); + // set forward and backward params + wf_ptr = w_ptr; + wb_ptr = w_ptr + cudnn_params; + std::cout<<"wf: "<cublasHandle, srcData, trans, dim.c, dim.h*dim.w*dim.l); + srcData = trans; + + // reposition in invered order + dnnType *srcBack; + checkCuda( cudaMalloc(&srcBack, dim.tot()*sizeof(dnnType))); + for(int i=0; icudnnHandle, + rnnDesc, + seqLen, // number of time steps (nT) + x_desc_vec_.data(), // input array of desc (nT*nC_in) + srcData, // input pointer + hx_desc_, // initial hidden state desc + hx_ptr, // initial hidden state pointer + cx_desc_, // initial cell state desc + cx_ptr, // initial cell state pointer + w_desc_, // weights desc + wf_ptr, // weights pointer + y_desc_vec_.data(), // output desc (nT*nC_out) + dstF, // output pointer + hy_desc_, // final hidden state desc + hy_ptr, // final hidden state pointer + cy_desc_, // final cell state desc + cy_ptr, // final cell state pointer + work_space_, // workspace pointer + workspace_byte_)); // workspace size + } + std::cout<<"OUTPUT F:\n"; + printDeviceVector(singleOutput.tot(), dstF); + + std::cout<<"INPUT:\n"; + printDeviceVector(input_dim.tot(), srcBack); + + // backward + { + // reset states + checkCuda( cudaMemset(hx_ptr, 0, stateDataDim*sizeof(float)) ); + checkCuda( cudaMemset(cx_ptr, 0, stateDataDim*sizeof(float)) ); + + checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle, + rnnDesc, + seqLen, // number of time steps (nT) + x_desc_vec_.data(), // input array of desc (nT*nC_in) + srcBack, // input pointer + hx_desc_, // initial hidden state desc + hx_ptr, // initial hidden state pointer + cx_desc_, // initial cell state desc + cx_ptr, // initial cell state pointer + w_desc_, // weights desc + wb_ptr, // weights pointer + y_desc_vec_.data(), // output desc (nT*nC_out) + dstB, // output pointer + hy_desc_, // final hidden state desc + hy_ptr, // final hidden state pointer + cy_desc_, // final cell state desc + cy_ptr, // final cell state pointer + work_space_, // workspace pointer + workspace_byte_)); // workspace size + } + + + // reposition in invered order + dnnType *dstBack; + checkCuda( cudaMalloc(&dstBack, singleOutput.tot()*sizeof(dnnType))); + for(int i=0; icublasHandle, dstF, trans, + singleOutput.h*singleOutput.w*singleOutput.l, singleOutput.c); + // backward transpose + matrixTranspose(net->cublasHandle, dstB, trans + singleOutput.tot(), + singleOutput.h*singleOutput.w*singleOutput.l, singleOutput.c); + dstData = trans; + } else { + // copy last of forward + checkCuda( cudaMemcpy(trans, dstF + singleOutput.tot() - singleOutput.c, singleOutput.c*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + // copy first of backward + checkCuda( cudaMemcpy(trans + singleOutput.c, dstB, singleOutput.c*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + dstData = trans; + } - checkCUDNN(cudnnRNNForwardInference(net->cudnnHandle, - rnnDesc, - seqLen, // number of time steps (nT) - x_desc_vec_.data(), // input array of desc (nT*nC_in) - srcData, // input pointer - hx_desc_, // initial hidden state desc - hx_ptr, // initial hidden state pointer - cx_desc_, // initial cell state desc - cx_ptr, // initial cell state pointer - w_desc_, // weights desc - w_ptr, // weights pointer - y_desc_vec_.data(), // output desc (nT*nC_out) - dstData, // output pointer - hy_desc_, // final hidden state desc - hy_ptr, // final hidden state pointer - cy_desc_, // final cell state desc - cy_ptr, // final cell state pointer - work_space_, // workspace pointer - workspace_byte_)); // workspace size dim = output_dim; return dstData; diff --git a/src/utils.cpp b/src/utils.cpp index 444318d..6789e8c 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -68,7 +68,7 @@ void printDeviceVector(int size, dnnType* vec_d, bool device) delete [] vec; } -int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { +int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int limit) { dnnType *data_h, *correct_h; const float eps = 0.02f; @@ -92,7 +92,7 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { diffs += 1; if(diffs == 1) std::cout<<"\n"; - if(diffs < 10) + if(diffs < limit) std::cout<<" | [ "< Date: Sun, 16 Feb 2020 17:08:19 +0100 Subject: [PATCH 29/30] structure ok, result wrong --- include/tkDNN/Layer.h | 9 +++- src/LSTM.cpp | 103 ++++++++++++++++---------------------- tests/imuodom/imuodom.cpp | 4 +- 3 files changed, 55 insertions(+), 61 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index f12027f..7e827d8 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -207,7 +207,9 @@ protected: /** Bidirectional LSTM layer - + ONLY BIDIRECTIONAL (TODO: more configurable) + currently implemented as 2 inferences: forward and backward (TODO: only 1 cudnn inference) + implementation info: https://github.com/jiangnanhugo/seq2seq_cuda/blob/e4dbdcfa0517c972bfd4beea9f11a5233954093c/src/rnn.cpp https://github.com/Jeffery-Song/mxnet-test/blob/aab666faad44011f7a67b527b5f6c960367d0422/src/operator/cudnn_rnn-inl.h @@ -261,6 +263,11 @@ protected: dnnType *w_ptr; dnnType *w_h; dnnType *wf_ptr, *wb_ptr; // params pointer forward and backward layer + + // used during inference + dataDim_t one_output_dim; // output dim of as single inference + dnnType *srcF, *srcB; // input of single inference + dnnType *dstF, *dstB_NR, *dstB; // output of single inference, dstB_NR = dstB not reversed }; diff --git a/src/LSTM.cpp b/src/LSTM.cpp index 00d8f76..44502fa 100644 --- a/src/LSTM.cpp +++ b/src/LSTM.cpp @@ -13,8 +13,6 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig seqLen = input_dim.w; stateSize = hiddensize; - std::cout<<"LSTM seqLen: "< x_vec(seqLen); std::vector y_vec(seqLen); @@ -110,7 +108,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig checkCUDNN(cudnnGetRNNParamsSize(net->cudnnHandle, rnnDesc,x_desc_vec_[0], &cudnn_param_size, net->dataType)); int cudnn_params = cudnn_param_size/sizeof(dnnType); - std::cout<<"LSTM params size: "<cudnnHandle, rnnDesc, i, x_desc_vec_[0], w_desc_, 0, j, m_desc, (void**)&p)); @@ -209,37 +217,27 @@ LSTM::~LSTM() { checkCuda(cudaFree(work_space_ )); checkCuda(cudaFree(dropout_states_)); + checkCuda(cudaFree(srcF)); + checkCuda(cudaFree(srcB)); + checkCuda(cudaFree(dstF)); + checkCuda(cudaFree(dstB_NR)); + checkCuda(cudaFree(dstB)); checkCuda(cudaFree(dstData)); } dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) { - std::cout<<"LSTM infer\n"; + // transpose input + matrixTranspose(net->cublasHandle, srcData, srcF, dim.c, dim.h*dim.w*dim.l); - dnnType *trans; - checkCuda( cudaMalloc(&trans, dim.tot()*sizeof(dnnType))); - matrixTranspose(net->cublasHandle, srcData, trans, dim.c, dim.h*dim.w*dim.l); - srcData = trans; - - // reposition in invered order - dnnType *srcBack; - checkCuda( cudaMalloc(&srcBack, dim.tot()*sizeof(dnnType))); + // build srcB as reversed srcF for(int i=0; i