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

Add Mobilenet2SSDLite test

Browse Source 
The new test works both with TensorRT and cuDNN. Preprocessing and
Postprocessing are missing. Add ClippedReLU (for ReLU6), groups for
Conv2d, additional bias for convolution.

Other minors:
-move the timer in the detector to measure all the
processing time for a given frame (both centernet and yolo);
-add int8 flag.

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
xavier
2020-02-21 10:45:46 +01:00
parent 40a4a55cd0
commit 808a84131c
17 changed files with 649 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/Activation.cpp
+3 -2
View File
@@ -5,10 +5,11 @@
namespace tk { namespace dnn {
Activation::Activation(Network *net, int act_mode) :
Activation::Activation(Network *net, int act_mode, const float ceiling) :
Layer(net) {
this->act_mode = act_mode;
this->ceiling = ceiling;
checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
if(int(act_mode) < 100) {
@@ -31,7 +32,7 @@ Activation::Activation(Network *net, int act_mode) :
checkCUDNN( cudnnSetActivationDescriptor(activDesc,
(cudnnActivationMode_t) act_mode,
CUDNN_PROPAGATE_NAN,
0.0) );
ceiling) );
}
}
src/CenternetDetection.cpp
+1 -2
View File
@@ -294,8 +294,6 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
netRT->infer(dim2, input_d);
TIMER_STOP
dim2.print();
stats.push_back(t_ns);
}
// checkResult(dim2.tot(), input_h, input);
step_t = std::chrono::steady_clock::now();
@@ -456,5 +454,6 @@ void CenternetDetection::update(cv::Mat &imageORIG) {
std::cout<<"TOTAL: \n";
TIMER_STOP
stats.push_back(t_ns);
}
}}
src/Conv2d.cpp
+23 -11
View File
@@ -100,7 +100,7 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) {
&beta, dstTensorDesc, dstData));
}
if(!batchnorm) {
if(!batchnorm && !additional_bias) { //CHECK WITH IF CORRECT
// bias
alpha = dnnType(1);
beta = dnnType(1);
@@ -108,23 +108,34 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) {
&alpha, biasTensorDesc, bias_d,
&beta, dstTensorDesc, dstData) );
} else {
alpha = dnnType(1);
beta = dnnType(0);
checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle,
CUDNN_BATCHNORM_SPATIAL, &alpha, &beta,
dstTensorDesc, dstData, dstTensorDesc,
dstData, biasTensorDesc, //same tensor descriptor as bias
scales_d, bias_d, mean_d, variance_d,
TKDNN_BN_MIN_EPSILON) );
if(additional_bias)
{
alpha = dnnType(1);
beta = dnnType(1);
checkCUDNN( cudnnAddTensor(net->cudnnHandle,
&alpha, biasTensorDesc, bias2_d,
&beta, dstTensorDesc, dstData) );
}
if(batchnorm)
{
alpha = dnnType(1);
beta = dnnType(0);
checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle,
CUDNN_BATCHNORM_SPATIAL, &alpha, &beta,
dstTensorDesc, dstData, dstTensorDesc,
dstData, biasTensorDesc, //same tensor descriptor as bias
scales_d, bias_d, mean_d, variance_d,
TKDNN_BN_MIN_EPSILON) );
}
}
}
Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
int strideH, int strideW, int paddingH, int paddingW,
std::string fname_weights, bool batchnorm, bool deConv, bool final, int groups) :
std::string fname_weights, bool batchnorm, bool deConv, bool final, int groups, bool additional_bias) :
LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1,
fname_weights, batchnorm, false, final, deConv, groups) {
fname_weights, batchnorm, additional_bias, final, deConv, groups) {
this->kernelH = kernelH;
this->kernelW = kernelW;
this->strideH = strideH;
@@ -133,6 +144,7 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
this->paddingW = paddingW;
this->deConv = deConv;
this->groups = groups;
this->additional_bias = additional_bias;
if(!deConv) {
output_dim.n = input_dim.n;
src/LayerWgs.cpp
+1 -5
View File
@@ -9,11 +9,7 @@ namespace tk { namespace dnn {
LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
int kh, int kw, int kl,
std::string fname_weights, bool batchnorm, bool additional_bias, bool final, bool deConv, int groups) : Layer(net, final) {
if(deConv)
inputs = inputs/groups;
else
outputs = outputs/groups;
inputs = inputs/groups;
this->inputs = inputs;
this->outputs = outputs;
src/Network.cpp
+12 -6
View File
@@ -22,19 +22,25 @@ Network::Network(dataDim_t input_dim) {
fp16 = false;
dla = false;
int8 = false;
if(const char* env_p = std::getenv("TKDNN_MODE")) {
if(strcmp(env_p, "FP16") == 0)
fp16 = true;
else if(strcmp(env_p, "DLA") == 0) {
dla = true;
fp16 = true;
}
else if(strcmp(env_p, "DLA") == 0) {
dla = true;
fp16 = true;
}
else if(strcmp(env_p, "INT8") == 0) {
int8 = true;
}
}
if(fp16)
std::cout<<COL_REDB<<"!! FP16 INERENCE ENABLED !!"<<COL_END<<"\n";
std::cout<<COL_REDB<<"!! FP16 INFERENCE ENABLED !!"<<COL_END<<"\n";
if(dla)
std::cout<<COL_GREENB<<"!! DLA INERENCE ENABLED !!"<<COL_END<<"\n";
std::cout<<COL_GREENB<<"!! DLA INFERENCE ENABLED !!"<<COL_END<<"\n";
if(int8)
std::cout<<COL_ORANGEB<<"!! INT8 INFERENCE ENABLED !!"<<COL_END<<"\n";
checkCUDNN( cudnnCreate(&cudnnHandle) );
src/NetworkRT.cpp
+22 -3
View File
@@ -9,6 +9,7 @@
#include "utils.h"
#include "NvInfer.h"
#include "NetworkRT.h"
// #include "calibrator.h"
using namespace nvinfer1;
@@ -58,6 +59,14 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
builderRT->setDefaultDeviceType(DeviceType::kDLA);
builderRT->setDLACore(0);
}
// if(net->int8 && builderRT->platformHasFastInt8())
// {
// dtRT = DataType::kINT8;
// builderRT->setInt8Mode(true);
// Int8EntropyCalibrator calibrator(1, "../demo/images.txt","../demo/yolov3-calibration.table", 416*416*3, 416, 416);
// builderRT->setInt8Calibrator((nvinfer1::IInt8Calibrator * )&calibrator);
// // builderRT->setStrictTypeConstraints(true);
// }
//add input layer
ITensor *input = networkRT->addInput("data", DataType::kFLOAT,
@@ -164,7 +173,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
return convert_layer(input, (Conv2d*) l);
if(type == LAYER_POOLING)
return convert_layer(input, (Pooling*) l);
if(type == LAYER_ACTIVATION)
if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY)
return convert_layer(input, (Activation*) l);
if(type == LAYER_SOFTMAX)
return convert_layer(input, (Softmax*) l);
@@ -337,7 +346,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kSIGMOID);
checkNULL(lRT);
return lRT;
}
else if(l->act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) {
IPlugin *plugin = new ActivationReLUCeiling(l->ceiling);
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
} else {
FatalError("this Activation mode is not yet implemented");
return NULL;
@@ -535,11 +549,16 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
std::string name(layerName);
std::cout<<name<<std::endl;
if(name.find("Activation") == 0) {
if(name.find("ActivationLeaky") == 0) {
ActivationLeakyRT *a = new ActivationLeakyRT();
a->size = readBUF<int>(buf);
return a;
}
if(name.find("ActivationCReLU") == 0) {
ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF<float>(buf));
a->size = readBUF<int>(buf);
return a;
}
if(name.find("Region") == 0) {
RegionRT *r = new RegionRT(readBUF<int>(buf), //classes
src/Yolo3Detection.cpp
+3 -2
View File
@@ -63,6 +63,7 @@ bool Yolo3Detection::init(std::string tensor_path) {
void Yolo3Detection::update(cv::Mat &imageORIG) {
TIMER_START
if(!imageORIG.data) {
std::cout<<"YOLO: NO IMAGE DATA\n";
return;
@@ -100,7 +101,6 @@ void Yolo3Detection::update(cv::Mat &imageORIG) {
stats.push_back(t_ns);
}
TIMER_START
// compute dets
ndets = 0;
for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
@@ -109,7 +109,6 @@ void Yolo3Detection::update(cv::Mat &imageORIG) {
yolo[i]->computeDetections(dets, ndets, netRT->input_dim.w, netRT->input_dim.h, thresh);
}
tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
TIMER_STOP
// fill detected
detected.clear();
@@ -148,6 +147,8 @@ void Yolo3Detection::update(cv::Mat &imageORIG) {
detected.push_back(res);
}
}
TIMER_STOP
stats.push_back(t_ns);
}
src/kernels/activation_relu_ceiling.cu
+33
View File
@@ -0,0 +1,33 @@
#include "kernels.h"
__global__
void activation_relu_ceiling(dnnType *input, dnnType *output, int size, const float ceiling) {
int i = blockDim.x*blockIdx.x + threadIdx.x;
if(i<size) {
if (input[i]>0)
{
if (input[i]>ceiling)
output[i] = ceiling;
else
output[i] = input[i];
}
else
output[i] = 0.0f;
}
}
/**
Relu ceiling activation function
*/
void activationReLUCeilingForward(dnnType* srcData, dnnType* dstData, int size, const float ceiling, cudaStream_t stream)
{
int blocks = (size+255)/256;
int threads = 256;
activation_relu_ceiling<<<blocks, threads, 0, stream>>>(srcData, dstData, size, ceiling);
}