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

added yolo4x and yolo4-csp from the github repo and download file corrections

Browse Source
This commit is contained in:
hchandirasekar
2021-03-26 12:18:46 +05:30
parent f3d1591430
commit f12ec3c935
12 changed files with 1444 additions and 40 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/Layer.h
+6 -1
View File
@@ -19,6 +19,7 @@ enum layerType_t {
LAYER_ACTIVATION_CRELU,
LAYER_ACTIVATION_LEAKY,
LAYER_ACTIVATION_MISH,
LAYER_ACTIVATION_LOGISTIC,
LAYER_FLATTEN,
LAYER_RESHAPE,
LAYER_MULADD,
@@ -68,6 +69,7 @@ public:
case LAYER_ACTIVATION_CRELU: return "ActivationCReLU";
case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky";
case LAYER_ACTIVATION_MISH: return "ActivationMish";
case LAYER_ACTIVATION_LOGISTIC: return "ActivationLogistic";
case LAYER_FLATTEN: return "Flatten";
case LAYER_RESHAPE: return "Reshape";
case LAYER_MULADD: return "MulAdd";
@@ -212,7 +214,8 @@ public:
typedef enum {
ACTIVATION_ELU = 100,
ACTIVATION_LEAKY = 101,
ACTIVATION_MISH = 102
ACTIVATION_MISH = 102,
ACTIVATION_LOGISTIC = 103
} tkdnnActivationMode_t;
/**
@@ -233,6 +236,8 @@ public:
return LAYER_ACTIVATION_LEAKY;
else if (act_mode == ACTIVATION_MISH)
return LAYER_ACTIVATION_MISH;
else if (act_mode == ACTIVATION_LOGISTIC)
return LAYER_ACTIVATION_LOGISTIC;
else
return LAYER_ACTIVATION;
};
include/tkDNN/NetworkRT.h
+1
View File
@@ -27,6 +27,7 @@ using namespace nvinfer1;
#include "pluginsRT/ActivationLeakyRT.h"
#include "pluginsRT/ActivationReLUCeilingRT.h"
#include "pluginsRT/ActivationMishRT.h"
#include "pluginsRT/ActivationLogisticRT.h"
#include "pluginsRT/ReorgRT.h"
#include "pluginsRT/RegionRT.h"
#include "pluginsRT/RouteRT.h"
include/tkDNN/pluginsRT/ActivationLogisticRT.h
+60
View File
@@ -0,0 +1,60 @@
#include<cassert>
#include "../kernels.h"
class ActivationLogisticRT : public IPlugin {
public:
ActivationLogisticRT() {
}
~ActivationLogisticRT(){
}
int getNbOutputs() const override {
return 1;
}
Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
return inputs[0];
}
void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
size = 1;
for(int i=0; i<outputDims[0].nbDims; i++)
size *= outputDims[0].d[i];
}
int initialize() override {
return 0;
}
virtual void terminate() override {
}
virtual size_t getWorkspaceSize(int maxBatchSize) const override {
return 0;
}
virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
activationLOGISTICForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]),
reinterpret_cast<dnnType*>(outputs[0]), batchSize*size, stream);
return 0;
}
virtual size_t getSerializationSize() override {
return 1*sizeof(int);
}
virtual void serialize(void* buffer) override {
char *buf = reinterpret_cast<char*>(buffer);
tk::dnn::writeBUF(buf, size);
}
int size;
};
include/tkDNN/pluginsRT/YoloRT.h
+16 -13
View File
@@ -64,20 +64,23 @@ public:
checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < n_masks; ++n){
int index = entry_index(b, n*w*h, 0);
if (new_coords == 1)
activationLOGISTICForward(srcData + index, dstData + index, 4*w*h, stream); //x,y,w,h
else
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
index = entry_index(b, n*w*h, 4);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
}
}
for (int b = 0; b < batchSize; ++b){
for(int n = 0; n < n_masks; ++n){
int index = entry_index(b, n*w*h, 0);
if (new_coords == 1){
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
}
else{
activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); //x,y
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * w*h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
index = entry_index(b, n*w*h, 4);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
}
}
}
//std::cout<<"YOLO END\n";
return 0;
src/Activation.cpp
+5 -1
View File
@@ -52,7 +52,11 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
else if(act_mode == ACTIVATION_MISH) {
activationMishForward(srcData, dstData, dim.tot());
} else {
}
else if(act_mode == ACTIVATION_LOGISTIC) {
activationLOGISTICForward(srcData, dstData, dim.tot());
}else {
dnnType alpha = dnnType(1);
dnnType beta = dnnType(0);
checkCUDNN( cudnnActivationForward(net->cudnnHandle,
src/DarknetParser.cpp
+1
View File
@@ -187,6 +187,7 @@ namespace tk { namespace dnn {
if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH;
else if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC;
else { FatalError("activation not supported: " + f.activation); }
netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);
};
src/NetworkRT.cpp
+12 -9
View File
@@ -229,7 +229,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
return convert_layer(input, (Conv2d*) l);
if(type == LAYER_POOLING)
return convert_layer(input, (Pooling*) l);
if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH)
if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_LOGISTIC)
return convert_layer(input, (Activation*) l);
if(type == LAYER_SOFTMAX)
return convert_layer(input, (Softmax*) l);
@@ -424,6 +424,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
checkNULL(lRT);
return lRT;
}
else if(l->act_mode == ACTIVATION_LOGISTIC) {
IPlugin *plugin = new ActivationLogisticRT();
IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
checkNULL(lRT);
return lRT;
}
else {
FatalError("this Activation mode is not yet implemented");
return NULL;
@@ -660,6 +666,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
assert(buf == bufCheck + serialLength);
return a;
}
if(name.find("ActivationLogistic") == 0) {
ActivationLogisticRT *a = new ActivationLogisticRT();
a->size = readBUF<int>(buf);
return a;
}
if(name.find("ActivationCReLU") == 0) {
float activationReluTemp = readBUF<float>(buf);
//ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF<float>(buf));
@@ -784,17 +795,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
float nms_thresh_temp = readBUF<float>(buf);
int nms_kind_temp = readBUF<int>(buf);
int new_coords_temp = readBUF<int>(buf);
std::cout << classes_temp << ":" << num_temp << ":" << ":" << n_masks_temp << ":" << nms_thresh_temp << ":" << nms_kind_temp << ":" << new_coords_temp << std::endl;
YoloRT *r = new YoloRT(classes_temp,num_temp,nullptr,n_masks_temp,scale_xy_temp,nms_thresh_temp,nms_kind_temp,new_coords_temp);
/* std::cout << "classes : " << r->classes;
std::cout << "num : " << r->num;
std::cout << "n_masks : " << r->n_masks;
std::cout << "scalexy : " << r->scaleXY;
std::cout << "nms_thresh : " << r->nms_thresh;
std::cout << "nms_kind : " << r->nms_kind;
std::cout << "new_coords : " << r->new_coords;*/
r->c = readBUF<int>(buf);
src/Yolo.cpp
+12 -9
View File
@@ -73,8 +73,8 @@ Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j,
b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
}
else{
b.x = (i + x[index + 0 * stride] * 2 - 0.5) / lw;
b.y = (j + x[index + 1 * stride] * 2 - 0.5) / lh;
b.x = (i + x[index + 0 * stride] ) / lw;
b.y = (j + x[index + 1 * stride] ) / lh;
b.w = x[index + 2 * stride] * x[index + 2 * stride] * 4 * biases[2 * n] / w;
b.h = x[index + 3 * stride] * x[index + 3 * stride] * 4 * biases[2 * n + 1] / h;
}
@@ -88,15 +88,18 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
for (int b = 0; b < dim.n; ++b){
for(int n = 0; n < n_masks; ++n){
int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim);
if (new_coords == 1)
activationLOGISTICForward(srcData + index, dstData + index, 4*dim.w*dim.h);
else
std::cout<<"new_coords"<<new_coords<<std::endl;
if (new_coords == 1){
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
}
else{
activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
}
}
}
tests/darknet/cfg/yolo4-csp.cfg
+1279
View File
File diff suppressed because it is too large Load Diff
tests/darknet/cfg/yolo4x.cfg
+15 -6
View File
@@ -5,8 +5,8 @@
# Training
batch=64
subdivisions=8
width=672
height=672
width=640
height=640
channels=3
momentum=0.949
decay=0.0005
@@ -15,7 +15,7 @@ saturation = 1.5
exposure = 1.5
hue=.1
learning_rate=0.00261
learning_rate=0.001
burn_in=1000
max_batches = 500500
policy=steps
@@ -26,6 +26,8 @@ mosaic=1
letter_box=1
#optimized_memory=1
[convolutional]
batch_normalize=1
filters=32
@@ -1131,6 +1133,7 @@ size=1
stride=1
pad=1
activation=mish
stopbackward=800
##########################
@@ -1147,7 +1150,7 @@ size=1
stride=1
pad=1
filters=255
activation=linear
activation=logistic
[yolo]
@@ -1156,6 +1159,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4
classes=80
num=9
jitter=.1
scale_x_y = 2.0
objectness_smooth=0
ignore_thresh = .7
truth_thresh = 1
@@ -1169,6 +1173,7 @@ iou_loss=ciou
nms_kind=diounms
beta_nms=0.6
new_coords=1
max_delta=5
[route]
layers = -4
@@ -1275,7 +1280,7 @@ size=1
stride=1
pad=1
filters=255
activation=linear
activation=logistic
[yolo]
@@ -1284,6 +1289,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4
classes=80
num=9
jitter=.1
scale_x_y = 2.0
objectness_smooth=1
ignore_thresh = .7
truth_thresh = 1
@@ -1297,6 +1303,7 @@ iou_loss=ciou
nms_kind=diounms
beta_nms=0.6
new_coords=1
max_delta=5
[route]
layers = -4
@@ -1403,7 +1410,7 @@ size=1
stride=1
pad=1
filters=255
activation=linear
activation=logistic
[yolo]
@@ -1412,6 +1419,7 @@ anchors = 12, 16, 19, 36, 40, 28, 36, 75, 76, 55, 72, 146, 142, 110, 192, 243, 4
classes=80
num=9
jitter=.1
scale_x_y = 2.0
objectness_smooth=1
ignore_thresh = .7
truth_thresh = 1
@@ -1425,3 +1433,4 @@ iou_loss=ciou
nms_kind=diounms
beta_nms=0.6
new_coords=1
max_delta=2
tests/darknet/yolo4-csp.cpp
+36
View File
@@ -0,0 +1,36 @@
#include<iostream>
#include<vector>
#include "tkdnn.h"
#include "test.h"
#include "DarknetParser.h"
int main() {
std::string bin_path = "yolo4-csp";
std::vector<std::string> input_bins = {
bin_path + "/layers/input.bin"
};
std::vector<std::string> output_bins = {
bin_path + "/debug/layer144_out.bin",
bin_path + "/debug/layer159_out.bin",
bin_path + "/debug/layer174_out.bin"
};
std::string wgs_path = bin_path + "/layers";
std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4-csp.cfg";
std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/AfzHE4BfTeEm2gH/download");
// parse darknet network
tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
net->print();
//convert network to tensorRT
tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str()));
int ret = testInference(input_bins, output_bins, net, netRT);
net->releaseLayers();
delete net;
delete netRT;
return ret;
}
tests/darknet/yolo4x.cpp
+1 -1
View File
@@ -17,7 +17,7 @@ int main() {
std::string wgs_path = bin_path + "/layers";
std::string cfg_path = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4x.cfg";
std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/BLPpiAigZJLorQD/download");
downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download");