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