diff --git a/.gitignore b/.gitignore index 463ed5e..858af8d 100644 --- a/.gitignore +++ b/.gitignore @@ -15,3 +15,5 @@ build/ *.tar.gz *.weights *.zip +.idea/ +*.hdf5 diff --git a/CMakeLists.txt b/CMakeLists.txt index dede4eb..570bcf4 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,50 +1,51 @@ -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_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC") +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) -set(BUILD_DEPS true CACHE BOOL "If true download deps") - -if( ${BUILD_DEPS} ) - 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) - message("Finished dowloading test weights") -endif() - +# project specific flags if(DEBUG) add_definitions(-DDEBUG) endif() -find_package(CUDA REQUIRED) + +#------------------------------------------------------------------------------- +# CUDA +#------------------------------------------------------------------------------- +find_package(CUDA 9.0 REQUIRED) +SET(CUDA_SEPARABLE_COMPILATION ON) +#set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'") + +find_package(CUDNN REQUIRED) + +# compile +file(GLOB tkdnn_CUSRC "src/kernels/*.cu") +cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) +cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) + + +#------------------------------------------------------------------------------- +# External Libraries +#------------------------------------------------------------------------------- find_package(OpenCV REQUIRED) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") include_directories(/usr/include/gdal) -# 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() - -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}) - +#------------------------------------------------------------------------------- +# Build Libraries +#------------------------------------------------------------------------------- file(GLOB tkdnn_SRC "src/*.cpp") -set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS} -lgdal) +set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS}) file(GLOB class_SRC "src/class_src/*.cpp") -set(class_LIBS ${OpenCV_LIBS} -lgdal yaml-cpp python2.7) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11 -O3") -include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES} "~/repos/cereal/include" ${CMAKE_CURRENT_SOURCE_DIR}/tracker_CLASS/c++/src) -include_directories( BEFORE ${MY_SOURCE_DIR}/src /usr/include/python2.7 ) +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES} "~/repos/cereal/include" ${CMAKE_CURRENT_SOURCE_DIR}/tracker_CLASS/c++/src /usr/include/python2.7) + +set(class_LIBS ${OpenCV_LIBS} -lgdal yaml-cpp python2.7) add_library(tkDNN SHARED ${tkdnn_SRC}) target_link_libraries(tkDNN ${tkdnn_LIBS}) @@ -100,6 +101,11 @@ target_link_libraries(test_yolo3_tetrapack_resize tkDNN) add_executable(test_yolo3_BCDS6 tests/yolo3_BCDS6/yolo3_BCDS6.cpp) target_link_libraries(test_yolo3_BCDS6 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) ################################################################################ @@ -117,15 +123,33 @@ target_link_libraries(yolo3_demo tkDNN CLASS) -#install + +#------------------------------------------------------------------------------- +# Install +#------------------------------------------------------------------------------- #if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT) # set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install" # 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 +install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/cmake/" # source directory + DESTINATION "share/tkDNN/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() + diff --git a/README.md b/README.md index cd06992..c76a4f8 100644 --- a/README.md +++ b/README.md @@ -11,7 +11,7 @@ this branch actually work on every NVIDIA GPU that support the dependencies: ## Dependencies ``` -sudo apt install libgdal-dev libeigen3-dev python-matplotlib libyaml-cpp-dev libcereal-dev +sudo apt install libgdal-dev libeigen3-dev python-matplotlib libyaml-cpp-dev libcereal-dev python2.7-dev ``` ## Workflow @@ -27,6 +27,7 @@ Build with cmake mkdir build cd build cmake .. +# use -DTEST_DATA=False to skip dataset download make ``` during the cmake configuration it will be dowloaded the weights needed for running @@ -54,5 +55,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](https://user-images.githubusercontent.com/11562617/72547657-540e7800-388d-11ea-83c6-49dfea2a0607.gif) diff --git a/cmake/FindCUDNN.cmake b/cmake/FindCUDNN.cmake new file mode 100644 index 0000000..f240fcb --- /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/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/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 new file mode 100644 index 0000000..4dbaf21 --- /dev/null +++ b/cmake/tkDNNConfig.cmake @@ -0,0 +1,24 @@ +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_package(CUDNN REQUIRED) + +set(tkDNN_INCLUDE_DIRS + ${CUDA_INCLUDE_DIRS} + ${OPENCV_INCLUDE_DIRS} + ${CUDNN_INCLUDE_DIRS} +) + +set(tkDNN_LIBRARIES + tkDNN + kernels + ${CUDA_LIBRARIES} + ${CUDA_CUBLAS_LIBRARIES} + ${CUDNN_LIBRARIES} + ${OpenCV_LIBS} +) + +set(tkDNN_FOUND true) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 79482bc..8e00a50 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -30,6 +30,7 @@ std::string obj_class[10]{"person", "car", "truck", "bus", "motor", "bike", "rid //mutex for some opencv operations std::mutex mutex_cv; Show_t updates; +bool SAVE_RESULT = false; void sig_handler(int signo) { @@ -302,6 +303,14 @@ void *computationTask(void *x_void_ptr) // float prob; cv::Scalar intensity; + + // 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::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h)); + // } + cv::Mat frame; cv::Mat frame_crop; cv::Mat dnn_input; @@ -546,6 +555,9 @@ int main(int argc, char *argv[]) { yolo[i].init(par.net); yolo[i].thresh = 0.25; + + // if(SAVE_RESULT) + // resultVideo << frame; } // tk::dnn::Yolo3Detection yolo; // yolo.init(net); diff --git a/include/Layer.h b/include/tkDNN/Layer.h similarity index 65% rename from include/Layer.h rename to include/tkDNN/Layer.h index 396288e..3c1945b 100644 --- a/include/Layer.h +++ b/include/tkDNN/Layer.h @@ -1,7 +1,8 @@ #ifndef LAYER_H #define LAYER_H -#include +#include +#include #include "utils.h" #include "Network.h" @@ -10,10 +11,11 @@ namespace tk namespace dnn { -enum layerType_t -{ +enum layerType_t { + LAYER_INPUT, LAYER_DENSE, LAYER_CONV2D, + LAYER_LSTM, LAYER_ACTIVATION, LAYER_FLATTEN, LAYER_MULADD, @@ -52,36 +54,23 @@ public: std::string getLayerName() { layerType_t type = getLayerType(); - switch (type) - { - case LAYER_DENSE: - return "Dense"; - case LAYER_CONV2D: - return "Conv2d"; - case LAYER_ACTIVATION: - return "Activation"; - case LAYER_FLATTEN: - return "Flatten"; - case LAYER_MULADD: - return "MulAdd"; - case LAYER_POOLING: - return "Pooling"; - case LAYER_SOFTMAX: - return "Softmax"; - case LAYER_ROUTE: - return "Route"; - case LAYER_REORG: - return "Reorg"; - case LAYER_SHORTCUT: - return "Shortcut"; - case LAYER_UPSAMPLE: - return "Upsample"; - case LAYER_REGION: - return "Region"; - case LAYER_YOLO: - return "Yolo"; - default: - return "unknown"; + 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"; + case LAYER_POOLING: return "Pooling"; + case LAYER_SOFTMAX: return "Softmax"; + case LAYER_ROUTE: return "Route"; + case LAYER_REORG: return "Reorg"; + case LAYER_SHORTCUT: return "Shortcut"; + case LAYER_UPSAMPLE: return "Upsample"; + case LAYER_REGION: return "Region"; + case LAYER_YOLO: return "Yolo"; + default: return "unknown"; } } @@ -97,8 +86,8 @@ 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); + LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, + std::string fname_weights, bool batchnorm = false); virtual ~LayerWgs(); int inputs, outputs; @@ -124,6 +113,27 @@ public: __half *variance16_h, *variance16_d; }; +/** + 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 */ @@ -131,7 +141,7 @@ 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; }; @@ -168,14 +178,22 @@ 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 { 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); + Conv2d( Network *net, int out_ch, int kernelH, int kernelW, + int strideH, int strideW, int paddingH, int paddingW, + std::string fname_weights, bool batchnorm = false); virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; @@ -193,6 +211,72 @@ protected: size_t ws_sizeInBytes; }; +/** + 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 + https://stackoverflow.com/a/38737941 + https://colah.github.io/posts/2015-08-Understanding-LSTMs/ + + PARAMS (numlayers*2): + layer0: + ( INCH, ? ) ??? + ( HIDDEN, ? ) ??? + ( HIDDEN * 8 ) ??? + layer2: + ( INCH, ? ) ??? + ( HIDDEN, ? ) ??? + ( HIDDEN * 8 ) ??? + + 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 { + +public: + 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 = 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 */ + int numLayers = 1; /**> number of internal layers */ + +protected: + cudnnRNNDescriptor_t rnnDesc; + cudnnDropoutDescriptor_t dropoutDesc; + dnnType *dropout_states_, *work_space_; + + size_t workspace_byte_, dropout_byte_; + int workspace_size_, dropout_size_; + + std::vector x_desc_vec_, y_desc_vec_; + 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; + 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 +}; + + /** Flatten layer is actually a matrix transposition @@ -384,13 +468,14 @@ 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; }; 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); @@ -422,8 +507,8 @@ class RegionInterpret { public: - RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, - int classes, int coords, int num, float thresh, const char *fname_weights); + RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, + int classes, int coords, int num, float thresh, std::string fname_weights); ~RegionInterpret(); dataDim_t input_dim, output_dim; 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 84% rename from include/Yolo3Detection.h rename to include/tkDNN/Yolo3Detection.h index 57f275c..2405a45 100644 --- a/include/Yolo3Detection.h +++ b/include/tkDNN/Yolo3Detection.h @@ -46,6 +46,9 @@ public: // this is filled with results std::vector detected; + // keep track of inference times (ms) + std::vector stats; + Yolo3Detection() {} virtual ~Yolo3Detection() {} @@ -58,6 +61,15 @@ public: bool init(std::string tensor_path); void addBorders(cv::Mat &imageORIG, cv::Mat &imageWBorders, int &top, int &left); void update(cv::Mat &frame); + + + tk::dnn::Yolo* getYoloLayer(int n=0) { + if(n<3) + return yolo[n]; + else + return nullptr; + } + }; } // namespace dnn 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/tkDNN/models/Yolo3.h b/include/tkDNN/models/Yolo3.h new file mode 100644 index 0000000..cd69b32 --- /dev/null +++ b/include/tkDNN/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/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 87% rename from include/pluginsRT/YoloRT.h rename to include/tkDNN/pluginsRT/YoloRT.h index 2e52142..dab7c18 100644 --- a/include/pluginsRT/YoloRT.h +++ b/include/tkDNN/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/include/tkdnn.h b/include/tkDNN/tkdnn.h similarity index 84% rename from include/tkdnn.h rename to include/tkDNN/tkdnn.h index dde3126..554daa1 100644 --- a/include/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/include/utils.h b/include/tkDNN/utils.h similarity index 95% rename from include/utils.h rename to include/tkDNN/utils.h index c6937e3..dd97c0c 100644 --- a/include/utils.h +++ b/include/tkDNN/utils.h @@ -104,9 +104,9 @@ 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); -int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true); -void printDeviceVector(int size, dnnType *vec_d, bool device = true); +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 limit = 10); +void printDeviceVector(int size, dnnType* vec_d, bool device = true); void resize(int size, dnnType **data); void matrixTranspose(cublasHandle_t handle, dnnType *srcData, dnnType *dstData, int rows, int cols); diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 9e3aeeb..4281d12 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -7,13 +7,12 @@ 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) : - - LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, - fname_weights, batchnorm) -{ +Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, + int strideH, int strideW, int paddingH, int paddingW, + std::string fname_weights, bool batchnorm) : + + LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, + fname_weights, batchnorm) { this->kernelH = kernelH; this->kernelW = kernelW; diff --git a/src/Dense.cpp b/src/Dense.cpp index f1a4da5..6306a3a 100644 --- a/src/Dense.cpp +++ b/src/Dense.cpp @@ -7,8 +7,8 @@ namespace tk namespace dnn { -Dense::Dense(Network *net, int out_ch, const char *fname_weights) : LayerWgs(net, net->getOutputDim().tot(), out_ch, 1, 1, 1, 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; output_dim.c = out_ch; diff --git a/src/LSTM.cpp b/src/LSTM.cpp new file mode 100644 index 0000000..6ecf4fd --- /dev/null +++ b/src/LSTM.cpp @@ -0,0 +1,327 @@ +#include + +#include "Layer.h" + +namespace tk { namespace dnn { + +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; + stateSize = hiddensize; + + // init Tensor Descriptors + std::vector x_vec(seqLen); + std::vector y_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])); + + dimA[0] = batchSize; + dimA[1] = inputSize; + dimA[2] = 1; + dimA[0] = batchSize; + dimA[1] = inputSize; + strideA[0] = dimA[2] * dimA[1]; + strideA[1] = dimA[2]; + strideA[2] = 1; + checkCUDNN(cudnnSetTensorNdDescriptor(x_vec[i], + net->dataType, 3, dimA, strideA)); + + dimA[0] = batchSize; + dimA[1] = 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)); + } + // apply tensordesc + x_desc_vec_ = x_vec; + y_desc_vec_ = y_vec; + + + // set the state tensors + dimA[0] = numLayers; + 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(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)); + // allocate dnnType *hx_ptr, *cx_ptr, *hy_ptr, *cy_ptr; + 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 ??? + 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, + //(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)); + + + // 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)); + + // load params + std::cout<<"Reading weights: PARAMS="<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-> "<cublasHandle, srcData, srcF, dim.c, dim.h*dim.w*dim.l); + + // build srcB as reversed srcF + for(int i=0; icudnnHandle, + rnnDesc, + seqLen, // number of time steps (nT) + x_desc_vec_.data(), // input array of desc (nT*nC_in) + srcF, // 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 + } + + // 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) + srcB, // 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_NR, // 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 + } + + + // reverse order of dstB + for(int i=0; icublasHandle, dstF, dstData, + one_output_dim.h* one_output_dim.w*one_output_dim.l, one_output_dim.c); + // backward transpose + matrixTranspose(net->cublasHandle, dstB, dstData + one_output_dim.tot(), + one_output_dim.h* one_output_dim.w*one_output_dim.l, one_output_dim.c); + } else { + // copy last of forward + checkCuda( cudaMemcpy(dstData, dstF + one_output_dim.tot() - one_output_dim.c, + one_output_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + // copy first of backward + checkCuda( cudaMemcpy(dstData + one_output_dim.c, dstB, + one_output_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + } + + dim = output_dim; + return dstData; +} + +}} diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index 869ef69..c9d2524 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -9,10 +9,9 @@ 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) -{ +LayerWgs::LayerWgs(Network *net, int inputs, int outputs, + int kh, int kw, int kl, + std::string fname_weights, bool batchnorm) : Layer(net) { this->inputs = inputs; this->outputs = outputs; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 2a73efa..f27f70d 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"; + +#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); @@ -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; } @@ -461,6 +475,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/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..babef38 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -11,20 +11,26 @@ 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; readBinaryFile(fname_weights, 3*num*2, &bias_h, &bias_d, seek); } + // init default classes name + classesNames.clear(); + for(int i=0; inum; // 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); 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); @@ -137,12 +138,12 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { 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]; @@ -155,6 +156,8 @@ void Yolo3Detection::update(cv::Mat &imageORIG) { netRT->infer(dim, input_d); TIMER_STOP dim.print(); + + stats.push_back(t_ns); } TIMER_START diff --git a/src/utils.cpp b/src/utils.cpp index 7955d15..6789e8c 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; @@ -39,7 +39,8 @@ void readBinaryFile(const char* 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 *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"; +const char *l0_bin = "../tests/imuodom/layers/bidirectional_3.bin"; +const char *l1_bin = "../tests/imuodom/layers/bidirectional_4.bin"; +const char *d0_bin = "../tests/imuodom/layers/dense_3.bin"; +const char *d1_bin = "../tests/imuodom/layers/dense_4.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, true, l0_bin); + tk::dnn::LSTM lstm1(&net, 128, false, l1_bin); + + tk::dnn::Dense d0 (&net, 3, d0_bin); + + tk::dnn::Layer *lstm1_l[1] = { &lstm1 }; + tk::dnn::Route lstm1_link (&net, lstm1_l, 1); + tk::dnn::Dense d1 (&net, 4, d1_bin); + net.print(); + + dnnType *data; + tk::dnn::dataDim_t dim; + + TIMER_START + // Inference + data = net.infer(dim, data); + TIMER_STOP + + // Print real test + std::cout<<"\n==== CHECK RESULT ====\n"; + dnnType *out0, *out1; + dnnType *out0_h, *out1_h; + readBinaryFile(o0_bin, d0.output_dim.tot(), &out0_h, &out0); + readBinaryFile(o1_bin, d1.output_dim.tot(), &out1_h, &out1); + d0.output_dim.print(); + checkResult(d0.output_dim.tot(), d0.dstData, out0); + d1.output_dim.print(); + checkResult(d1.output_dim.tot(), d1.dstData, out1); + return 0; +} diff --git a/tests/imuodom/infer.py b/tests/imuodom/infer.py new file mode 100644 index 0000000..6c197f4 --- /dev/null +++ b/tests/imuodom/infer.py @@ -0,0 +1,78 @@ +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() + + np.random.seed(2) + 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 = 'dense_4' + #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("Aggregate:") + #print(intermediate_output.tolist()) + + 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..554cee0 100644 --- a/tests/simple/test_model.py +++ b/tests/simple/test_model.py @@ -1,43 +1,54 @@ 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 import Bidirectional, CuDNNLSTM 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((3, 8), name='x1') + conv = Conv1D(4, 2)(x1) + lstm = Bidirectional(CuDNNLSTM(5, return_sequences=True))(conv) + lstm2 = Bidirectional(CuDNNLSTM(5, return_sequences=False))(lstm) + model = Model([x1], [lstm2]) + 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 + np.random.seed(2) + x = np.random.rand(1,1,3,8) + 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)) + print("output: ", r.tolist()) + + 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..f1f2b3a 100644 --- a/tests/simple/test_simple.cpp +++ b/tests/simple/test_simple.cpp @@ -2,23 +2,21 @@ #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 *l1_bin = "../tests/simple/layers/bidirectional_1.bin"; +const char *l2_bin = "../tests/simple/layers/bidirectional_2.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, 8, 1, 3); 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); + tk::dnn::LSTM l1(&net, 5, true, l1_bin); + tk::dnn::LSTM l2(&net, 5, false, l2_bin); + + net.print(); // Load input dnnType *data; diff --git a/tests/weights_exporter.py b/tests/weights_exporter.py index df8bab3..266c374 100644 --- a/tests/weights_exporter.py +++ b/tests/weights_exporter.py @@ -5,96 +5,89 @@ 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) + elif(weights.ndim == 2): + weights = weights.transpose(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 +def export_bidir(name, params, paramsb): + print ("######## EXPORT", name, "LAYER ########") + + f = open(name + ".bin", mode='wb') - 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)) + print("FORWARD") + ker = params[0] + rec_ker = params[1] + bias = params[2] + print ("export kernels: ", np.shape(ker)) + units = np.shape(ker)[1] // 4 + bin_write(f, ker[:,:units]) + bin_write(f, ker[:,units:units*2]) + bin_write(f, ker[:,units*2:units*3]) + bin_write(f, ker[:,units*3:]) + print ("export recurrent kernels: ", np.shape(rec_ker)) + bin_write(f, rec_ker[:,:units]) + bin_write(f, rec_ker[:,units:units*2]) + bin_write(f, rec_ker[:,units*2:units*3]) + bin_write(f, rec_ker[:,units*3:]) + print ("export kernels: ", np.shape(ker)) + bin_write(f, bias) + print("WEIGHTS saved\n") + print("BACKWARD") + ker = paramsb[0] + rec_ker = paramsb[1] + bias = paramsb[2] + print ("export kernels: ", np.shape(ker)) + units = np.shape(ker)[1] // 4 + bin_write(f, ker[:,:units]) + bin_write(f, ker[:,units:units*2]) + bin_write(f, ker[:,units*2:units*3]) + bin_write(f, ker[:,units*3:]) + print ("export recurrent kernels: ", np.shape(rec_ker)) + bin_write(f, rec_ker[:,:units]) + bin_write(f, rec_ker[:,units:units*2]) + bin_write(f, rec_ker[:,units*2:units*3]) + bin_write(f, rec_ker[:,units*3:]) + print ("export kernels: ", np.shape(ker)) + bin_write(f, bias) + print("WEIGHTS saved\n") #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 +96,43 @@ 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"): + wgs = l.forward_layer.get_weights() + export_bidir(args.output + "/" + name, l.forward_layer.get_weights(), l.backward_layer.get_weights()) + else: + print ("skip:", name, "has no weights") + continue + diff --git a/tests/yolo3_berkeley/yolo3_berkeley.cpp b/tests/yolo3_berkeley/yolo3_berkeley.cpp index 3a69e32..70e8b60 100644 --- a/tests/yolo3_berkeley/yolo3_berkeley.cpp +++ b/tests/yolo3_berkeley/yolo3_berkeley.cpp @@ -1,295 +1,23 @@ #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); + // create yolo3 model + std::string bin_path = "../tests/yolo3_berkeley"; + int classes = 10; + tk::dnn::Yolo *yolo [3]; + #include "models/Yolo3.h" - 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); + // 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; @@ -304,9 +32,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 +43,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<