mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Merge with master, add yolov4tiny_512

Browse Source 
Signed-off-by: tk <micaelaverucchi@gmail.com>
This commit is contained in:
tk
2020-07-06 17:57:50 +02:00
parent 8d08f5aade 7c2155decf
commit a68f45cb4e
38 changed files with 1328 additions and 361 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+3
View File
@@ -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")
README.md
+30 -2
View File
@@ -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)<sup>4</sup> | [COCO 2017](http://cocodataset.org/) | 80 | 512x512 | [weights](https://cloud.hipert.unimore.it/s/5BTjHMWBcJk8g3i/download) |
| csresnext50-panet-spp | Cross Stage Partial Network <sup>7</sup> | [COCO 2014](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/Kcs4xBozwY4wFx8/download) |
| yolo4 | Yolov4 <sup>8</sup> | [COCO 2017](http://cocodataset.org/) | 80 | 416x416 | [weights](https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download) |
| yolo4_berkeley | Yolov4 <sup>8</sup> | [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
cmake/FindCUDNN.cmake
+62 -29
View File
@@ -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)
include/tkDNN/DarknetParser.h
+15 -260
View File
@@ -1,6 +1,6 @@
#pragma once
#include <iostream>
#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<int> 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<int> fromStringToIntVec(const std::string& line, const char delimiter){
std::stringstream linestream(line);
std::string value;
std::vector<int> 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: "<<line<<std::endl;
return true;
}
tk::dnn::Network *darknetAddNet(darknetFields_t &fields) {
//std::cout<<"Add Net: "<<fields.type<<"\n";
dataDim_t dim(1, fields.channels, fields.height, fields.width);
return new tk::dnn::Network(dim);
}
void darknetAddLayer(tk::dnn::Network *net, darknetFields_t &f, std::string wgs_path, std::vector<tk::dnn::Layer*> &netLayers, const std::vector<std::string>& 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: "<<f.type<<"\n";
if(f.type == "convolutional") {
std::string wgs = wgs_path + "/c" + std::to_string(netLayers.size()) + ".bin";
//printf("%d (%d,%d) (%d,%d) (%d,%d) %s %d %d\n", f.filters, f.size_x, f.size_y, f.stride_x, f.stride_y, f.padding_x, f.padding_y, wgs.c_str(), f.batch_normalize, f.groups);
tk::dnn::Conv2d *l= new tk::dnn::Conv2d(net, f.filters, f.size_x, f.size_y, f.stride_x,
f.stride_y, f.padding_x, f.padding_y, wgs, f.batch_normalize, false, f.groups);
netLayers.push_back(l);
} else if(f.type == "maxpool") {
if(f.stride_x == 1 && f.stride_y == 1)
netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y,
f.padding_x, f.padding_y, tk::dnn::POOLING_MAX_FIXEDSIZE));
else
netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y,
f.padding_x, f.padding_y, tk::dnn::POOLING_MAX));
} else if(f.type == "avgpool") {
netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y,
f.padding_x, f.padding_y, tk::dnn::POOLING_AVERAGE));
} else if(f.type == "shortcut") {
if(f.layers.size() != 1) FatalError("no layers to shortcut\n");
int layerIdx = f.layers[0];
if(layerIdx < 0)
layerIdx = netLayers.size() + layerIdx;
if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to shortcut\n");
//std::cout<<"shortcut to "<<layerIdx<<" "<<netLayers[layerIdx]->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<tk::dnn::Layer*> layers;
for(int i=0; i<f.layers.size(); i++) {
int layerIdx = f.layers[i];
if(layerIdx < 0)
layerIdx = netLayers.size() + layerIdx;
if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to route\n");
//std::cout<<"Route to "<<layerIdx<<" "<<netLayers[layerIdx]->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<std::string> 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<std::string> 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<tk::dnn::Layer*> netLayers;
std::ifstream if_cfg(cfg_file);
if(!if_cfg.is_open())
FatalError("cloud not open cfg file: " + cfg_file);
std::vector<std::string> 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: "<<type<<"\n";
fields = darknetFields_t(); // reset to default
fields.type = type;
continue;
}
if(darknetParseFields(line, fields)) {
// already parsed do nothing
} else {
FatalError("could not parse line: " + line);
}
}
// end of filled type
if(fields.type != "") {
darknetAddLayer(net, fields, wgs_path, netLayers, names);
}
if(net == nullptr) {
FatalError("net not found\n");
}
return net;
}
std::string darknetParseType(const std::string& line);
bool divideNameAndValue(const std::string& line, std::string&name, std::string& value);
std::vector<int> 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<tk::dnn::Layer*> &netLayers, const std::vector<std::string>& names);
std::vector<std::string> 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);
}}
include/tkDNN/Layer.h
+3 -1
View File
@@ -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;
};
include/tkDNN/NetworkRT.h
+1 -1
View File
@@ -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"
include/tkDNN/NetworkViz.h
+12
View File
@@ -0,0 +1,12 @@
#pragma once
#include <iostream>
#include <opencv2/core/types.hpp>
#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);
}}
include/tkDNN/pluginsRT/RouteRT.h
+19 -10
View File
@@ -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<nbInputDims; i++) out_c += inputs[i].d[0];
return DimsCHW{out_c, inputs[0].d[1], inputs[0].d[2]};
return DimsCHW{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
@@ -34,6 +36,7 @@ public:
}
h = inputDims[0].d[1];
w = inputDims[0].d[2];
c /= groups;
}
int initialize() override {
@@ -49,15 +52,18 @@ public:
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
int offset = 0;
for(int i=0; i<in; i++) {
dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(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<batchSize; b++) {
int offset = 0;
for(int i=0; i<in; i++) {
dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(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<char*>(buffer);
tk::dnn::writeBUF(buf, groups);
tk::dnn::writeBUF(buf, group_id);
tk::dnn::writeBUF(buf, in);
for(int i=0; i<MAX_INPUTS; i++)
tk::dnn::writeBUF(buf, c_in[i]);
@@ -83,4 +91,5 @@ public:
int in;
int c_in[MAX_INPUTS];
int c, h, w;
int groups, group_id;
};
include/tkDNN/test.h
+8 -4
View File
@@ -1,7 +1,7 @@
#include <tkdnn.h>
int testInference(std::vector<std::string> input_bins, std::vector<std::string> output_bins,
tk::dnn::Network *net, tk::dnn::NetworkRT *netRT = nullptr) {
tk::dnn::Network *net, tk::dnn::NetworkRT *netRT = nullptr) {
std::vector<tk::dnn::Layer*> outputs;
for(int i=0; i<net->num_layers; i++) {
@@ -67,7 +67,11 @@ int testInference(std::vector<std::string> input_bins, std::vector<std::string>
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;
}
delete [] out_h;
checkCuda( cudaFree(out) );
}
delete [] input_h;
checkCuda( cudaFree(data) );
return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
}
include/tkDNN/utils.h
+4
View File
@@ -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
scripts/test_all_tests.sh
+1
View File
@@ -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
src/DarknetParser.cpp
+263
View File
@@ -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<int> fromStringToIntVec(const std::string& line, const char delimiter){
std::stringstream linestream(line);
std::string value;
std::vector<int> 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: "<<line<<std::endl;
return true;
}
tk::dnn::Network *darknetAddNet(darknetFields_t &fields) {
//std::cout<<"Add Net: "<<fields.type<<"\n";
dataDim_t dim(1, fields.channels, fields.height, fields.width);
return new tk::dnn::Network(dim);
}
void darknetAddLayer(tk::dnn::Network *net, darknetFields_t &f, std::string wgs_path, std::vector<tk::dnn::Layer*> &netLayers, const std::vector<std::string>& 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: "<<f.type<<"\n";
if(f.type == "convolutional") {
std::string wgs = wgs_path + "/c" + std::to_string(netLayers.size()) + ".bin";
//printf("%d (%d,%d) (%d,%d) (%d,%d) %s %d %d\n", f.filters, f.size_x, f.size_y, f.stride_x, f.stride_y, f.padding_x, f.padding_y, wgs.c_str(), f.batch_normalize, f.groups);
tk::dnn::Conv2d *l= new tk::dnn::Conv2d(net, f.filters, f.size_x, f.size_y, f.stride_x,
f.stride_y, f.padding_x, f.padding_y, wgs, f.batch_normalize, false, f.groups);
netLayers.push_back(l);
} else if(f.type == "maxpool") {
if(f.stride_x == 1 && f.stride_y == 1)
netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y,
f.padding_x, f.padding_y, tk::dnn::POOLING_MAX_FIXEDSIZE));
else
netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y,
f.padding_x, f.padding_y, tk::dnn::POOLING_MAX));
} else if(f.type == "avgpool") {
netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y,
f.padding_x, f.padding_y, tk::dnn::POOLING_AVERAGE));
} else if(f.type == "shortcut") {
if(f.layers.size() != 1) FatalError("no layers to shortcut\n");
int layerIdx = f.layers[0];
if(layerIdx < 0)
layerIdx = netLayers.size() + layerIdx;
if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to shortcut\n");
//std::cout<<"shortcut to "<<layerIdx<<" "<<netLayers[layerIdx]->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<tk::dnn::Layer*> layers;
for(int i=0; i<f.layers.size(); i++) {
int layerIdx = f.layers[i];
if(layerIdx < 0)
layerIdx = netLayers.size() + layerIdx;
if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to route\n");
//std::cout<<"Route to "<<layerIdx<<" "<<netLayers[layerIdx]->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<std::string> 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<std::string> 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<tk::dnn::Layer*> netLayers;
std::ifstream if_cfg(cfg_file);
if(!if_cfg.is_open())
FatalError("cloud not open cfg file: " + cfg_file);
std::vector<std::string> 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: "<<type<<"\n";
fields = darknetFields_t(); // reset to default
fields.type = type;
continue;
}
if(darknetParseFields(line, fields)) {
// already parsed do nothing
} else {
FatalError("could not parse line: " + line);
}
}
// end of filled type
if(fields.type != "") {
darknetAddLayer(net, fields, wgs_path, netLayers, names);
}
if(net == nullptr) {
FatalError("net not found\n");
}
return net;
}
}}
src/Int8BatchStream.cpp
+6 -13
View File
@@ -132,21 +132,14 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector<float>& res
void BatchStream::readLabels(std::string inputFileName, std::vector<float>& 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;
}
}
}
src/NetworkRT.cpp
+12 -9
View File
@@ -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<int>(buf);
return r;
}
/*
if(name.find("Route") == 0) {
RouteRT *r = new RouteRT();
RouteRT *r = new RouteRT(readBUF<int>(buf),readBUF<int>(buf));
r->in = readBUF<int>(buf);
for(int i=0; i<RouteRT::MAX_INPUTS; i++)
r->c_in[i] = readBUF<int>(buf);
@@ -777,7 +780,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
r->w = readBUF<int>(buf);
return r;
}
*/
if(name.find("Deformable") == 0) {
DeformableConvRT *r = new DeformableConvRT(readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
src/NetworkViz.cpp
+69
View File
@@ -0,0 +1,69 @@
#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/videoio.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#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<dim.c;i++) {
cv::Mat raw = vizFloat2colorMap(cv::Mat(cv::Size(dim.w, dim.h),CV_32FC1, data + dim.w*dim.h*i));
int r = i / gridDim;
int c = i - r * gridDim;
raw.copyTo(grid.rowRange(r*dim.h, r*dim.h + dim.h).colRange(c*dim.w, c*dim.w + dim.w));
}
float ar = float(dim.w)/dim.h;
cv::Size vdim(ar*imgdim, imgdim);
cv::Mat viz;
cv::resize(grid, viz, vdim, 0, 0, 0);
// free memory
if(isCudaPointer(dataInput)) {
delete [] data;
}
return viz;
}
cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim) {
if(layer >= 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);
}
}}
src/Route.cpp
+7 -3
View File
@@ -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; i<layers_n; i++) {
dnnType *input = layers[i]->dstData;
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
tests/centernet/dla34_cnet/dla34_cnet.cpp
+12
View File
@@ -479,6 +479,18 @@ int main()
//print network model
net.print();
// for(int i=0; i<net.num_layers; i++) {
// if(net.layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net.layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net.layers[i]->dstData != nullptr) {
// cudaFree(net.layers[i]->dstData);
// net.layers[i]->dstData = nullptr;
// }
// }
//convert network to tensorRT
tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("dla34_cnet"));
tests/centernet/resnet101_cnet/resnet101_cnet.cpp
+12
View File
@@ -353,6 +353,18 @@ int main()
//print network model
net.print();
// for(int i=0; i<net.num_layers; i++) {
// if(net.layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net.layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net.layers[i]->dstData != nullptr) {
// cudaFree(net.layers[i]->dstData);
// net.layers[i]->dstData = nullptr;
// }
// }
//convert network to tensorRT
tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("resnet101_cnet"));
tests/darknet/cfg/yolo4tiny.cfg
+281
View File
@@ -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
tests/darknet/cfg/yolo4tiny_512.cfg
+281
View File
@@ -0,0 +1,281 @@
[net]
# Testing
#batch=1
#subdivisions=1
# Training
batch=64
subdivisions=1
width=512
height=512
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
tests/darknet/csresnext50-panet-spp.cpp
+2 -2
View File
@@ -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
tests/darknet/csresnext50-panet-spp_berkeley.cpp
+3 -4
View File
@@ -15,10 +15,9 @@ 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";
// FIXME: wrong weights
// downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s//download");
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";
downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/q82qHAtqpoaFYo5/download");
// parse darknet network
tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
tests/darknet/viz_yolo3.cpp
+70
View File
@@ -0,0 +1,70 @@
#include<iostream>
#include<vector>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#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; i<net->input_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; i<net->num_layers; i++) {
std::string output_png = output_viz + "/layer" + std::to_string(i) + ".png";
std::cout<<"saving "<<output_png<<"\n";
cv::Mat viz = vizLayer2Mat(net, i);
cv::imwrite(output_png, viz);
//cv::imshow("layer", viz);
//cv::waitKey(0);
}
checkCuda(cudaFree(input_d));
net->releaseLayers();
delete net;
return 0;
}
tests/darknet/yolo2.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo2_voc.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo2tiny.cpp
+2 -2
View File
@@ -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");
tests/darknet/yolo3.cpp
+3 -3
View File
@@ -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;
}
}
tests/darknet/yolo3_512.cpp
+14 -2
View File
@@ -15,14 +15,26 @@ 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
tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
net->print();
// for(int i=0; i<net->num_layers; i++) {
// if(net->layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net->layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net->layers[i]->dstData != nullptr) {
// cudaFree(net->layers[i]->dstData);
// net->layers[i]->dstData = nullptr;
// }
// }
//convert network to tensorRT
tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str()));
tests/darknet/yolo3_berkeley.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo3_coco4.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo3_flir.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo3tiny.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo3tiny_512.cpp
+15 -2
View File
@@ -14,14 +14,27 @@ 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
tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
net->print();
// for(int i=0; i<net->num_layers; i++) {
// if(net->layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net->layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net->layers[i]->dstData != nullptr) {
// cudaFree(net->layers[i]->dstData);
// net->layers[i]->dstData = nullptr;
// }
// }
//convert network to tensorRT
tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str()));
tests/darknet/yolo4_512.cpp
+13
View File
@@ -23,6 +23,19 @@ int main() {
tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
net->print();
// for(int i=0; i<net->num_layers; i++) {
// if(net->layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net->layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net->layers[i]->dstData != nullptr) {
// cudaFree(net->layers[i]->dstData);
// net->layers[i]->dstData = nullptr;
// }
// }
//convert network to tensorRT
tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str()));
tests/darknet/yolo4_berkeley.cpp
+2 -2
View File
@@ -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
tests/darknet/yolo4tiny.cpp
+33
View File
@@ -0,0 +1,33 @@
#include<iostream>
#include<vector>
#include "tkdnn.h"
#include "test.h"
#include "DarknetParser.h"
int main() {
std::string bin_path = "yolo4tiny";
std::vector<std::string> input_bins = {
bin_path + "/layers/input.bin"
};
std::vector<std::string> 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;
}
tests/darknet/yolo4tiny_512.cpp
+45
View File
@@ -0,0 +1,45 @@
#include<iostream>
#include<vector>
#include "tkdnn.h"
#include "test.h"
#include "DarknetParser.h"
int main() {
std::string bin_path = "yolo4tiny_512";
std::vector<std::string> input_bins = {
bin_path + "/layers/input.bin"
};
std::vector<std::string> 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_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/qa2ws4GXg7mS5nN/download");
// parse darknet network
tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
net->print();
// for(int i=0; i<net->num_layers; i++) {
// if(net->layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net->layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net->layers[i]->dstData != nullptr) {
// cudaFree(net->layers[i]->dstData);
// net->layers[i]->dstData = nullptr;
// }
// }
//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;
}
tests/mobilenet/mobilenetv2ssd512/mobilenetv2ssd512.cpp
+13
View File
@@ -469,6 +469,19 @@ int main()
//print network model
net.print();
// for(int i=0; i<net.num_layers; i++) {
// if(net.layers[i]->getLayerType() == tk::dnn::LAYER_CONV2D) {
// tk::dnn::Conv2d *c = (tk::dnn::Conv2d*) net.layers[i];
// c->releaseDevice();
// c->releaseHost(true, false);
// }
// if(net.layers[i]->dstData != nullptr) {
// cudaFree(net.layers[i]->dstData);
// net.layers[i]->dstData = nullptr;
// }
// }
// convert network to tensorRT
tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("mobilenetv2ssd512"));