From 56feb54377c0678e42077fabbf84ce2fc138f4c5 Mon Sep 17 00:00:00 2001 From: perseusdg Date: Thu, 21 Jan 2021 00:22:15 +0400 Subject: [PATCH 01/22] able to build kernels as shared object file(dll),and minor changes to lstm.cpp and utils.cpp to overcome minor msvc build errors --- CMakeLists.txt | 11 +++++++++-- include/tkDNN/DetectionNN.h | 5 +++++ include/tkDNN/ImuOdom.h | 6 ++++++ include/tkDNN/Int8BatchStream.h | 7 ++++++- include/tkDNN/utils.h | 8 ++++++++ src/LSTM.cpp | 13 +++++++++---- src/utils.cpp | 7 ++++++- 7 files changed, 49 insertions(+), 8 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 8c8619d..03e26c5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,7 +2,13 @@ cmake_minimum_required(VERSION 3.5) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) +if(LINUX) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") +endif() +if(WIN32) +set(CMAKE_CXX_STANDARD 14) +set(CMAKE_CXX_FLAGS "/O2 /FS ") +endif() include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) # project specific flags @@ -18,7 +24,7 @@ add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}") find_package(CUDA 9.0 REQUIRED) SET(CUDA_SEPARABLE_COMPILATION ON) #set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'") -set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32) +set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32 -arch=sm_61 ) find_package(CUDNN REQUIRED) include_directories(${CUDNN_INCLUDE_DIR}) @@ -28,6 +34,7 @@ include_directories(${CUDNN_INCLUDE_DIR}) file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) +target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES}) #------------------------------------------------------------------------------- @@ -48,7 +55,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") file(GLOB tkdnn_SRC "src/*.cpp") set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS} yaml-cpp) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}") include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES}) add_library(tkDNN SHARED ${tkdnn_SRC}) target_link_libraries(tkDNN ${tkdnn_LIBS}) diff --git a/include/tkDNN/DetectionNN.h b/include/tkDNN/DetectionNN.h index 0498d41..9ce33c1 100644 --- a/include/tkDNN/DetectionNN.h +++ b/include/tkDNN/DetectionNN.h @@ -4,7 +4,12 @@ #include #include #include +#ifdef __linux__ #include +#elif _WIN32 +#include +#endif + #include #include "utils.h" diff --git a/include/tkDNN/ImuOdom.h b/include/tkDNN/ImuOdom.h index 6d8d4cb..aace012 100644 --- a/include/tkDNN/ImuOdom.h +++ b/include/tkDNN/ImuOdom.h @@ -1,7 +1,13 @@ #include #include #include /* srand, rand */ + +#ifdef __linux__ #include +#elif _WIN32 +#include +#endif + #include #include #include "utils.h" diff --git a/include/tkDNN/Int8BatchStream.h b/include/tkDNN/Int8BatchStream.h index 4349c1f..7d2cef5 100644 --- a/include/tkDNN/Int8BatchStream.h +++ b/include/tkDNN/Int8BatchStream.h @@ -11,8 +11,13 @@ #include #include #include -#include +#include +#ifdef __linux__ #include +#elif _WIN32 +#include +#endif + #include #include "NvInfer.h" diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index 538a3f3..cb3c18f 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -12,7 +12,12 @@ #include #include +#ifdef __linux__ #include +#elif _WIN32 +#include +#endif + #include @@ -39,6 +44,7 @@ #define TKDNN_VERBOSE 0 // Simple Timer +#ifdef __linux__ #define TKDNN_TSTART timespec start, end; \ clock_gettime(CLOCK_MONOTONIC, &start); @@ -48,6 +54,8 @@ if(show) std::cout< 7 - checkCUDNN(cudnnSetRNNDescriptor_v6(net->cudnnHandle, + checkCUDNN(cudnnSetRNNDescriptor_v6(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)); #else - checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, -#endif - rnnDesc, stateSize, numLayers, dropoutDesc, + 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)); +#endif // Get temp space sizes diff --git a/src/utils.cpp b/src/utils.cpp index 65030f0..e143cce 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -170,6 +170,7 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){ using std::ios_base; using std::ifstream; using std::string; + SYSTEM_INFO sysInfo; vm_usage_kb = 0.0; resident_set_kb = 0.0; @@ -191,8 +192,12 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){ >> O >> itrealvalue >> starttime >> vsize >> rss; stat_stream.close(); - +#ifdef __linux__ long page_size_kb = sysconf(_SC_PAGE_SIZE) / 1024; // in case x86-64 is configured to use 2MB pages +#elif _WIN32 + long page_size_kb = sysInfo.dwPageSize/1024; +#endif + vm_usage_kb = vsize / 1024.0; resident_set_kb = rss * page_size_kb; } -- 2.52.0 From 512acd8cba99c7ec21d542b32a2881671da39cba Mon Sep 17 00:00:00 2001 From: perseusdg Date: Thu, 21 Jan 2021 09:29:27 +0400 Subject: [PATCH 02/22] minor fixes --- Issues.md | 1 + demo/demo/map.cpp | 5 +++++ include/tkDNN/utils.h | 14 +++++++++++--- src/Yolo.cpp | 9 +++++---- 4 files changed, 22 insertions(+), 7 deletions(-) create mode 100644 Issues.md diff --git a/Issues.md b/Issues.md new file mode 100644 index 0000000..4875b13 --- /dev/null +++ b/Issues.md @@ -0,0 +1 @@ +1)error C2131 @ Yolo3Detection.cpp(97) -> expression doesnt evaluate to a constant caused to read of variable outside its lifetime \ No newline at end of file diff --git a/demo/demo/map.cpp b/demo/demo/map.cpp index 356e35a..8e363ee 100644 --- a/demo/demo/map.cpp +++ b/demo/demo/map.cpp @@ -2,7 +2,12 @@ #include #include #include /* srand, rand */ +#ifdef __linux__ #include +#elif _WIN32 +#include +#endif + #include #include "utils.h" diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index cb3c18f..1404509 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -15,10 +15,12 @@ #ifdef __linux__ #include #elif _WIN32 -#include -#endif +#define NOMINMAX +#include +#endif #include +#include #define dnnType float @@ -55,7 +57,13 @@ #define TKDNN_TSTOP TKDNN_TSTOP_C(COL_CYANB, TKDNN_VERBOSE) #elif _WIN32 -#endif +#define TKDNN_TSTART auto start = std::chrono::high_resolution_clock::now(); +#define TKDNN_TSTOP auto stop = std::chrono::high_resolution_clock::now(); \ +std::chrono::duration duration = stop -start; \ +auto time_ms = std::chrono::duration_cast(duration);\ +double t_ns = time_ms.count(); +#endif + /******************************************************** * Prints the error message, and exits diff --git a/src/Yolo.cpp b/src/Yolo.cpp index 9737e74..6d0b546 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -9,6 +9,7 @@ #include "Layer.h" #include "kernels.h" + namespace tk { namespace dnn { Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks, float scale_xy, double nms_thresh, nmsKind_t nsm_kind, int new_coords) : @@ -209,10 +210,10 @@ float yolo_box_iou(Yolo::box a, Yolo::box b) } void box_c(const Yolo::box a, const Yolo::box b, float& top, float& bot, float& left, float& right) { - top = std::min(a.y - a.h / 2, b.y - b.h / 2); - bot = std::max(a.y + a.h / 2, b.y + b.h / 2); - left = std::min(a.x - a.w / 2, b.x - b.w / 2); - right = std::max(a.x + a.w / 2, b.x + b.w / 2); + top = (std::min)(a.y - a.h / 2, b.y - b.h / 2); + bot = (std::max)(a.y + a.h / 2, b.y + b.h / 2); + left = (std::min)(a.x - a.w / 2, b.x - b.w / 2); + right = (std::max)(a.x + a.w / 2, b.x + b.w / 2); } // https://github.com/Zzh-tju/DIoU-darknet -- 2.52.0 From adac8576b0faf515ad3f459b1f50fd16cef6d64d Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Fri, 22 Jan 2021 17:57:42 +0100 Subject: [PATCH 03/22] Add support to Scaled-YOLO4, update Yolov4x-mish (tested) Signed-off-by: Micaela Verucchi --- README.md | 4 +- include/tkDNN/Layer.h | 7 +- include/tkDNN/NetworkRT.h | 1 + .../tkDNN/pluginsRT/ActivationLogisticRT.h | 60 + include/tkDNN/pluginsRT/YoloRT.h | 14 +- scripts/test_all_tests.sh | 5 +- src/Activation.cpp | 4 + src/DarknetParser.cpp | 1 + src/NetworkRT.cpp | 13 +- src/Yolo.cpp | 21 +- tests/darknet/cfg/yolo4-csp.cfg | 1279 +++++++++++++++++ tests/darknet/cfg/yolo4x.cfg | 21 +- tests/darknet/yolo4-csp.cpp | 36 + tests/darknet/yolo4x.cpp | 2 +- 14 files changed, 1441 insertions(+), 27 deletions(-) create mode 100644 include/tkDNN/pluginsRT/ActivationLogisticRT.h create mode 100644 tests/darknet/cfg/yolo4-csp.cfg create mode 100644 tests/darknet/yolo4-csp.cpp diff --git a/README.md b/README.md index 84e0037..d9927c0 100644 --- a/README.md +++ b/README.md @@ -353,7 +353,8 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing | yolo4 | Yolov4 8 | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download) | | yolo4_berkeley | Yolov4 8 | [BDD100K ](https://bair.berkeley.edu/blog/2018/05/30/bdd/) | 10 | 540x320 | [weights](https://cloud.hipert.unimore.it/s/nkWFa5fgb4NTdnB/download) | | yolo4tiny | Yolov4 tiny 9 | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download) | -| yolo4x | Yolov4x-mish 9 | [COCO 2017](http://cocodataset.org/) | 80 | 672x672 | [weights](https://cloud.hipert.unimore.it/s/BLPpiAigZJLorQD/download) | +| yolo4x | Yolov4x-mish 9 | [COCO 2017](http://cocodataset.org/) | 80 | 640x640 | [weights](https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download) | +| yolo4x-cps | Scaled Yolov4 10 | [COCO 2017](http://cocodataset.org/) | 80 | 512x512 | [weights](https://cloud.hipert.unimore.it/s/AfzHE4BfTeEm2gH/download) | ## References @@ -367,3 +368,4 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing 7. Wang, Chien-Yao, et al. "CSPNet: A New Backbone that can Enhance Learning Capability of CNN." arXiv preprint arXiv:1911.11929 (2019). 8. Bochkovskiy, Alexey, Chien-Yao Wang, and Hong-Yuan Mark Liao. "YOLOv4: Optimal Speed and Accuracy of Object Detection." arXiv preprint arXiv:2004.10934 (2020). 9. Bochkovskiy, Alexey, "Yolo v4, v3 and v2 for Windows and Linux" (https://github.com/AlexeyAB/darknet) +10. Wang, Chien-Yao, Alexey Bochkovskiy, and Hong-Yuan Mark Liao. "Scaled-YOLOv4: Scaling Cross Stage Partial Network." arXiv preprint arXiv:2011.08036 (2020). diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 25c4565..e097372 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -19,6 +19,7 @@ enum layerType_t { LAYER_ACTIVATION_CRELU, LAYER_ACTIVATION_LEAKY, LAYER_ACTIVATION_MISH, + LAYER_ACTIVATION_LOGISTIC, LAYER_FLATTEN, LAYER_RESHAPE, LAYER_MULADD, @@ -68,6 +69,7 @@ public: case LAYER_ACTIVATION_CRELU: return "ActivationCReLU"; case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky"; case LAYER_ACTIVATION_MISH: return "ActivationMish"; + case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic"; case LAYER_FLATTEN: return "Flatten"; case LAYER_RESHAPE: return "Reshape"; case LAYER_MULADD: return "MulAdd"; @@ -212,7 +214,8 @@ public: typedef enum { ACTIVATION_ELU = 100, ACTIVATION_LEAKY = 101, - ACTIVATION_MISH = 102 + ACTIVATION_MISH = 102, + ACTIVATION_LOGISTIC = 103 } tkdnnActivationMode_t; /** @@ -233,6 +236,8 @@ public: return LAYER_ACTIVATION_LEAKY; else if (act_mode == ACTIVATION_MISH) return LAYER_ACTIVATION_MISH; + else if (act_mode == ACTIVATION_LOGISTIC) + return LAYER_ACTIVATION_LOGISTIC; else return LAYER_ACTIVATION; }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 4c6c816..4fe2e0e 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -24,6 +24,7 @@ template T readBUF(const char*& buffer) using namespace nvinfer1; #include "pluginsRT/ActivationLeakyRT.h" +#include "pluginsRT/ActivationLogisticRT.h" #include "pluginsRT/ActivationReLUCeilingRT.h" #include "pluginsRT/ActivationMishRT.h" #include "pluginsRT/ReorgRT.h" diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h new file mode 100644 index 0000000..a1ceb6b --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h @@ -0,0 +1,60 @@ +#include +#include "../kernels.h" + +class ActivationLogisticRT : public IPlugin { + +public: + ActivationLogisticRT() { + + + } + + ~ActivationLogisticRT(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, stream); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 1*sizeof(int); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, size); + } + + int size; +}; diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 9af8587..f2bdaa1 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -67,15 +67,17 @@ public: for (int b = 0; b < batchSize; ++b){ for(int n = 0; n < n_masks; ++n){ int index = entry_index(b, n*w*h, 0); - if (new_coords == 1) - activationLOGISTICForward(srcData + index, dstData + index, 4*w*h, stream); //x,y,w,h - else + if (new_coords == 1){ + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + } + else{ activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y - if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); - index = entry_index(b, n*w*h, 4); - activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream); + index = entry_index(b, n*w*h, 4); + activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream); + } } } diff --git a/scripts/test_all_tests.sh b/scripts/test_all_tests.sh index af04aff..6aab775 100644 --- a/scripts/test_all_tests.sh +++ b/scripts/test_all_tests.sh @@ -69,10 +69,11 @@ do echo -e "${ORANGE}Batch $TKDNN_BATCHSIZE ${NC}" test_net mnist - ./test_imuodom &>> $out_file - print_output $? imuodom + # ./test_imuodom &>> $out_file + # print_output $? imuodom test_net yolo4 + test_net yolo4-csp test_net yolo4x test_net yolo4_berkeley test_net yolo4tiny diff --git a/src/Activation.cpp b/src/Activation.cpp index 28c7624..a8642e7 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -52,6 +52,10 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { else if(act_mode == ACTIVATION_MISH) { activationMishForward(srcData, dstData, dim.tot()); + } + else if(act_mode == ACTIVATION_LOGISTIC) { + activationLOGISTICForward(srcData, dstData, dim.tot()); + } else { dnnType alpha = dnnType(1); dnnType beta = dnnType(0); diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 7b5410c..69b6b29 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -187,6 +187,7 @@ namespace tk { namespace dnn { if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; + else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC; else { FatalError("activation not supported: " + f.activation); } netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); }; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 501ade4..c915ba5 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -226,7 +226,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Conv2d*) l); if(type == LAYER_POOLING) return convert_layer(input, (Pooling*) l); - if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -421,6 +421,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { checkNULL(lRT); return lRT; } + else if(l->act_mode == ACTIVATION_LOGISTIC) { + IPlugin *plugin = new ActivationLogisticRT(); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + return lRT; + } else { FatalError("this Activation mode is not yet implemented"); return NULL; @@ -653,6 +659,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa a->size = readBUF(buf); return a; } + if(name.find("ActivationLogistic") == 0) { + ActivationLogisticRT *a = new ActivationLogisticRT(); + a->size = readBUF(buf); + return a; + } if(name.find("ActivationCReLU") == 0) { ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); a->size = readBUF(buf); diff --git a/src/Yolo.cpp b/src/Yolo.cpp index 9737e74..61ed3cf 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -72,8 +72,8 @@ Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h; } else{ - b.x = (i + x[index + 0 * stride] * 2 - 0.5) / lw; - b.y = (j + x[index + 1 * stride] * 2 - 0.5) / lh; + b.x = (i + x[index + 0 * stride] ) / lw; + b.y = (j + x[index + 1 * stride] ) / lh; b.w = x[index + 2 * stride] * x[index + 2 * stride] * 4 * biases[2 * n] / w; b.h = x[index + 3 * stride] * x[index + 3 * stride] * 4 * biases[2 * n + 1] / h; } @@ -87,15 +87,18 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) { for (int b = 0; b < dim.n; ++b){ for(int n = 0; n < n_masks; ++n){ int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim); - if (new_coords == 1) - activationLOGISTICForward(srcData + index, dstData + index, 4*dim.w*dim.h); - else + std::cout<<"new_coords"<scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + } + else{ activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h); - if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); - - index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim); - activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h); + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + + index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim); + activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h); + } } } diff --git a/tests/darknet/cfg/yolo4-csp.cfg b/tests/darknet/cfg/yolo4-csp.cfg new file mode 100644 index 0000000..691ec03 --- /dev/null +++ b/tests/darknet/cfg/yolo4-csp.cfg @@ -0,0 +1,1279 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +subdivisions=8 +width=512 +height=512 +channels=3 +momentum=0.949 +decay=0.0005 +angle=0 +saturation = 1.5 +exposure = 1.5 +hue=.1 + +learning_rate=0.001 +burn_in=1000 +max_batches = 500500 +policy=steps +steps=400000,450000 +scales=.1,.1 + +mosaic=1 + +letter_box=1 + +ema_alpha=0.9998 + +#optimized_memory=1 + +#23:104x104 54:52x52 85:26x26 104:13x13 for 416 + + + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=2 +pad=1 +activation=mish + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +#[route] +#layers = -2 + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +[convolutional] +batch_normalize=1 +filters=32 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +#[route] +#layers = -1,-7 + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-10 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-28 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-28 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-16 + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=mish + +########################## + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +### SPP ### +[maxpool] +stride=1 +size=5 + +[route] +layers=-2 + +[maxpool] +stride=1 +size=9 + +[route] +layers=-4 + +[maxpool] +stride=1 +size=13 + +[route] +layers=-1,-3,-5,-6 +### End SPP ### + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[route] +layers = -1, -13 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[upsample] +stride=2 + +[route] +layers = 79 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[route] +layers = -1, -6 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[upsample] +stride=2 + +[route] +layers = 48 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=mish + +[route] +layers = -1, -6 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +########################## + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + + +[yolo] +mask = 0,1,2 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=0 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=4.0 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=5 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=256 +activation=mish + +[route] +layers = -1, -20 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[route] +layers = -1,-6 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + + +[yolo] +mask = 3,4,5 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=1.0 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=5 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=512 +activation=mish + +[route] +layers = -1, -49 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[route] +layers = -1,-6 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=mish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + + +[yolo] +mask = 6,7,8 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 diff --git a/tests/darknet/cfg/yolo4x.cfg b/tests/darknet/cfg/yolo4x.cfg index 89f2564..2ff854f 100644 --- a/tests/darknet/cfg/yolo4x.cfg +++ b/tests/darknet/cfg/yolo4x.cfg @@ -5,8 +5,8 @@ # Training batch=64 subdivisions=8 -width=672 -height=672 +width=640 +height=640 channels=3 momentum=0.949 decay=0.0005 @@ -15,7 +15,7 @@ saturation = 1.5 exposure = 1.5 hue=.1 -learning_rate=0.00261 +learning_rate=0.001 burn_in=1000 max_batches = 500500 policy=steps @@ -26,6 +26,8 @@ mosaic=1 letter_box=1 +#optimized_memory=1 + [convolutional] batch_normalize=1 filters=32 @@ -1131,6 +1133,7 @@ size=1 stride=1 pad=1 activation=mish +stopbackward=800 ########################## @@ -1147,7 +1150,7 @@ size=1 stride=1 pad=1 filters=255 -activation=linear +activation=logistic [yolo] @@ -1156,6 +1159,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4 classes=80 num=9 jitter=.1 +scale_x_y = 2.0 objectness_smooth=0 ignore_thresh = .7 truth_thresh = 1 @@ -1169,6 +1173,7 @@ iou_loss=ciou nms_kind=diounms beta_nms=0.6 new_coords=1 +max_delta=5 [route] layers = -4 @@ -1275,7 +1280,7 @@ size=1 stride=1 pad=1 filters=255 -activation=linear +activation=logistic [yolo] @@ -1284,6 +1289,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4 classes=80 num=9 jitter=.1 +scale_x_y = 2.0 objectness_smooth=1 ignore_thresh = .7 truth_thresh = 1 @@ -1297,6 +1303,7 @@ iou_loss=ciou nms_kind=diounms beta_nms=0.6 new_coords=1 +max_delta=5 [route] layers = -4 @@ -1403,7 +1410,7 @@ size=1 stride=1 pad=1 filters=255 -activation=linear +activation=logistic [yolo] @@ -1412,6 +1419,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4 classes=80 num=9 jitter=.1 +scale_x_y = 2.0 objectness_smooth=1 ignore_thresh = .7 truth_thresh = 1 @@ -1425,3 +1433,4 @@ iou_loss=ciou nms_kind=diounms beta_nms=0.6 new_coords=1 +max_delta=2 diff --git a/tests/darknet/yolo4-csp.cpp b/tests/darknet/yolo4-csp.cpp new file mode 100644 index 0000000..af8a7fc --- /dev/null +++ b/tests/darknet/yolo4-csp.cpp @@ -0,0 +1,36 @@ +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "yolo4-csp"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector output_bins = { + bin_path + "/debug/layer144_out.bin", + bin_path + "/debug/layer159_out.bin", + bin_path + "/debug/layer174_out.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4-csp.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/AfzHE4BfTeEm2gH/download"); + + + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + //convert network to tensorRT + tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + + int ret = testInference(input_bins, output_bins, net, netRT); + net->releaseLayers(); + delete net; + delete netRT; + return ret; +} diff --git a/tests/darknet/yolo4x.cpp b/tests/darknet/yolo4x.cpp index b9ad003..8df1aef 100644 --- a/tests/darknet/yolo4x.cpp +++ b/tests/darknet/yolo4x.cpp @@ -17,7 +17,7 @@ int main() { std::string wgs_path = bin_path + "/layers"; std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4x.cfg"; std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; - downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/BLPpiAigZJLorQD/download"); + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download"); -- 2.52.0 From fb52444cdc06b6a3868dd13cdd2cc0d0601cc6c0 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Mon, 25 Jan 2021 04:12:05 +0530 Subject: [PATCH 04/22] Replaced dynamic arrays with std::vector ,works on linux ..needs to be tested on windows after clearing up the lnk2019 error --- CMakeLists.txt | 4 ++-- src/Yolo3Detection.cpp | 7 ++++--- src/utils.cpp | 2 +- 3 files changed, 7 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 03e26c5..c11b56f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,13 +2,13 @@ cmake_minimum_required(VERSION 3.5) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) -if(LINUX) +if(UNIX) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") endif() if(WIN32) set(CMAKE_CXX_STANDARD 14) set(CMAKE_CXX_FLAGS "/O2 /FS ") -endif() +endif(WIN32) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) # project specific flags diff --git a/src/Yolo3Detection.cpp b/src/Yolo3Detection.cpp index b94eea9..0c638e6 100644 --- a/src/Yolo3Detection.cpp +++ b/src/Yolo3Detection.cpp @@ -94,9 +94,10 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){ void Yolo3Detection::postprocess(const int bi, const bool mAP){ //get yolo outputs - dnnType *rt_out[netRT->pluginFactory->n_yolos]; - for(int i=0; ipluginFactory->n_yolos; i++) - rt_out[i] = (dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi; + std::vector rt_out; + //dnnType *rt_out[netRT->pluginFactory->n_yolos]; + for(int i=0; ipluginFactory->n_yolos; i++) + rt_out.push_back((dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi); float x_ratio = float(originalSize[bi].width) / float(netRT->input_dim.w); float y_ratio = float(originalSize[bi].height) / float(netRT->input_dim.h); diff --git a/src/utils.cpp b/src/utils.cpp index e143cce..bbdf516 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -170,7 +170,6 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){ using std::ios_base; using std::ifstream; using std::string; - SYSTEM_INFO sysInfo; vm_usage_kb = 0.0; resident_set_kb = 0.0; @@ -195,6 +194,7 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){ #ifdef __linux__ long page_size_kb = sysconf(_SC_PAGE_SIZE) / 1024; // in case x86-64 is configured to use 2MB pages #elif _WIN32 +SYSTEM_INFO sysInfo; long page_size_kb = sysInfo.dwPageSize/1024; #endif -- 2.52.0 From 2d4dececb683ffa28f7f8aaf72a2e645e8c45adb Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Tue, 26 Jan 2021 20:17:43 +0400 Subject: [PATCH 05/22] Builds on windows successfully,issues with deserialization and downloading weights --- .gitignore | 3 ++- CMakeLists.txt | 9 +++++---- demo/demo/demo.cpp | 2 +- include/tkDNN/test.h | 3 ++- 4 files changed, 10 insertions(+), 7 deletions(-) diff --git a/.gitignore b/.gitignore index b56526f..5be5a73 100644 --- a/.gitignore +++ b/.gitignore @@ -12,5 +12,6 @@ build/ *.hdf5 *.pk *.table +cmake-build-release/ demo/COCO_val2017 -demo/BDD100K_val \ No newline at end of file +demo/BDD100K_val diff --git a/CMakeLists.txt b/CMakeLists.txt index c11b56f..4904e5d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,13 +1,14 @@ cmake_minimum_required(VERSION 3.5) -project (tkDNN) +project (tkDNN CUDA) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) if(UNIX) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") endif() if(WIN32) set(CMAKE_CXX_STANDARD 14) -set(CMAKE_CXX_FLAGS "/O2 /FS ") +set(CMAKE_CXX_FLAGS "/O2 ") +set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) endif(WIN32) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) @@ -33,7 +34,7 @@ include_directories(${CUDNN_INCLUDE_DIR}) # compile file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) -cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) +add_library(kernels SHARED ${tkdnn_CUSRC}) target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES}) @@ -47,7 +48,7 @@ find_package(OpenCV REQUIRED) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV") # gives problems in cross-compiling, probably malformed cmake config -#find_package(yaml-cpp REQUIRED) +find_package(yaml-cpp REQUIRED) #------------------------------------------------------------------------------- # Build Libraries diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 9f50d0b..609affe 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -1,7 +1,7 @@ #include #include #include /* srand, rand */ -#include +//#include #include #include "CenternetDetection.h" diff --git a/include/tkDNN/test.h b/include/tkDNN/test.h index 13c943e..f1d44bd 100644 --- a/include/tkDNN/test.h +++ b/include/tkDNN/test.h @@ -29,7 +29,8 @@ int testInference(std::vector input_bins, std::vector readBinaryFile(input_bins[0], net->input_dim.tot(), &input_h, &data); // outputs - dnnType *cudnn_out[outputs.size()], *rt_out[outputs.size()]; + //dnnType *cudnn_out[outputs.size()], *rt_out[outputs.size()]; + std::vector cudnn_out,rt_out; tk::dnn::dataDim_t dim1 = net->input_dim; //input dim printCenteredTitle(" CUDNN inference ", '=', 30); { -- 2.52.0 From 4a9031433399b6dbf5d08a6c963f8e3f3a829b72 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Tue, 26 Jan 2021 23:11:58 +0530 Subject: [PATCH 06/22] minor fixes in test.h --- CMakeLists.txt | 4 ++-- include/tkDNN/test.h | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 4904e5d..c173eca 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,6 +1,6 @@ cmake_minimum_required(VERSION 3.5) -project (tkDNN CUDA) +project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) if(UNIX) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") @@ -34,7 +34,7 @@ include_directories(${CUDNN_INCLUDE_DIR}) # compile file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu") cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS}) -add_library(kernels SHARED ${tkdnn_CUSRC}) +cuda_add_library(kernels SHARED ${tkdnn_CUSRC}) target_link_libraries(kernels ${CUDA_CUBLAS_LIBRARIES}) diff --git a/include/tkDNN/test.h b/include/tkDNN/test.h index f1d44bd..e842269 100644 --- a/include/tkDNN/test.h +++ b/include/tkDNN/test.h @@ -40,7 +40,7 @@ int testInference(std::vector input_bins, std::vector TKDNN_TSTOP dim1.print(); } - for(int i=0; idstData; + for(int i=0; idstData); if(netRT != nullptr) { tk::dnn::dataDim_t dim2 = net->input_dim; @@ -51,7 +51,7 @@ int testInference(std::vector input_bins, std::vector TKDNN_TSTOP dim2.print(); } - for(int i=0; ibuffersRT[i+1]; + for(int i=0; ibuffersRT[i+1]); } int ret_cudnn = 0, ret_tensorrt = 0, ret_cudnn_tensorrt = 0; -- 2.52.0 From 6aa8666be54ce7726654afd3062704555a836161 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Wed, 10 Mar 2021 12:59:04 +0530 Subject: [PATCH 07/22] Commits for msvc 16.9 --- CMakeLists.txt | 2 +- demo/demo/demo.cpp | 2 +- src/kernels/deformable_conv.cu | 2 +- src/utils.cpp | 11 ++++++++++- 4 files changed, 13 insertions(+), 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c173eca..77425d3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,7 +7,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declara endif() if(WIN32) set(CMAKE_CXX_STANDARD 14) -set(CMAKE_CXX_FLAGS "/O2 ") +set(CMAKE_CXX_FLAGS "/O2 /FS ") set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) endif(WIN32) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 609affe..c622f2b 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -22,7 +22,7 @@ int main(int argc, char *argv[]) { signal(SIGINT, sig_handler); - std::string net = "yolo3_berkeley.rt"; + std::string net = "yolo4tiny_fp32.rt"; if(argc > 1) net = argv[1]; std::string input = "../demo/yolo_test.mp4"; diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 592c538..4dbc552 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -18,7 +18,7 @@ inline int GET_BLOCKS(const int N) } -__device__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width, +__device__ __host__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width, const int height, const int width, float h, float w) { int h_low = floor(h); int w_low = floor(w); diff --git a/src/utils.cpp b/src/utils.cpp index bbdf516..52775df 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -23,14 +23,23 @@ bool fileExist(const char *fname) { void downloadWeightsifDoNotExist(const std::string& input_bin, const std::string& test_folder, const std::string& weights_url){ if(!fileExist(input_bin.c_str())){ std::string mkdir_cmd = "mkdir " + test_folder; - std::string wget_cmd = "wget " + weights_url + " -O " + test_folder + "/weights.zip"; + std::string wget_cmd = "curl " + weights_url + " --output " + test_folder + "/weights.zip"; +#ifdef __linux__ std::string unzip_cmd = "unzip " + test_folder + "/weights.zip -d" + test_folder; std::string rm_cmd = "rm " + test_folder + "/weights.zip"; + +#elif _WIN32 + + std::string unzip_cmd = "7z x " + test_folder + "/weights.zip -o" + test_folder; +#endif int err = 0; err = system(mkdir_cmd.c_str()); err = system(wget_cmd.c_str()); err = system(unzip_cmd.c_str()); +#ifdef __linux__ err = system(rm_cmd.c_str()); +#endif + } } -- 2.52.0 From 304ab49897938beb3ea6619fb5720bc1aae1280d Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar Date: Wed, 17 Mar 2021 22:26:26 +0530 Subject: [PATCH 08/22] shared_ptr migrations --- CMakeLists.txt | 2 +- include/tkDNN/NetworkRT.h | 7 ++++++- include/tkDNN/utils.h | 11 +++++++++++ src/NetworkRT.cpp | 11 +++++++---- 4 files changed, 25 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 77425d3..f478cab 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,7 +3,7 @@ cmake_minimum_required(VERSION 3.5) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) if(UNIX) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") endif() if(WIN32) set(CMAKE_CXX_STANDARD 14) diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 4c6c816..ca012fb 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -6,6 +6,7 @@ #include "Network.h" #include "Layer.h" #include "NvInfer.h" +#include namespace tk { namespace dnn { @@ -59,7 +60,8 @@ public: #if NV_TENSORRT_MAJOR >= 6 nvinfer1::IBuilderConfig *configRT; #endif - nvinfer1::ICudaEngine *engineRT; + std::shared_ptr engineRT; + //nvinfer1::ICudaEngine *engineRT; nvinfer1::IExecutionContext *contextRT; const static int MAX_BUFFERS_RT = 10; @@ -114,6 +116,9 @@ public: bool serialize(const char *filename); bool deserialize(const char *filename); + + + }; }} diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index 1404509..edc770e 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -110,6 +110,17 @@ double t_ns = time_ms.count(); FatalError(_error.str()); \ } \ } +struct InferDeleter +{ + template + void operator()(T* obj) const + { + if (obj) + { + obj->destroy(); + } + } +}; typedef enum { ERROR_CUDNN = 2, diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 501ade4..4006caf 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -137,9 +137,11 @@ NetworkRT::NetworkRT(Network *net, const char *name) { printCudaMemUsage(); std::cout<<"Building tensorRT cuda engine...\n"; #if NV_TENSORRT_MAJOR >= 6 - engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); + //engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); + engineRT = std::shared_ptr(builderRT->buildEngineWithConfig(*networkRT,*configRT),InferDeleter()); #else - engineRT = builderRT->buildCudaEngine(*networkRT); + //engineRT = builderRT->buildCudaEngine(*networkRT); + engineRT = std::shared_ptr(builderRT->buildCudaEngine(*networkRT)); #endif if(engineRT == nullptr) FatalError("cloud not build cuda engine") @@ -561,7 +563,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); checkNULL(lRT); lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() ); - delete(inputs); + delete[](inputs); // batchnorm void *bias_b, *power_b, *mean_b, *variance_b, *scales_b; if(dtRT == DataType::kHALF) { @@ -629,7 +631,8 @@ bool NetworkRT::deserialize(const char *filename) { pluginFactory = new PluginFactory(); runtimeRT = createInferRuntime(loggerRT); - engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); + //engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); + engineRT = std::shared_ptr(runtimeRT->deserializeCudaEngine(gieModelStream,size,(IPluginFactory*)pluginFactory),InferDeleter()); //if (gieModelStream) delete [] gieModelStream; return true; -- 2.52.0 From 06787a931fd8897ce8400d62c60393b8ef5fdc85 Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar Date: Wed, 17 Mar 2021 23:06:50 +0530 Subject: [PATCH 09/22] minor migrations --- include/tkDNN/NetworkRT.h | 6 ++++-- src/NetworkRT.cpp | 6 ++++-- 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index ca012fb..a50cd8c 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -54,8 +54,10 @@ class NetworkRT { public: nvinfer1::DataType dtRT; - nvinfer1::IBuilder *builderRT; - nvinfer1::IRuntime *runtimeRT; + //nvinfer1::IBuilder *builderRT; + std::unique_ptr builderRT; + //nvinfer1::IRuntime *runtimeRT; + std::unique_ptr runtimeRT; nvinfer1::INetworkDefinition *networkRT; #if NV_TENSORRT_MAJOR >= 6 nvinfer1::IBuilderConfig *configRT; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 4006caf..e6eb778 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -33,7 +33,8 @@ NetworkRT::NetworkRT(Network *net, const char *name) { float(NV_TENSORRT_PATCH)/100; std::cout<<"New NetworkRT (TensorRT v"<(createInferBuilder(loggerRT)); + //builderRT = createInferBuilder(loggerRT); std::cout<<"Float16 support: "<platformHasFastFp16()<<"\n"; std::cout<<"Int8 support: "<platformHasFastInt8()<<"\n"; #if NV_TENSORRT_MAJOR >= 5 @@ -630,7 +631,8 @@ bool NetworkRT::deserialize(const char *filename) { } pluginFactory = new PluginFactory(); - runtimeRT = createInferRuntime(loggerRT); + //runtimeRT = createInferRuntime(loggerRT); + runtimeRT = std::unique_ptr(createInferRuntime(loggerRT)); //engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); engineRT = std::shared_ptr(runtimeRT->deserializeCudaEngine(gieModelStream,size,(IPluginFactory*)pluginFactory),InferDeleter()); //if (gieModelStream) delete [] gieModelStream; -- 2.52.0 From e94e1f7622d5bc4479036b39e012213acf9c1fdb Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Wed, 24 Mar 2021 22:26:46 +0530 Subject: [PATCH 10/22] tkdnn first patch for windows --- .gitignore | 1 + demo/demo/demo.cpp | 2 +- include/tkDNN/DetectionNN.h | 1 - include/tkDNN/ImuOdom.h | 2 + include/tkDNN/pluginsRT/ActivationLeakyRT.h | 3 +- include/tkDNN/pluginsRT/ActivationMishRT.h | 3 +- .../tkDNN/pluginsRT/ActivationReLUCeilingRT.h | 3 +- include/tkDNN/pluginsRT/ActivationSigmoidRT.h | 3 +- include/tkDNN/pluginsRT/DeformableConvRT.h | 3 +- include/tkDNN/pluginsRT/FlattenConcatRT.h | 3 +- .../tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 3 +- include/tkDNN/pluginsRT/RegionRT.h | 3 +- include/tkDNN/pluginsRT/ReorgRT.h | 3 +- include/tkDNN/pluginsRT/ReshapeRT.h | 3 +- include/tkDNN/pluginsRT/ResizeLayerRT.h | 3 +- include/tkDNN/pluginsRT/RouteRT.h | 3 +- include/tkDNN/pluginsRT/ShortcutRT.h | 3 +- include/tkDNN/pluginsRT/UpsampleRT.h | 5 +- include/tkDNN/pluginsRT/YoloRT.h | 36 +++--- src/NetworkRT.cpp | 119 ++++++++++++++---- 20 files changed, 151 insertions(+), 54 deletions(-) diff --git a/.gitignore b/.gitignore index 5be5a73..7e5c4ca 100644 --- a/.gitignore +++ b/.gitignore @@ -15,3 +15,4 @@ build/ cmake-build-release/ demo/COCO_val2017 demo/BDD100K_val +/.vs diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index c622f2b..f97ead9 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -131,7 +131,7 @@ int main(int argc, char *argv[]) { double mean = 0; std::cout<stats.begin(), detNN->stats.end())/n_batch<<" ms\n"; + std::cout<<"Min: "<<*std::min_element(detNN->stats.begin(), detNN->stats.end())/n_batch<<" ms\n"; std::cout<<"Max: "<<*std::max_element(detNN->stats.begin(), detNN->stats.end())/n_batch<<" ms\n"; for(int i=0; istats.size(); i++) mean += detNN->stats[i]; mean /= detNN->stats.size(); std::cout<<"Avg: "< #elif _WIN32 +#define _USE_MATH_DEFINES +#include #include #endif diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index d3f66fb..9e26b2b 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -52,8 +52,9 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); } int size; diff --git a/include/tkDNN/pluginsRT/ActivationMishRT.h b/include/tkDNN/pluginsRT/ActivationMishRT.h index 1744ab0..5d660af 100644 --- a/include/tkDNN/pluginsRT/ActivationMishRT.h +++ b/include/tkDNN/pluginsRT/ActivationMishRT.h @@ -52,8 +52,9 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); } int size; diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h index 286f22e..50ceb81 100644 --- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -51,9 +51,10 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, ceiling); tk::dnn::writeBUF(buf, size); + assert(buf = a + getSerializationSize()); } diff --git a/include/tkDNN/pluginsRT/ActivationSigmoidRT.h b/include/tkDNN/pluginsRT/ActivationSigmoidRT.h index 1d47136..bcc58c7 100644 --- a/include/tkDNN/pluginsRT/ActivationSigmoidRT.h +++ b/include/tkDNN/pluginsRT/ActivationSigmoidRT.h @@ -52,8 +52,9 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); } int size; diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 225a24e..5cb2bab 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -116,7 +116,7 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, chunk_dim); tk::dnn::writeBUF(buf, kh); tk::dnn::writeBUF(buf, kw); @@ -163,6 +163,7 @@ public: for(int i=0; i(buffer); + char *buf = reinterpret_cast(buffer),*a = buf; tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); tk::dnn::writeBUF(buf, rows); tk::dnn::writeBUF(buf, cols); + assert(buf == a + getSerializationSize()); } int c, h, w; diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 911fca2..0899a34 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -55,7 +55,7 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, this->c); tk::dnn::writeBUF(buf, this->h); @@ -65,6 +65,7 @@ public: tk::dnn::writeBUF(buf, this->stride_W); tk::dnn::writeBUF(buf, this->winSize); tk::dnn::writeBUF(buf, this->padding); + assert(buf == a + getSerializationSize()); } int n, c, h, w; diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index f0d127e..8487652 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -73,13 +73,14 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, classes); tk::dnn::writeBUF(buf, coords); tk::dnn::writeBUF(buf, num); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + assert(buf == a + getSerializationSize()); } int c, h, w; diff --git a/include/tkDNN/pluginsRT/ReorgRT.h b/include/tkDNN/pluginsRT/ReorgRT.h index ee85718..c1b529a 100644 --- a/include/tkDNN/pluginsRT/ReorgRT.h +++ b/include/tkDNN/pluginsRT/ReorgRT.h @@ -52,11 +52,12 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, stride); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + assert(buf == a + getSerializationSize()); } int c, h, w, stride; diff --git a/include/tkDNN/pluginsRT/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h index 97030db..37017c7 100644 --- a/include/tkDNN/pluginsRT/ReshapeRT.h +++ b/include/tkDNN/pluginsRT/ReshapeRT.h @@ -50,11 +50,12 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a = buf; tk::dnn::writeBUF(buf, n); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + assert(buf == a + getSerializationSize()); } int n, c, h, w; diff --git a/include/tkDNN/pluginsRT/ResizeLayerRT.h b/include/tkDNN/pluginsRT/ResizeLayerRT.h index ae87dbf..cde52bf 100644 --- a/include/tkDNN/pluginsRT/ResizeLayerRT.h +++ b/include/tkDNN/pluginsRT/ResizeLayerRT.h @@ -52,7 +52,7 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, o_c); tk::dnn::writeBUF(buf, o_h); @@ -61,6 +61,7 @@ public: tk::dnn::writeBUF(buf, i_c); tk::dnn::writeBUF(buf, i_h); tk::dnn::writeBUF(buf, i_w); + assert(buf == a + getSerializationSize()); } int i_c, i_h, i_w, o_c, o_h, o_w; diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 23f30b7..5a8c170 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -75,7 +75,7 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, groups); tk::dnn::writeBUF(buf, group_id); tk::dnn::writeBUF(buf, in); @@ -85,6 +85,7 @@ public: tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + assert(buf == a + getSerializationSize()); } static const int MAX_INPUTS = 4; diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index 3eadd3f..17f050f 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -59,13 +59,14 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, bc); tk::dnn::writeBUF(buf, bh); tk::dnn::writeBUF(buf, bw); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + assert(buf == a + getSerializationSize()); } diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h index 7a62abc..5350b7e 100644 --- a/include/tkDNN/pluginsRT/UpsampleRT.h +++ b/include/tkDNN/pluginsRT/UpsampleRT.h @@ -54,11 +54,14 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); + char *buf = reinterpret_cast(buffer),*a=buf; tk::dnn::writeBUF(buf, stride); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); + std::cout << "Upsample Serialization SIze" << getSerializationSize() << std::endl; + + assert(buf == a + getSerializationSize()); } int c, h, w, stride; diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 9af8587..451d99f 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -89,21 +89,25 @@ public: } virtual void serialize(void* buffer) override { - char *buf = reinterpret_cast(buffer); - tk::dnn::writeBUF(buf, classes); - tk::dnn::writeBUF(buf, num); - tk::dnn::writeBUF(buf, n_masks); - tk::dnn::writeBUF(buf, scaleXY); - tk::dnn::writeBUF(buf, nms_thresh); - tk::dnn::writeBUF(buf, nms_kind); - tk::dnn::writeBUF(buf, new_coords); - tk::dnn::writeBUF(buf, c); - tk::dnn::writeBUF(buf, h); - tk::dnn::writeBUF(buf, w); - for(int i=0; i(buffer),*a=buf; + tk::dnn::writeBUF(buf, classes); std::cout << "Classes :" << classes << std::endl; + tk::dnn::writeBUF(buf, num); std::cout << "Num : " << num << std::endl; + tk::dnn::writeBUF(buf, n_masks); std::cout << "N_Masks" << n_masks << std::endl; + tk::dnn::writeBUF(buf, scaleXY); std::cout << "ScaleXY :" << scaleXY << std::endl; + tk::dnn::writeBUF(buf, nms_thresh); std::cout << "nms_thresh :" << nms_thresh << std::endl; + tk::dnn::writeBUF(buf, nms_kind); std::cout << "nms_kind : " << nms_kind << std::endl; + tk::dnn::writeBUF(buf, new_coords); std::cout << "new_coords : " << new_coords << std::endl; + tk::dnn::writeBUF(buf, c); std::cout << "C : " << c << std::endl; + tk::dnn::writeBUF(buf, h); std::cout << "H : " << h << std::endl; + tk::dnn::writeBUF(buf, w); std::cout << "C : " << c << std::endl; + for (int i = 0; i < n_masks; i++) + { + tk::dnn::writeBUF(buf, mask[i]); std::cout << "mask[i] : " << mask[i] << std::endl; + } + for (int i = 0; i < n_masks * 2 * num; i++) + { + tk::dnn::writeBUF(buf, bias[i]); std::cout << "bias[i] : " << bias[i] << std::endl; + } // save classes names for(int i=0; i(serialData); + const char * buf = reinterpret_cast(serialData),*bufCheck = buf; std::string name(layerName); - //std::cout<size = readBUF(buf); + assert(buf == bufCheck + serialLength); return a; } if(name.find("ActivationMish") == 0) { ActivationMishRT *a = new ActivationMishRT(); a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); return a; } if(name.find("ActivationCReLU") == 0) { - ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); + float activationReluTemp = readBUF(buf); + //ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); + ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp); a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); return a; } if(name.find("Region") == 0) { - RegionRT *r = new RegionRT(readBUF(buf), //classes + int classesTemp = readBUF(buf); + int coordsTemp = readBUF(buf); + int numTemp = readBUF(buf); + /*RegionRT *r = new RegionRT(readBUF(buf), //classes readBUF(buf), //coords - readBUF(buf)); //num + readBUF(buf)); //num8*/ + RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); + assert(buf == bufCheck + serialLength); return r; } if(name.find("Reorg") == 0) { - ReorgRT *r = new ReorgRT(readBUF(buf)); //stride + int strideTemp = readBUF(buf); + //ReorgRT *r = new ReorgRT(readBUF(buf)); //stride + ReorgRT *r = new ReorgRT(strideTemp); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); + assert(buf == bufCheck + serialLength); return r; } @@ -695,27 +708,46 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->h = readBUF(buf); r->w = readBUF(buf); return r; + assert(buf == bufCheck + serialLength); } if(name.find("Pooling") == 0) { - MaxPoolFixedSizeRT *r = new MaxPoolFixedSizeRT( readBUF(buf), //c + /* MaxPoolFixedSizeRT *r = new MaxPoolFixedSizeRT( readBUF(buf), //c readBUF(buf), //h readBUF(buf), //w readBUF(buf), //n readBUF(buf), //strideH readBUF(buf), //strideW readBUF(buf), //winSize - readBUF(buf)); //padding + readBUF(buf)); //padding*/ + + int cTemp = readBUF(buf); + int hTemp = readBUF(buf); + int wTemp = readBUF(buf); + int nTemp = readBUF(buf); + int strideHTemp = readBUF(buf); + int strideWTemp = readBUF(buf); + int winSizeTemp = readBUF(buf); + int paddingTemp = readBUF(buf); + + MaxPoolFixedSizeRT* r = new MaxPoolFixedSizeRT(cTemp, hTemp, wTemp, nTemp, strideHTemp, strideWTemp, winSizeTemp, paddingTemp); + assert(buf == bufCheck + serialLength); return r; } if(name.find("Resize") == 0) { - ResizeLayerRT *r = new ResizeLayerRT(readBUF(buf), //o_c + /*ResizeLayerRT *r = new ResizeLayerRT(readBUF(buf), //o_c readBUF(buf), //o_h - readBUF(buf)); //o_w + readBUF(buf)); //o_w*/ + int o_cTemp = readBUF(buf); + int o_hTemp = readBUF(buf); + int o_wTemp = readBUF(buf); + ResizeLayerRT* r = new ResizeLayerRT(o_cTemp, o_hTemp, o_wTemp); + r->i_c = readBUF(buf); r->i_h = readBUF(buf); r->i_w = readBUF(buf); + assert(buf == bufCheck + serialLength); return r; } @@ -726,6 +758,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa r->w = readBUF(buf); r->rows = readBUF(buf); r->cols = readBUF(buf); + assert(buf == bufCheck + serialLength); return r; } @@ -737,20 +770,33 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa new_dim.h = readBUF(buf); new_dim.w = readBUF(buf); ReshapeRT *r = new ReshapeRT(new_dim); + assert(buf == bufCheck + serialLength); return r; } if(name.find("Yolo") == 0) { - YoloRT *r = new YoloRT(readBUF(buf), //classes - readBUF(buf), //num - nullptr, //yolo - readBUF(buf), //n_masks - readBUF(buf), //scale_xy - readBUF(buf), //nms_thresh - readBUF(buf), //nms_kind - readBUF(buf) //new_coords - ); + + int classes_temp = readBUF(buf); + int num_temp = readBUF(buf); + int n_masks_temp = readBUF(buf); + float scale_xy_temp = readBUF(buf); + float nms_thresh_temp = readBUF(buf); + int nms_kind_temp = readBUF(buf); + int new_coords_temp = readBUF(buf); + std::cout << classes_temp << ":" << num_temp << ":" << ":" << n_masks_temp << ":" << nms_thresh_temp << ":" << nms_kind_temp << ":" << new_coords_temp << std::endl; + + YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp); + + /* std::cout << "classes : " << r->classes; + std::cout << "num : " << r->num; + std::cout << "n_masks : " << r->n_masks; + std::cout << "scalexy : " << r->scaleXY; + std::cout << "nms_thresh : " << r->nms_thresh; + std::cout << "nms_kind : " << r->nms_kind; + std::cout << "new_coords : " << r->new_coords;*/ + + r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); @@ -767,36 +813,62 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa tmp[j] = readBUF(buf); r->classesNames[i] = std::string(tmp); } + assert(buf == bufCheck + serialLength); yolos[n_yolos++] = r; return r; } if(name.find("Upsample") == 0) { - UpsampleRT *r = new UpsampleRT(readBUF(buf)); //stride + //UpsampleRT *r = new UpsampleRT(readBUF(buf)); //stride + int strideTemp = readBUF(buf); + UpsampleRT* r = new UpsampleRT(strideTemp); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); + assert(buf == bufCheck + serialLength); return r; } if(name.find("Route") == 0) { - RouteRT *r = new RouteRT(readBUF(buf),readBUF(buf)); + //RouteRT *r = new RouteRT(readBUF(buf),readBUF(buf)); + int groupsTemp = readBUF(buf); + int group_idTemp = readBUF(buf); + RouteRT* r = new RouteRT(groupsTemp, group_idTemp); r->in = readBUF(buf); for(int i=0; ic_in[i] = readBUF(buf); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); + assert(buf == bufCheck + serialLength); return r; } if(name.find("Deformable") == 0) { - DeformableConvRT *r = new DeformableConvRT(readBUF(buf), readBUF(buf), readBUF(buf), + /*DeformableConvRT *r = new DeformableConvRT(readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf), readBUF(buf),readBUF(buf),readBUF(buf),readBUF(buf), readBUF(buf),readBUF(buf),readBUF(buf),readBUF(buf), - nullptr); + nullptr); */ + int chuck_dimTemp = readBUF(buf); + int khTemp = readBUF(buf); + int kwTemp = readBUF(buf); + int shTemp = readBUF(buf); + int swTemp = readBUF(buf); + int phTemp = readBUF(buf); + int pwTemp = readBUF(buf); + int deformableGroupTemp = readBUF(buf); + int i_nTemp = readBUF(buf); + int i_cTemp = readBUF(buf); + int i_hTemp = readBUF(buf); + int i_wTemp = readBUF(buf); + int o_nTemp = readBUF(buf); + int o_cTemp = readBUF(buf); + int o_hTemp = readBUF(buf); + int o_wTemp = readBUF(buf); + + DeformableConvRT* r = new DeformableConvRT(chuck_dimTemp, khTemp, kwTemp, shTemp, swTemp, phTemp, pwTemp, deformableGroupTemp, i_nTemp, i_cTemp, i_hTemp, i_wTemp, o_nTemp, o_cTemp, o_hTemp, o_wTemp, nullptr); dnnType *aus = new dnnType[r->chunk_dim*2]; for(int i=0; ichunk_dim*2; i++) aus[i] = readBUF(buf); @@ -827,6 +899,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa aus[i] = readBUF(buf); checkCuda( cudaMemcpy(r->ones_d2, aus, sizeof(dnnType)*r->dim_ones, cudaMemcpyHostToDevice) ); free(aus); + assert(buf == bufCheck + serialLength); return r; } -- 2.52.0 From 78859fe19109f265489090911ee57629bf31ff67 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Thu, 25 Mar 2021 16:38:53 +0530 Subject: [PATCH 11/22] timer fix --- .gitignore | 4 ++++ CMakeLists.txt | 2 +- include/tkDNN/utils.h | 2 +- 3 files changed, 6 insertions(+), 2 deletions(-) diff --git a/.gitignore b/.gitignore index 7e5c4ca..c1d362c 100644 --- a/.gitignore +++ b/.gitignore @@ -16,3 +16,7 @@ cmake-build-release/ demo/COCO_val2017 demo/BDD100K_val /.vs +cmake-build-minsizerel/* +scripts/COCO_val2017/* +scripts/COCO_val2017.zip +scripts/all_labels.txt \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index f478cab..579a8bc 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,7 +7,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declara endif() if(WIN32) set(CMAKE_CXX_STANDARD 14) -set(CMAKE_CXX_FLAGS "/O2 /FS ") +set(CMAKE_CXX_FLAGS "/O1 /FS") set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) endif(WIN32) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index edc770e..ce4a617 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -60,7 +60,7 @@ #define TKDNN_TSTART auto start = std::chrono::high_resolution_clock::now(); #define TKDNN_TSTOP auto stop = std::chrono::high_resolution_clock::now(); \ std::chrono::duration duration = stop -start; \ -auto time_ms = std::chrono::duration_cast(duration);\ +auto time_ms = std::chrono::duration_cast(duration);\ double t_ns = time_ms.count(); #endif -- 2.52.0 From 44b71ae6f33acbaf09fd8a4064d15b8395768cfb Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Thu, 25 Mar 2021 20:44:43 +0530 Subject: [PATCH 12/22] ReadMe.md windows changes --- CMakeLists.txt | 2 +- README.md | 85 ++++++++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 86 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 579a8bc..1b7ed63 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,7 +7,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declara endif() if(WIN32) set(CMAKE_CXX_STANDARD 14) -set(CMAKE_CXX_FLAGS "/O1 /FS") +set(CMAKE_CXX_FLAGS "/O1 /FS /EHsc") set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) endif(WIN32) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) diff --git a/README.md b/README.md index 84e0037..242220a 100644 --- a/README.md +++ b/README.md @@ -80,6 +80,13 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001 - [mAP demo](#map-demo) - [Existing tests and supported networks](#existing-tests-and-supported-networks) - [References](#references) + - [tkDNN on Windows 10 (experimental)](#tkdnn-on-windows) + - [Dependencies](#dependencies) + - [Compiling tkDNN on Windows](#tkdnn-windows-compile) + - [Run the demo on Windows](#run-the-demo-on-windows) + - [FP16 interference windows](#fp16-windows) + - [INT8 interference windows](#int8-windows) + @@ -355,6 +362,84 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing | yolo4tiny | Yolov4 tiny 9 | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download) | | yolo4x | Yolov4x-mish 9 | [COCO 2017](http://cocodataset.org/) | 80 | 672x672 | [weights](https://cloud.hipert.unimore.it/s/BLPpiAigZJLorQD/download) | +##tkDNN on Windows 10 (experimental) + +### Dependencies +This branch should work on every NVIDIA GPU supported in windows with the following dependencies: + +* WINDOWS 10 1803 or HIGHER +* CUDA 10.0 (Recommended CUDA 11.0 +) +* CUDNN 7.6 (Recommended CUDNN 8.0.0 +) +* TENSORRT 6.0.1 (Recommended TENSORRT 7.1 +) +* OPENCV 3.4 (Recommended OPENCV 4.2.0 +) +* MSVC 16.7 (Recommended MSVC 16.8/16.9) +* YAML-CPP 0.5.2 +* EIGEN3 +* 7ZIP (ADD TO PATH) +* NINJA 1.10 + +All the above mentioned dependencies except 7ZIP can be installed using Microsoft's [VCPKG](https://github.com/microsoft/vcpkg.git) . +After bootstrapping VCPKG the dependencies can be built and installed using the following command : + +```vcpkg.exe install opencv4[tbb,jpeg,tiff,opengl,openmp,png,ffmpeg]:x64-windows yaml-cpp:x64-windows eigen3:x64-windows --x-install-root=C:\opt --x-buildtrees-root=C:\temp_vcpkg_build``` + +After VCPKG finishes building and installing all the packages delete C:\temp_vcpkg_build and add C:\opt\x64-windows\bin and C:\opt\x64-windows\debug\bin to path + +### Compiling tkDNN on Windows + +tkDNN is built with cmake(3.15+) on windows along with ninja.Msbuild and NMake Makefiles are drastically slower when compiling the library compared to windows +``` +git clone https://git.hipert.unimore.it/research-cv-chandirasekar/tkdnn-windows.git +cd tkdnn-windows +mkdir build +cd build +cmake -DCMAKE_BUILD_TYPE=Release -G"Ninja" .. +ninja -j4 +``` + +### Run the demo on Windows + +This example uses yolo4_tiny.\ +To run the object detection file create .rt file bu running: +``` +.\test_yolo4tiny.exe +``` + +Once the rt file has been successfully create,run the demo using the following command: +``` +.\demo.exe yolo4tiny_fp32.rt ..\demo\yolo_test.mp4 y +``` + For general info on more demo paramters,check Run the demo section on top + +### FP16 interference windows + +This is an untested feature on windows.To run the object detection demo with FP16 interference follow the below steps(example with yolo4tiny): +``` +set TKDNN_MODE=FP16 +del /f yolo4tiny_fp16.rt +.\test_yolo4tiny.exe +.\demo.exe yolo4tiny_fp16.rt ..\demo\yolo_test.mp4 +``` + +### INT8 interference windows +To run object detection demo with INT8 (example with yolo4tiny): +``` +set TKDNN_MODE=INT8 +set TKDNN_CALIB_LABEL_PATH=..\demo\COCO_val2017\all_labels.txt +set TKDNN_CALIB_IMG_PATH=..\demo\COCO_val2017\all_images.txt +del /f yolo4tiny_int8.rt # be sure to delete(or move) old tensorRT files +.\test_yolo4tiny.exe # run the yolo test (is slow) +.\demo.exe yolo4tiny_int8.rt ..\demo\yolo_test.mp4 y + +``` + + + + + + + + ## References -- 2.52.0 From f3d159143025672d5cd1bc91a89370dac1d902c2 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Thu, 25 Mar 2021 20:46:12 +0530 Subject: [PATCH 13/22] ReadMe.md windows changes --- README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 242220a..f685b58 100644 --- a/README.md +++ b/README.md @@ -81,7 +81,7 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001 - [Existing tests and supported networks](#existing-tests-and-supported-networks) - [References](#references) - [tkDNN on Windows 10 (experimental)](#tkdnn-on-windows) - - [Dependencies](#dependencies) + - [Dependencies-Windows](#dependencies-windows) - [Compiling tkDNN on Windows](#tkdnn-windows-compile) - [Run the demo on Windows](#run-the-demo-on-windows) - [FP16 interference windows](#fp16-windows) @@ -364,7 +364,7 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing ##tkDNN on Windows 10 (experimental) -### Dependencies +### Dependencies-Windows This branch should work on every NVIDIA GPU supported in windows with the following dependencies: * WINDOWS 10 1803 or HIGHER -- 2.52.0 From f12ec3c935ce4fdcb45c966e9b4f753454d630b9 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Fri, 26 Mar 2021 12:18:46 +0530 Subject: [PATCH 14/22] added yolo4x and yolo4-csp from the github repo and download file corrections --- include/tkDNN/Layer.h | 7 +- include/tkDNN/NetworkRT.h | 1 + .../tkDNN/pluginsRT/ActivationLogisticRT.h | 60 + include/tkDNN/pluginsRT/YoloRT.h | 29 +- src/Activation.cpp | 6 +- src/DarknetParser.cpp | 1 + src/NetworkRT.cpp | 21 +- src/Yolo.cpp | 21 +- tests/darknet/cfg/yolo4-csp.cfg | 1279 +++++++++++++++++ tests/darknet/cfg/yolo4x.cfg | 21 +- tests/darknet/yolo4-csp.cpp | 36 + tests/darknet/yolo4x.cpp | 2 +- 12 files changed, 1444 insertions(+), 40 deletions(-) create mode 100644 include/tkDNN/pluginsRT/ActivationLogisticRT.h create mode 100644 tests/darknet/cfg/yolo4-csp.cfg create mode 100644 tests/darknet/yolo4-csp.cpp diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 25c4565..e097372 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -19,6 +19,7 @@ enum layerType_t { LAYER_ACTIVATION_CRELU, LAYER_ACTIVATION_LEAKY, LAYER_ACTIVATION_MISH, + LAYER_ACTIVATION_LOGISTIC, LAYER_FLATTEN, LAYER_RESHAPE, LAYER_MULADD, @@ -68,6 +69,7 @@ public: case LAYER_ACTIVATION_CRELU: return "ActivationCReLU"; case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky"; case LAYER_ACTIVATION_MISH: return "ActivationMish"; + case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic"; case LAYER_FLATTEN: return "Flatten"; case LAYER_RESHAPE: return "Reshape"; case LAYER_MULADD: return "MulAdd"; @@ -212,7 +214,8 @@ public: typedef enum { ACTIVATION_ELU = 100, ACTIVATION_LEAKY = 101, - ACTIVATION_MISH = 102 + ACTIVATION_MISH = 102, + ACTIVATION_LOGISTIC = 103 } tkdnnActivationMode_t; /** @@ -233,6 +236,8 @@ public: return LAYER_ACTIVATION_LEAKY; else if (act_mode == ACTIVATION_MISH) return LAYER_ACTIVATION_MISH; + else if (act_mode == ACTIVATION_LOGISTIC) + return LAYER_ACTIVATION_LOGISTIC; else return LAYER_ACTIVATION; }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index a50cd8c..b39360c 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -27,6 +27,7 @@ using namespace nvinfer1; #include "pluginsRT/ActivationLeakyRT.h" #include "pluginsRT/ActivationReLUCeilingRT.h" #include "pluginsRT/ActivationMishRT.h" +#include "pluginsRT/ActivationLogisticRT.h" #include "pluginsRT/ReorgRT.h" #include "pluginsRT/RegionRT.h" #include "pluginsRT/RouteRT.h" diff --git a/include/tkDNN/pluginsRT/ActivationLogisticRT.h b/include/tkDNN/pluginsRT/ActivationLogisticRT.h new file mode 100644 index 0000000..83f62ff --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationLogisticRT.h @@ -0,0 +1,60 @@ +#include +#include "../kernels.h" + +class ActivationLogisticRT : public IPlugin { + +public: + ActivationLogisticRT() { + + + } + + ~ActivationLogisticRT(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, stream); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 1*sizeof(int); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, size); + } + + int size; +}; \ No newline at end of file diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 451d99f..0dd26e1 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -64,20 +64,23 @@ public: checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - for (int b = 0; b < batchSize; ++b){ - for(int n = 0; n < n_masks; ++n){ - int index = entry_index(b, n*w*h, 0); - if (new_coords == 1) - activationLOGISTICForward(srcData + index, dstData + index, 4*w*h, stream); //x,y,w,h - else - activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y - if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); - - index = entry_index(b, n*w*h, 4); - activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream); - } - } + for (int b = 0; b < batchSize; ++b){ + for(int n = 0; n < n_masks; ++n){ + int index = entry_index(b, n*w*h, 0); + if (new_coords == 1){ + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + } + else{ + activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y + + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + + index = entry_index(b, n*w*h, 4); + activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream); + } + } + } //std::cout<<"YOLO END\n"; return 0; diff --git a/src/Activation.cpp b/src/Activation.cpp index 28c7624..4219271 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -52,7 +52,11 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { else if(act_mode == ACTIVATION_MISH) { activationMishForward(srcData, dstData, dim.tot()); - } else { + } + else if(act_mode == ACTIVATION_LOGISTIC) { + activationLOGISTICForward(srcData, dstData, dim.tot()); + + }else { dnnType alpha = dnnType(1); dnnType beta = dnnType(0); checkCUDNN( cudnnActivationForward(net->cudnnHandle, diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 7b5410c..69b6b29 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -187,6 +187,7 @@ namespace tk { namespace dnn { if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; + else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC; else { FatalError("activation not supported: " + f.activation); } netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); }; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 29df411..1e0b062 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -229,7 +229,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Conv2d*) l); if(type == LAYER_POOLING) return convert_layer(input, (Pooling*) l); - if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -424,6 +424,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { checkNULL(lRT); return lRT; } + else if(l->act_mode == ACTIVATION_LOGISTIC) { + IPlugin *plugin = new ActivationLogisticRT(); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + return lRT; + } else { FatalError("this Activation mode is not yet implemented"); return NULL; @@ -660,6 +666,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa assert(buf == bufCheck + serialLength); return a; } + if(name.find("ActivationLogistic") == 0) { + ActivationLogisticRT *a = new ActivationLogisticRT(); + a->size = readBUF(buf); + return a; + } if(name.find("ActivationCReLU") == 0) { float activationReluTemp = readBUF(buf); //ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); @@ -784,17 +795,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa float nms_thresh_temp = readBUF(buf); int nms_kind_temp = readBUF(buf); int new_coords_temp = readBUF(buf); - std::cout << classes_temp << ":" << num_temp << ":" << ":" << n_masks_temp << ":" << nms_thresh_temp << ":" << nms_kind_temp << ":" << new_coords_temp << std::endl; YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp); - /* std::cout << "classes : " << r->classes; - std::cout << "num : " << r->num; - std::cout << "n_masks : " << r->n_masks; - std::cout << "scalexy : " << r->scaleXY; - std::cout << "nms_thresh : " << r->nms_thresh; - std::cout << "nms_kind : " << r->nms_kind; - std::cout << "new_coords : " << r->new_coords;*/ r->c = readBUF(buf); diff --git a/src/Yolo.cpp b/src/Yolo.cpp index 6d0b546..1eb7843 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -73,8 +73,8 @@ Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h; } else{ - b.x = (i + x[index + 0 * stride] * 2 - 0.5) / lw; - b.y = (j + x[index + 1 * stride] * 2 - 0.5) / lh; + b.x = (i + x[index + 0 * stride] ) / lw; + b.y = (j + x[index + 1 * stride] ) / lh; b.w = x[index + 2 * stride] * x[index + 2 * stride] * 4 * biases[2 * n] / w; b.h = x[index + 3 * stride] * x[index + 3 * stride] * 4 * biases[2 * n + 1] / h; } @@ -88,15 +88,18 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) { for (int b = 0; b < dim.n; ++b){ for(int n = 0; n < n_masks; ++n){ int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim); - if (new_coords == 1) - activationLOGISTICForward(srcData + index, dstData + index, 4*dim.w*dim.h); - else + std::cout<<"new_coords"<scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + } + else{ activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h); - if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); - - index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim); - activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h); + if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1); + + index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim); + activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h); + } } } diff --git a/tests/darknet/cfg/yolo4-csp.cfg b/tests/darknet/cfg/yolo4-csp.cfg new file mode 100644 index 0000000..887898e --- /dev/null +++ b/tests/darknet/cfg/yolo4-csp.cfg @@ -0,0 +1,1279 @@ +[net] +# Testing +#batch=1 +#subdivisions=1 +# Training +batch=64 +subdivisions=8 +width=512 +height=512 +channels=3 +momentum=0.949 +decay=0.0005 +angle=0 +saturation = 1.5 +exposure = 1.5 +hue=.1 + +learning_rate=0.001 +burn_in=1000 +max_batches = 500500 +policy=steps +steps=400000,450000 +scales=.1,.1 + +mosaic=1 + +letter_box=1 + +ema_alpha=0.9998 + +#optimized_memory=1 + +#23:104x104 54:52x52 85:26x26 104:13x13 for 416 + + + +[convolutional] +batch_normalize=1 +filters=32 +size=3 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=2 +pad=1 +activation=mish + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +#[route] +#layers = -2 + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +[convolutional] +batch_normalize=1 +filters=32 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +#[route] +#layers = -1,-7 + +#[convolutional] +#batch_normalize=1 +#filters=64 +#size=1 +#stride=1 +#pad=1 +#activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=64 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=64 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-10 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-28 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-28 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +# Downsample + +[convolutional] +batch_normalize=1 +filters=1024 +size=3 +stride=2 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=3 +stride=1 +pad=1 +activation=mish + +[shortcut] +from=-3 +activation=linear + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1,-16 + +[convolutional] +batch_normalize=1 +filters=1024 +size=1 +stride=1 +pad=1 +activation=mish + +########################## + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +### SPP ### +[maxpool] +stride=1 +size=5 + +[route] +layers=-2 + +[maxpool] +stride=1 +size=9 + +[route] +layers=-4 + +[maxpool] +stride=1 +size=13 + +[route] +layers=-1,-3,-5,-6 +### End SPP ### + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[route] +layers = -1, -13 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[upsample] +stride=2 + +[route] +layers = 79 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[route] +layers = -1, -6 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[upsample] +stride=2 + +[route] +layers = 48 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -1, -3 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=mish + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=128 +activation=mish + +[route] +layers = -1, -6 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=mish + +########################## + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + + +[yolo] +mask = 0,1,2 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=0 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=4.0 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=5 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=256 +activation=mish + +[route] +layers = -1, -20 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=256 +activation=mish + +[route] +layers = -1,-6 + +[convolutional] +batch_normalize=1 +filters=256 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + + +[yolo] +mask = 3,4,5 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=1.0 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=5 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +size=3 +stride=2 +pad=1 +filters=512 +activation=mish + +[route] +layers = -1, -49 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[route] +layers = -2 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=512 +activation=mish + +[route] +layers = -1,-6 + +[convolutional] +batch_normalize=1 +filters=512 +size=1 +stride=1 +pad=1 +activation=mish + +[convolutional] +batch_normalize=1 +size=3 +stride=1 +pad=1 +filters=1024 +activation=mish + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=logistic + + +[yolo] +mask = 6,7,8 +anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 459, 401 +classes=80 +num=9 +jitter=.1 +scale_x_y = 2.0 +objectness_smooth=1 +ignore_thresh = .7 +truth_thresh = 1 +#random=1 +resize=1.5 +iou_thresh=0.2 +iou_normalizer=0.05 +cls_normalizer=0.5 +obj_normalizer=0.4 +iou_loss=ciou +nms_kind=diounms +beta_nms=0.6 +new_coords=1 +max_delta=2 \ No newline at end of file diff --git a/tests/darknet/cfg/yolo4x.cfg b/tests/darknet/cfg/yolo4x.cfg index 89f2564..f6604f6 100644 --- a/tests/darknet/cfg/yolo4x.cfg +++ b/tests/darknet/cfg/yolo4x.cfg @@ -5,8 +5,8 @@ # Training batch=64 subdivisions=8 -width=672 -height=672 +width=640 +height=640 channels=3 momentum=0.949 decay=0.0005 @@ -15,7 +15,7 @@ saturation = 1.5 exposure = 1.5 hue=.1 -learning_rate=0.00261 +learning_rate=0.001 burn_in=1000 max_batches = 500500 policy=steps @@ -26,6 +26,8 @@ mosaic=1 letter_box=1 +#optimized_memory=1 + [convolutional] batch_normalize=1 filters=32 @@ -1131,6 +1133,7 @@ size=1 stride=1 pad=1 activation=mish +stopbackward=800 ########################## @@ -1147,7 +1150,7 @@ size=1 stride=1 pad=1 filters=255 -activation=linear +activation=logistic [yolo] @@ -1156,6 +1159,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4 classes=80 num=9 jitter=.1 +scale_x_y = 2.0 objectness_smooth=0 ignore_thresh = .7 truth_thresh = 1 @@ -1169,6 +1173,7 @@ iou_loss=ciou nms_kind=diounms beta_nms=0.6 new_coords=1 +max_delta=5 [route] layers = -4 @@ -1275,7 +1280,7 @@ size=1 stride=1 pad=1 filters=255 -activation=linear +activation=logistic [yolo] @@ -1284,6 +1289,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4 classes=80 num=9 jitter=.1 +scale_x_y = 2.0 objectness_smooth=1 ignore_thresh = .7 truth_thresh = 1 @@ -1297,6 +1303,7 @@ iou_loss=ciou nms_kind=diounms beta_nms=0.6 new_coords=1 +max_delta=5 [route] layers = -4 @@ -1403,7 +1410,7 @@ size=1 stride=1 pad=1 filters=255 -activation=linear +activation=logistic [yolo] @@ -1412,6 +1419,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4 classes=80 num=9 jitter=.1 +scale_x_y = 2.0 objectness_smooth=1 ignore_thresh = .7 truth_thresh = 1 @@ -1425,3 +1433,4 @@ iou_loss=ciou nms_kind=diounms beta_nms=0.6 new_coords=1 +max_delta=2 \ No newline at end of file diff --git a/tests/darknet/yolo4-csp.cpp b/tests/darknet/yolo4-csp.cpp new file mode 100644 index 0000000..3802a9a --- /dev/null +++ b/tests/darknet/yolo4-csp.cpp @@ -0,0 +1,36 @@ +#include +#include +#include "tkdnn.h" +#include "test.h" +#include "DarknetParser.h" + +int main() { + std::string bin_path = "yolo4-csp"; + std::vector input_bins = { + bin_path + "/layers/input.bin" + }; + std::vector output_bins = { + bin_path + "/debug/layer144_out.bin", + bin_path + "/debug/layer159_out.bin", + bin_path + "/debug/layer174_out.bin" + }; + std::string wgs_path = bin_path + "/layers"; + std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4-csp.cfg"; + std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/AfzHE4BfTeEm2gH/download"); + + + + // parse darknet network + tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path); + net->print(); + + //convert network to tensorRT + tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str())); + + int ret = testInference(input_bins, output_bins, net, netRT); + net->releaseLayers(); + delete net; + delete netRT; + return ret; +} \ No newline at end of file diff --git a/tests/darknet/yolo4x.cpp b/tests/darknet/yolo4x.cpp index b9ad003..8df1aef 100644 --- a/tests/darknet/yolo4x.cpp +++ b/tests/darknet/yolo4x.cpp @@ -17,7 +17,7 @@ int main() { std::string wgs_path = bin_path + "/layers"; std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4x.cfg"; std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names"; - downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/BLPpiAigZJLorQD/download"); + downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download"); -- 2.52.0 From 7018d163ed2741bc1f095f9f129368beecdee567 Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Fri, 26 Mar 2021 19:37:44 +0530 Subject: [PATCH 15/22] python download file --- scripts/download_validation.py | 40 ++++++++++++++++++++++++++++++++++ 1 file changed, 40 insertions(+) create mode 100644 scripts/download_validation.py diff --git a/scripts/download_validation.py b/scripts/download_validation.py new file mode 100644 index 0000000..e49f2ab --- /dev/null +++ b/scripts/download_validation.py @@ -0,0 +1,40 @@ +import os +from pathlib import Path +import urllib.request as dowReq +import zipfile + +val = input("Enter BDD or COCO :") +if(val == "COCO"): + url = "https://cloud.hipert.unimore.it/s/LNxBDk4wzqXPL8c/download" + lib = "..\demo\COCO_val2017" + lib_zip = "COCO_val2017.zip" +elif(val == "BDD"): + url = "https://cloud.hipert.unimore.it/s/bikqk3FzCq2tg4D/download" + lib = "..\demo\BDD100k_val" + lib_zip = "BDD100k_val.zip" + +dowReq.urlretrieve(url,lib_zip) + +with zipfile.ZipFile(lib_zip,'r') as zip_ref: + zip_ref.extractall(lib) + +labelFolder = lib + "\labels" +imageFolder = lib + "\images" + +file1 = open(".\\..\\demo\\all_labels.txt","a") +path1 = os.path.realpath(labelFolder) +for file in os.listdir(labelFolder): + valTemp = path1 + "\\" + file + valTemp = valTemp + " \n" + file1.write(valTemp) +file1.close() + +file2 = open(".\\..\\demo\\all_images.txt","a") +path2 = os.path.realpath(imageFolder) +for file in os.listdir(imageFolder): + pathtemp = path2 + "\\" + file + pathtemp = pathtemp + " \n" + file2.write(pathtemp) +file2.close() + +print("Completed") \ No newline at end of file -- 2.52.0 From 5f3ab1472c8c02be43039abc273bcd6afc9f823d Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar Date: Fri, 26 Mar 2021 15:09:28 +0100 Subject: [PATCH 16/22] Update download_validation.py --- scripts/download_validation.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/scripts/download_validation.py b/scripts/download_validation.py index e49f2ab..e531b4f 100644 --- a/scripts/download_validation.py +++ b/scripts/download_validation.py @@ -1,5 +1,4 @@ import os -from pathlib import Path import urllib.request as dowReq import zipfile @@ -37,4 +36,4 @@ for file in os.listdir(imageFolder): file2.write(pathtemp) file2.close() -print("Completed") \ No newline at end of file +print("Completed") -- 2.52.0 From 2e92944f1da276ec31eeb08bd99cf7e78679b23c Mon Sep 17 00:00:00 2001 From: hchandirasekar Date: Wed, 31 Mar 2021 03:44:51 -0700 Subject: [PATCH 17/22] Opencv cuda fix --- demo/demo/demo.cpp | 2 +- include/tkDNN/DetectionNN.h | 4 +--- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index f97ead9..59bb1a2 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -25,7 +25,7 @@ int main(int argc, char *argv[]) { std::string net = "yolo4tiny_fp32.rt"; if(argc > 1) net = argv[1]; - std::string input = "../demo/yolo_test.mp4"; + std::string input = "..\..\..\demo\yolo_test.mp4"; if(argc > 2) input = argv[2]; char ntype = 'y'; diff --git a/include/tkDNN/DetectionNN.h b/include/tkDNN/DetectionNN.h index fb29be1..b1266e0 100644 --- a/include/tkDNN/DetectionNN.h +++ b/include/tkDNN/DetectionNN.h @@ -6,8 +6,6 @@ #include #ifdef __linux__ #include -#elif _WIN32 -#include #endif #include @@ -19,7 +17,7 @@ #include "tkdnn.h" -// #define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib. +#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib. #ifdef OPENCV_CUDACONTRIB #include -- 2.52.0 From cc594f09efdbf0e7c5a33e4e4c40ea24ac2e7048 Mon Sep 17 00:00:00 2001 From: perseusdg Date: Thu, 8 Apr 2021 00:28:00 +0530 Subject: [PATCH 18/22] merge from gitlab --- include/tkDNN/NetworkRT.h | 10 ++++------ include/tkDNN/utils.h | 14 -------------- src/NetworkRT.cpp | 37 ++++--------------------------------- src/utils.cpp | 3 +-- 4 files changed, 9 insertions(+), 55 deletions(-) diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index b39360c..66892f5 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -55,16 +55,14 @@ class NetworkRT { public: nvinfer1::DataType dtRT; - //nvinfer1::IBuilder *builderRT; - std::unique_ptr builderRT; - //nvinfer1::IRuntime *runtimeRT; - std::unique_ptr runtimeRT; + nvinfer1::IBuilder *builderRT; + nvinfer1::IRuntime *runtimeRT; nvinfer1::INetworkDefinition *networkRT; #if NV_TENSORRT_MAJOR >= 6 nvinfer1::IBuilderConfig *configRT; #endif - std::shared_ptr engineRT; - //nvinfer1::ICudaEngine *engineRT; + + nvinfer1::ICudaEngine *engineRT; nvinfer1::IExecutionContext *contextRT; const static int MAX_BUFFERS_RT = 10; diff --git a/include/tkDNN/utils.h b/include/tkDNN/utils.h index ce4a617..eeef3c2 100644 --- a/include/tkDNN/utils.h +++ b/include/tkDNN/utils.h @@ -14,9 +14,6 @@ #ifdef __linux__ #include -#elif _WIN32 -#define NOMINMAX -#include #endif #include @@ -110,17 +107,6 @@ double t_ns = time_ms.count(); FatalError(_error.str()); \ } \ } -struct InferDeleter -{ - template - void operator()(T* obj) const - { - if (obj) - { - obj->destroy(); - } - } -}; typedef enum { ERROR_CUDNN = 2, diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 1e0b062..b5005db 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -33,8 +33,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { float(NV_TENSORRT_PATCH)/100; std::cout<<"New NetworkRT (TensorRT v"<(createInferBuilder(loggerRT)); - //builderRT = createInferBuilder(loggerRT); + builderRT = createInferBuilder(loggerRT); std::cout<<"Float16 support: "<platformHasFastFp16()<<"\n"; std::cout<<"Int8 support: "<platformHasFastInt8()<<"\n"; #if NV_TENSORRT_MAJOR >= 5 @@ -138,8 +137,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { printCudaMemUsage(); std::cout<<"Building tensorRT cuda engine...\n"; #if NV_TENSORRT_MAJOR >= 6 - //engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); - engineRT = std::shared_ptr(builderRT->buildEngineWithConfig(*networkRT,*configRT),InferDeleter()); + engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); #else //engineRT = builderRT->buildCudaEngine(*networkRT); engineRT = std::shared_ptr(builderRT->buildCudaEngine(*networkRT)); @@ -637,10 +635,8 @@ bool NetworkRT::deserialize(const char *filename) { } pluginFactory = new PluginFactory(); - //runtimeRT = createInferRuntime(loggerRT); - runtimeRT = std::unique_ptr(createInferRuntime(loggerRT)); - //engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); - engineRT = std::shared_ptr(runtimeRT->deserializeCudaEngine(gieModelStream,size,(IPluginFactory*)pluginFactory),InferDeleter()); + runtimeRT = createInferRuntime(loggerRT); + engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory); //if (gieModelStream) delete [] gieModelStream; return true; @@ -673,7 +669,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa } if(name.find("ActivationCReLU") == 0) { float activationReluTemp = readBUF(buf); - //ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); ActivationReLUCeiling* a = new ActivationReLUCeiling(activationReluTemp); a->size = readBUF(buf); assert(buf == bufCheck + serialLength); @@ -684,9 +679,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa int classesTemp = readBUF(buf); int coordsTemp = readBUF(buf); int numTemp = readBUF(buf); - /*RegionRT *r = new RegionRT(readBUF(buf), //classes - readBUF(buf), //coords - readBUF(buf)); //num8*/ RegionRT* r = new RegionRT(classesTemp, coordsTemp, numTemp); r->c = readBUF(buf); @@ -698,7 +690,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa if(name.find("Reorg") == 0) { int strideTemp = readBUF(buf); - //ReorgRT *r = new ReorgRT(readBUF(buf)); //stride ReorgRT *r = new ReorgRT(strideTemp); r->c = readBUF(buf); r->h = readBUF(buf); @@ -723,15 +714,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa } if(name.find("Pooling") == 0) { - /* MaxPoolFixedSizeRT *r = new MaxPoolFixedSizeRT( readBUF(buf), //c - readBUF(buf), //h - readBUF(buf), //w - readBUF(buf), //n - readBUF(buf), //strideH - readBUF(buf), //strideW - readBUF(buf), //winSize - readBUF(buf)); //padding*/ - int cTemp = readBUF(buf); int hTemp = readBUF(buf); int wTemp = readBUF(buf); @@ -747,9 +729,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa } if(name.find("Resize") == 0) { - /*ResizeLayerRT *r = new ResizeLayerRT(readBUF(buf), //o_c - readBUF(buf), //o_h - readBUF(buf)); //o_w*/ int o_cTemp = readBUF(buf); int o_hTemp = readBUF(buf); int o_wTemp = readBUF(buf); @@ -822,7 +801,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa return r; } if(name.find("Upsample") == 0) { - //UpsampleRT *r = new UpsampleRT(readBUF(buf)); //stride int strideTemp = readBUF(buf); UpsampleRT* r = new UpsampleRT(strideTemp); r->c = readBUF(buf); @@ -833,7 +811,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa } if(name.find("Route") == 0) { - //RouteRT *r = new RouteRT(readBUF(buf),readBUF(buf)); int groupsTemp = readBUF(buf); int group_idTemp = readBUF(buf); RouteRT* r = new RouteRT(groupsTemp, group_idTemp); @@ -848,12 +825,6 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa } if(name.find("Deformable") == 0) { - /*DeformableConvRT *r = new DeformableConvRT(readBUF(buf), readBUF(buf), readBUF(buf), - readBUF(buf), readBUF(buf), readBUF(buf), - readBUF(buf), readBUF(buf), - readBUF(buf),readBUF(buf),readBUF(buf),readBUF(buf), - readBUF(buf),readBUF(buf),readBUF(buf),readBUF(buf), - nullptr); */ int chuck_dimTemp = readBUF(buf); int khTemp = readBUF(buf); int kwTemp = readBUF(buf); diff --git a/src/utils.cpp b/src/utils.cpp index 52775df..fa6458f 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -203,8 +203,7 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){ #ifdef __linux__ long page_size_kb = sysconf(_SC_PAGE_SIZE) / 1024; // in case x86-64 is configured to use 2MB pages #elif _WIN32 -SYSTEM_INFO sysInfo; - long page_size_kb = sysInfo.dwPageSize/1024; + long page_size_kb = 4096/1024; #endif vm_usage_kb = vsize / 1024.0; -- 2.52.0 From 37b2a5bd9856b8a6becb8ac696c95f70eaef20e9 Mon Sep 17 00:00:00 2001 From: perseusdg Date: Thu, 8 Apr 2021 00:57:58 +0530 Subject: [PATCH 19/22] minor modifications --- demo/demo/map.cpp | 2 -- include/tkDNN/ImuOdom.h | 1 - include/tkDNN/Int8BatchStream.h | 2 -- 3 files changed, 5 deletions(-) diff --git a/demo/demo/map.cpp b/demo/demo/map.cpp index 8e363ee..58490e5 100644 --- a/demo/demo/map.cpp +++ b/demo/demo/map.cpp @@ -4,8 +4,6 @@ #include /* srand, rand */ #ifdef __linux__ #include -#elif _WIN32 -#include #endif #include diff --git a/include/tkDNN/ImuOdom.h b/include/tkDNN/ImuOdom.h index fa870f3..d5429a8 100644 --- a/include/tkDNN/ImuOdom.h +++ b/include/tkDNN/ImuOdom.h @@ -7,7 +7,6 @@ #elif _WIN32 #define _USE_MATH_DEFINES #include -#include #endif #include diff --git a/include/tkDNN/Int8BatchStream.h b/include/tkDNN/Int8BatchStream.h index 7d2cef5..c39a11c 100644 --- a/include/tkDNN/Int8BatchStream.h +++ b/include/tkDNN/Int8BatchStream.h @@ -14,8 +14,6 @@ #include #ifdef __linux__ #include -#elif _WIN32 -#include #endif #include -- 2.52.0 From 1de804f98dd67e66e6893e0c64749befdb9b6371 Mon Sep 17 00:00:00 2001 From: perseusdg Date: Fri, 9 Apr 2021 13:17:41 +0530 Subject: [PATCH 20/22] Code cleanup and readme fixes --- CMakeLists.txt | 6 ++--- README.md | 39 +++++++++++++++++----------- demo/demo/demo.cpp | 7 ++++- include/tkDNN/DetectionNN.h | 2 +- include/tkDNN/pluginsRT/UpsampleRT.h | 2 -- include/tkDNN/pluginsRT/YoloRT.h | 1 - scripts/download_validation.py | 4 +-- src/NetworkRT.cpp | 2 +- 8 files changed, 37 insertions(+), 26 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 1b7ed63..c27e519 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,11 +3,11 @@ cmake_minimum_required(VERSION 3.5) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) if(UNIX) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declarations -Wno-unused-variable") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable -g ") endif() if(WIN32) -set(CMAKE_CXX_STANDARD 14) -set(CMAKE_CXX_FLAGS "/O1 /FS /EHsc") +set(CMAKE_CXX_STANDARD 11) +set(CMAKE_CXX_FLAGS "/O2 /FS /EHsc") set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) endif(WIN32) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN) diff --git a/README.md b/README.md index f685b58..cbc65b0 100644 --- a/README.md +++ b/README.md @@ -80,12 +80,13 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001 - [mAP demo](#map-demo) - [Existing tests and supported networks](#existing-tests-and-supported-networks) - [References](#references) - - [tkDNN on Windows 10 (experimental)](#tkdnn-on-windows) + - [tkDNN on Windows 10 (experimental)](#tkdnn-on-windows-10-experimental) - [Dependencies-Windows](#dependencies-windows) - - [Compiling tkDNN on Windows](#tkdnn-windows-compile) + - [Compiling tkDNN on Windows](#compiling-tkdnn-on-windows) - [Run the demo on Windows](#run-the-demo-on-windows) - - [FP16 interference windows](#fp16-windows) - - [INT8 interference windows](#int8-windows) + - [FP16 inference windows](#fp16-inference-windows) + - [INT8 inference windows](#int8-inference-windows) + - [Known issues with tkDNN on Windows](#known-issues-with-tkdnn-on-windows) @@ -362,26 +363,31 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing | yolo4tiny | Yolov4 tiny 9 | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download) | | yolo4x | Yolov4x-mish 9 | [COCO 2017](http://cocodataset.org/) | 80 | 672x672 | [weights](https://cloud.hipert.unimore.it/s/BLPpiAigZJLorQD/download) | -##tkDNN on Windows 10 (experimental) +### tkDNN on Windows 10 (experimental) ### Dependencies-Windows This branch should work on every NVIDIA GPU supported in windows with the following dependencies: * WINDOWS 10 1803 or HIGHER -* CUDA 10.0 (Recommended CUDA 11.0 +) -* CUDNN 7.6 (Recommended CUDNN 8.0.0 +) -* TENSORRT 6.0.1 (Recommended TENSORRT 7.1 +) -* OPENCV 3.4 (Recommended OPENCV 4.2.0 +) -* MSVC 16.7 (Recommended MSVC 16.8/16.9) -* YAML-CPP 0.5.2 +* CUDA 10.0 (Recommended CUDA 11.2 ) +* CUDNN 7.6 (Recommended CUDNN 8.1.1 ) +* TENSORRT 6.0.1 (Recommended TENSORRT 7.2.3.4 ) +* OPENCV 3.4 (Recommended OPENCV 4.2.0 ) +* MSVC 16.7 +* YAML-CPP * EIGEN3 * 7ZIP (ADD TO PATH) * NINJA 1.10 + All the above mentioned dependencies except 7ZIP can be installed using Microsoft's [VCPKG](https://github.com/microsoft/vcpkg.git) . After bootstrapping VCPKG the dependencies can be built and installed using the following command : -```vcpkg.exe install opencv4[tbb,jpeg,tiff,opengl,openmp,png,ffmpeg]:x64-windows yaml-cpp:x64-windows eigen3:x64-windows --x-install-root=C:\opt --x-buildtrees-root=C:\temp_vcpkg_build``` +``` +opencv4(normal) - vcpkg.exe install opencv4[tbb,jpeg,tiff,opengl,openmp,png,ffmpeg,eigen]:x64-windows yaml-cpp:x64-windows eigen3:x64-windows --x-install-root=C:\opt --x-buildtrees-root=C:\temp_vcpkg_build + +opencv4(cuda) - vcpkg.exe install opencv4[cuda,nonfree,contrib,eigen,tbb,jpeg,tiff,opengl,openmp,png,ffmpeg]:x64-windows yaml-cpp:x64-windows eigen3:x64-windows --x-install-root=C:\opt --x-buildtrees-root=C:\temp_vcpkg_build +``` After VCPKG finishes building and installing all the packages delete C:\temp_vcpkg_build and add C:\opt\x64-windows\bin and C:\opt\x64-windows\debug\bin to path @@ -411,7 +417,7 @@ Once the rt file has been successfully create,run the demo using the following c ``` For general info on more demo paramters,check Run the demo section on top -### FP16 interference windows +### FP16 inference windows This is an untested feature on windows.To run the object detection demo with FP16 interference follow the below steps(example with yolo4tiny): ``` @@ -421,7 +427,7 @@ del /f yolo4tiny_fp16.rt .\demo.exe yolo4tiny_fp16.rt ..\demo\yolo_test.mp4 ``` -### INT8 interference windows +### INT8 inference windows To run object detection demo with INT8 (example with yolo4tiny): ``` set TKDNN_MODE=INT8 @@ -433,10 +439,13 @@ del /f yolo4tiny_int8.rt # be sure to delete(or move) old tensorRT files ``` +### Known issues with tkDNN on Windows +Mobilenet and Centernet demos work properly only when built with msvc 16.7 in Release Mode,when built in debug mode for the mentioned networks one might encounter opencv assert errors +All Darknet models work properly with demo using MSVC version(16.7-16.9) - +It is recommended to use Nvidia Driver(465+),Cuda unknown errors have been observed when using older drivers on pascal(SM 61) devices. diff --git a/demo/demo/demo.cpp b/demo/demo/demo.cpp index 59bb1a2..317a574 100644 --- a/demo/demo/demo.cpp +++ b/demo/demo/demo.cpp @@ -25,7 +25,12 @@ int main(int argc, char *argv[]) { std::string net = "yolo4tiny_fp32.rt"; if(argc > 1) net = argv[1]; - std::string input = "..\..\..\demo\yolo_test.mp4"; + #ifdef __linux__ + std::string input = "../demo/yolo_test.mp4"; + #elif _WIN32 + std::string input = "..\\..\\..\\demo\\yolo_test.mp4"; + #endif + if(argc > 2) input = argv[2]; char ntype = 'y'; diff --git a/include/tkDNN/DetectionNN.h b/include/tkDNN/DetectionNN.h index b1266e0..a8c81f7 100644 --- a/include/tkDNN/DetectionNN.h +++ b/include/tkDNN/DetectionNN.h @@ -17,7 +17,7 @@ #include "tkdnn.h" -#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib. +//#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib. #ifdef OPENCV_CUDACONTRIB #include diff --git a/include/tkDNN/pluginsRT/UpsampleRT.h b/include/tkDNN/pluginsRT/UpsampleRT.h index 5350b7e..a11d7b4 100644 --- a/include/tkDNN/pluginsRT/UpsampleRT.h +++ b/include/tkDNN/pluginsRT/UpsampleRT.h @@ -59,8 +59,6 @@ public: tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); - std::cout << "Upsample Serialization SIze" << getSerializationSize() << std::endl; - assert(buf == a + getSerializationSize()); } diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index 0dd26e1..2911869 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -120,7 +120,6 @@ public: tk::dnn::writeBUF(buf, tmp[j]); } } - std::cout << getSerializationSize() << std::endl; assert(buf == a + getSerializationSize()); } diff --git a/scripts/download_validation.py b/scripts/download_validation.py index e531b4f..0e3b1d4 100644 --- a/scripts/download_validation.py +++ b/scripts/download_validation.py @@ -24,7 +24,7 @@ file1 = open(".\\..\\demo\\all_labels.txt","a") path1 = os.path.realpath(labelFolder) for file in os.listdir(labelFolder): valTemp = path1 + "\\" + file - valTemp = valTemp + " \n" + valTemp = valTemp + '\n' file1.write(valTemp) file1.close() @@ -32,7 +32,7 @@ file2 = open(".\\..\\demo\\all_images.txt","a") path2 = os.path.realpath(imageFolder) for file in os.listdir(imageFolder): pathtemp = path2 + "\\" + file - pathtemp = pathtemp + " \n" + pathtemp = pathtemp + '\n' file2.write(pathtemp) file2.close() diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index b5005db..5dc8ee0 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -648,7 +648,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa const char * buf = reinterpret_cast(serialData),*bufCheck = buf; std::string name(layerName); - std::cout< Date: Fri, 9 Apr 2021 13:20:44 +0530 Subject: [PATCH 21/22] Update CMakeLists.txt --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c27e519..d3a89f5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,7 +3,7 @@ cmake_minimum_required(VERSION 3.5) project (tkDNN) set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake) if(UNIX) -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable -g ") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable ") endif() if(WIN32) set(CMAKE_CXX_STANDARD 11) -- 2.52.0 From 39323ca8d3ac02a8a1e74ede6c851e57546f18ab Mon Sep 17 00:00:00 2001 From: Harshvardhan Chandirasekar <43143075+perseusdg@users.noreply.github.com> Date: Wed, 14 Apr 2021 17:22:26 +0530 Subject: [PATCH 22/22] Update README.md --- README.md | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index cbc65b0..7fb389b 100644 --- a/README.md +++ b/README.md @@ -388,6 +388,7 @@ opencv4(normal) - vcpkg.exe install opencv4[tbb,jpeg,tiff,opengl,openmp,png,ffmp opencv4(cuda) - vcpkg.exe install opencv4[cuda,nonfree,contrib,eigen,tbb,jpeg,tiff,opengl,openmp,png,ffmpeg]:x64-windows yaml-cpp:x64-windows eigen3:x64-windows --x-install-root=C:\opt --x-buildtrees-root=C:\temp_vcpkg_build ``` +To build opencv4 with cuda and cudnn version corresponding to your cuda version,vcpkg's cudnn portfile needs to be modified by adding ```$ENV{CUDA_PATH}``` at lines 16 and 17 in the portfile.cmake After VCPKG finishes building and installing all the packages delete C:\temp_vcpkg_build and add C:\opt\x64-windows\bin and C:\opt\x64-windows\debug\bin to path @@ -395,7 +396,7 @@ After VCPKG finishes building and installing all the packages delete C:\temp_vcp tkDNN is built with cmake(3.15+) on windows along with ninja.Msbuild and NMake Makefiles are drastically slower when compiling the library compared to windows ``` -git clone https://git.hipert.unimore.it/research-cv-chandirasekar/tkdnn-windows.git +git clone https://github.com/ceccocats/tkDNN.git cd tkdnn-windows mkdir build cd build @@ -416,6 +417,7 @@ Once the rt file has been successfully create,run the demo using the following c .\demo.exe yolo4tiny_fp32.rt ..\demo\yolo_test.mp4 y ``` For general info on more demo paramters,check Run the demo section on top + To run the test_all_tests.sh on windows,use git bash or msys2 ### FP16 inference windows -- 2.52.0