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

Merge remote-tracking branch 'origin/master' into cnet

Browse Source 
Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
Davide Sapienza
2020-12-07 17:45:44 +01:00
parent 3d2405323b b8855b9599
commit 7f65ee0b2a
138 changed files with 6667 additions and 5411 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/CenternetDetection.cpp
+20 -17
View File
@@ -3,10 +3,12 @@
namespace tk { namespace dnn {
bool CenternetDetection::init(const std::string& tensor_path, const int n_classes){
bool CenternetDetection::init(const std::string& tensor_path, const int n_classes, const int n_batches, const float conf_thresh){
std::cout<<(tensor_path).c_str()<<"\n";
netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str() );
classes = n_classes;
nBatches = n_batches;
confThreshold = conf_thresh;
dim = netRT->input_dim;
@@ -41,7 +43,7 @@ bool CenternetDetection::init(const std::string& tensor_path, const int n_classe
trans = cv::Mat(cv::Size(3,2), CV_32F);
trans2 = cv::Mat(cv::Size(3,2), CV_32F);
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot() * nBatches));
dim_hm = tk::dnn::dataDim_t(1, 80, 128, 128, 1);
dim_wh = tk::dnn::dataDim_t(1, 2, 128, 128, 1);
@@ -98,7 +100,7 @@ bool CenternetDetection::init(const std::string& tensor_path, const int n_classe
checkCuda(cudaMemcpy(mean_d, mean, 3*sizeof(float), cudaMemcpyHostToDevice));
checkCuda(cudaMemcpy(stddev_d, stddev, 3*sizeof(float), cudaMemcpyHostToDevice));
#else
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()* nBatches));
mean << 0.408, 0.447, 0.47;
stddev << 0.289, 0.274, 0.278;
#endif
@@ -120,13 +122,13 @@ bool CenternetDetection::init(const std::string& tensor_path, const int n_classe
}
void CenternetDetection::preprocess(cv::Mat &frame){
void CenternetDetection::preprocess(cv::Mat &frame, const int bi){
// -----------------------------------pre-process ------------------------------------------
// auto start_t = std::chrono::steady_clock::now();
// auto step_t = std::chrono::steady_clock::now();
// auto end_t = std::chrono::steady_clock::now();
cv::Size sz = originalSize;
cv::Size sz = originalSize[bi];
// std::cout<<"image: "<<sz.width<<", "<<sz.height<<std::endl;
cv::Size sz_old;
float scale = 1.0;
@@ -212,7 +214,7 @@ void CenternetDetection::preprocess(cv::Mat &frame){
// std::cout << " TIME normalize: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
// step_t = end_t;
checkCuda(cudaMemcpy(input_d, d_ptrs, dim2.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda(cudaMemcpy(input_d+ netRT->input_dim.tot()*bi, d_ptrs, dim2.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// end_t = std::chrono::steady_clock::now();
// std::cout << " TIME Memcpy to input_d: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
@@ -254,18 +256,18 @@ void CenternetDetection::preprocess(cv::Mat &frame){
int idx = i*imageF.rows*imageF.cols;
int ch = dim2.c-3 +i;
// std::cout<<"i: "<<i<<", idx: "<<idx<<", ch: "<<ch<<std::endl;
memcpy((void*)&input[idx], (void*)bgr[ch].data, imageF.rows*imageF.cols*sizeof(dnnType));
memcpy((void*)&input[idx+ netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imageF.rows*imageF.cols*sizeof(dnnType));
}
checkCuda(cudaMemcpyAsync(input_d, input, dim2.tot()*sizeof(dnnType), cudaMemcpyHostToDevice));
checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input+ netRT->input_dim.tot()*bi, dim2.tot()*sizeof(dnnType), cudaMemcpyHostToDevice));
#endif
}
void CenternetDetection::postprocess(){
void CenternetDetection::postprocess(const int bi, const bool mAP){
dnnType *rt_out[4];
rt_out[0] = (dnnType *)netRT->buffersRT[1];
rt_out[1] = (dnnType *)netRT->buffersRT[2];
rt_out[2] = (dnnType *)netRT->buffersRT[3];
rt_out[3] = (dnnType *)netRT->buffersRT[4];
rt_out[0] = (dnnType *)netRT->buffersRT[1]+ netRT->buffersDIM[1].tot()*bi;
rt_out[1] = (dnnType *)netRT->buffersRT[2]+ netRT->buffersDIM[2].tot()*bi;
rt_out[2] = (dnnType *)netRT->buffersRT[3]+ netRT->buffersDIM[3].tot()*bi;
rt_out[3] = (dnnType *)netRT->buffersRT[4]+ netRT->buffersDIM[4].tot()*bi;
// auto start_t = std::chrono::steady_clock::now();
// auto step_t = std::chrono::steady_clock::now();
@@ -370,10 +372,10 @@ void CenternetDetection::postprocess(){
// std::cout<<"th: "<<scores[j]<<" - cl: "<<clses[j]<<" i: "<<i<<std::endl;
//add coco bbox
//det[0:4], i, det[4]
int x0 = target_coords[j*4];
int y0 = target_coords[j*4+1];
int x1 = target_coords[j*4+2];
int y1 = target_coords[j*4+3];
float x0 = target_coords[j*4];
float y0 = target_coords[j*4+1];
float x1 = target_coords[j*4+2];
float y1 = target_coords[j*4+3];
int obj_class = clses[j];
float prob = scores[j];
// std::cout<<"("<<x0<<", "<<y0<<"),("<<x1<<", "<<y1<<")"<<std::endl;
@@ -389,6 +391,7 @@ void CenternetDetection::postprocess(){
}
}
batchDetected.push_back(detected);
// end_t = std::chrono::steady_clock::now();
// std::cout << " TIME detections: " << std::chrono::duration_cast<std::chrono:: microseconds>(end_t - step_t).count() << " us" << std::endl;
// step_t = end_t;
src/Conv2d.cpp
+18 -13
View File
@@ -62,25 +62,30 @@ void Conv2d::initCUDNN(bool back) {
// init workspace
workSpace = NULL;
ws_sizeInBytes = 0;
int algo_count = 0;
if(back) {
checkCUDNN( cudnnGetConvolutionBackwardDataAlgorithm(net->cudnnHandle,
filterDesc, dstTensor, convDesc, srcTensor,
CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST, 0, &bwAlgo) );
checkCUDNN( cudnnGetConvolutionBackwardDataAlgorithm_v7(net->cudnnHandle,
filterDesc, dstTensor, convDesc, srcTensor, 1, &algo_count, &bwAlgo) );
checkCUDNN(cudnnGetConvolutionBackwardDataWorkspaceSize(net->cudnnHandle,
filterDesc, dstTensor, convDesc, srcTensor,
bwAlgo, &ws_sizeInBytes));
filterDesc, dstTensor, convDesc, srcTensor,
bwAlgo.algo, &ws_sizeInBytes));
// invert tensors
srcTensorDesc = dstTensor;
dstTensorDesc = srcTensor;
} else {
checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle,
srcTensor, filterDesc, convDesc, dstTensor,
CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo) );
checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle,
srcTensor, filterDesc, convDesc, dstTensor,
algo, &ws_sizeInBytes));
checkCUDNN( cudnnGetConvolutionForwardAlgorithm_v7(net->cudnnHandle,
srcTensor, filterDesc, convDesc, dstTensor,
1, &algo_count, &algo) );
checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle,
srcTensor, filterDesc, convDesc, dstTensor,
algo.algo, &ws_sizeInBytes));
}
if(algo_count < 1)
FatalError("Cannot retrieve convolutional algo");
}
void Conv2d::inferCUDNN(dnnType* srcData, bool back) {
@@ -91,12 +96,12 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) {
checkCUDNN(cudnnConvolutionBackwardData(net->cudnnHandle,
&alpha, filterDesc, data_d,
srcTensorDesc, srcData,
convDesc, bwAlgo, workSpace, ws_sizeInBytes,
convDesc, bwAlgo.algo, workSpace, ws_sizeInBytes,
&beta, dstTensorDesc, dstData));
} else {
checkCUDNN(cudnnConvolutionForward(net->cudnnHandle,
&alpha, srcTensorDesc, srcData, filterDesc,
data_d, convDesc, algo, workSpace, ws_sizeInBytes,
data_d, convDesc, algo.algo, workSpace, ws_sizeInBytes,
&beta, dstTensorDesc, dstData));
}
src/DarknetParser.cpp
+273
View File
@@ -0,0 +1,273 @@
#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("new_coords") != std::string::npos)
fields.new_coords = std::stoi(value);
else 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("beta_nms") != std::string::npos)
fields.nms_thresh = std::stof(value);
else if(name.find("nms_kind") != std::string::npos){
if(value == "greedynms") fields.nms_kind = 0;
else if(value == "diounms") fields.nms_kind = 1;
else std::cout<<"Not supported nms_kind "<<value<<", setting to greedynms"<<std::endl;
}
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, f.nms_thresh, (tk::dnn::Yolo::nmsKind_t) f.nms_kind, f.new_coords);
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 retrieve 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/DeformConv2d.cpp
+1 -1
View File
@@ -95,7 +95,7 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
// split conv2d outputs into offset and mask
checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// kernel sigmoide
// kernel sigmoid
activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution
src/Dense.cpp
+1 -1
View File
@@ -37,7 +37,7 @@ dnnType* Dense::infer(dataDim_t &dim, dnnType* srcData) {
// place bias into dstData
checkCuda( cudaMemcpy(dstData, bias_d, dim_y*sizeof(dnnType), cudaMemcpyDeviceToDevice) );
//do matrix moltiplication
//do matrix multiplication
checkERROR( cublasSgemv(net->cublasHandle, CUBLAS_OP_T,
dim_x, dim_y,
&alpha,
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/LSTM.cpp
+10 -6
View File
@@ -86,7 +86,11 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig
// RNN descriptors
checkCUDNN(cudnnCreateRNNDescriptor(&rnnDesc));
checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle,
#if CUDNN_MAJOR > 7
checkCUDNN(cudnnSetRNNDescriptor_v6(net->cudnnHandle,
#else
checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle,
#endif
rnnDesc, stateSize, numLayers, dropoutDesc,
cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT,
//(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL),
@@ -129,7 +133,7 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig
output_dim = input_dim;
output_dim.c = stateSize*(bidirectional ? 2 : 1);
// if retunseq is disabled only the last timestep is returned
// if retunseq is disabled only the last timestamp is returned
if(!returnSeq) {
output_dim.h = 1;
output_dim.w = 1;
@@ -250,7 +254,7 @@ dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) {
rnnDesc,
seqLen, // number of time steps (nT)
x_desc_vec_.data(), // input array of desc (nT*nC_in)
srcF, // input pointer
srcF, // input pointer
hx_desc_, // initial hidden state desc
hx_ptr, // initial hidden state pointer
cx_desc_, // initial cell state desc
@@ -277,7 +281,7 @@ dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) {
rnnDesc,
seqLen, // number of time steps (nT)
x_desc_vec_.data(), // input array of desc (nT*nC_in)
srcB, // input pointer
srcB, // input pointer
hx_desc_, // initial hidden state desc
hx_ptr, // initial hidden state pointer
cx_desc_, // initial cell state desc
@@ -285,7 +289,7 @@ dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) {
w_desc_, // weights desc
wb_ptr, // weights pointer
y_desc_vec_.data(), // output desc (nT*nC_out)
dstB_NR, // output pointer
dstB_NR, // output pointer
hy_desc_, // final hidden state desc
hy_ptr, // final hidden state pointer
cy_desc_, // final cell state desc
@@ -303,7 +307,7 @@ dnnType* LSTM::infer(dataDim_t &dim, dnnType* srcData) {
one_output_dim.c*sizeof(dnnType), cudaMemcpyDeviceToDevice));
}
// if retunseq is disabled only the last timestep is returned
// if retunseq is disabled only the last timestamp is returned
if(returnSeq) {
// forward transpose
matrixTranspose(net->cublasHandle, dstF, dstData,
src/Layer.cpp
+5
View File
@@ -24,6 +24,11 @@ Layer::~Layer() {
checkCUDNN( cudnnDestroyTensorDescriptor(srcTensorDesc) );
checkCUDNN( cudnnDestroyTensorDescriptor(dstTensorDesc) );
if(dstData != nullptr) {
cudaFree(dstData);
dstData = nullptr;
}
}
}}
src/LayerWgs.cpp
+5 -16
View File
@@ -95,7 +95,6 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
cudaMemcpy(power16_h, power16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
//mean array
cudaMemcpy(tmp_d, mean_h, b_size*sizeof(float), cudaMemcpyHostToDevice);
float2half(tmp_d, mean16_d, b_size);
cudaMemcpy(mean16_h, mean16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
@@ -106,27 +105,17 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
float2half(tmp_d, variance16_d, b_size);
cudaMemcpy(variance16_h, variance16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
//conver scales
//convert scales
float2half(scales_d, scales16_d, b_size);
cudaMemcpy(scales16_h, scales16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
cudaFree(tmp_d);
}
}
LayerWgs::~LayerWgs() {
delete [] data_h;
delete [] bias_h;
checkCuda( cudaFree(data_d) );
checkCuda( cudaFree(bias_d) );
if(batchnorm) {
delete [] scales_h;
delete [] mean_h;
delete [] variance_h;
checkCuda( cudaFree(scales_d) );
checkCuda( cudaFree(mean_d) );
checkCuda( cudaFree(variance_d) );
}
releaseHost();
releaseDevice();
}
}}
src/MobilenetDetection.cpp
+19 -12
View File
@@ -126,11 +126,13 @@ float MobilenetDetection::iou(const tk::dnn::box &a, const tk::dnn::box &b){
return iou;
}
bool MobilenetDetection::init(const std::string& tensor_path, const int n_classes){
bool MobilenetDetection::init(const std::string& tensor_path, const int n_classes, const int n_batches, const float conf_thresh){
std::cout<<(tensor_path).c_str()<<"\n";
netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str());
imageSize = netRT->input_dim.h;
classes = n_classes;
nBatches = n_batches;
confThreshold = conf_thresh;
SSDSpec specs[N_SSDSPEC];
@@ -157,9 +159,9 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
generate_ssd_priors(specs, N_SSDSPEC);
#ifndef OPENCV_CUDACONTRIB
checkCuda(cudaMallocHost(&input, sizeof(dnnType) * netRT->input_dim.tot()));
checkCuda(cudaMallocHost(&input, sizeof(dnnType) * netRT->input_dim.tot() * nBatches));
#endif
checkCuda(cudaMalloc(&input_d, sizeof(dnnType) * netRT->input_dim.tot()));
checkCuda(cudaMalloc(&input_d, sizeof(dnnType) * netRT->input_dim.tot() * nBatches));
locations_h = (float *)malloc(N_COORDS * nPriors * sizeof(float));
confidences_h = (float *)malloc(nPriors * classes * sizeof(float));
@@ -208,7 +210,7 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
return 1;
}
void MobilenetDetection::preprocess(cv::Mat &frame){
void MobilenetDetection::preprocess(cv::Mat &frame, const int bi){
#ifdef OPENCV_CUDACONTRIB
//move original image on GPU
cv::cuda::GpuMat orig_img, frame_nomean;
@@ -224,7 +226,7 @@ void MobilenetDetection::preprocess(cv::Mat &frame){
for(int i=0; i < netRT->input_dim.c; i++){
int idx = i * imagePreproc.rows * imagePreproc.cols;
checkCuda( cudaMemcpy((void *)&input_d[idx], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols* sizeof(float), cudaMemcpyDeviceToDevice) );
checkCuda( cudaMemcpy((void *)&input_d[idx + netRT->input_dim.tot()*bi], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols* sizeof(float), cudaMemcpyDeviceToDevice) );
}
#else
//resize image, remove mean, divide by std
@@ -237,17 +239,17 @@ void MobilenetDetection::preprocess(cv::Mat &frame){
cv::split(imagePreproc, bgr);
for (int i = 0; i < netRT->input_dim.c; i++){
int idx = i * imagePreproc.rows * imagePreproc.cols;
memcpy((void *)&input[idx], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols * sizeof(dnnType));
memcpy((void *)&input[idx + netRT->input_dim.tot()*bi], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols * sizeof(dnnType));
}
checkCuda(cudaMemcpyAsync(input_d, input, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
#endif
}
void MobilenetDetection::postprocess(){
void MobilenetDetection::postprocess(const int bi, const bool mAP){
//get confidences and locations_h
dnnType *rt_out[2];
rt_out[0] = (dnnType *)netRT->buffersRT[3];
rt_out[1] = (dnnType *)netRT->buffersRT[4];
rt_out[0] = (dnnType *)netRT->buffersRT[3]+ netRT->buffersDIM[3].tot()*bi;
rt_out[1] = (dnnType *)netRT->buffersRT[4]+ netRT->buffersDIM[4].tot()*bi;
detected.clear();
@@ -255,8 +257,8 @@ void MobilenetDetection::postprocess(){
checkCuda(cudaMemcpy(locations_h, rt_out[1], N_COORDS * nPriors * sizeof(float), cudaMemcpyDeviceToHost));
convert_locatios_to_boxes_and_center();
int width = originalSize.width;
int height = originalSize.height;
int width = originalSize[bi].width;
int height = originalSize[bi].height;
float *conf_per_class;
for (int i = 1; i < classes; i++){
@@ -273,6 +275,10 @@ void MobilenetDetection::postprocess(){
b.w = locations_h[j * N_COORDS + 2];
b.h = locations_h[j * N_COORDS + 3];
if(mAP)
for(int c=1; c<classes; c++)
b.probs.push_back(confidences_h[c * nPriors + j]);
boxes.push_back(b);
}
}
@@ -298,6 +304,7 @@ void MobilenetDetection::postprocess(){
boxes = remaining;
}
}
batchDetected.push_back(detected);
}
src/MulAdd.cpp
+1 -1
View File
@@ -12,7 +12,7 @@ MulAdd::MulAdd(Network *net, dnnType mul, dnnType add) : Layer(net) {
int size = input_dim.tot();
// create a vector with all value setted to add
// create a vector with all value set to add
dnnType *add_vector_h = new dnnType[size];
for(int i=0; i<size; i++)
add_vector_h[i] = add;
src/Network.cpp
+7 -1
View File
@@ -59,11 +59,16 @@ Network::Network(dataDim_t input_dim) {
}
Network::~Network() {
checkCUDNN( cudnnDestroy(cudnnHandle) );
checkERROR( cublasDestroy(cublasHandle) );
}
void Network::releaseLayers() {
for(int i=0; i<num_layers; i++)
delete layers[i];
num_layers = 0;
}
dnnType* Network::infer(dataDim_t &dim, dnnType* data) {
//do infer for every layer
@@ -123,6 +128,7 @@ void Network::print() {
}
printCenteredTitle("", '=', 60);
std::cout<<"\n";
printCudaMemUsage();
}
const char *Network::getNetworkRTName(const char *network_name){
networkName = network_name;
src/NetworkRT.cpp
+23 -15
View File
@@ -122,7 +122,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
input = Ilay->getOutput(0);
input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() );
if(l->getLayerType() == LAYER_YOLO || l->final)
if(l->final)
networkRT->markOutput(*input);
tensors[l] = input;
}
@@ -134,6 +134,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
networkRT->markOutput(*input);
std::cout<<"Selected maxBatchSize: "<<builderRT->getMaxBatchSize()<<"\n";
printCudaMemUsage();
std::cout<<"Building tensorRT cuda engine...\n";
#if NV_TENSORRT_MAJOR >= 6
engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT);
@@ -162,7 +163,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
// note that indices are guaranteed to be less than IEngine::getNbBindings()
buf_input_idx = engineRT->getBindingIndex("data");
buf_output_idx = engineRT->getBindingIndex("out");
std::cout<<"input idex = "<<buf_input_idx<<" -> output index = "<<buf_output_idx<<"\n";
std::cout<<"input index = "<<buf_input_idx<<" -> output index = "<<buf_output_idx<<"\n";
Dims iDim = engineRT->getBindingDimensions(buf_input_idx);
@@ -448,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;
}
@@ -525,7 +529,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
//std::cout<<"convert Yolo\n";
//std::cout<<"New plugin YOLO\n";
IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY);
IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY, l->nms_thresh, l->nsm_kind, l->new_coords);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
@@ -594,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;
@@ -735,12 +739,16 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
if(name.find("Yolo") == 0) {
YoloRT *r = new YoloRT(readBUF<int>(buf), //classes
readBUF<int>(buf), //num
nullptr,
readBUF<int>(buf)); //n_masks
nullptr, //yolo
readBUF<int>(buf), //n_masks
readBUF<float>(buf), //scale_xy
readBUF<float>(buf), //nms_thresh
readBUF<int>(buf), //nms_kind
readBUF<int>(buf) //new_coords
);
r->c = readBUF<int>(buf);
r->h = readBUF<int>(buf);
r->w = readBUF<int>(buf);
r->scaleXY = readBUF<float>(buf);
for(int i=0; i<r->n_masks; i++)
r->mask[i] = readBUF<dnnType>(buf);
for(int i=0; i<r->n_masks*2*r->num; i++)
@@ -765,9 +773,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);
@@ -776,7 +784,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/Region.cpp
+2 -3
View File
@@ -12,8 +12,7 @@
namespace tk { namespace dnn {
Region::Region(Network *net, int classes, int coords, int num) :
Layer(net) {
Layer(net) {
this->classes = classes;
this->coords = coords;
this->num = num;
@@ -64,7 +63,7 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) {
}
/* Intepret class */
/* Interpret class */
RegionInterpret::RegionInterpret(dataDim_t input_dim, dataDim_t output_dim,
int classes, int coords, int num, float thresh, std::string fname_weights) {
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
src/Shortcut.cpp
+1 -1
View File
@@ -13,7 +13,7 @@ Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) {
if( /*backLayer->output_dim.c != input_dim.c ||*/
backLayer->output_dim.w != input_dim.w ||
backLayer->output_dim.h != input_dim.h )
FatalError("Shortcut dim missmatch");
FatalError("Shortcut dim mismatch");
}
Shortcut::~Shortcut() {
src/Yolo.cpp
+56 -15
View File
@@ -11,13 +11,17 @@
namespace tk { namespace dnn {
Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks, float scale_xy) :
Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks, float scale_xy, double nms_thresh, nmsKind_t nsm_kind, int new_coords) :
Layer(net) {
this->final = true;
this->classes = classes;
this->num = num;
this->n_masks = n_masks;
this->scaleXY = scale_xy;
this->nms_thresh = nms_thresh;
this->nsm_kind = nsm_kind;
this->new_coords = new_coords;
// load anchors
if(fname_weights != "") {
@@ -58,12 +62,21 @@ int entry_index(int batch, int location, int entry,
entry*input_dim.w*input_dim.h + loc;
}
Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, int lw, int lh, int w, int h, int stride) {
Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, int lw, int lh, int w, int h, int stride, int new_coords) {
Yolo::box b;
b.x = (i + x[index + 0*stride]) / lw;
b.y = (j + x[index + 1*stride]) / lh;
b.w = exp(x[index + 2*stride]) * biases[2*n] / w;
b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
if(new_coords == 0){
b.x = (i + x[index + 0*stride]) / lw;
b.y = (j + x[index + 1*stride]) / lh;
b.w = exp(x[index + 2*stride]) * biases[2*n] / w;
b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
}
else{
b.x = (i + x[index + 0 * stride] * 2 - 0.5) / lw;
b.y = (j + x[index + 1 * stride] * 2 - 0.5) / lh;
b.w = x[index + 2 * stride] * x[index + 2 * stride] * 4 * biases[2 * n] / w;
b.h = x[index + 3 * stride] * x[index + 3 * stride] * 4 * biases[2 * n + 1] / h;
}
return b;
}
@@ -74,7 +87,10 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
for (int b = 0; b < dim.n; ++b){
for(int n = 0; n < n_masks; ++n){
int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
if (new_coords == 1)
activationLOGISTICForward(srcData + index, dstData + index, 4*dim.w*dim.h);
else
activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
@@ -115,7 +131,7 @@ void correct_yolo_boxes(Yolo::detection *dets, int n, int w, int h, int netw, in
}
}
int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh) {
int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh, int new_coords) {
if(predictions == nullptr)
predictions = new dnnType[output_dim.tot()];
@@ -139,7 +155,7 @@ int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int net
if(objectness <= thresh) continue;
int box_index = entry_index(0, n*lw*lh + i, 0, classes, input_dim, output_dim);
dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh);
dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh, new_coords);
dets[count].objectness = objectness;
dets[count].classes = classes;
for(j = 0; j < classes; ++j){
@@ -192,6 +208,32 @@ float yolo_box_iou(Yolo::box a, Yolo::box b)
return yolo_box_intersection(a, b)/yolo_box_union(a, b);
}
void box_c(const Yolo::box a, const Yolo::box b, float& top, float& bot, float& left, float& right) {
top = std::min(a.y - a.h / 2, b.y - b.h / 2);
bot = std::max(a.y + a.h / 2, b.y + b.h / 2);
left = std::min(a.x - a.w / 2, b.x - b.w / 2);
right = std::max(a.x + a.w / 2, b.x + b.w / 2);
}
// https://github.com/Zzh-tju/DIoU-darknet
// https://arxiv.org/abs/1911.08287
float yolo_box_diou(const Yolo::box a, const Yolo::box b, const float nms_thresh=0.6)
{
float top, bot, left, right;
box_c(a, b, top, bot, left, right);
float w = right - left;
float h = bot - top;
float c = w * w + h * h;
float iou = yolo_box_iou(a, b);
if (c == 0)
return iou;
float d = (a.x - b.x) * (a.x - b.x) + (a.y - b.y) * (a.y - b.y);
float u = pow(d / c, nms_thresh);
float diou_term = u;
return iou - diou_term;
}
int yolo_nms_comparator(const void *pa, const void *pb)
{
Yolo::detection a = *(Yolo::detection *)pa;
@@ -218,8 +260,7 @@ Yolo::detection *Yolo::allocateDetections(int nboxes, int classes) {
return dets;
}
void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes) {
double nms_thresh = 0.45;
void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes, double nms_thresh, nmsKind_t nsm_kind) {
int total = ndets;
int i, j, k;
@@ -245,13 +286,13 @@ void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes) {
box a = dets[i].bbox;
for(j = i+1; j < total; ++j){
box b = dets[j].bbox;
if (yolo_box_iou(a, b) > nms_thresh){
if (nsm_kind == GREEDY_NMS && yolo_box_iou(a, b) > nms_thresh)
dets[j].prob[k] = 0;
else if (nsm_kind == DIOU_NMS && yolo_box_diou(a, b, nms_thresh) > nms_thresh)
dets[j].prob[k] = 0;
}
}
}
}
}
}}
src/Yolo3Detection.cpp
+53 -40
View File
@@ -3,12 +3,17 @@
namespace tk { namespace dnn {
bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes) {
bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, const int n_batches, const float conf_thresh) {
//convert network to tensorRT
std::cout<<(tensor_path).c_str()<<"\n";
netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str() );
nBatches = n_batches;
confThreshold = conf_thresh;
tk::dnn::dataDim_t idim = netRT->input_dim;
idim.n = nBatches;
if(netRT->pluginFactory->n_yolos < 2 ) {
FatalError("this is not yolo3");
}
@@ -19,7 +24,7 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes) {
num = yRT->num;
nMasks = yRT->n_masks;
// make a yolo layer for interpret predictions
// make a yolo layer to interpret predictions
yolo[i] = new tk::dnn::Yolo(nullptr, classes, nMasks, ""); // yolo without input and bias
yolo[i]->mask_h = new dnnType[nMasks];
yolo[i]->bias_h = new dnnType[num*nMasks*2];
@@ -27,13 +32,16 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes) {
memcpy(yolo[i]->bias_h, yRT->bias, sizeof(dnnType)*num*nMasks*2);
yolo[i]->input_dim = yolo[i]->output_dim = tk::dnn::dataDim_t(1, yRT->c, yRT->h, yRT->w);
yolo[i]->classesNames = yRT->classesNames;
yolo[i]->nms_thresh = yRT->nms_thresh;
yolo[i]->nsm_kind = (tk::dnn::Yolo::nmsKind_t) yRT->nms_kind;
yolo[i]->new_coords = yRT->new_coords;
}
dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
#ifndef OPENCV_CUDACONTRIB
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMallocHost(&input, sizeof(dnnType)*idim.tot()));
#endif
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*idim.tot()));
// class colors precompute
for(int c=0; c<classes; c++) {
@@ -48,7 +56,7 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes) {
return true;
}
void Yolo3Detection::preprocess(cv::Mat &frame){
void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
#ifdef OPENCV_CUDACONTRIB
cv::cuda::GpuMat orig_img, img_resized;
orig_img = cv::cuda::GpuMat(frame);
@@ -64,7 +72,7 @@ void Yolo3Detection::preprocess(cv::Mat &frame){
int size = imagePreproc.rows * imagePreproc.cols;
int ch = netRT->input_dim.c-1 -i;
bgr[ch].download(bgr_h); //TODO: don't copy back on CPU
checkCuda( cudaMemcpy(input_d + i*size, (float*)bgr_h.data, size*sizeof(dnnType), cudaMemcpyHostToDevice));
checkCuda( cudaMemcpy(input_d + i*size + netRT->input_dim.tot()*bi, (float*)bgr_h.data, size*sizeof(dnnType), cudaMemcpyHostToDevice));
}
#else
cv::resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
@@ -77,64 +85,69 @@ void Yolo3Detection::preprocess(cv::Mat &frame){
for(int i=0; i<netRT->input_dim.c; i++) {
int idx = i*imagePreproc.rows*imagePreproc.cols;
int ch = netRT->input_dim.c-1 -i;
memcpy((void*)&input[idx], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType));
memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType));
}
checkCuda(cudaMemcpyAsync(input_d, input, netRT->input_dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
checkCuda(cudaMemcpyAsync(input_d + netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
#endif
}
void Yolo3Detection::postprocess(){
void Yolo3Detection::postprocess(const int bi, const bool mAP){
//get yolo outputs
dnnType *rt_out[netRT->pluginFactory->n_yolos];
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
rt_out[i] = (dnnType*)netRT->buffersRT[i+1];
}
for(int i=0; i<netRT->pluginFactory->n_yolos; i++)
rt_out[i] = (dnnType*)netRT->buffersRT[i+1] + netRT->buffersDIM[i+1].tot()*bi;
float x_ratio = float(originalSize.width) / float(netRT->input_dim.w);
float y_ratio = float(originalSize.height) / float(netRT->input_dim.h);
float x_ratio = float(originalSize[bi].width) / float(netRT->input_dim.w);
float y_ratio = float(originalSize[bi].height) / float(netRT->input_dim.h);
// compute dets
nDets = 0;
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
yolo[i]->dstData = rt_out[i];
yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold);
yolo[i]->computeDetections(dets, nDets, netRT->input_dim.w, netRT->input_dim.h, confThreshold, yolo[i]->new_coords);
}
tk::dnn::Yolo::mergeDetections(dets, nDets, classes);
tk::dnn::Yolo::mergeDetections(dets, nDets, classes, yolo[0]->nms_thresh, yolo[0]->nsm_kind);
// fill detected
detected.clear();
for(int j=0; j<nDets; j++) {
tk::dnn::Yolo::box b = dets[j].bbox;
int x0 = (b.x-b.w/2.);
int x1 = (b.x+b.w/2.);
int y0 = (b.y-b.h/2.);
int y1 = (b.y+b.h/2.);
int obj_class = -1;
float prob = 0;
float x0 = (b.x-b.w/2.);
float x1 = (b.x+b.w/2.);
float y0 = (b.y-b.h/2.);
float y1 = (b.y+b.h/2.);
// convert to image coords
x0 = x_ratio*x0;
x1 = x_ratio*x1;
y0 = y_ratio*y0;
y1 = y_ratio*y1;
for(int c=0; c<classes; c++) {
if(dets[j].prob[c] >= confThreshold) {
obj_class = c;
prob = dets[j].prob[c];
int obj_class = c;
float prob = dets[j].prob[c];
tk::dnn::box res;
res.cl = obj_class;
res.prob = prob;
res.x = x0;
res.y = y0;
res.w = x1 - x0;
res.h = y1 - y0;
// FIXME: this shuld be useless
// if(mAP)
// for(int c=0; c<classes; c++)
// res.probs.push_back(dets[j].prob[c]);
detected.push_back(res);
}
}
if(obj_class >= 0) {
// convert to image coords
x0 = x_ratio*x0;
x1 = x_ratio*x1;
y0 = y_ratio*y0;
y1 = y_ratio*y1;
tk::dnn::box res;
res.cl = obj_class;
res.prob = prob;
res.x = x0;
res.y = y0;
res.w = x1 - x0;
res.h = y1 - y0;
detected.push_back(res);
}
}
batchDetected.push_back(detected);
}
src/evaluation.cpp
+45 -4
View File
@@ -63,7 +63,7 @@ double computeMap( std::vector<Frame> &images,const int classes,
int gt_checked = 0;
// for each detection comput IoU with groundtruth and match detetcion and
// for each detection compute IoU with groundtruth and match detetcion and
// groundtruth with IoU greater than IoU_thresh
for(auto &img:images){
for(size_t i=0; i<img.det.size(); i++){
@@ -153,7 +153,7 @@ double computeMap( std::vector<Frame> &images,const int classes,
}
}
//compute average precision for each class. Two methods are avaible,
//compute average precision for each class. Two methods are available,
//based on map_points required
double mean_average_precision = 0;
double last_recall, last_precision, delta_recall;
@@ -287,7 +287,7 @@ void computeTPFPFN( std::vector<Frame> &images,const int classes,
}
}
//count all TP, FP, FN and compute precsion, recall and f1-score
//count all TP, FP, FN and compute precision, recall and f1-score
double avg_precision = 0, avg_recall = 0, f1_score = 0;
int TP = 0, FP = 0, FN = 0;
for(size_t i=0; i<classes; i++){
@@ -314,5 +314,46 @@ void computeTPFPFN( std::vector<Frame> &images,const int classes,
std::cout<<"avg precision: "<<avg_precision<<"\tavg recall: "<<avg_recall<<"\tavg f1 score:"<<f1_score<<std::endl;
}
void printJsonCOCOFormat(std::ofstream *out_file, const std::string image_path, std::vector<tk::dnn::box> bbox, const int classes, const int w, const int h)
{
int coco_ids[] = { 1,2,3,4,5,6,7,8,9,10,11,13,14,15,16,17,18,19,20,21,22,23,24,25,27,28,31,32,33,34,35,36,37,38,39,40,41,42,43,44,46,47,48,49,50,51,52,53,54,55,56,57,58,59,60,61,62,63,64,65,67,70,72,73,74,75,76,77,78,79,80,81,82,84,85,86,87,88,89,90 };
std::string id = image_path.substr(image_path.find("images/")+7, image_path.find(".jpg") - image_path.find("images/") -7);
int image_id = std::stoi(id);
for (int i = 0; i < bbox.size(); ++i) {
float xmin = bbox[i].x ;
float xmax = bbox[i].x + float(bbox[i].w);
float ymin = bbox[i].y;
float ymax = bbox[i].y + float(bbox[i].h);
//limit to image borders
if (xmin < 0) xmin = 0;
if (ymin < 0) ymin = 0;
if (xmax > w) xmax = w;
if (ymax > h) ymax = h;
float bx = xmin;
float by = ymin;
float bw = xmax - xmin;
float bh = ymax - ymin;
if(bbox[i].probs.size() == classes)
for (int j = 0; j < classes; ++j) {
//min threshold confidence is set in DetectionNN.h
if (bbox[i].probs[j] > 0) {
*out_file << "{\"image_id\":" << image_id <<
", \"category_id\":" << coco_ids[j] <<
", \"bbox\":[" << bx << ", " << by << ", " << bw << ", " << bh <<
"], \"score\":" << bbox[i].probs[j] << "},\n";
}
}
else
*out_file << "{\"image_id\":" << image_id <<
", \"category_id\":" << coco_ids[bbox[i].cl] <<
", \"bbox\":[" << bx << ", " << by << ", " << bw << ", " << bh <<
"], \"score\":" << bbox[i].prob << "},\n";
}
}
}}
src/kernels/activation.mish.cu → src/kernels/activation_mish.cu
+22 -3
View File
@@ -3,20 +3,39 @@
#define MISH_THRESHOLD 20
__device__ float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
__device__ float softplus_kernel(float x, float threshold = 20) {
__device__
float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
__device__
float softplus_kernel(float x, float threshold = 20) {
if (x > threshold) return x; // too large
else if (x < -threshold) return expf(x); // too small
return logf(expf(x) + 1);
}
__device__
float mish_yashas(float x) {
float e = __expf(x);
if (x <= -18.0f)
return x * e;
float n = e * e + 2 * e;
if (x <= -5.0f)
return x * __fdividef(n, n + 2);
return x - 2 * __fdividef(x, n + 2);
}
// https://github.com/digantamisra98/Mish
// https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu
__global__
void activation_mish(dnnType *input, dnnType *output, int size) {
int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
if (i < size)
output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));
// output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));
output[i] = mish_yashas(input[i]);
}
/**
src/utils.cpp
+11 -4
View File
@@ -26,10 +26,11 @@ void downloadWeightsifDoNotExist(const std::string& input_bin, const std::string
std::string wget_cmd = "wget " + weights_url + " -O " + test_folder + "/weights.zip";
std::string unzip_cmd = "unzip " + test_folder + "/weights.zip -d" + test_folder;
std::string rm_cmd = "rm " + test_folder + "/weights.zip";
system(mkdir_cmd.c_str());
system(wget_cmd.c_str());
system(unzip_cmd.c_str());
system(rm_cmd.c_str());
int err = 0;
err = system(mkdir_cmd.c_str());
err = system(wget_cmd.c_str());
err = system(unzip_cmd.c_str());
err = system(rm_cmd.c_str());
}
}
@@ -196,6 +197,12 @@ void getMemUsage(double& vm_usage_kb, double& resident_set_kb){
resident_set_kb = rss * page_size_kb;
}
void printCudaMemUsage() {
size_t free, total;
checkCuda( cudaMemGetInfo(&free, &total) );
std::cout<<"GPU free memory: "<<double(free)/1e6<<" mb.\n";
}
void removePathAndExtension(const std::string &full_string, std::string &name){
name = full_string;
std::string tmp_str = full_string;